From 90e7799bf41b78fba20c31dca99607d2e815b61a Mon Sep 17 00:00:00 2001 From: Alex Voicu Date: Thu, 31 May 2018 15:27:12 +0100 Subject: [PATCH] Add missing interop with volatile. Fix unit tests. [ROCm/hip commit: 0108819e2aa0c66145b59118c325d7ee67c1cfd7] --- .../hip/include/hip/hcc_detail/hip_fp16.h | 50 +++++++++- .../hip/tests/src/deviceLib/hipTestHalf.cpp | 25 +++-- .../tests/src/deviceLib/hipTestNativeHalf.cpp | 94 +++++++++++++------ 3 files changed, 124 insertions(+), 45 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/hip_fp16.h b/projects/hip/include/hip/hcc_detail/hip_fp16.h index 50cbb20709..8657bc30a3 100644 --- a/projects/hip/include/hip/hcc_detail/hip_fp16.h +++ b/projects/hip/include/hip/hcc_detail/hip_fp16.h @@ -58,6 +58,11 @@ THE SOFTWARE. #include "hip_vector_types.h" #include "host_defines.h" + namespace std + { + template<> struct is_floating_point<_Float16> : std::true_type {}; + } + template using Enable_if_t = typename std::enable_if::type; @@ -111,6 +116,32 @@ THE SOFTWARE. data = x.data; return *this; } + __host__ __device__ + volatile __half& operator=(const __half_raw& x) volatile + { + data = x.data; + return *this; + } + volatile __half& operator=(const volatile __half_raw& x) volatile + { + data = x.data; + return *this; + } + __half& operator=(__half_raw&& x) + { + data = x.data; + return *this; + } + volatile __half& operator=(__half_raw&& x) volatile + { + data = x.data; + return *this; + } + volatile __half& operator=(volatile __half_raw&& x) volatile + { + data = x.data; + return *this; + } #if !defined(__HIP_NO_HALF_CONVERSIONS__) template< typename T, @@ -182,18 +213,27 @@ THE SOFTWARE. // ACCESSORS #if !defined(__HIP_NO_HALF_CONVERSIONS__) - __host__ __device__ - operator decltype(data)() const { return data; } - __host__ __device__ - operator float() const { return data; } + template< + typename T, + Enable_if_t< + std::is_floating_point{} && + !std::is_same{}>* = nullptr> + operator T() const { return data; } #endif __host__ __device__ operator __half_raw() const { return __half_raw{data}; } + __host__ __device__ + operator volatile __half_raw() const volatile + { + return __half_raw{data}; + } // ACCESSORS - DEVICE ONLY #if !defined(__HIP_NO_HALF_CONVERSIONS__) + template< + typename T, Enable_if_t{}>* = nullptr> __device__ - operator bool() const { return data; } + operator T() const { return data; } #endif #if !defined(__HIP_NO_HALF_OPERATORS__) diff --git a/projects/hip/tests/src/deviceLib/hipTestHalf.cpp b/projects/hip/tests/src/deviceLib/hipTestHalf.cpp index 6b2ee5e29e..125edd76eb 100644 --- a/projects/hip/tests/src/deviceLib/hipTestHalf.cpp +++ b/projects/hip/tests/src/deviceLib/hipTestHalf.cpp @@ -31,10 +31,7 @@ THE SOFTWARE. #if __HIP_ARCH_GFX803__ || __HIP_ARCH_GFX900__ || __HIP_ARCH_GFX906__ __global__ -__attribute__((optnone)) -void __halfMath(bool* result) { - __half a{1}; - +void __halfMath(bool* result, __half a) { result[0] = __heq(__hadd(a, __half{1}), __half{2}); result[0] = __heq(__hadd_sat(a, __half{1}), __half{1}) && result[0]; result[0] = __heq(__hfma(a, __half{2}, __half{3}), __half{5}) && result[0]; @@ -56,11 +53,9 @@ bool to_bool(const __half2& x) } __global__ -__attribute__((optnone)) -void __half2Math(bool* result) { - __half2 a{1, 1}; - - result[0] = to_bool(__heq2(__hadd2(a, __half2{1, 1}), __half2{2, 2})); +void __half2Math(bool* result, __half2 a) { + result[0] = + to_bool(__heq2(__hadd2(a, __half2{1, 1}), __half2{2, 2})); result[0] = to_bool(__heq2(__hadd2_sat(a, __half2{1, 1}), __half2{1, 1})) && result[0]; result[0] = to_bool(__heq2( @@ -79,12 +74,14 @@ void __half2Math(bool* result) { result[0]; } -__global__ void kernel_hisnan(__half* input, int* output) { +__global__ +void kernel_hisnan(__half* input, int* output) { int tx = threadIdx.x; output[tx] = __hisnan(input[tx]); } -__global__ void kernel_hisinf(__half* input, int* output) { +__global__ +void kernel_hisinf(__half* input, int* output) { int tx = threadIdx.x; output[tx] = __hisinf(input[tx]); } @@ -235,13 +232,15 @@ int main() { hipHostMalloc(&result, sizeof(result)); result[0] = false; - hipLaunchKernelGGL(__halfMath, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, result); + hipLaunchKernelGGL( + __halfMath, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, result, __half{1}); hipDeviceSynchronize(); if (!result[0]) { failed("Failed __half tests."); } result[0] = false; - hipLaunchKernelGGL(__half2Math, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, result); + hipLaunchKernelGGL( + __half2Math, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, result, __half2{1, 1}); hipDeviceSynchronize(); if (!result[0]) { failed("Failed __half2 tests."); } diff --git a/projects/hip/tests/src/deviceLib/hipTestNativeHalf.cpp b/projects/hip/tests/src/deviceLib/hipTestNativeHalf.cpp index 54a19f42bf..6e618618f6 100644 --- a/projects/hip/tests/src/deviceLib/hipTestNativeHalf.cpp +++ b/projects/hip/tests/src/deviceLib/hipTestNativeHalf.cpp @@ -28,28 +28,60 @@ THE SOFTWARE. #include "test_common.h" +#include + +using namespace std; + #if __HIP_ARCH_GFX803__ || __HIP_ARCH_GFX900__ || __HIP_ARCH_GFX906__ __global__ -__attribute__((optnone)) -void __halfTest(bool* result) { +void __halfTest(bool* result, __half a) { // Construction - __half a{1}; result[0] = __heq(a, 1); - a = __half{1.0f}; result[0] = __heq(a, 1) && result[0]; - a = __half{1.0}; result[0] = __heq(a, 1) && result[0]; - a = __half{static_cast(1)}; - result[0] = __heq(a, 1) && result[0]; - a = __half{static_cast(1)}; result[0] = __heq(a, 1) && result[0]; - a = __half{1u}; result[0] = __heq(a, 1) && result[0]; - a = __half{1ul}; result[0] = __heq(a, 1) && result[0]; - a = __half{1l}; result[0] = __heq(a, 1) && result[0]; - a = __half{1ll}; result[0] = __heq(a, 1) && result[0]; - a = __half{1ull}; result[0] = __heq(a, 1) && result[0]; + static_assert(is_default_constructible<__half>{}, ""); + static_assert(is_copy_constructible<__half>{}, ""); + static_assert(is_move_constructible<__half>{}, ""); + static_assert(is_constructible<__half, float>{}, ""); + static_assert(is_constructible<__half, double>{}, ""); + static_assert(is_constructible<__half, unsigned short>{}, ""); + static_assert(is_constructible<__half, short>{}, ""); + static_assert(is_constructible<__half, unsigned int>{}, ""); + static_assert(is_constructible<__half, int>{}, ""); + static_assert(is_constructible<__half, unsigned long>{}, ""); + static_assert(is_constructible<__half, long>{}, ""); + static_assert(is_constructible<__half, long long>{}, ""); + static_assert(is_constructible<__half, unsigned long long>{}, ""); + static_assert(is_constructible<__half, __half_raw>{}, ""); // Assignment - a = 0.0f; result[0] = __heq(a, 0) && result[0]; - a = 1.0; result[0] = __heq(a, 1) && result[0]; - a = __half_raw{2}; result[0] = __heq(a, 2) && result[0]; + static_assert(is_copy_assignable<__half>{}, ""); + static_assert(is_move_assignable<__half>{}, ""); + static_assert(is_assignable<__half, float>{}, ""); + static_assert(is_assignable<__half, double>{}, ""); + static_assert(is_assignable<__half, unsigned short>{}, ""); + static_assert(is_assignable<__half, short>{}, ""); + static_assert(is_assignable<__half, unsigned int>{}, ""); + static_assert(is_assignable<__half, int>{}, ""); + static_assert(is_assignable<__half, unsigned long>{}, ""); + static_assert(is_assignable<__half, long>{}, ""); + static_assert(is_assignable<__half, long long>{}, ""); + static_assert(is_assignable<__half, unsigned long long>{}, ""); + static_assert(is_assignable<__half, __half_raw>{}, ""); + static_assert(is_assignable<__half, volatile __half_raw&>{}, ""); + static_assert(is_assignable<__half, volatile __half_raw&&>{}, ""); + + // Conversion + static_assert(is_convertible<__half, float>{}, ""); + static_assert(is_convertible<__half, unsigned short>{}, ""); + static_assert(is_convertible<__half, short>{}, ""); + static_assert(is_convertible<__half, unsigned int>{}, ""); + static_assert(is_convertible<__half, int>{}, ""); + static_assert(is_convertible<__half, unsigned long>{}, ""); + static_assert(is_convertible<__half, long>{}, ""); + static_assert(is_convertible<__half, long long>{}, ""); + static_assert(is_convertible<__half, bool>{}, ""); + static_assert(is_convertible<__half, unsigned long long>{}, ""); + static_assert(is_convertible<__half, __half_raw>{}, ""); + static_assert(is_convertible<__half, volatile __half_raw>{}, ""); // Nullary result[0] = __heq(a, +a) && result[0]; @@ -83,17 +115,23 @@ bool to_bool(const __half2& x) return r.data.x != 0 && r.data.y != 0; } + __global__ -__attribute__((optnone)) -void __half2Test(bool* result) { +void __half2Test(bool* result, __half2 a) { // Construction - __half2 a{1}; - result[0] = to_bool(__heq2(a, 1)); - a = __half2{__half{1}, __half{1}}; - result[0] = to_bool(__heq2(a, {1, 1})) && result[0]; + static_assert(is_default_constructible<__half2>{}, ""); + static_assert(is_copy_constructible<__half2>{}, ""); + static_assert(is_move_constructible<__half2>{}, ""); + static_assert(is_constructible<__half2, __half, __half>{}, ""); + static_assert(is_constructible<__half2, __half2_raw>{}, ""); // Assignment - a = __half2_raw{2}; result[0] = to_bool(__heq2(a, {2, 2})) && result[0]; + static_assert(is_copy_assignable<__half2>{}, ""); + static_assert(is_move_assignable<__half2>{}, ""); + static_assert(is_assignable<__half2, __half2_raw>{}, ""); + + // Conversion + static_assert(is_convertible<__half2, __half2_raw>{}, ""); // Nullary result[0] = to_bool(__heq2(a, +a)) && result[0]; @@ -126,14 +164,16 @@ int main() { bool* result{nullptr}; hipHostMalloc(&result, 1); - result[0] = false; - hipLaunchKernelGGL(__halfTest, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, result); + result[0] = true; + hipLaunchKernelGGL( + __halfTest, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, result, __half{1}); hipDeviceSynchronize(); if (!result[0]) { failed("Failed __half tests."); } - result[0] = false; - hipLaunchKernelGGL(__half2Test, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, result); + result[0] = true; + hipLaunchKernelGGL( + __half2Test, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, result, __half2{1, 1}); hipDeviceSynchronize(); if (!result[0]) { failed("Failed __half2 tests."); }