From 00d735cdc95376151fcc65ede8bef9183fd0721a Mon Sep 17 00:00:00 2001 From: jglaser Date: Mon, 9 Dec 2019 23:10:15 -0500 Subject: [PATCH] fix linking of vector types with gcc (#1690) * fix linking of vector types when linking hipcc objects with gcc * use __atribute__((vector_size)) with both clang and gcc and reinstate nonaligned n=3 vector type * use implicit conversion to value and ext_vector_type when available * Alternate formulation for GCC compatibility * Built-in arrays don't mix well with placement new * Fix typo * Add conversions to enum * Fix Scalar_accessor assignment. * Update hip_vector_types.h * stir up the underlying_type hideous mess This fixes the HIP build issue "error: only enumeration types have underlying types". --- include/hip/hcc_detail/hip_vector_types.h | 114 +++++++++++++++++++--- 1 file changed, 98 insertions(+), 16 deletions(-) diff --git a/include/hip/hcc_detail/hip_vector_types.h b/include/hip/hcc_detail/hip_vector_types.h index 582a359fbd..0e8dad595a 100644 --- a/include/hip/hcc_detail/hip_vector_types.h +++ b/include/hip/hcc_detail/hip_vector_types.h @@ -34,15 +34,31 @@ THE SOFTWARE. #include "hip/hcc_detail/host_defines.h" -#if !defined(_MSC_VER) || __clang__ -#if defined(__clang__) - #define __NATIVE_VECTOR__(n, ...) __attribute__((ext_vector_type(n))) +#if !defined(_MSC_VER) + #if __has_attribute(ext_vector_type) + #define __NATIVE_VECTOR__(n, ...) __attribute__((ext_vector_type(n))) + #else + #define __NATIVE_VECTOR__(n, ...) [n] #endif -#if defined(__cplusplus) && defined(__clang__) +#if defined(__cplusplus) + #include #include #include + namespace hip_impl { + template struct Scalar_accessor; + } // Namespace hip_impl. + + namespace std { + template + struct is_integral> + : is_integral {}; + template + struct is_floating_point> + : is_floating_point {}; + } // Namespace std. + namespace hip_impl { template struct Scalar_accessor { @@ -93,6 +109,27 @@ THE SOFTWARE. operator T() const noexcept { return data[idx]; } __host__ __device__ operator T() const volatile noexcept { return data[idx]; } + + // The conversions to enum are fairly ghastly, but unfortunately used in + // some pre-existing, difficult to modify, code. + template< + typename U, + typename std::enable_if< + !std::is_same{} && + std::is_enum{} && + std::is_convertible< + T, typename std::enable_if::value, std::underlying_type>::type::type>{}>::type* = nullptr> + __host__ __device__ + operator U() const noexcept { return static_cast(data[idx]); } + template< + typename U, + typename std::enable_if< + !std::is_same{} && + std::is_enum{} && + std::is_convertible< + T, typename std::enable_if::value, std::underlying_type>::type::type>{}>::type* = nullptr> + __host__ __device__ + operator U() const volatile noexcept { return static_cast(data[idx]); } __host__ __device__ operator T&() noexcept { @@ -107,7 +144,13 @@ THE SOFTWARE. __host__ __device__ Address operator&() const noexcept { return Address{this}; } + + __host__ __device__ + Scalar_accessor& operator=(const Scalar_accessor& x) noexcept { + data[idx] = x.data[idx]; + return *this; + } __host__ __device__ Scalar_accessor& operator=(T x) noexcept { data[idx] = x; @@ -179,7 +222,7 @@ THE SOFTWARE. typename std::enable_if< std::is_convertible{}>::type* = nullptr> __host__ __device__ - Scalar_accessor& operator/=(T x) noexcept { + Scalar_accessor& operator/=(U x) noexcept { data[idx] /= x; return *this; } @@ -245,7 +288,7 @@ THE SOFTWARE. template struct HIP_vector_base { - typedef T Native_vec_ __NATIVE_VECTOR__(1, T); + using Native_vec_ = T __NATIVE_VECTOR__(1, T); union { Native_vec_ data; @@ -253,11 +296,22 @@ THE SOFTWARE. }; using value_type = T; + + __host__ __device__ + HIP_vector_base& operator=(const HIP_vector_base& x) noexcept { + #if __has_attribute(ext_vector_type) + data = x.data; + #else + data[0] = x.data[0]; + #endif + + return *this; + } }; template struct HIP_vector_base { - typedef T Native_vec_ __NATIVE_VECTOR__(2, T); + using Native_vec_ = T __NATIVE_VECTOR__(2, T); union { Native_vec_ data; @@ -266,6 +320,18 @@ THE SOFTWARE. }; using value_type = T; + + __host__ __device__ + HIP_vector_base& operator=(const HIP_vector_base& x) noexcept { + #if __has_attribute(ext_vector_type) + data = x.data; + #else + data[0] = x.data[0]; + data[1] = x.data[1]; + #endif + + return *this; + } }; template @@ -404,15 +470,11 @@ THE SOFTWARE. return *this; } - using Vec3_cmp = int __NATIVE_VECTOR__(3, int); + using Vec3_cmp = int __attribute__((vector_size(4 * sizeof(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; + return Vec3_cmp{d[0] == x.d[0], d[1] == x.d[1], d[2] == x.d[2]}; } }; @@ -430,7 +492,7 @@ THE SOFTWARE. template struct HIP_vector_base { - typedef T Native_vec_ __NATIVE_VECTOR__(4, T); + using Native_vec_ = T __NATIVE_VECTOR__(4, T); union { Native_vec_ data; @@ -441,6 +503,20 @@ THE SOFTWARE. }; using value_type = T; + + __host__ __device__ + HIP_vector_base& operator=(const HIP_vector_base& x) noexcept { + #if __has_attribute(ext_vector_type) + data = x.data; + #else + data[0] = x.data[0]; + data[1] = x.data[1]; + data[2] = x.data[2]; + data[3] = x.data[3]; + #endif + + return *this; + } }; template @@ -464,7 +540,14 @@ THE SOFTWARE. typename std::enable_if< (rank > 1) && sizeof...(Us) == rank>::type* = nullptr> inline __host__ __device__ - HIP_vector_type(Us... xs) noexcept { data = Native_vec_{static_cast(xs)...}; } + HIP_vector_type(Us... xs) noexcept + { + #if __has_attribute(ext_vector_type) + new (&data) Native_vec_{static_cast(xs)...}; + #else + new (&data) std::array{static_cast(xs)...}; + #endif + } inline __host__ __device__ HIP_vector_type(const HIP_vector_type&) = default; inline __host__ __device__ @@ -651,7 +734,6 @@ THE SOFTWARE. } }; - template inline __host__ __device__ HIP_vector_type operator+(