diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h index f80745038a..b203d942a8 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_vector_types.h @@ -44,18 +44,35 @@ THE SOFTWARE. __attribute__((vector_size(__ROUND_UP_TO_NEXT_POT__(n) * sizeof(T)))) #endif -#if defined(__cplusplus) +#if defined(__cplusplus) && defined(__clang__) #include namespace hip_impl { template struct Scalar_accessor { + struct Address { + const Scalar_accessor* p; + + __host__ __device__ + operator const T*() const noexcept { + return &reinterpret_cast(p)[idx]; + } + __host__ __device__ + operator T*() noexcept { + return &reinterpret_cast( + const_cast(p))[idx]; + } + }; + // Idea from https://t0rakka.silvrback.com/simd-scalar-accessor Vector data; __host__ __device__ operator T() const noexcept { return data[idx]; } + __host__ __device__ + Address operator&() const noexcept { return Address{this}; } + __host__ __device__ Scalar_accessor& operator=(T x) noexcept { data[idx] = x; @@ -63,6 +80,29 @@ THE SOFTWARE. return *this; } + __host__ __device__ + Scalar_accessor& operator++() noexcept { + ++data[idx]; + return *this; + } + __host__ __device__ + T operator++(int) noexcept { + auto r{data[idx]}; + ++data[idx]; + return *this; + } + __host__ __device__ + Scalar_accessor& operator--() noexcept { + --data[idx]; + return *this; + } + __host__ __device__ + T operator--(int) noexcept { + auto r{data[idx]}; + --data[idx]; + return *this; + } + __host__ __device__ Scalar_accessor& operator+=(T x) noexcept { data[idx] += x;