From e4744692134c55231f32eb4b82c2bce01ce56708 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Mon, 24 Jun 2019 20:02:09 -0500 Subject: [PATCH] Put 3-wide vector types on a ketogenic diet. (#1180) * Put 3-wide vector types on a ketogenic diet. * Remove needless include. * Do not be narrow-minded. * Do not be narrow-minded. * Put the C people on a diet too. [ROCm/hip commit: 67abac1365c26fdf499e288ac7f3eec88b52c816] --- .../include/hip/hcc_detail/hip_vector_types.h | 148 +++++++++++++++++- .../tests/src/deviceLib/hipVectorTypes.cpp | 2 +- 2 files changed, 147 insertions(+), 3 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_vector_types.h b/projects/hip/include/hip/hcc_detail/hip_vector_types.h index 49df41f5a7..0d424965ca 100644 --- a/projects/hip/include/hip/hcc_detail/hip_vector_types.h +++ b/projects/hip/include/hip/hcc_detail/hip_vector_types.h @@ -76,7 +76,151 @@ THE SOFTWARE. template struct HIP_vector_base { - typedef T Native_vec_ __NATIVE_VECTOR__(3, T); + struct Native_vec_ { + T d[3]; + + __host__ __device__ + constexpr + Native_vec_() = default; + __host__ __device__ + explicit + constexpr + Native_vec_(T x) noexcept : d{x, x, x} {} + __host__ __device__ + constexpr + Native_vec_(T x, T y, T z) noexcept : d{x, y, z} {} + __host__ __device__ + constexpr + Native_vec_(const Native_vec_&) = default; + __host__ __device__ + constexpr + Native_vec_(Native_vec_&&) = default; + __host__ __device__ + ~Native_vec_() = default; + + __host__ __device__ + Native_vec_& operator=(const Native_vec_&) = default; + __host__ __device__ + Native_vec_& operator=(Native_vec_&&) = default; + + __host__ __device__ + T& operator[](unsigned int idx) noexcept { return d[idx]; } + __host__ __device__ + T operator[](unsigned int idx) const noexcept { return d[idx]; } + + __host__ __device__ + Native_vec_& operator+=(const Native_vec_& x) noexcept + { + for (auto i = 0u; i != 3u; ++i) d[i] += x.d[i]; + return *this; + } + __host__ __device__ + Native_vec_& operator-=(const Native_vec_& x) noexcept + { + for (auto i = 0u; i != 3u; ++i) d[i] -= x.d[i]; + return *this; + } + + __host__ __device__ + Native_vec_& operator*=(const Native_vec_& x) noexcept + { + for (auto i = 0u; i != 3u; ++i) d[i] *= x.d[i]; + return *this; + } + __host__ __device__ + Native_vec_& operator/=(const Native_vec_& x) noexcept + { + for (auto i = 0u; i != 3u; ++i) d[i] /= x.d[i]; + return *this; + } + + template< + typename U = T, + typename std::enable_if{}>::type* = nullptr> + __host__ __device__ + Native_vec_ operator-() const noexcept + { + auto r{*this}; + for (auto&& x : r.d) x = -x; + return r; + } + + template< + typename U = T, + typename std::enable_if{}>::type* = nullptr> + __host__ __device__ + Native_vec_ operator~() const noexcept + { + auto r{*this}; + for (auto&& x : r.d) x = ~x; + return r; + } + template< + typename U = T, + typename std::enable_if{}>::type* = nullptr> + __host__ __device__ + Native_vec_& operator%=(const Native_vec_& x) noexcept + { + for (auto i = 0u; i != 3u; ++i) d[i] %= x.d[i]; + return *this; + } + template< + typename U = T, + typename std::enable_if{}>::type* = nullptr> + __host__ __device__ + Native_vec_& operator^=(const Native_vec_& x) noexcept + { + for (auto i = 0u; i != 3u; ++i) d[i] ^= x.d[i]; + return *this; + } + template< + typename U = T, + typename std::enable_if{}>::type* = nullptr> + __host__ __device__ + Native_vec_& operator|=(const Native_vec_& x) noexcept + { + for (auto i = 0u; i != 3u; ++i) d[i] |= x.d[i]; + return *this; + } + template< + typename U = T, + typename std::enable_if{}>::type* = nullptr> + __host__ __device__ + Native_vec_& operator&=(const Native_vec_& x) noexcept + { + for (auto i = 0u; i != 3u; ++i) d[i] &= x.d[i]; + return *this; + } + template< + typename U = T, + typename std::enable_if{}>::type* = nullptr> + __host__ __device__ + Native_vec_& operator>>=(const Native_vec_& x) noexcept + { + for (auto i = 0u; i != 3u; ++i) d[i] >>= x.d[i]; + return *this; + } + template< + typename U = T, + typename std::enable_if{}>::type* = nullptr> + __host__ __device__ + Native_vec_& operator<<=(const Native_vec_& x) noexcept + { + for (auto i = 0u; i != 3u; ++i) d[i] <<= x.d[i]; + return *this; + } + + using Vec3_cmp = T __NATIVE_VECTOR__(3, int); + __host__ __device__ + Vec3_cmp operator==(const Native_vec_& x) const noexcept + { + Vec3_cmp r; + r[0] = d[0] == x.d[0]; + r[1] = d[1] == x.d[1]; + r[2] = d[2] == x.d[2]; + return -r; + } + }; union { Native_vec_ data; @@ -632,7 +776,7 @@ THE SOFTWARE. } CUDA_name##2;\ typedef struct {\ union {\ - CUDA_name##_impl3 data;\ + T data[3];\ struct {\ T x;\ T y;\ diff --git a/projects/hip/tests/src/deviceLib/hipVectorTypes.cpp b/projects/hip/tests/src/deviceLib/hipVectorTypes.cpp index f05e6372d8..14479881ff 100644 --- a/projects/hip/tests/src/deviceLib/hipVectorTypes.cpp +++ b/projects/hip/tests/src/deviceLib/hipVectorTypes.cpp @@ -190,7 +190,7 @@ bool CheckVectorTypes() { int main() { static_assert(sizeof(float1) == 4, ""); static_assert(sizeof(float2) >= 8, ""); - static_assert(sizeof(float3) >= 12, ""); + static_assert(sizeof(float3) == 12, ""); static_assert(sizeof(float4) >= 16, ""); if (CheckVectorTypes()) {