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 a03a46b8cf..f80745038a 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 @@ -47,6 +47,95 @@ THE SOFTWARE. #if defined(__cplusplus) #include + namespace hip_impl { + template + struct Scalar_accessor { + // Idea from https://t0rakka.silvrback.com/simd-scalar-accessor + Vector data; + + __host__ __device__ + operator T() const noexcept { return data[idx]; } + + __host__ __device__ + Scalar_accessor& operator=(T x) noexcept { + data[idx] = x; + + return *this; + } + + __host__ __device__ + Scalar_accessor& operator+=(T x) noexcept { + data[idx] += x; + return *this; + } + __host__ __device__ + Scalar_accessor& operator-=(T x) noexcept { + data[idx] -= x; + return *this; + } + + __host__ __device__ + Scalar_accessor& operator*=(T x) noexcept { + data[idx] *= x; + return *this; + } + __host__ __device__ + Scalar_accessor& operator/=(T x) noexcept { + data[idx] /= x; + return *this; + } + template< + typename U = T, + typename std::enable_if{}>::type* = nullptr> + __host__ __device__ + Scalar_accessor& operator%=(T x) noexcept { + data[idx] %= x; + return *this; + } + + template< + typename U = T, + typename std::enable_if{}>::type* = nullptr> + __host__ __device__ + Scalar_accessor& operator>>=(T x) noexcept { + data[idx] >>= x; + return *this; + } + template< + typename U = T, + typename std::enable_if{}>::type* = nullptr> + __host__ __device__ + Scalar_accessor& operator<<=(T x) noexcept { + data[idx] <<= x; + return *this; + } + template< + typename U = T, + typename std::enable_if{}>::type* = nullptr> + __host__ __device__ + Scalar_accessor& operator&=(T x) noexcept { + data[idx] &= x; + return *this; + } + template< + typename U = T, + typename std::enable_if{}>::type* = nullptr> + __host__ __device__ + Scalar_accessor& operator|=(T x) noexcept { + data[idx] |= x; + return *this; + } + template< + typename U = T, + typename std::enable_if{}>::type* = nullptr> + __host__ __device__ + Scalar_accessor& operator^=(T x) noexcept { + data[idx] ^= x; + return *this; + } + }; + } // Namespace hip_impl. + template struct HIP_vector_base; template @@ -55,9 +144,7 @@ THE SOFTWARE. union { Native_vec_ data; - struct { - T x; - }; + hip_impl::Scalar_accessor x; }; }; @@ -67,10 +154,8 @@ THE SOFTWARE. union { Native_vec_ data; - struct { - T x; - T y; - }; + hip_impl::Scalar_accessor x; + hip_impl::Scalar_accessor y; }; }; @@ -238,12 +323,10 @@ THE SOFTWARE. union { Native_vec_ data; - struct { - T x; - T y; - T z; - T w; - }; + hip_impl::Scalar_accessor x; + hip_impl::Scalar_accessor y; + hip_impl::Scalar_accessor z; + hip_impl::Scalar_accessor w; }; };