From 6a09344dba91a1a9816cb6bcdcc6d8bc6ea564c3 Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary <51944368+cjatin@users.noreply.github.com> Date: Mon, 6 Sep 2021 12:21:38 +0530 Subject: [PATCH] SWDEV-296918 - Adding unsafe atomics test for gfx90a (#2331) Change-Id: I3936e0b30e15aff154bdab3efe3f6da2ccce3578 --- tests/catch/unit/deviceLib/CMakeLists.txt | 1 + tests/catch/unit/deviceLib/unsafeAtomicAdd.cc | 133 ++++++++++++++++++ 2 files changed, 134 insertions(+) create mode 100644 tests/catch/unit/deviceLib/unsafeAtomicAdd.cc diff --git a/tests/catch/unit/deviceLib/CMakeLists.txt b/tests/catch/unit/deviceLib/CMakeLists.txt index ce11801786..191553ca99 100644 --- a/tests/catch/unit/deviceLib/CMakeLists.txt +++ b/tests/catch/unit/deviceLib/CMakeLists.txt @@ -17,6 +17,7 @@ set(TEST_SRC # AMD only tests set(AMD_TEST_SRC + unsafeAtomicAdd.cc vectorTypesDevice.cc mbcnt.cc bitExtract.cc diff --git a/tests/catch/unit/deviceLib/unsafeAtomicAdd.cc b/tests/catch/unit/deviceLib/unsafeAtomicAdd.cc new file mode 100644 index 0000000000..24e1329c7a --- /dev/null +++ b/tests/catch/unit/deviceLib/unsafeAtomicAdd.cc @@ -0,0 +1,133 @@ +/* +Copyright (c) 2021 - 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. +*/ + +#include + +#include +#include + +#include + +static constexpr auto kernel{ + R"( +extern "C" +__global__ +void unsafeAdd_f(float *p, float v) +{ + auto val = unsafeAtomicAdd(p, v); +} + +extern "C" +__global__ +void unsafeAdd_d(double *p, double v) +{ + auto val = unsafeAtomicAdd(p, v); +} +)"}; + + +TEST_CASE("Unit_unsafeAtomicAdd") { + using namespace std; + int device = 0; + hipDeviceProp_t props; + HIP_CHECK(hipGetDeviceProperties(&props, device)); + std::string gfxName(props.gcnArchName); + + if (gfxName == "gfx90a" || gfxName.find("gfx90a:") == 0) { + hiprtcProgram prog; + hiprtcCreateProgram(&prog, // prog + kernel, // buffer + "kernel.cu", // name + 0, nullptr, nullptr); + std::string sarg = std::string("--gpu-architecture=") + props.gcnArchName; + const char* options[] = {sarg.c_str()}; + hiprtcResult compileResult{hiprtcCompileProgram(prog, 1, options)}; + + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + INFO(log); + } + + REQUIRE(compileResult == HIPRTC_SUCCESS); + size_t codeSize; + HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize)); + + vector code(codeSize); + HIPRTC_CHECK(hiprtcGetCode(prog, code.data())); + + HIPRTC_CHECK(hiprtcDestroyProgram(&prog)); + + float* fX; + double* dX; + HIP_CHECK(hipMalloc(&fX, sizeof(float))); + HIP_CHECK(hipMalloc(&dX, sizeof(double))); + + hipModule_t module; + hipFunction_t f_kernel, d_kernel; + HIP_CHECK(hipModuleLoadData(&module, code.data())); + HIP_CHECK(hipModuleGetFunction(&f_kernel, module, "unsafeAdd_f")); + HIP_CHECK(hipModuleGetFunction(&d_kernel, module, "unsafeAdd_d")); + + float f_val = 10.1f; + double d_val = 10.1; + + HIP_CHECK(hipMemcpy(fX, &f_val, sizeof(float), hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(dX, &d_val, sizeof(double), hipMemcpyHostToDevice)); + + struct { + float* p; + float val; + } args_f{fX, f_val}; + + struct { + double* p; + double val; + } args_d{dX, d_val}; + + auto size_f = sizeof(args_f); + void* config_f[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f, HIP_LAUNCH_PARAM_BUFFER_SIZE, + &size_f, HIP_LAUNCH_PARAM_END}; + + auto size_d = sizeof(args_d); + void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_d, HIP_LAUNCH_PARAM_BUFFER_SIZE, + &size_d, HIP_LAUNCH_PARAM_END}; + + hipModuleLaunchKernel(f_kernel, 10, 1, 1, 100, 1, 1, 0, nullptr, nullptr, config_f); + hipModuleLaunchKernel(d_kernel, 10, 1, 1, 100, 1, 1, 0, nullptr, nullptr, config_d); + + float res_f = 0.0f; + double res_d = 0.0; + HIP_CHECK(hipMemcpy(&res_f, fX, sizeof(float), hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(&res_d, dX, sizeof(double), hipMemcpyDeviceToHost)); + + hipFree(dX); + hipFree(fX); + + HIP_CHECK(hipModuleUnload(module)); + + REQUIRE(fabs(res_f - (1000 * f_val)) <= 0.2f); + REQUIRE(fabs(res_d - (1000 * d_val)) <= 0.2); + } +}