From 9d91b802a5880eddba2273c07413d64b66b448b2 Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Mon, 25 Jun 2018 17:49:50 +0100 Subject: [PATCH] Let's try this again... --- include/hip/hcc_detail/hip_vector_types.h | 571 +++++++++++++++++-- src/device_functions.cpp | 36 +- tests/src/deviceLib/hipVectorTypes.cpp | 104 ++-- tests/src/deviceLib/hipVectorTypesDevice.cpp | 107 ++-- tests/src/deviceLib/vector_test_common.h | 36 -- 5 files changed, 646 insertions(+), 208 deletions(-) diff --git a/include/hip/hcc_detail/hip_vector_types.h b/include/hip/hcc_detail/hip_vector_types.h index 59b9c247e3..704ff7e550 100644 --- a/include/hip/hcc_detail/hip_vector_types.h +++ b/include/hip/hcc_detail/hip_vector_types.h @@ -34,6 +34,8 @@ THE SOFTWARE. #include "hip/hcc_detail/host_defines.h" +#include + #if defined(__clang__) #define __NATIVE_VECTOR__(n, ...) __attribute__((ext_vector_type(n))) #elif defined(__GNUC__) // N.B.: GCC does not support .xyzw syntax. @@ -43,65 +45,538 @@ THE SOFTWARE. __attribute__((vector_size(__ROUND_UP_TO_NEXT_POT__(n) * sizeof(T)))) #endif -typedef unsigned char uchar1 __NATIVE_VECTOR__(1, unsigned char); -typedef unsigned char uchar2 __NATIVE_VECTOR__(2, unsigned char); -typedef unsigned char uchar3 __NATIVE_VECTOR__(3, unsigned char); -typedef unsigned char uchar4 __NATIVE_VECTOR__(4, unsigned char); +#if defined(__cplusplus) + template struct HIP_vector_base; -typedef char char1 __NATIVE_VECTOR__(1, char); -typedef char char2 __NATIVE_VECTOR__(2, char); -typedef char char3 __NATIVE_VECTOR__(3, char); -typedef char char4 __NATIVE_VECTOR__(4, char); + template + struct HIP_vector_base { + typedef T Native_vec_ __NATIVE_VECTOR__(1, T); -typedef unsigned short ushort1 __NATIVE_VECTOR__(1, unsigned short); -typedef unsigned short ushort2 __NATIVE_VECTOR__(2, unsigned short); -typedef unsigned short ushort3 __NATIVE_VECTOR__(3, unsigned short); -typedef unsigned short ushort4 __NATIVE_VECTOR__(4, unsigned short); + union { + Native_vec_ data; + struct { + typename std::decay< + decltype(std::declval().x)>::type x; + }; + }; + }; -typedef short short1 __NATIVE_VECTOR__(1, short); -typedef short short2 __NATIVE_VECTOR__(2, short); -typedef short short3 __NATIVE_VECTOR__(3, short); -typedef short short4 __NATIVE_VECTOR__(4, short); + template + struct HIP_vector_base { + typedef T Native_vec_ __NATIVE_VECTOR__(2, T); -typedef unsigned int uint1 __NATIVE_VECTOR__(1, unsigned int); -typedef unsigned int uint2 __NATIVE_VECTOR__(2, unsigned int); -typedef unsigned int uint3 __NATIVE_VECTOR__(3, unsigned int); -typedef unsigned int uint4 __NATIVE_VECTOR__(4, unsigned int); + union { + Native_vec_ data; + struct { + typename std::decay< + decltype(std::declval().x)>::type x; + typename std::decay< + decltype(std::declval().y)>::type y; + }; + }; + }; -typedef int int1 __NATIVE_VECTOR__(1, int); -typedef int int2 __NATIVE_VECTOR__(2, int); -typedef int int3 __NATIVE_VECTOR__(3, int); -typedef int int4 __NATIVE_VECTOR__(4, int); + template + struct HIP_vector_base { + typedef T Native_vec_ __NATIVE_VECTOR__(3, T); -typedef unsigned long ulong1 __NATIVE_VECTOR__(1, unsigned long); -typedef unsigned long ulong2 __NATIVE_VECTOR__(2, unsigned long); -typedef unsigned long ulong3 __NATIVE_VECTOR__(3, unsigned long); -typedef unsigned long ulong4 __NATIVE_VECTOR__(4, unsigned long); + union { + Native_vec_ data; + struct { + typename std::decay< + decltype(std::declval().x)>::type x; + typename std::decay< + decltype(std::declval().y)>::type y; + typename std::decay< + decltype(std::declval().z)>::type z; + }; + }; + }; -typedef long long1 __NATIVE_VECTOR__(1, long); -typedef long long2 __NATIVE_VECTOR__(2, long); -typedef long long3 __NATIVE_VECTOR__(3, long); -typedef long long4 __NATIVE_VECTOR__(4, long); + template + struct HIP_vector_base { + typedef T Native_vec_ __NATIVE_VECTOR__(4, T); -typedef unsigned long long ulonglong1 __NATIVE_VECTOR__(1, unsigned long long); -typedef unsigned long long ulonglong2 __NATIVE_VECTOR__(2, unsigned long long); -typedef unsigned long long ulonglong3 __NATIVE_VECTOR__(3, unsigned long long); -typedef unsigned long long ulonglong4 __NATIVE_VECTOR__(4, unsigned long long); + union { + Native_vec_ data; + struct { + typename std::decay< + decltype(std::declval().x)>::type x; + typename std::decay< + decltype(std::declval().y)>::type y; + typename std::decay< + decltype(std::declval().z)>::type z; + typename std::decay< + decltype(std::declval().w)>::type w; + }; + }; + }; -typedef long long longlong1 __NATIVE_VECTOR__(1, long long); -typedef long long longlong2 __NATIVE_VECTOR__(2, long long); -typedef long long longlong3 __NATIVE_VECTOR__(3, long long); -typedef long long longlong4 __NATIVE_VECTOR__(4, long long); + template + struct HIP_vector_type : public HIP_vector_base { + using HIP_vector_base::data; + using typename HIP_vector_base::Native_vec_; -typedef float float1 __NATIVE_VECTOR__(1, float); -typedef float float2 __NATIVE_VECTOR__(2, float); -typedef float float3 __NATIVE_VECTOR__(3, float); -typedef float float4 __NATIVE_VECTOR__(4, float); + __host__ __device__ + HIP_vector_type() = default; + template< + typename U, + typename std::enable_if< + std::is_convertible{}>::type* = nullptr> + __host__ __device__ + explicit + HIP_vector_type(U x) noexcept { data = Native_vec_(x); } + template< // TODO: constrain based on type as well. + typename... Us, + typename std::enable_if::type* = nullptr> + __host__ __device__ + HIP_vector_type(Us... xs) noexcept { data = Native_vec_{xs...}; } + __host__ __device__ + HIP_vector_type(const HIP_vector_type&) = default; + __host__ __device__ + HIP_vector_type(HIP_vector_type&&) = default; + __host__ __device__ + ~HIP_vector_type() = default; -typedef double double1 __NATIVE_VECTOR__(1, double); -typedef double double2 __NATIVE_VECTOR__(2, double); -typedef double double3 __NATIVE_VECTOR__(3, double); -typedef double double4 __NATIVE_VECTOR__(4, double); + __host__ __device__ + HIP_vector_type& operator=(const HIP_vector_type&) = default; + __host__ __device__ + HIP_vector_type& operator=(HIP_vector_type&&) = default; + + // Operators + __host__ __device__ + HIP_vector_type& operator++() noexcept + { + data += Native_vec_(1); + return *this; + } + __host__ __device__ + HIP_vector_type operator++(int) noexcept + { + auto tmp(*this); + ++*this; + return tmp; + } + __host__ __device__ + HIP_vector_type& operator--() noexcept + { + data -= Native_vec_(1); + return *this; + } + __host__ __device__ + HIP_vector_type operator--(int) noexcept + { + auto tmp(*this); + --*this; + return tmp; + } + __host__ __device__ + HIP_vector_type& operator+=(const HIP_vector_type& x) noexcept + { + data += x.data; + return *this; + } + __host__ __device__ + HIP_vector_type& operator-=(const HIP_vector_type& x) noexcept + { + data -= x.data; + return *this; + } + __host__ __device__ + HIP_vector_type& operator*=(const HIP_vector_type& x) noexcept + { + data *= x.data; + return *this; + } + __host__ __device__ + HIP_vector_type& operator/=(const HIP_vector_type& x) noexcept + { + data /= x.data; + return *this; + } + + template< + typename U = T, + typename std::enable_if{}>::type* = nullptr> + __host__ __device__ + HIP_vector_type operator-() noexcept + { + auto tmp(*this); + tmp.data = -tmp.data; + return tmp; + } + + template< + typename U = T, + typename std::enable_if{}>::type* = nullptr> + __host__ __device__ + HIP_vector_type operator~() noexcept + { + HIP_vector_type r{*this}; + r.data = ~r.data; + return r; + } + template< + typename U = T, + typename std::enable_if{}>::type* = nullptr> + __host__ __device__ + HIP_vector_type& operator%=(const HIP_vector_type& x) noexcept + { + data %= x.data; + return *this; + } + template< + typename U = T, + typename std::enable_if{}>::type* = nullptr> + __host__ __device__ + HIP_vector_type& operator^=(const HIP_vector_type& x) noexcept + { + data ^= x.data; + return *this; + } + template< + typename U = T, + typename std::enable_if{}>::type* = nullptr> + __host__ __device__ + HIP_vector_type& operator|=(const HIP_vector_type& x) noexcept + { + data |= x.data; + return *this; + } + template< + typename U = T, + typename std::enable_if{}>::type* = nullptr> + __host__ __device__ + HIP_vector_type& operator&=(const HIP_vector_type& x) noexcept + { + data &= x.data; + return *this; + } + template< + typename U = T, + typename std::enable_if{}>::type* = nullptr> + __host__ __device__ + HIP_vector_type& operator>>=(const HIP_vector_type& x) noexcept + { + data >>= x.data; + return *this; + } + template< + typename U = T, + typename std::enable_if{}>::type* = nullptr> + __host__ __device__ + HIP_vector_type& operator<<=(const HIP_vector_type& x) noexcept + { + data <<= x.data; + return *this; + } + }; + + + template + __host__ __device__ + inline + HIP_vector_type operator+( + const HIP_vector_type& x, const HIP_vector_type& y) noexcept + { + return HIP_vector_type{x} += y; + } + + template + __host__ __device__ + inline + HIP_vector_type operator-( + const HIP_vector_type& x, const HIP_vector_type& y) noexcept + { + return HIP_vector_type{x} -= y; + } + + template + __host__ __device__ + inline + HIP_vector_type operator*( + const HIP_vector_type& x, const HIP_vector_type& y) noexcept + { + return HIP_vector_type{x} *= y; + } + + template + __host__ __device__ + inline + HIP_vector_type operator/( + const HIP_vector_type& x, const HIP_vector_type& y) noexcept + { + return HIP_vector_type{x} /= y; + } + + template + __host__ __device__ + inline + bool operator==( + const HIP_vector_type& x, const HIP_vector_type& y) noexcept + { + auto tmp = x.data == y.data; + for (auto i = 0u; i != n; ++i) if (tmp[i] == 0) return false; + return true; + } + + template + __host__ __device__ + inline + bool operator!=( + const HIP_vector_type& x, const HIP_vector_type& y) noexcept + { + return !(x == y); + } + + template + __host__ __device__ + inline + bool operator<( + const HIP_vector_type& x, const HIP_vector_type& y) noexcept + { + auto tmp = x.data < y.data; + for (auto i = 0u; i != n; ++i) if (tmp[i] == 0) return false; + return true; + } + + template + __host__ __device__ + inline + bool operator>( + const HIP_vector_type& x, const HIP_vector_type& y) noexcept + { + return y < x; + } + + template + __host__ __device__ + inline + bool operator<=( + const HIP_vector_type& x, const HIP_vector_type& y) noexcept + { + return !(y < x); + } + + template + __host__ __device__ + inline + bool operator>=( + const HIP_vector_type& x, const HIP_vector_type& y) noexcept + { + return !(x < y); + } + + template< + typename T, + unsigned int n, + typename std::enable_if{}>* = nullptr> + inline + HIP_vector_type operator%( + const HIP_vector_type& x, const HIP_vector_type& y) noexcept + { + return HIP_vector_type{x} %= y; + } + + template< + typename T, + unsigned int n, + typename std::enable_if{}>* = nullptr> + inline + HIP_vector_type operator^( + const HIP_vector_type& x, const HIP_vector_type& y) noexcept + { + return HIP_vector_type{x} ^= y; + } + + template< + typename T, + unsigned int n, + typename std::enable_if{}>* = nullptr> + inline + HIP_vector_type operator|( + const HIP_vector_type& x, const HIP_vector_type& y) noexcept + { + return HIP_vector_type{x} |= y; + } + + template< + typename T, + unsigned int n, + typename std::enable_if{}>* = nullptr> + inline + HIP_vector_type operator&( + const HIP_vector_type& x, const HIP_vector_type& y) noexcept + { + return HIP_vector_type{x} &= y; + } + + template< + typename T, + unsigned int n, + typename std::enable_if{}>* = nullptr> + inline + HIP_vector_type operator>>( + const HIP_vector_type& x, const HIP_vector_type& y) noexcept + { + return HIP_vector_type{x} >>= y; + } + + template< + typename T, + unsigned int n, + typename std::enable_if{}>* = nullptr> + inline + HIP_vector_type operator<<( + const HIP_vector_type& x, const HIP_vector_type& y) noexcept + { + return HIP_vector_type{x} <<= y; + } + + // TODO: the following are rather dubious in terms of general utility. + template + inline + bool operator||( + const HIP_vector_type& x, const HIP_vector_type& y) noexcept + { + auto tmp = x.data || y.data; + for (auto i = 0u; i != n; ++i) if (tmp[i] == 0) return false; + return true; + } + + template + inline + bool operator&&( + const HIP_vector_type& x, const HIP_vector_type& y) noexcept + { + auto tmp = x.data && y.data; + for (auto i = 0u; i != n; ++i) if (tmp[i] == 0) return false; + return true; + } + + #define __MAKE_VECTOR_TYPE__(CUDA_name, T, n) \ + using CUDA_name = HIP_vector_type; +#else + typedef unsigned char uchar1 __NATIVE_VECTOR__(1, unsigned char); + typedef unsigned char uchar2 __NATIVE_VECTOR__(2, unsigned char); + typedef unsigned char uchar3 __NATIVE_VECTOR__(3, unsigned char); + typedef unsigned char uchar4 __NATIVE_VECTOR__(4, unsigned char); + + typedef char char1 __NATIVE_VECTOR__(1, char); + typedef char char2 __NATIVE_VECTOR__(2, char); + typedef char char3 __NATIVE_VECTOR__(3, char); + typedef char char4 __NATIVE_VECTOR__(4, char); + + typedef unsigned short ushort1 __NATIVE_VECTOR__(1, unsigned short); + typedef unsigned short ushort2 __NATIVE_VECTOR__(2, unsigned short); + typedef unsigned short ushort3 __NATIVE_VECTOR__(3, unsigned short); + typedef unsigned short ushort4 __NATIVE_VECTOR__(4, unsigned short); + + typedef short short1 __NATIVE_VECTOR__(1, short); + typedef short short2 __NATIVE_VECTOR__(2, short); + typedef short short3 __NATIVE_VECTOR__(3, short); + typedef short short4 __NATIVE_VECTOR__(4, short); + + typedef unsigned int uint1 __NATIVE_VECTOR__(1, unsigned int); + typedef unsigned int uint2 __NATIVE_VECTOR__(2, unsigned int); + typedef unsigned int uint3 __NATIVE_VECTOR__(3, unsigned int); + typedef unsigned int uint4 __NATIVE_VECTOR__(4, unsigned int); + + typedef int int1 __NATIVE_VECTOR__(1, int); + typedef int int2 __NATIVE_VECTOR__(2, int); + typedef int int3 __NATIVE_VECTOR__(3, int); + typedef int int4 __NATIVE_VECTOR__(4, int); + + typedef unsigned long ulong1 __NATIVE_VECTOR__(1, unsigned long); + typedef unsigned long ulong2 __NATIVE_VECTOR__(2, unsigned long); + typedef unsigned long ulong3 __NATIVE_VECTOR__(3, unsigned long); + typedef unsigned long ulong4 __NATIVE_VECTOR__(4, unsigned long); + + typedef long long1 __NATIVE_VECTOR__(1, long); + typedef long long2 __NATIVE_VECTOR__(2, long); + typedef long long3 __NATIVE_VECTOR__(3, long); + typedef long long4 __NATIVE_VECTOR__(4, long); + + typedef unsigned long long ulonglong1 __NATIVE_VECTOR__(1, unsigned long long); + typedef unsigned long long ulonglong2 __NATIVE_VECTOR__(2, unsigned long long); + typedef unsigned long long ulonglong3 __NATIVE_VECTOR__(3, unsigned long long); + typedef unsigned long long ulonglong4 __NATIVE_VECTOR__(4, unsigned long long); + + typedef long long longlong1 __NATIVE_VECTOR__(1, long long); + typedef long long longlong2 __NATIVE_VECTOR__(2, long long); + typedef long long longlong3 __NATIVE_VECTOR__(3, long long); + typedef long long longlong4 __NATIVE_VECTOR__(4, long long); + + typedef float float1 __NATIVE_VECTOR__(1, float); + typedef float float2 __NATIVE_VECTOR__(2, float); + typedef float float3 __NATIVE_VECTOR__(3, float); + typedef float float4 __NATIVE_VECTOR__(4, float); + + typedef double double1 __NATIVE_VECTOR__(1, double); + typedef double double2 __NATIVE_VECTOR__(2, double); + typedef double double3 __NATIVE_VECTOR__(3, double); + typedef double double4 __NATIVE_VECTOR__(4, double); +#endif + +__MAKE_VECTOR_TYPE__(uchar1, unsigned char, 1); +__MAKE_VECTOR_TYPE__(uchar2, unsigned char, 2); +__MAKE_VECTOR_TYPE__(uchar3, unsigned char, 3); +__MAKE_VECTOR_TYPE__(uchar4, unsigned char, 4); + +__MAKE_VECTOR_TYPE__(char1, char, 1); +__MAKE_VECTOR_TYPE__(char2, char, 2); +__MAKE_VECTOR_TYPE__(char3, char, 3); +__MAKE_VECTOR_TYPE__(char4, char, 4); + +__MAKE_VECTOR_TYPE__(ushort1, unsigned short, 1); +__MAKE_VECTOR_TYPE__(ushort2, unsigned short, 2); +__MAKE_VECTOR_TYPE__(ushort3, unsigned short, 3); +__MAKE_VECTOR_TYPE__(ushort4, unsigned short, 4); + +__MAKE_VECTOR_TYPE__(short1, short, 1); +__MAKE_VECTOR_TYPE__(short2, short, 2); +__MAKE_VECTOR_TYPE__(short3, short, 3); +__MAKE_VECTOR_TYPE__(short4, short, 4); + +__MAKE_VECTOR_TYPE__(uint1, unsigned int, 1); +__MAKE_VECTOR_TYPE__(uint2, unsigned int, 2); +__MAKE_VECTOR_TYPE__(uint3, unsigned int, 3); +__MAKE_VECTOR_TYPE__(uint4, unsigned int, 4); + +__MAKE_VECTOR_TYPE__(int1, int, 1); +__MAKE_VECTOR_TYPE__(int2, int, 2); +__MAKE_VECTOR_TYPE__(int3, int, 3); +__MAKE_VECTOR_TYPE__(int4, int, 4); + +__MAKE_VECTOR_TYPE__(ulong1, unsigned long, 1); +__MAKE_VECTOR_TYPE__(ulong2, unsigned long, 2); +__MAKE_VECTOR_TYPE__(ulong3, unsigned long, 3); +__MAKE_VECTOR_TYPE__(ulong4, unsigned long, 4); + +__MAKE_VECTOR_TYPE__(long1, long, 1); +__MAKE_VECTOR_TYPE__(long2, long, 2); +__MAKE_VECTOR_TYPE__(long3, long, 3); +__MAKE_VECTOR_TYPE__(long4, long, 4); + +__MAKE_VECTOR_TYPE__(ulonglong1, unsigned long long, 1); +__MAKE_VECTOR_TYPE__(ulonglong2, unsigned long long, 2); +__MAKE_VECTOR_TYPE__(ulonglong3, unsigned long long, 3); +__MAKE_VECTOR_TYPE__(ulonglong4, unsigned long long, 4); + +__MAKE_VECTOR_TYPE__(longlong1, long long, 1); +__MAKE_VECTOR_TYPE__(longlong2, long long, 2); +__MAKE_VECTOR_TYPE__(longlong3, long long, 3); +__MAKE_VECTOR_TYPE__(longlong4, long long, 4); + +__MAKE_VECTOR_TYPE__(float1, float, 1); +__MAKE_VECTOR_TYPE__(float2, float, 2); +__MAKE_VECTOR_TYPE__(float3, float, 3); +__MAKE_VECTOR_TYPE__(float4, float, 4); + +__MAKE_VECTOR_TYPE__(double1, double, 1); +__MAKE_VECTOR_TYPE__(double2, double, 2); +__MAKE_VECTOR_TYPE__(double3, double, 3); +__MAKE_VECTOR_TYPE__(double4, double, 4); #define DECLOP_MAKE_ONE_COMPONENT(comp, type) \ __device__ __host__ \ diff --git a/src/device_functions.cpp b/src/device_functions.cpp index 8ef19bab3f..8dfd5a07c4 100644 --- a/src/device_functions.cpp +++ b/src/device_functions.cpp @@ -355,33 +355,33 @@ __device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask __device__ char4 __hip_hc_add8pk(char4 in1, char4 in2) { char4 out; - unsigned one1 = in1.a & MASK1; - unsigned one2 = in2.a & MASK1; - out.a = (one1 + one2) & MASK1; - one1 = in1.a & MASK2; - one2 = in2.a & MASK2; - out.a = out.a | ((one1 + one2) & MASK2); + unsigned one1 = in1.w & MASK1; + unsigned one2 = in2.w & MASK1; + out.w = (one1 + one2) & MASK1; + one1 = in1.w & MASK2; + one2 = in2.w & MASK2; + out.w = out.w | ((one1 + one2) & MASK2); return out; } __device__ char4 __hip_hc_sub8pk(char4 in1, char4 in2) { char4 out; - unsigned one1 = in1.a & MASK1; - unsigned one2 = in2.a & MASK1; - out.a = (one1 - one2) & MASK1; - one1 = in1.a & MASK2; - one2 = in2.a & MASK2; - out.a = out.a | ((one1 - one2) & MASK2); + unsigned one1 = in1.w & MASK1; + unsigned one2 = in2.w & MASK1; + out.w = (one1 - one2) & MASK1; + one1 = in1.w & MASK2; + one2 = in2.w & MASK2; + out.w = out.w | ((one1 - one2) & MASK2); return out; } __device__ char4 __hip_hc_mul8pk(char4 in1, char4 in2) { char4 out; - unsigned one1 = in1.a & MASK1; - unsigned one2 = in2.a & MASK1; - out.a = (one1 * one2) & MASK1; - one1 = in1.a & MASK2; - one2 = in2.a & MASK2; - out.a = out.a | ((one1 * one2) & MASK2); + unsigned one1 = in1.w & MASK1; + unsigned one2 = in2.w & MASK1; + out.w = (one1 * one2) & MASK1; + one1 = in1.w & MASK2; + one2 = in2.w & MASK2; + out.w = out.w | ((one1 * one2) & MASK2); return out; } diff --git a/tests/src/deviceLib/hipVectorTypes.cpp b/tests/src/deviceLib/hipVectorTypes.cpp index 3c36fb5d2e..bc7e4eb356 100644 --- a/tests/src/deviceLib/hipVectorTypes.cpp +++ b/tests/src/deviceLib/hipVectorTypes.cpp @@ -52,20 +52,20 @@ template< __device__ bool integer_unary_tests(V& f1, V& f2) { f1 %= f2; - if (!cmp(f1, 0)) return false; + if (f1 != V{0}) return false; f1 &= f2; - if (!cmp(f1, 0)) return false; + if (f1 != V{0}) return false; f1 |= f2; - if (!cmp(f1, 1)) return false; + if (f1 != V{1}) return false; f1 ^= f2; - if (!cmp(f1, 0)) return false; - f1.x = 1; + if (f1 != V{0}) return false; + f1 = V{1}; f1 <<= f2; - if (!cmp(f1, 2)) return false; + if (f1 != V{2}) return false; f1 >>= f2; - if (!cmp(f1, 1)) return false; + if (f1 != V{1}) return false; f2 = ~f1; - return cmp(f2, ~1); + return f2 == V{~1}; } template< @@ -74,17 +74,17 @@ template< __device__ bool integer_binary_tests(V& f1, V& f2, V& f3) { f3 = f1 % f2; - if (!cmp(f3, 0)) return false; + if (f3 != V{0}) return false; f1 = f3 & f2; - if (!cmp(f1, 0)) return false; + if (f1 != V{0}) return false; f2 = f1 ^ f3; - if (!cmp(f2, 0)) return false; - f1.x = 1; - f2.x = 2; + if (f2 != V{0}) return false; + f1 = V{1}; + f2 = V{2}; f3 = f1 << f2; - if (!cmp(f3, 4)) return false; + if (f3 != V{4}) return false; f2 = f3 >> f1; - if (!cmp(f2, 2)) return false; + return f2 == V{2}; } template @@ -107,60 +107,58 @@ bool constructor_tests() { template bool TestVectorType() { - V f1(1); - V f2(1); + V f1{1}; + V f2{1}; V f3 = f1 + f2; - if (!cmp(f3, 2)) return false; + if (f3 != V{2}) return false; f2 = f3 - f1; - if (!cmp(f2, 1)) return false; + if (f2 != V{1}) return false; f1 = f2 * f3; - if (!cmp(f1, 2)) return false; + if (f1 != V{2}) return false; f2 = f1 / f3; - if (!cmp(f2, 2 / 2)) return false; + if (f2 != V{1}) return false; if (!integer_binary_tests(f1, f2, f3)) return false; - f1 = V(2); - f2 = V(1); + f1 = V{2}; + f2 = V{1}; f1 += f2; - if (!cmp(f1, 3)) return false; + if (f1 != V{3}) return false; f1 -= f2; - if (!cmp(f1, 2)) return false; + if (f1 != V{2}) return false; f1 *= f2; - if (!cmp(f1, 2)) return false; + if (f1 != V{2}) return false; f1 /= f2; - if (!cmp(f1, 2)) return false; + if (f1 != V{2}) return false; if (!integer_unary_tests(f1, f2)) return false; - #if false // We do not enable nullary increment / decrement yet. - f1 = V(2); - f2 = f1++; - if (!cmp(f1, 3)) return false; - if (!cmp(f2, 2)) return false; - f2 = f1--; - if (!cmp(f2, 3)) return false; - if (!cmp(f1, 2)) return false; - f2 = ++f1; - if (!cmp(f1, 3)) return false; - if (!cmp(f2, 3)) return false; - f2 = --f1; - if (!cmp(f1, 2)) return false; - if (!cmp(f2, 2)) return false; - #endif + f1 = V{2}; + f2 = f1++; + if (f1 != V{3}) return false; + if (f2 != V{2}) return false; + f2 = f1--; + if (f2 != V{3}) return false; + if (f1 != V{2}) return false; + f2 = ++f1; + if (f1 != V{3}) return false; + if (f2 != V{3}) return false; + f2 = --f1; + if (f1 != V{2}) return false; + if (f2 != V{2}) return false; if (!constructor_tests()) return false; - f1 = V(3); - f2 = V(4); - f3 = V(3); - if (cmp(f1 == f2, true)) return false; - if (cmp(f1 != f2, false)) return false; - if (cmp(f1 < f2, false)) return false; - if (cmp(f2 > f1, false)) return false; - if (cmp(f1 >= f3, false)) return false; - if (cmp(f1 <= f3, false)) return false; + f1 = V{3}; + f2 = V{4}; + f3 = V{3}; + if (f1 == f2) return false; + if (!(f1 != f2)) return false; + if (!(f1 < f2)) return false; + if (!(f2 > f1)) return false; + if (!(f1 >= f3)) return false; + if (!(f1 <= f3)) return false; - if (cmp(f1 && f2, false)) return false; - if (cmp(f1 || f2, false)) return false; + if (!(f1 && f2)) return false; + if (!(f1 || f2)) return false; return true; } diff --git a/tests/src/deviceLib/hipVectorTypesDevice.cpp b/tests/src/deviceLib/hipVectorTypesDevice.cpp index edb817ced1..21dc1f1a75 100644 --- a/tests/src/deviceLib/hipVectorTypesDevice.cpp +++ b/tests/src/deviceLib/hipVectorTypesDevice.cpp @@ -51,20 +51,23 @@ template< __device__ bool integer_unary_tests(V& f1, V& f2) { f1 %= f2; - if (!cmp(f1, 0)) return false; + if (f1 != V{0}) return false; + f1 &= f2; - if (!cmp(f1, 0)) return false; + if (f1 != V{0}) return false; f1 |= f2; - if (!cmp(f1, 1)) return false; + if (f1 != V{1}) return false; f1 ^= f2; - if (!cmp(f1, 0)) return false; - f1.x = 1; + if (f1 != V{0}) return false; + f1 = V{1}; f1 <<= f2; - if (!cmp(f1, 2)) return false; + if (f1 != V{2}) return false; f1 >>= f2; - if (!cmp(f1, 1)) return false; + if (f1 != V{1}) return false; f2 = ~f1; - return cmp(f2, ~1); + return f2 == V{~1}; + + return true; } template< @@ -81,74 +84,72 @@ template< __device__ bool integer_binary_tests(V& f1, V& f2, V& f3) { f3 = f1 % f2; - if (!cmp(f3, 0)) return false; + if (f3 != V{0}) return false; f1 = f3 & f2; - if (!cmp(f1, 0)) return false; + if (f1 != V{0}) return false; f2 = f1 ^ f3; - if (!cmp(f2, 0)) return false; - f1.x = 1; - f2.x = 2; + if (f2 != V{0}) return false; + f1 = V{1}; + f2 = V{2}; f3 = f1 << f2; - if (!cmp(f3, 4)) return false; + if (f3 != V{4}) return false; f2 = f3 >> f1; - if (!cmp(f2, 2)) return false; + return f2 == V{2}; } template __device__ bool TestVectorType() { - V f1(1); - V f2(1); + V f1{1}; + V f2{1}; V f3 = f1 + f2; - if (!cmp(f3, 2)) return false; + if (f3 != V{2}) return false; f2 = f3 - f1; - if (!cmp(f2, 1)) return false; + if (f2 != V{1}) return false; f1 = f2 * f3; - if (!cmp(f1, 2)) return false; + if (f1 != V{2}) return false; f2 = f1 / f3; - if (!cmp(f2, 2 / 2)) return false; + if (f2 != V{1}) return false; if (!integer_binary_tests(f1, f2, f3)) return false; - f1 = V(2); - f2 = V(1); + f1 = V{2}; + f2 = V{1}; f1 += f2; - if (!cmp(f1, 3)) return false; + if (f1 != V{3}) return false; f1 -= f2; - if (!cmp(f1, 2)) return false; + if (f1 != V{2}) return false; f1 *= f2; - if (!cmp(f1, 2)) return false; + if (f1 != V{2}) return false; f1 /= f2; - if (!cmp(f1, 2)) return false; + if (f1 != V{2}) return false; if (!integer_unary_tests(f1, f2)) return false; - #if false // We do not enable nullary increment / decrement yet. - f1 = V(2); - f2 = f1++; - if (!cmp(f1, 3)) return false; - if (!cmp(f2, 2)) return false; - f2 = f1--; - if (!cmp(f2, 3)) return false; - if (!cmp(f1, 2)) return false; - f2 = ++f1; - if (!cmp(f1, 3)) return false; - if (!cmp(f2, 3)) return false; - f2 = --f1; - if (!cmp(f1, 2)) return false; - if (!cmp(f2, 2)) return false; - #endif + f1 = V{2}; + f2 = f1++; + if (f1 != V{3}) return false; + if (f2 != V{2}) return false; + f2 = f1--; + if (f2 != V{3}) return false; + if (f1 != V{2}) return false; + f2 = ++f1; + if (f1 != V{3}) return false; + if (f2 != V{3}) return false; + f2 = --f1; + if (f1 != V{2}) return false; + if (f2 != V{2}) return false; - f1 = V(3); - f2 = V(4); - f3 = V(3); - if (cmp(f1 == f2, true)) return false; - if (cmp(f1 != f2, false)) return false; - if (cmp(f1 < f2, false)) return false; - if (cmp(f2 > f1, false)) return false; - if (cmp(f1 >= f3, false)) return false; - if (cmp(f1 <= f3, false)) return false; + f1 = V{3}; + f2 = V{4}; + f3 = V{3}; + if (f1 == f2) return false; + if (!(f1 != f2)) return false; + if (!(f1 < f2)) return false; + if (!(f2 > f1)) return false; + if (!(f1 >= f3)) return false; + if (!(f1 <= f3)) return false; - if (cmp(f1 && f2, false)) return false; - if (cmp(f1 || f2, false)) return false; + if (!(f1 && f2)) return false; + if (!(f1 || f2)) return false; return true; } diff --git a/tests/src/deviceLib/vector_test_common.h b/tests/src/deviceLib/vector_test_common.h index d5bc4c57a2..fac5ab84a1 100644 --- a/tests/src/deviceLib/vector_test_common.h +++ b/tests/src/deviceLib/vector_test_common.h @@ -66,40 +66,4 @@ bool is_vec() { ((dimension == 2) ? decltype(is_vec2(std::declval())){} : ((dimension == 3) ? decltype(is_vec3(std::declval())){} : decltype(is_vec4(std::declval())){})); -} - -template()>* = nullptr> -__host__ __device__ -inline -bool cmp(const T& x, U expected) { - const auto r = x == T(expected); - - return r.x != 0; -} - -template()>* = nullptr> -__host__ __device__ -inline -bool cmp(const T& x, U expected) { - const auto r = x == T(expected); - - return r.x != 0 && r.y != 0; -} - -template()>* = nullptr> -__host__ __device__ -inline -bool cmp(const T& x, U expected) { - const auto r = x == T(expected); - - return r.x != 0 && r.y != 0 && r.z != 0; -} - -template()>* = nullptr> -__host__ __device__ -inline -bool cmp(const T& x, U expected) { - const auto r = x == T(expected); - - return r.x != 0 && r.y != 0 && r.z != 0 && r.w != 0; } \ No newline at end of file