From 2d4040f75ddfd85329b11e5de31eda757f30eeb6 Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Fri, 10 Aug 2018 11:12:53 -0400 Subject: [PATCH] Add fma function with float and _Float16 arguments [ROCm/hip commit: bd622a4b4ab4679c8d902f797b3cd60b95f13d9c] --- .../include/hip/hcc_detail/math_functions.h | 10 ++ .../hip/tests/src/deviceLib/hipTestFMA.cpp | 142 ++++++++++++++++++ 2 files changed, 152 insertions(+) create mode 100644 projects/hip/tests/src/deviceLib/hipTestFMA.cpp diff --git a/projects/hip/include/hip/hcc_detail/math_functions.h b/projects/hip/include/hip/hcc_detail/math_functions.h index 702c120b86..c1adef68fd 100644 --- a/projects/hip/include/hip/hcc_detail/math_functions.h +++ b/projects/hip/include/hip/hcc_detail/math_functions.h @@ -1166,6 +1166,16 @@ long long llabs(long long x) #endif // END INTEGER +__DEVICE__ +inline _Float16 fma(_Float16 x, _Float16 y, _Float16 z) { + return __ocml_fma_f16(x, y, z); +} + +__DEVICE__ +inline float fma(float x, float y, float z) { + return fmaf(x, y, z); +} + #pragma push_macro("__DEF_FLOAT_FUN") #pragma push_macro("__DEF_FLOAT_FUN2") #pragma push_macro("__DEF_FLOAT_FUN2I") diff --git a/projects/hip/tests/src/deviceLib/hipTestFMA.cpp b/projects/hip/tests/src/deviceLib/hipTestFMA.cpp new file mode 100644 index 0000000000..5e1913a5c7 --- /dev/null +++ b/projects/hip/tests/src/deviceLib/hipTestFMA.cpp @@ -0,0 +1,142 @@ +/* +Copyright (c) 2015-2016 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 + * RUN: %t + * HIT_END + */ + +#include "test_common.h" +#include +#include +#include + +#define HIP_ASSERT(status) assert(status == hipSuccess) + +#define LEN 50 +#define SIZE (LEN * sizeof(bool)) + +struct TestFMA { + static __global__ void kernel(bool *Ad) { + float f = 1.0f / 3.0f; + double d = f; + int i = 0; + auto Check = [&](bool Cond) { Ad[i++] = Cond; }; + // f * f + 3.0f will be different if promoted to double. + float floatResult = fma(f, f, 3.0f); + double doubleResult = fma(d, d, 3.0); + Check(floatResult != doubleResult); + + // check promote to float. + Check(fma(f, f, 3) == floatResult); + Check(fma(f, f, (char)3) == floatResult); + Check(fma(f, f, (unsigned char)3) == floatResult); + Check(fma(f, f, (short)3) == floatResult); + Check(fma(f, f, (unsigned short)3) == floatResult); + Check(fma(f, f, (int)3) == floatResult); + Check(fma(f, f, (unsigned int)3) == floatResult); + Check(fma(f, f, (long)3) == floatResult); + Check(fma(f, f, (unsigned long)3) == floatResult); + Check(fma(f, f, true) == fma(f, f, 1.0f)); + + // check promote to double. + Check(fma(d, (double)f, 3) == doubleResult); + Check(fma(d, (double)f, (char)3) == doubleResult); + Check(fma(d, (double)f, (unsigned char)3) == doubleResult); + Check(fma(d, (double)f, (short)3) == doubleResult); + Check(fma(d, (double)f, (unsigned short)3) == doubleResult); + Check(fma(d, (double)f, (int)3) == doubleResult); + Check(fma(d, (double)f, (unsigned int)3) == doubleResult); + Check(fma(d, (double)f, (long)3) == doubleResult); + Check(fma(d, (double)f, (unsigned long)3) == doubleResult); + Check(fma(d, (double)f, true) == fma((double)f, (double)f, 1.0)); + + while (i < LEN) + Check(true); + } + void run() { + bool *Ad; + bool A[LEN]; + for (unsigned i = 0; i < LEN; i++) { + A[i] = 0; + } + + HIP_ASSERT(hipMalloc((void **)&Ad, SIZE)); + hipLaunchKernelGGL(kernel, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, Ad); + HIP_ASSERT(hipMemcpy(A, Ad, SIZE, hipMemcpyDeviceToHost)); + + for (unsigned i = 0; i < LEN; i++) { + assert(A[i]); + } + } +}; + +struct TestHalfFMA { + static __global__ void kernel(bool *Ad) { + _Float16 h = (_Float16)(1.0f/3.0f); + float f = h; + double d = f; + int i = 0; + auto Check = [&](bool Cond) { Ad[i++] = Cond; }; + // h * h + 3 will be different if promoted to float. + _Float16 halfResult = fma(h, h, (_Float16)3); + float floatResult = fma(f, f, 3.0f); + double doubleResult = fma(d, d, 3.0); + Check(halfResult != floatResult); + Check(halfResult != doubleResult); + + // check promote to half. + Check(fma(h, h, 3) == halfResult); + Check(fma(h, h, (char)3) == halfResult); + Check(fma(h, h, (unsigned char)3) == halfResult); + Check(fma(h, h, (short)3) == halfResult); + Check(fma(h, h, (unsigned short)3) == halfResult); + Check(fma(h, h, (int)3) == halfResult); + Check(fma(h, h, (unsigned int)3) == halfResult); + Check(fma(h, h, (long)3) == halfResult); + Check(fma(h, h, (unsigned long)3) == halfResult); + Check(fma(h, h, true) == fma(h, h, (_Float16)1)); + + while (i < LEN) + Check(true); + } + + void run() { + bool *Ad; + bool A[LEN]; + for (unsigned i = 0; i < LEN; i++) { + A[i] = 0; + } + + HIP_ASSERT(hipMalloc((void **)&Ad, SIZE)); + hipLaunchKernelGGL(kernel, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0, Ad); + HIP_ASSERT(hipMemcpy(A, Ad, SIZE, hipMemcpyDeviceToHost)); + + for (unsigned i = 0; i < LEN; i++) { + assert(A[i]); + } + } +}; + +int main() { + TestFMA().run(); + TestHalfFMA().run(); + passed(); +}