From bae9739c15c00486322e3c48aae65d64a6908558 Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary Date: Fri, 14 Jul 2023 16:08:47 +0100 Subject: [PATCH] SWDEV-411156 - Add basic tests for fp16 Recreation of github PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/384 Change-Id: Ie4f00ebd198c8689d34757a3c27aaddb85d15370 [ROCm/hip-tests commit: 51c5b7b612246377a9dab3c1cd400be616e00993] --- .../catch/unit/deviceLib/CMakeLists.txt | 2 + .../catch/unit/deviceLib/fp16_ops.cc | 188 ++++++++++++++++++ 2 files changed, 190 insertions(+) create mode 100644 projects/hip-tests/catch/unit/deviceLib/fp16_ops.cc diff --git a/projects/hip-tests/catch/unit/deviceLib/CMakeLists.txt b/projects/hip-tests/catch/unit/deviceLib/CMakeLists.txt index 67dd75e094..0645defeb6 100644 --- a/projects/hip-tests/catch/unit/deviceLib/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/deviceLib/CMakeLists.txt @@ -87,7 +87,9 @@ set(AMD_TEST_SRC hip_test_make_type.cc bfloat16.cc AtomicsWithRandomActiveLanesInWavefront.cc + fp16_ops.cc ) + set(AMD_ARCH_SPEC_TEST_SRC AtomicAdd_Coherent_withunsafeflag.cc AtomicAdd_Coherent_withoutflag.cc diff --git a/projects/hip-tests/catch/unit/deviceLib/fp16_ops.cc b/projects/hip-tests/catch/unit/deviceLib/fp16_ops.cc new file mode 100644 index 0000000000..841e17f1b4 --- /dev/null +++ b/projects/hip-tests/catch/unit/deviceLib/fp16_ops.cc @@ -0,0 +1,188 @@ +/* +Copyright (c) 2023 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. +*/ + +#include + +#include +#include +#include +#include + +__global__ void fp16_arith_gpu(float* a, float* b, float* c) { + c[0] = __half2float(__hadd(__float2half_rn(a[0]), __float2half_rn(b[0]))); + c[1] = __half2float(__hsub(__float2half_rn(a[1]), __float2half_rn(b[1]))); + c[2] = __half2float(__hmul(__float2half_rn(a[2]), __float2half_rn(b[2]))); + c[3] = __half2float(__hdiv(__float2half_rn(a[3]), __float2half_rn(b[3]))); + c[4] = __half2float(hfloor(__float2half_rn(a[4]))); + c[5] = __half2float(htrunc(__float2half_rn(a[5]))); + c[6] = __half2float(hceil(__float2half_rn(a[6]))); + c[7] = __half2float(hrint(__float2half_rn(a[7]))); + c[8] = __half2float(hsin(__float2half_rn(a[8]))); + c[9] = __half2float(hcos(__float2half_rn(a[9]))); + c[10] = __half2float(hexp(__float2half_rn(a[10]))); + c[11] = __half2float(hexp2(__float2half_rn(a[11]))); + c[12] = __half2float(hlog2(__float2half_rn(a[12]))); + c[13] = __half2float(hlog(__float2half_rn(a[13]))); + c[14] = __half2float(hlog10(__float2half_rn(a[14]))); + c[15] = __half2float(hsqrt(__float2half_rn(a[15]))); + c[16] = __half2float(__hneg(__float2half_rn(a[16]))); + c[17] = __half2float(hrcp(__float2half_rn(a[17]))); +} + +void fp16_arith_cpu(const std::vector& a, const std::vector& b, + std::vector& c) { + c[0] = a[0] + b[0]; + c[1] = a[0] - b[0]; + c[2] = a[0] * b[0]; + c[3] = a[0] / b[0]; + c[4] = std::floorf(a[4]); + c[5] = std::truncf(a[5]); + c[6] = std::ceilf(a[6]); + c[7] = std::rintf(a[7]); + c[8] = std::sinf(a[8]); + c[9] = std::cosf(a[9]); + c[10] = std::expf(a[10]); + c[11] = std::exp2f(a[11]); + c[12] = std::log2f(a[12]); + c[13] = std::logf(a[13]); + c[14] = std::log10f(a[14]); + c[15] = std::sqrtf(a[15]); + c[16] = -a[16]; + c[17] = 1.0f / a[17]; +} + +TEST_CASE("Unit_fp16_arith") { + constexpr size_t num_of_ops = 18; + constexpr size_t iters = 100; + Catch::Generators::RandomFloatingGenerator input1_gen(2.2f, 10.f); + constexpr float input2 = 1.1f; + for (size_t iter = 0; iter < iters; iter++) { + auto input1 = input1_gen.get(); + + std::vector in1(num_of_ops, input1); + std::vector in2(num_of_ops, input2); + float *din1, *din2, *dout; + HIP_CHECK(hipMalloc(&dout, sizeof(float) * num_of_ops)); + HIP_CHECK(hipMalloc(&din1, sizeof(float) * num_of_ops)); + HIP_CHECK(hipMalloc(&din2, sizeof(float) * num_of_ops)); + + HIP_CHECK(hipMemcpy(din1, in1.data(), sizeof(float) * in1.size(), hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(din2, in2.data(), sizeof(float) * in2.size(), hipMemcpyHostToDevice)); + + fp16_arith_gpu<<<1, 1>>>(din1, din2, dout); + std::vector cpuout(num_of_ops, 0.0f); + fp16_arith_cpu(in1, in2, cpuout); + + std::vector out(num_of_ops, 0.0f); + HIP_CHECK(hipMemcpy(out.data(), dout, sizeof(float) * out.size(), hipMemcpyDeviceToHost)); + + for (size_t i = 0; i < out.size(); i++) { + INFO("Iter: " << i << " In1: " << in1[i] << " CPU res: " << cpuout[i] + << " GPU res: " << out[i]); + REQUIRE(out[i] == Approx(cpuout[i]).epsilon(0.1)); + } + HIP_CHECK(hipFree(dout)); + HIP_CHECK(hipFree(din1)); + HIP_CHECK(hipFree(din2)); + } +} + +__device__ void fp162_arith_impl(float2* a, float2* b, float2* c) { + c[0] = __half22float2(__hadd2(__float22half2_rn(a[0]), __float22half2_rn(b[0]))); + c[1] = __half22float2(__hsub2(__float22half2_rn(a[1]), __float22half2_rn(b[1]))); + c[2] = __half22float2(__hmul2(__float22half2_rn(a[2]), __float22half2_rn(b[2]))); + c[3] = __half22float2(__h2div(__float22half2_rn(a[3]), __float22half2_rn(b[3]))); + c[4] = __half22float2(h2floor(__float22half2_rn(a[4]))); + c[5] = __half22float2(h2trunc(__float22half2_rn(a[5]))); + c[6] = __half22float2(h2ceil(__float22half2_rn(a[6]))); + c[7] = __half22float2(h2rint(__float22half2_rn(a[7]))); + c[8] = __half22float2(h2sin(__float22half2_rn(a[8]))); + c[9] = __half22float2(h2cos(__float22half2_rn(a[9]))); + c[10] = __half22float2(h2exp(__float22half2_rn(a[10]))); + c[11] = __half22float2(h2exp2(__float22half2_rn(a[11]))); + c[12] = __half22float2(h2log2(__float22half2_rn(a[12]))); + c[13] = __half22float2(h2log(__float22half2_rn(a[13]))); + c[14] = __half22float2(h2log10(__float22half2_rn(a[14]))); + c[15] = __half22float2(h2sqrt(__float22half2_rn(a[15]))); + c[16] = __half22float2(__hneg2(__float22half2_rn(a[16]))); + c[17] = __half22float2(h2rcp(__float22half2_rn(a[17]))); +} + +__global__ void fp162_arith_gpu(float2* a, float2* b, float2* c) { fp162_arith_impl(a, b, c); } + +void fp162_arith_cpu(std::vector& a, std::vector& b, std::vector& c) { + c[0] = a[0] + b[0]; + c[1] = a[1] - b[1]; + c[2] = a[2] * b[2]; + c[3] = a[3] / b[3]; + c[4] = {std::floorf(a[4].x), std::floorf(a[4].y)}; + c[5] = {std::truncf(a[5].x), std::truncf(a[5].y)}; + c[6] = {std::ceilf(a[6].x), std::ceilf(a[6].y)}; + c[7] = {std::rintf(a[7].x), std::rintf(a[7].y)}; + c[8] = {std::sinf(a[8].x), std::sinf(a[8].y)}; + c[9] = {std::cosf(a[9].x), std::cosf(a[9].y)}; + c[10] = {std::expf(a[10].x), std::expf(a[10].y)}; + c[11] = {std::exp2f(a[11].x), std::exp2f(a[11].y)}; + c[12] = {std::log2f(a[12].x), std::log2f(a[12].y)}; + c[13] = {std::logf(a[13].x), std::logf(a[13].y)}; + c[14] = {std::log10f(a[14].x), std::log10f(a[14].y)}; + c[15] = {std::sqrtf(a[15].x), std::sqrtf(a[15].y)}; + c[16] = {-a[16].x, -a[16].y}; + c[17] = {1.0f / a[17].x, 1.0f / a[17].y}; +} + +TEST_CASE("Unit_fp162_arith") { + constexpr size_t num_of_ops = 18; + constexpr size_t iters = 100; + Catch::Generators::RandomFloatingGenerator input1_gen(2.2f, 10.f); + for (size_t iter = 0; iter < iters; iter++) { + auto input1 = input1_gen.get(); + auto input2 = input1_gen.get(); + + std::vector in1(num_of_ops, float2{input1, input2}); + std::vector in2(num_of_ops, float2{input1_gen.get(), input1_gen.get()}); + float2 *din1, *din2, *dout; + HIP_CHECK(hipMalloc(&dout, sizeof(float2) * num_of_ops)); + HIP_CHECK(hipMalloc(&din1, sizeof(float2) * num_of_ops)); + HIP_CHECK(hipMalloc(&din2, sizeof(float2) * num_of_ops)); + + HIP_CHECK(hipMemcpy(din1, in1.data(), sizeof(float2) * in1.size(), hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(din2, in2.data(), sizeof(float2) * in2.size(), hipMemcpyHostToDevice)); + + fp162_arith_gpu<<<1, 1>>>(din1, din2, dout); + std::vector cpuout(num_of_ops, float2{0.0f, 0.0f}); + fp162_arith_cpu(in1, in2, cpuout); + + std::vector out(num_of_ops, float2{0.0f, 0.0f}); + HIP_CHECK(hipMemcpy(out.data(), dout, sizeof(float2) * out.size(), hipMemcpyDeviceToHost)); + + for (size_t i = 0; i < out.size(); i++) { + INFO("Iter: " << i << " In1: " << in1[i].x << " - " << in1[i].y << " CPU res: " << cpuout[i].x + << " - " << cpuout[i].y << " GPU res: " << out[i].x << " - " << out[i].y); + REQUIRE(out[i].x == Approx(cpuout[i].x).epsilon(0.1)); + REQUIRE(out[i].y == Approx(cpuout[i].y).epsilon(0.1)); + } + HIP_CHECK(hipFree(dout)); + HIP_CHECK(hipFree(din1)); + HIP_CHECK(hipFree(din2)); + } +}