Re-factor half support to match CUDA whilst exploiting native support.
[ROCm/clr commit: 36e805cf76]
This commit is contained in:
Alex Voicu
@@ -195,7 +195,6 @@ if(HIP_PLATFORM STREQUAL "hcc")
|
||||
set(SOURCE_FILES_DEVICE
|
||||
src/device_util.cpp
|
||||
src/hip_ldg.cpp
|
||||
src/hip_fp16.cpp
|
||||
src/device_functions.cpp
|
||||
src/math_functions.cpp)
|
||||
|
||||
|
||||
تفاوت فایلی نمایش داده نمی شود زیرا این فایل بسیار بزرگ است
Diff را بارگزاری کن
@@ -0,0 +1,252 @@
|
||||
#pragma once
|
||||
|
||||
#if defined(__cplusplus)
|
||||
#include <cstring>
|
||||
#endif
|
||||
|
||||
struct __half_raw {
|
||||
unsigned short x;
|
||||
};
|
||||
|
||||
struct __half2_raw {
|
||||
unsigned short x;
|
||||
unsigned short y;
|
||||
};
|
||||
|
||||
#if defined(__cplusplus)
|
||||
struct __half;
|
||||
|
||||
__half __float2half(float);
|
||||
float __half2float(__half);
|
||||
|
||||
// BEGIN STRUCT __HALF
|
||||
struct __half {
|
||||
protected:
|
||||
unsigned short __x;
|
||||
public:
|
||||
// CREATORS
|
||||
__half() = default;
|
||||
__half(const __half_raw& x) : __x{x.x} {}
|
||||
#if !defined(__HIP_NO_HALF_CONVERSIONS__)
|
||||
__half(float x) : __x{__float2half(x).__x} {}
|
||||
__half(double x) : __x{__float2half(x).__x} {}
|
||||
#endif
|
||||
__half(const __half&) = default;
|
||||
__half(__half&&) = default;
|
||||
~__half() = default;
|
||||
|
||||
// MANIPULATORS
|
||||
__half& operator=(const __half&) = default;
|
||||
__half& operator=(__half&&) = default;
|
||||
__half& operator=(const __half_raw& x) { __x = x.x; return *this; }
|
||||
#if !defined(__HIP_NO_HALF_CONVERSIONS__)
|
||||
__half& operator=(float x)
|
||||
{
|
||||
__x = __float2half(x).__x;
|
||||
return *this;
|
||||
}
|
||||
__half& operator=(double x)
|
||||
{
|
||||
return *this = static_cast<float>(x);
|
||||
}
|
||||
#endif
|
||||
|
||||
// ACCESSORS
|
||||
operator float() const { return __half2float(*this); }
|
||||
operator __half_raw() const { return __half_raw{__x}; }
|
||||
};
|
||||
// END STRUCT __HALF
|
||||
|
||||
// BEGIN STRUCT __HALF2
|
||||
struct __half2 {
|
||||
protected:
|
||||
__half x;
|
||||
__half y;
|
||||
public:
|
||||
// CREATORS
|
||||
__half2() = default;
|
||||
__half2(const __half2_raw& ix)
|
||||
:
|
||||
x{reinterpret_cast<const __half&>(ix.x)},
|
||||
y{reinterpret_cast<const __half&>(ix.y)}
|
||||
{}
|
||||
__half2(const __half& ix, const __half& iy) : x{ix}, y{iy} {}
|
||||
__half2(const __half2&) = default;
|
||||
__half2(__half2&&) = default;
|
||||
~__half2() = default;
|
||||
|
||||
// MANIPULATORS
|
||||
__half2& operator=(const __half2&) = default;
|
||||
__half2& operator=(__half2&&) = default;
|
||||
__half2& operator=(const __half2_raw& ix)
|
||||
{
|
||||
x = reinterpret_cast<const __half_raw&>(ix.x);
|
||||
y = reinterpret_cast<const __half_raw&>(ix.y);
|
||||
return *this;
|
||||
}
|
||||
|
||||
// ACCESSORS
|
||||
operator __half2_raw() const
|
||||
{
|
||||
return __half2_raw{
|
||||
reinterpret_cast<const unsigned short&>(x),
|
||||
reinterpret_cast<const unsigned short&>(y)};
|
||||
}
|
||||
};
|
||||
// END STRUCT __HALF2
|
||||
|
||||
namespace
|
||||
{
|
||||
inline
|
||||
unsigned short __internal_float2half(
|
||||
float flt, unsigned int& sgn, unsigned int& rem)
|
||||
{
|
||||
unsigned int x{};
|
||||
std::memcpy(&x, &flt, sizeof(flt));
|
||||
|
||||
unsigned int u = (x & 0x7fffffffU);
|
||||
sgn = ((x >> 16) & 0x8000U);
|
||||
|
||||
// NaN/+Inf/-Inf
|
||||
if (u >= 0x7f800000U) {
|
||||
rem = 0;
|
||||
return static_cast<unsigned short>(
|
||||
(u == 0x7f800000U) ? (sgn | 0x7c00U) : 0x7fffU);
|
||||
}
|
||||
// Overflows
|
||||
if (u > 0x477fefffU) {
|
||||
rem = 0x80000000U;
|
||||
return static_cast<unsigned short>(sgn | 0x7bffU);
|
||||
}
|
||||
// Normal numbers
|
||||
if (u >= 0x38800000U) {
|
||||
rem = u << 19;
|
||||
u -= 0x38000000U;
|
||||
return static_cast<unsigned short>(sgn | (u >> 13));
|
||||
}
|
||||
// +0/-0
|
||||
if (u < 0x33000001U) {
|
||||
rem = u;
|
||||
return static_cast<unsigned short>(sgn);
|
||||
}
|
||||
// Denormal numbers
|
||||
unsigned int exponent = u >> 23;
|
||||
unsigned int mantissa = (u & 0x7fffffU);
|
||||
unsigned int shift = 0x7eU - exponent;
|
||||
mantissa |= 0x800000U;
|
||||
rem = mantissa << (32 - shift);
|
||||
return static_cast<unsigned short>(sgn | (mantissa >> shift));
|
||||
}
|
||||
|
||||
inline
|
||||
__half __float2half(float x)
|
||||
{
|
||||
__half_raw r;
|
||||
unsigned int sgn{};
|
||||
unsigned int rem{};
|
||||
r.x = __internal_float2half(x, sgn, rem);
|
||||
if (rem > 0x80000000U || (rem == 0x80000000U && (r.x & 0x1))) ++r.x;
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
inline
|
||||
__half __float2half_rn(float x) { return __float2half(x); }
|
||||
|
||||
inline
|
||||
__half __float2half_rz(float x)
|
||||
{
|
||||
__half_raw r;
|
||||
unsigned int sgn{};
|
||||
unsigned int rem{};
|
||||
r.x = __internal_float2half(x, sgn, rem);
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
inline
|
||||
__half __float2half_rd(float x)
|
||||
{
|
||||
__half_raw r;
|
||||
unsigned int sgn{};
|
||||
unsigned int rem{};
|
||||
r.x = __internal_float2half(x, sgn, rem);
|
||||
if (rem && sgn) ++r.x;
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
inline
|
||||
__half __float2half_ru(float x)
|
||||
{
|
||||
__half_raw r;
|
||||
unsigned int sgn{};
|
||||
unsigned int rem{};
|
||||
r.x = __internal_float2half(x, sgn, rem);
|
||||
if (rem && !sgn) ++r.x;
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
inline
|
||||
__half2 __float2half2_rn(float x)
|
||||
{
|
||||
return __half2{__float2half_rn(x), __float2half_rn(x)};
|
||||
}
|
||||
|
||||
inline
|
||||
__half2 __floats2half2_rn(float x, float y)
|
||||
{
|
||||
return __half2{__float2half_rn(x), __float2half_rn(y)};
|
||||
}
|
||||
|
||||
inline
|
||||
float __internal_half2float(unsigned short x)
|
||||
{
|
||||
unsigned int sign = ((x >> 15) & 1);
|
||||
unsigned int exponent = ((x >> 10) & 0x1f);
|
||||
unsigned int mantissa = ((x & 0x3ff) << 13);
|
||||
|
||||
if (exponent == 0x1fU) { /* NaN or Inf */
|
||||
mantissa = (mantissa ? (sign = 0, 0x7fffffU) : 0);
|
||||
exponent = 0xffU;
|
||||
} else if (!exponent) { /* Denorm or Zero */
|
||||
if (mantissa) {
|
||||
unsigned int msb;
|
||||
exponent = 0x71U;
|
||||
do {
|
||||
msb = (mantissa & 0x400000U);
|
||||
mantissa <<= 1; /* normalize */
|
||||
--exponent;
|
||||
} while (!msb);
|
||||
mantissa &= 0x7fffffU; /* 1.mantissa is implicit */
|
||||
}
|
||||
} else {
|
||||
exponent += 0x70U;
|
||||
}
|
||||
unsigned int u = ((sign << 31) | (exponent << 23) | mantissa);
|
||||
float f;
|
||||
memcpy(&f, &u, sizeof(u));
|
||||
|
||||
return f;
|
||||
}
|
||||
|
||||
inline
|
||||
float __half2float(__half x)
|
||||
{
|
||||
return __internal_half2float(static_cast<__half_raw>(x).x);
|
||||
}
|
||||
|
||||
inline
|
||||
float __low2float(__half2 x)
|
||||
{
|
||||
return __internal_half2float(static_cast<__half2_raw>(x).x);
|
||||
}
|
||||
|
||||
inline
|
||||
float __high2float(__half2 x)
|
||||
{
|
||||
return __internal_half2float(static_cast<__half2_raw>(x).y);
|
||||
}
|
||||
}
|
||||
#endif // defined(__cplusplus)
|
||||
@@ -0,0 +1,76 @@
|
||||
/*
|
||||
Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#pragma once
|
||||
|
||||
// /*
|
||||
// Half Math Functions
|
||||
// */
|
||||
|
||||
extern "C"
|
||||
{
|
||||
__attribute__((const)) _Float16 __ocml_ceil_f16(_Float16);
|
||||
_Float16 __ocml_cos_f16(_Float16);
|
||||
__attribute__((pure)) _Float16 __ocml_exp_f16(_Float16);
|
||||
__attribute__((pure)) _Float16 __ocml_exp10_f16(_Float16);
|
||||
__attribute__((pure)) _Float16 __ocml_exp2_f16(_Float16);
|
||||
__attribute__((const)) _Float16 __ocml_floor_f16(_Float16);
|
||||
__attribute__((const))
|
||||
_Float16 __ocml_fma_f16(_Float16, _Float16, _Float16);
|
||||
__attribute__((const)) int __ocml_isinf_f16(_Float16);
|
||||
__attribute__((const)) int __ocml_isnan_f16(_Float16);
|
||||
__attribute__((pure)) _Float16 __ocml_log_f16(_Float16);
|
||||
__attribute__((pure)) _Float16 __ocml_log10_f16(_Float16);
|
||||
__attribute__((pure)) _Float16 __ocml_log2_f16(_Float16);
|
||||
__attribute__((const)) _Float16 __llvm_amdgcn_rcp_f16(_Float16);
|
||||
__attribute__((const)) _Float16 __ocml_rint_f16(_Float16);
|
||||
__attribute__((const)) _Float16 __ocml_rsqrt_f16(_Float16);
|
||||
_Float16 __ocml_sin_f16(_Float16);
|
||||
__attribute__((const)) _Float16 __ocml_sqrt_f16(_Float16);
|
||||
__attribute__((const)) _Float16 __ocml_trunc_f16(_Float16);
|
||||
|
||||
typedef _Float16 __2f16 __attribute__((ext_vector_type(2)));
|
||||
typedef short __2i16 __attribute__((ext_vector_type(2)));
|
||||
|
||||
__attribute__((const)) __2f16 __ocml_ceil_2f16(__2f16);
|
||||
__2f16 __ocml_cos_2f16(__2f16);
|
||||
__attribute__((pure)) __2f16 __ocml_exp_2f16(__2f16);
|
||||
__attribute__((pure)) __2f16 __ocml_exp10_2f16(__2f16);
|
||||
__attribute__((pure)) __2f16 __ocml_exp2_2f16(__2f16);
|
||||
__attribute__((const)) __2f16 __ocml_floor_2f16(__2f16);
|
||||
__attribute__((const)) __2f16 __ocml_fma_2f16(__2f16, __2f16, __2f16);
|
||||
__attribute__((const)) __2i16 __ocml_isinf_2f16(__2f16);
|
||||
__attribute__((const)) __2i16 __ocml_isnan_2f16(__2f16);
|
||||
__attribute__((pure)) __2f16 __ocml_log_2f16(__2f16);
|
||||
__attribute__((pure)) __2f16 __ocml_log10_2f16(__2f16);
|
||||
__attribute__((pure)) __2f16 __ocml_log2_2f16(__2f16);
|
||||
inline
|
||||
__2f16 __llvm_amdgcn_rcp_2f16(__2f16 x) // Not currently exposed by ROCDL.
|
||||
{
|
||||
return __2f16{__llvm_amdgcn_rcp_f16(x.x), __llvm_amdgcn_rcp_f16(x.y)};
|
||||
}
|
||||
__attribute__((const)) __2f16 __ocml_rint_2f16(__2f16);
|
||||
__attribute__((const)) __2f16 __ocml_rsqrt_2f16(__2f16);
|
||||
__2f16 __ocml_sin_2f16(__2f16);
|
||||
__attribute__((const)) __2f16 __ocml_sqrt_2f16(__2f16);
|
||||
__attribute__((const)) __2f16 __ocml_trunc_2f16(__2f16);
|
||||
}
|
||||
@@ -34,7 +34,7 @@ THE SOFTWARE.
|
||||
|
||||
#if __HIP_ARCH_GFX803__ || __HIP_ARCH_GFX900__ || __HIP_ARCH_GFX906__
|
||||
|
||||
__global__ void __halfMath(hipLaunchParm lp, __half* A, __half* B, __half* C) {
|
||||
__global__ void __halfMath(__half* A, __half* B, __half* C) {
|
||||
int tx = threadIdx.x;
|
||||
__half a = A[tx];
|
||||
__half b = B[tx];
|
||||
@@ -47,10 +47,10 @@ __global__ void __halfMath(hipLaunchParm lp, __half* A, __half* B, __half* C) {
|
||||
c = __hsub_sat(b, c);
|
||||
c = __hmul(a, c);
|
||||
c = __hmul_sat(b, c);
|
||||
c = hdiv(a, c);
|
||||
c = __hdiv(a, c);
|
||||
}
|
||||
|
||||
__global__ void __half2Math(hipLaunchParm lp, __half2* A, __half2* B, __half2* C) {
|
||||
__global__ void __half2Math(__half2* A, __half2* B, __half2* C) {
|
||||
int tx = threadIdx.x;
|
||||
__half2 a = A[tx];
|
||||
__half2 b = B[tx];
|
||||
@@ -65,12 +65,12 @@ __global__ void __half2Math(hipLaunchParm lp, __half2* A, __half2* B, __half2* C
|
||||
c = __hmul2_sat(b, c);
|
||||
}
|
||||
|
||||
__global__ void kernel_hisnan(hipLaunchParm lp, __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(hipLaunchParm lp, __half* input, int* output) {
|
||||
__global__ void kernel_hisinf(__half* input, int* output) {
|
||||
int tx = threadIdx.x;
|
||||
output[tx] = __hisinf(input[tx]);
|
||||
}
|
||||
@@ -93,7 +93,8 @@ void check_hisnan(int NUM_INPUTS, __half* inputCPU, __half* inputGPU) {
|
||||
hipMalloc((void**)&outputGPU, memsize);
|
||||
|
||||
// launch the kernel
|
||||
hipLaunchKernel(kernel_hisnan, dim3(1), dim3(NUM_INPUTS), 0, 0, inputGPU, outputGPU);
|
||||
hipLaunchKernelGGL(
|
||||
kernel_hisnan, dim3(1), dim3(NUM_INPUTS), 0, 0, inputGPU, outputGPU);
|
||||
|
||||
// copy output from device
|
||||
int* outputCPU = (int*) malloc(memsize);
|
||||
@@ -103,12 +104,18 @@ void check_hisnan(int NUM_INPUTS, __half* inputCPU, __half* inputGPU) {
|
||||
for (int i=0; i<NUM_INPUTS; i++) {
|
||||
if ((2 <= i) && (i <= 5)) { // inputs are nan, output should be true
|
||||
if (outputCPU[i] == 0) {
|
||||
failed("__hisnan() returned false for %f (input idx = %d)\n", inputCPU[i], i);
|
||||
failed(
|
||||
"__hisnan() returned false for %f (input idx = %d)\n",
|
||||
static_cast<float>(inputCPU[i]),
|
||||
i);
|
||||
}
|
||||
}
|
||||
else { // inputs are NOT nan, output should be false
|
||||
if (outputCPU[i] != 0) {
|
||||
failed("__hisnan() returned true for %f (input idx = %d)\n", inputCPU[i], i);
|
||||
failed(
|
||||
"__hisnan() returned true for %f (input idx = %d)\n",
|
||||
static_cast<float>(inputCPU[i]),
|
||||
i);
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -129,7 +136,8 @@ void check_hisinf(int NUM_INPUTS, __half* inputCPU, __half* inputGPU) {
|
||||
hipMalloc((void**)&outputGPU, memsize);
|
||||
|
||||
// launch the kernel
|
||||
hipLaunchKernel(kernel_hisinf, dim3(1), dim3(NUM_INPUTS), 0, 0, inputGPU, outputGPU);
|
||||
hipLaunchKernelGGL(
|
||||
kernel_hisinf, dim3(1), dim3(NUM_INPUTS), 0, 0, inputGPU, outputGPU);
|
||||
|
||||
// copy output from device
|
||||
int* outputCPU = (int*) malloc(memsize);
|
||||
@@ -139,12 +147,18 @@ void check_hisinf(int NUM_INPUTS, __half* inputCPU, __half* inputGPU) {
|
||||
for (int i=0; i<NUM_INPUTS; i++) {
|
||||
if ((0 <= i) && (i <= 1)) { // inputs are inf, output should be true
|
||||
if (outputCPU[i] == 0) {
|
||||
failed("__hisinf() returned false for %f (input idx = %d)\n", inputCPU[i], i);
|
||||
failed(
|
||||
"__hisinf() returned false for %f (input idx = %d)\n",
|
||||
static_cast<float>(inputCPU[i]),
|
||||
i);
|
||||
}
|
||||
}
|
||||
else { // inputs are NOT inf, output should be false
|
||||
if (outputCPU[i] != 0) {
|
||||
failed("__hisinf() returned true for %f (input idx = %d)\n", inputCPU[i], i);
|
||||
failed(
|
||||
"__hisinf() returned true for %f (input idx = %d)\n",
|
||||
static_cast<float>(inputCPU[i]),
|
||||
i);
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -160,11 +174,11 @@ void check_hisinf(int NUM_INPUTS, __half* inputCPU, __half* inputGPU) {
|
||||
|
||||
void checkFunctional() {
|
||||
|
||||
// allocate memory
|
||||
// allocate memory
|
||||
const int NUM_INPUTS = 16;
|
||||
auto memsize = NUM_INPUTS * sizeof(__half);
|
||||
__half* inputCPU = (__half*) malloc(memsize);
|
||||
|
||||
|
||||
// populate inputs
|
||||
inputCPU[0] = host_ushort_as_half(0x7c00); // inf
|
||||
inputCPU[1] = host_ushort_as_half(0xfc00); // -inf
|
||||
@@ -207,7 +221,8 @@ int main() {
|
||||
hipMalloc(&A, HALF_SIZE);
|
||||
hipMalloc(&B, HALF_SIZE);
|
||||
hipMalloc(&C, HALF_SIZE);
|
||||
hipLaunchKernel(__halfMath, dim3(1, 1, 1), dim3(LEN, 1, 1), 0, 0, A, B, C);
|
||||
hipLaunchKernelGGL(
|
||||
__halfMath, dim3(1, 1, 1), dim3(LEN, 1, 1), 0, 0, A, B, C);
|
||||
hipFree(A);
|
||||
hipFree(B);
|
||||
hipFree(C);
|
||||
@@ -215,13 +230,14 @@ int main() {
|
||||
hipMalloc(&A2, HALF2_SIZE);
|
||||
hipMalloc(&B2, HALF2_SIZE);
|
||||
hipMalloc(&C2, HALF2_SIZE);
|
||||
hipLaunchKernel(__half2Math, dim3(1, 1, 1), dim3(LEN, 1, 1), 0, 0, A2, B2, C2);
|
||||
hipLaunchKernelGGL(
|
||||
__half2Math, dim3(1, 1, 1), dim3(LEN, 1, 1), 0, 0, A2, B2, C2);
|
||||
hipFree(A2);
|
||||
hipFree(B2);
|
||||
hipFree(C2);
|
||||
|
||||
// run some functional checks
|
||||
checkFunctional();
|
||||
|
||||
|
||||
passed();
|
||||
}
|
||||
|
||||
@@ -0,0 +1,142 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
|
||||
* RUN: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include <hip/hip_fp16.h>
|
||||
#include "hip/hip_runtime.h"
|
||||
|
||||
#include "test_common.h"
|
||||
|
||||
#if __HIP_ARCH_GFX803__ || __HIP_ARCH_GFX900__ || __HIP_ARCH_GFX906__
|
||||
|
||||
__global__
|
||||
__attribute__((optnone))
|
||||
void __halfTest(bool* result) {
|
||||
// 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<unsigned short>(1)};
|
||||
result[0] = __heq(a, 1) && result[0];
|
||||
a = __half{static_cast<short>(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];
|
||||
|
||||
// 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];
|
||||
|
||||
// Nullary
|
||||
result[0] = __heq(a, +a) && result[0];
|
||||
result[0] = __heq(__hneg(a), -a) && result[0];
|
||||
|
||||
// Unary arithmetic
|
||||
result[0] = __heq(a += 0, a) && result[0];
|
||||
result[0] = __heq(a -= 0, a) && result[0];
|
||||
result[0] = __heq(a *= 1, a) && result[0];
|
||||
result[0] = __heq(a /= 1, a) && result[0];
|
||||
|
||||
// Binary arithmetic
|
||||
result[0] = __heq((a + a), __hadd(a, a)) && result[0];
|
||||
result[0] = __heq((a - a), __hsub(a, a)) && result[0];
|
||||
result[0] = __heq((a * a), __hmul(a, a)) && result[0];
|
||||
result[0] = __heq((a / a), __hdiv(a, a)) && result[0];
|
||||
|
||||
// Relations
|
||||
result[0] = (a == a) && result[0];
|
||||
result[0] = !(a != a) && result[0];
|
||||
result[0] = (a <= a) && result[0];
|
||||
result[0] = (a >= a) && result[0];
|
||||
result[0] = !(a < a) && result[0];
|
||||
result[0] = !(a > a) && result[0];
|
||||
}
|
||||
|
||||
__device__
|
||||
bool to_bool(const __half2& x)
|
||||
{
|
||||
auto r = static_cast<const __half2_raw&>(x);
|
||||
|
||||
return r.data.x != 0 && r.data.y != 0;
|
||||
}
|
||||
__global__
|
||||
__attribute__((optnone))
|
||||
void __half2Test(bool* result) {
|
||||
// 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];
|
||||
|
||||
// Assignment
|
||||
a = __half2_raw{2}; result[0] = to_bool(__heq2(a, {2, 2})) && result[0];
|
||||
|
||||
// Nullary
|
||||
result[0] = to_bool(__heq2(a, +a)) && result[0];
|
||||
result[0] = to_bool(__heq2(__hneg2(a), -a)) && result[0];
|
||||
|
||||
// Unary arithmetic
|
||||
result[0] = to_bool(__heq2(a += 0, a)) && result[0];
|
||||
result[0] = to_bool(__heq2(a -= 0, a)) && result[0];
|
||||
result[0] = to_bool(__heq2(a *= 1, a)) && result[0];
|
||||
result[0] = to_bool(__heq2(a /= 1, a)) && result[0];
|
||||
|
||||
// Binary arithmetic
|
||||
result[0] = to_bool(__heq2((a + a), __hadd2(a, a))) && result[0];
|
||||
result[0] = to_bool(__heq2((a - a), __hsub2(a, a))) && result[0];
|
||||
result[0] = to_bool(__heq2((a * a), __hmul2(a, a))) && result[0];
|
||||
result[0] = to_bool(__heq2((a / a), __h2div(a, a))) && result[0];
|
||||
|
||||
// Relations
|
||||
result[0] = (a == a) && result[0];
|
||||
result[0] = !(a != a) && result[0];
|
||||
result[0] = (a <= a) && result[0];
|
||||
result[0] = (a >= a) && result[0];
|
||||
result[0] = !(a < a) && result[0];
|
||||
result[0] = !(a > a) && result[0];
|
||||
}
|
||||
|
||||
#endif
|
||||
|
||||
int main() {
|
||||
bool* result{nullptr};
|
||||
hipHostMalloc(&result, 1);
|
||||
|
||||
result[0] = false;
|
||||
hipLaunchKernelGGL(__halfTest, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, result);
|
||||
hipDeviceSynchronize();
|
||||
|
||||
if (!result[0]) { failed("Failed __half tests."); }
|
||||
|
||||
result[0] = false;
|
||||
hipLaunchKernelGGL(__half2Test, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, result);
|
||||
hipDeviceSynchronize();
|
||||
|
||||
if (!result[0]) { failed("Failed __half2 tests."); }
|
||||
|
||||
passed();
|
||||
}
|
||||
مرجع در شماره جدید
Block a user