diff --git a/projects/hip-tests/catch/unit/deviceLib/CMakeLists.txt b/projects/hip-tests/catch/unit/deviceLib/CMakeLists.txt index 2ee700545c..c54cd68c08 100644 --- a/projects/hip-tests/catch/unit/deviceLib/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/deviceLib/CMakeLists.txt @@ -27,6 +27,7 @@ set(AMD_TEST_SRC bitExtract.cc bitInsert.cc floatTM.cc + hipMathFunctions.cc ) set(AMD_ARCH_SPEC_TEST_SRC AtomicAdd_Coherent_withunsafeflag.cc @@ -76,6 +77,7 @@ if(${ARCH_CHECK} GREATER_EQUAL 0) set_source_files_properties(unsafeAtomicAdd_NonCoherent_withunsafeflag.cc PROPERTIES COMPILE_OPTIONS "-munsafe-fp-atomics") set_source_files_properties(unsafeAtomicAdd_Coherent_withnounsafeflag.cc PROPERTIES COMPILE_OPTIONS "-mno-unsafe-fp-atomics") set_source_files_properties(unsafeAtomicAdd_NonCoherent_withnounsafeflag.cc PROPERTIES COMPILE_OPTIONS "-mno-unsafe-fp-atomics") + set_source_files_properties(hipMathFunctions.cc PROPERTIES COMPILE_FLAGS "-Xclang -fallow-half-arguments-and-returns") file(GLOB AtomicAdd_files *AtomicAdd_*_*.cc) set_property(SOURCE ${AtomicAdd_files} PROPERTY COMPILE_FLAGS --save-temps) file(GLOB unsafeAtomicAdd_files *unsafeAtomicAdd_*_*.cc) diff --git a/projects/hip-tests/catch/unit/deviceLib/hipMathFunctions.cc b/projects/hip-tests/catch/unit/deviceLib/hipMathFunctions.cc new file mode 100644 index 0000000000..543bacc647 --- /dev/null +++ b/projects/hip-tests/catch/unit/deviceLib/hipMathFunctions.cc @@ -0,0 +1,136 @@ +/* +Copyright (c) 2022 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. +*/ + +/* + Test Scenarios : + 1) Verification of absolute int64 operation performed at device. + 2) Verification of __fp16 operation performed at device. + 3) Verification of pow operations performed at device. +*/ + +#include + +__global__ void kernel_abs_int64(long long *input, long long *output) { // NOLINT + int tx = threadIdx.x; + output[tx] = abs(input[tx]); +} + + +#define CHECK_ABS_INT64(IN, OUT, EXP) \ + { \ + if (OUT != EXP) { \ + INFO("check_abs_int64 failed on " << IN << ", output " << OUT << \ + ", expected " << EXP); \ + REQUIRE(false); \ + } \ + } + +template +__global__ void kernel_simple(F f, T *out) { + *out = f(); +} + +template +void check_simple(F f, T expected, const char* file, unsigned line) { + auto memsize = sizeof(T); + T *outputCPU = reinterpret_cast(malloc(memsize)); + T *outputGPU = nullptr; + REQUIRE(outputCPU != nullptr); + HIP_CHECK(hipMalloc(&outputGPU, memsize)); + hipLaunchKernelGGL(kernel_simple, 1, 1, 0, 0, f, outputGPU); + HIP_CHECK(hipMemcpy(outputCPU, outputGPU, memsize, hipMemcpyDeviceToHost)); + if (*outputCPU != expected) { + INFO("File " << file << ", line " << line << " check failed." << + " output = " << static_cast(*outputCPU) << " expected " + << static_cast(expected)); + REQUIRE(false); + } + HIP_CHECK(hipFree(outputGPU)); + free(outputCPU); +} + +#define CHECK_SIMPLE(lambda, expected) \ + check_simple(lambda, expected, __FILE__, __LINE__); + + +/** + Verification of absolute int64 operation performed at device. + */ +TEST_CASE("Unit_abs_int64_Verification") { + using datatype_t = long long; // NOLINT + + datatype_t *inputCPU{}, *outputCPU{}; + datatype_t *inputGPU{}, *outputGPU{}; + const int NUM_INPUTS = 8; + auto memsize = NUM_INPUTS * sizeof(datatype_t); + + // allocate memories + inputCPU = reinterpret_cast(malloc(memsize)); + outputCPU = reinterpret_cast(malloc(memsize)); + REQUIRE(inputCPU != nullptr); + REQUIRE(outputCPU != nullptr); + HIP_CHECK(hipMalloc(&inputGPU, memsize)); + HIP_CHECK(hipMalloc(&outputGPU, memsize)); + + // populate input with constants + inputCPU[0] = -81985529216486895ll; + inputCPU[1] = 81985529216486895ll; + inputCPU[2] = -1250999896491ll; + inputCPU[3] = 1250999896491ll; + inputCPU[4] = -19088743ll; + inputCPU[5] = 19088743ll; + inputCPU[6] = -291ll; + inputCPU[7] = 291ll; + + // copy inputs to device + HIP_CHECK(hipMemcpy(inputGPU, inputCPU, memsize, hipMemcpyHostToDevice)); + + // launch kernel + hipLaunchKernelGGL(kernel_abs_int64, dim3(1), dim3(NUM_INPUTS), 0, 0, + inputGPU, outputGPU); + // copy outputs from device + HIP_CHECK(hipMemcpy(outputCPU, outputGPU, memsize, hipMemcpyDeviceToHost)); + + // check outputs + CHECK_ABS_INT64(inputCPU[0], outputCPU[0], outputCPU[1]); + CHECK_ABS_INT64(inputCPU[1], outputCPU[1], outputCPU[1]); + CHECK_ABS_INT64(inputCPU[2], outputCPU[2], outputCPU[3]); + CHECK_ABS_INT64(inputCPU[3], outputCPU[3], outputCPU[3]); + CHECK_ABS_INT64(inputCPU[4], outputCPU[4], outputCPU[5]); + CHECK_ABS_INT64(inputCPU[5], outputCPU[5], outputCPU[5]); + CHECK_ABS_INT64(inputCPU[6], outputCPU[6], outputCPU[7]); + CHECK_ABS_INT64(inputCPU[7], outputCPU[7], outputCPU[7]); + + // free memories + HIP_CHECK(hipFree(inputGPU)); + HIP_CHECK(hipFree(outputGPU)); + free(inputCPU); + free(outputCPU); +} + +/** + Verification of pow operations performed at device. + */ +TEST_CASE("Unit_pown_Verification") { + CHECK_SIMPLE([]__device__(){ return powif(2.0f, 2); }, 4.0f); + CHECK_SIMPLE([]__device__(){ return powi(2.0, 2); }, 4.0); + CHECK_SIMPLE([]__device__(){ return pow(2.0f, 2); }, 4.0f); + CHECK_SIMPLE([]__device__(){ return pow(2.0, 2); }, 4.0); + CHECK_SIMPLE([]__device__(){ return pow(2.0f16, 2); }, 4.0f16); +}