From 3dae157ab51141b64bc4b06ce22ad4efa87018e6 Mon Sep 17 00:00:00 2001 From: ROCm CI Service Account <66695075+rocm-ci@users.noreply.github.com> Date: Thu, 31 Mar 2022 21:33:17 +0530 Subject: [PATCH] SWDEV-293806 - Added tests for unsafeAtomicAdd and builtin APIs (#2597) Added tests for unsafeAtomicAdd and Builtin APIs for coherent/non-coherent memory with RTC and without RTC Change-Id: I8b243ac82d0f14a38f6b5553107066248aafb41d [ROCm/hip-tests commit: 34bcf0398fac02664ef069d817c7619ae9ac99e6] --- .../catch/unit/deviceLib/BuiltIns_fadd.cc | 292 +++++++++ .../catch/unit/deviceLib/BuiltIns_fmax.cc | 354 +++++++++++ .../catch/unit/deviceLib/BuiltIns_fmin.cc | 357 +++++++++++ .../catch/unit/deviceLib/CMakeLists.txt | 16 + ...safeAtomicAdd_Coherent_withnounsafeflag.cc | 99 +++ .../unsafeAtomicAdd_Coherent_withoutflag.cc | 99 +++ ...unsafeAtomicAdd_Coherent_withunsafeflag.cc | 100 +++ ...eAtomicAdd_NonCoherent_withnounsafeflag.cc | 97 +++ ...unsafeAtomicAdd_NonCoherent_withoutflag.cc | 97 +++ ...afeAtomicAdd_NonCoherent_withunsafeflag.cc | 97 +++ .../unit/deviceLib/unsafeAtomicAdd_RTC.cc | 583 ++++++++++++++++++ 11 files changed, 2191 insertions(+) create mode 100644 projects/hip-tests/catch/unit/deviceLib/BuiltIns_fadd.cc create mode 100644 projects/hip-tests/catch/unit/deviceLib/BuiltIns_fmax.cc create mode 100644 projects/hip-tests/catch/unit/deviceLib/BuiltIns_fmin.cc create mode 100644 projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withnounsafeflag.cc create mode 100644 projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withoutflag.cc create mode 100644 projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withunsafeflag.cc create mode 100644 projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_NonCoherent_withnounsafeflag.cc create mode 100644 projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_NonCoherent_withoutflag.cc create mode 100644 projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_NonCoherent_withunsafeflag.cc create mode 100644 projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_RTC.cc diff --git a/projects/hip-tests/catch/unit/deviceLib/BuiltIns_fadd.cc b/projects/hip-tests/catch/unit/deviceLib/BuiltIns_fadd.cc new file mode 100644 index 0000000000..472e89491c --- /dev/null +++ b/projects/hip-tests/catch/unit/deviceLib/BuiltIns_fadd.cc @@ -0,0 +1,292 @@ +/* + Copyright (c) 2021 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. + */ +/* +This testfile verifies __builtin_amdgcn_global_atomic_fadd_f64 API scenarios +1. AtomicAdd on Coherent Memory +2. AtomicAdd on Non-Coherent Memory +3. AtomicAdd on Coherent Memory with RTC +4. AtomicAdd on Non-Coherent Memory with RTC +*/ + +#include +#include +#include + +#define INC_VAL 10 +#define INITIAL_VAL 5 +__global__ void AtomicAdd_GlobalMem(double* addr, double* result) { + double inc_val = 10; + *result = __builtin_amdgcn_global_atomic_fadd_f64(addr, inc_val); +} +static constexpr auto AtomicAddGlobalMem{ +R"( +extern "C" +__global__ void AtomicAdd_GlobalMem(double* addr, double* result) { + double inc_val = 10; + *result = __builtin_amdgcn_global_atomic_fadd_f64(addr, inc_val); +} +)"}; +/* +This test verifies the built in atomic add API on Coherent Memory +Input: A_h with INITIAL_VAL +Output: A_h will not get updated with Coherent Memory + A_h will be INITIAL_VAL + ret value would be 0, B_h would be 0 +*/ +TEST_CASE("Unit_BuiltInAtomicAdd_CoherentGlobalMem") { + hipDeviceProp_t prop; + int device; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + std::string gfxName(prop.gcnArchName); + if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) { + if (prop.canMapHostMemory != 1) { + SUCCEED("Does support HostPinned Memory"); + } else { + double *A_h, *result_h, *result; + double *A_d; + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_h), sizeof(double), + hipHostMallocCoherent)); + HIP_CHECK(hipHostMalloc(reinterpret_cast(&result_h), + sizeof(double), hipHostMallocCoherent)); + A_h[0] = INITIAL_VAL; + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&A_d), + A_h, 0)); + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&result), + result_h, 0)); + std::cout << "test" << std::endl; + hipLaunchKernelGGL(AtomicAdd_GlobalMem, dim3(1), dim3(1), + 0, 0, A_d, + result); + std::cout << "test 1" << std::endl; + HIP_CHECK(hipDeviceSynchronize()); + REQUIRE(A_h[0] == INITIAL_VAL); + REQUIRE(*result_h == 0); + HIP_CHECK(hipHostFree(A_h)); + HIP_CHECK(hipFree(result)); + } + } else { + SUCCEED("Memory model feature is only supported for gfx90a Hence" + "skipping the testcase for this GPU " << device); + } +} + +/* +This test verifies the built in atomic add API on Non-Coherent Memory +Input: A_h with INITIAL_VAL +Output: A_h will not get updated with Coherent Memory + A_h will be INITIAL_VAL+INC_VAL + B_h would be initial value of A_h, B_h would be INITIAL_VAL +*/ +TEST_CASE("Unit_BuiltInAtomicAdd_NonCoherentGlobalMem") { + hipDeviceProp_t prop; + int device; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + std::string gfxName(prop.gcnArchName); + if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) { + if (prop.canMapHostMemory != 1) { + SUCCEED("Does not support HostPinned Memory"); + } else { + double *A_h, *result, *B_h; + double *A_d; + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_h), sizeof(double), + hipHostMallocNonCoherent)); + B_h = reinterpret_cast(malloc(sizeof(double))); + HIP_CHECK(hipMalloc(reinterpret_cast(&result), sizeof(double))); + A_h[0] = INITIAL_VAL; + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&A_d), + A_h, 0)); + hipLaunchKernelGGL(AtomicAdd_GlobalMem, dim3(1), dim3(1), + 0, 0, static_cast(A_d), + static_cast(result)); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipMemcpy(B_h, result, sizeof(double), hipMemcpyDeviceToHost)); + REQUIRE(A_h[0] == INITIAL_VAL + INC_VAL); + REQUIRE(*B_h == INITIAL_VAL); + HIP_CHECK(hipHostFree(A_h)); + HIP_CHECK(hipFree(result)); + free(B_h); + } + } else { + SUCCEED("Memory model feature is only supported for gfx90a" + "Hence skipping the testcase for GPU-0"); + } +} +/* +This test verifies the built in atomic add API on Coherent Memory with RTC +Input: A_h with INITIAL_VAL +Output: A_h will not get updated with Coherent Memory + A_h will be INITIAL_VAL + ret value would be 0, B_h would be 0 +*/ +TEST_CASE("Unit_BuiltInAtomicAdd_CoherentGlobalMemWithRtc") { + hipDeviceProp_t prop; + int device; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + std::string gfxName(prop.gcnArchName); + if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) { + if (prop.canMapHostMemory != 1) { + SUCCEED("Does not support HostPinned Memory"); + } else { + hiprtcProgram prog; + hiprtcCreateProgram(&prog, // prog + AtomicAddGlobalMem, // buffer + "kernel.cu", // name + 0, nullptr, nullptr); + std::string sarg = std::string("--gpu-architecture=") + prop.gcnArchName; + const char* options[] = {sarg.c_str()}; + hiprtcResult compileResult{hiprtcCompileProgram(prog, 1, options)}; + + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + INFO(log); + } + + REQUIRE(compileResult == HIPRTC_SUCCESS); + size_t codeSize; + HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize)); + + std::vector code(codeSize); + HIPRTC_CHECK(hiprtcGetCode(prog, code.data())); + HIPRTC_CHECK(hiprtcDestroyProgram(&prog)); + + hipModule_t module; + hipFunction_t fmaxkernel; + HIP_CHECK(hipModuleLoadData(&module, code.data())); + HIP_CHECK(hipModuleGetFunction(&fmaxkernel, module, + "AtomicAdd_GlobalMem")); + double *A_h, *result, *B_h; + double *A_d; + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_h), sizeof(double), + hipHostMallocCoherent)); + B_h = reinterpret_cast(malloc(sizeof(double))); + HIP_CHECK(hipMalloc(reinterpret_cast(&result), sizeof(double))); + A_h[0] = INITIAL_VAL; + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&A_d), + A_h, 0)); + struct { + double* p; + double* res; + } args_f{A_d, result}; + auto size = sizeof(args_f); + void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f, + HIP_LAUNCH_PARAM_BUFFER_SIZE, + &size, HIP_LAUNCH_PARAM_END}; + hipModuleLaunchKernel(fmaxkernel, 1, 1, 1, 1, 1, 1, 0, + nullptr, nullptr, config_d); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipMemcpy(B_h, result, sizeof(double), hipMemcpyDeviceToHost)); + REQUIRE(A_h[0] == INITIAL_VAL); + REQUIRE(*B_h == 0); + HIP_CHECK(hipHostFree(A_h)); + HIP_CHECK(hipFree(result)); + free(B_h); + } + } else { + SUCCEED("Memory model feature is only supported for gfx90a, Hence" + "skipping the testcase for this GPU " << device); + } +} + +/* +This test verifies the built in atomic add API on Non-Coherent Memory +Input: A_h with INITIAL_VAL +Output: A_h will not get updated with Coherent Memory + A_h will be INITIAL_VAL+INC_VAL + B_h would be initial value of A_h, B_h would be INITIAL_VAL +*/ +TEST_CASE("Unit_BuiltInAtomicAdd_NonCoherentGlobalMemWithRtc") { + hipDeviceProp_t prop; + int device; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + std::string gfxName(prop.gcnArchName); + if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) { + if (prop.canMapHostMemory != 1) { + SUCCEED("Does support HostPinned Memory"); + } else { + hiprtcProgram prog; + hiprtcCreateProgram(&prog, // prog + AtomicAddGlobalMem, // buffer + "kernel.cu", // name + 0, nullptr, nullptr); + std::string sarg = std::string("--gpu-architecture=") + prop.gcnArchName; + const char* options[] = {sarg.c_str()}; + hiprtcResult compileResult{hiprtcCompileProgram(prog, 1, options)}; + + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + WARN(log); + } + + REQUIRE(compileResult == HIPRTC_SUCCESS); + size_t codeSize; + HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize)); + + std::vector code(codeSize); + HIPRTC_CHECK(hiprtcGetCode(prog, code.data())); + HIPRTC_CHECK(hiprtcDestroyProgram(&prog)); + + hipModule_t module; + hipFunction_t fmaxkernel; + HIP_CHECK(hipModuleLoadData(&module, code.data())); + HIP_CHECK(hipModuleGetFunction(&fmaxkernel, module, + "AtomicAdd_GlobalMem")); + double *A_h, *result, *B_h; + double *A_d; + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_h), sizeof(double), + hipHostMallocNonCoherent)); + B_h = reinterpret_cast(malloc(sizeof(double))); + HIP_CHECK(hipMalloc(reinterpret_cast(&result), sizeof(double))); + A_h[0] = INITIAL_VAL; + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&A_d), + A_h, 0)); + struct { + double* p; + double* res; + } args_f{A_d, result}; + auto size = sizeof(args_f); + void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f, + HIP_LAUNCH_PARAM_BUFFER_SIZE, + &size, HIP_LAUNCH_PARAM_END}; + hipModuleLaunchKernel(fmaxkernel, 1, 1, 1, 1, 1, 1, 0, + nullptr, nullptr, config_d); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipMemcpy(B_h, result, sizeof(double), hipMemcpyDeviceToHost)); + REQUIRE(A_h[0] == INITIAL_VAL + INC_VAL); + REQUIRE(*B_h == INITIAL_VAL); + HIP_CHECK(hipHostFree(A_h)); + HIP_CHECK(hipFree(result)); + free(B_h); + } + } else { + SUCCEED("Memory model feature is only supported for gfx90a, Hence" + "skipping the testcase for this GPU " << device); + } +} diff --git a/projects/hip-tests/catch/unit/deviceLib/BuiltIns_fmax.cc b/projects/hip-tests/catch/unit/deviceLib/BuiltIns_fmax.cc new file mode 100644 index 0000000000..40f4c52397 --- /dev/null +++ b/projects/hip-tests/catch/unit/deviceLib/BuiltIns_fmax.cc @@ -0,0 +1,354 @@ +/* + Copyright (c) 2021 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. + */ +/* +This testfile verifies Built fmax API scenarios +1. Builtin fmax on Coherent Memory with memory type as global +2. Builtin fmax on Non-Coherent Memory with memory type as global +3. Builtin fmax with memory type as flat +4. Builtin fmax on Coherent Memory with RTC and memory type as global +5. Builtin fmax on Non-Coherent Memory with RTC and memory type as global +6. Builtin fmax with RTC and memory type as flat +*/ + + +#include +#include +#include + +#define INITIAL_VAL 5 + +__global__ void unsafeAtomicMax_FlatMem(double* addr, double* result) { + __shared__ double int_val; + int_val = 5; + double comp = 10; + if (__builtin_amdgcn_is_shared( + (const __attribute__((address_space(0))) void*)(&int_val))) + *result = __builtin_amdgcn_flat_atomic_fmax_f64(&int_val, comp); + else + *result = __builtin_amdgcn_global_atomic_fmax_f64(&int_val, comp); + *addr = int_val; +} +__global__ void unsafeAtomicMax_GlobalMem(double* addr, double* result) { + double comp = 10; + if (__builtin_amdgcn_is_shared( + (const __attribute__((address_space(0))) void*)(addr))) + *result = __builtin_amdgcn_flat_atomic_fmax_f64(addr, comp); + else + *result = __builtin_amdgcn_global_atomic_fmax_f64(addr, comp); +} +static constexpr auto fmaxFlatMem { +R"( +extern "C" +__global__ void unsafeAtomicMax_FlatMem(double* addr, double* result) { + __shared__ double int_val; + int_val = 5; + double comp = 10; + if (__builtin_amdgcn_is_shared( + (const __attribute__((address_space(0))) void*)(&int_val))) + *result = __builtin_amdgcn_flat_atomic_fmax_f64(&int_val, comp); + else + *result = __builtin_amdgcn_global_atomic_fmax_f64(&int_val, comp); + *addr = int_val; +} +)"}; + +static constexpr auto fmaxGlobalMem { +R"( +extern "C" +__global__ void unsafeAtomicMax_GlobalMem(double* addr, double* result) { + double comp = 10; + if (__builtin_amdgcn_is_shared( + (const __attribute__((address_space(0))) void*)(addr))) + *result = __builtin_amdgcn_flat_atomic_fmax_f64(addr, comp); + else + *result = __builtin_amdgcn_global_atomic_fmax_f64(addr, comp); +} +)"}; + +/* +This testcase verifies the builtinAtomic fmax API on Coherent memory +with memory type as global +Input: A_h with INITIAL_VAL +Output: Return val would be 0 and the input value to API will not + get updated. A_h would be INITIAL_VAL, B_h is 0 +*/ +TEST_CASE("Unit_BuiltinAtomics_fmaxCoherentGlobalMem") { + hipDeviceProp_t prop; + int device; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + std::string gfxName(prop.gcnArchName); + if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) { + if (prop.canMapHostMemory != 1) { + SUCCEED("Does not support HostPinned Memory"); + } else { + double *A_h, *B_h; + double *A_d; + double *result; + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_h), sizeof(double), + hipHostMallocCoherent)); + A_h[0] = INITIAL_VAL; + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&A_d), + A_h, 0)); + B_h = reinterpret_cast(malloc(sizeof(double))); + HIP_CHECK(hipMalloc(reinterpret_cast(&result), sizeof(double))); + hipLaunchKernelGGL(unsafeAtomicMax_GlobalMem, dim3(1), dim3(1), + 0, 0, static_cast(A_d), result); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipMemcpy(B_h, result, sizeof(double), hipMemcpyDeviceToHost)); + REQUIRE(*B_h == 0); + REQUIRE(A_h[0] == INITIAL_VAL); + HIP_CHECK(hipHostFree(A_h)); + HIP_CHECK(hipFree(result)); + free(B_h); + } + } else { + SUCCEED("Memory model feature is only supported for gfx90a, Hence" + "skipping the testcase for this GPU " << device); + } +} + +/* +This testcase verifies the builtinAtomic fmax API +1. Non Coherent memory with memory type as global +2. Memory type as flat +Input: A_h with INITIAL_VAL +Output: Return val would be initial val of A_h and the input value of + API would be updated with the max value + A_h would be 10, B_h would be INITIAL_VAL +*/ +TEST_CASE("Unit_BuiltinAtomics_fmaxNonCoherentGlobalFlatMem") { + int mem_type = GENERATE(0, 1); + hipDeviceProp_t prop; + int device; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + std::string gfxName(prop.gcnArchName); + if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) { + if (prop.canMapHostMemory != 1) { + SUCCEED("Does not support HostPinned Memory"); + } else { + double *A_h, *B_h; + double *A_d; + double *result; + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_h), sizeof(double), + hipHostMallocNonCoherent)); + A_h[0] = INITIAL_VAL; + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&A_d), + A_h, 0)); + B_h = reinterpret_cast(malloc(sizeof(double))); + HIP_CHECK(hipMalloc(reinterpret_cast(&result), sizeof(double))); + if (mem_type) { + hipLaunchKernelGGL(unsafeAtomicMax_GlobalMem, dim3(1), dim3(1), + 0, 0, static_cast(A_d), result); + } else { + hipLaunchKernelGGL(unsafeAtomicMax_FlatMem, dim3(1), dim3(1), + 0, 0, static_cast(A_d), result); + } + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipMemcpy(B_h, result, sizeof(double), hipMemcpyDeviceToHost)); + REQUIRE(*B_h == INITIAL_VAL); + REQUIRE(A_h[0] == 10); + HIP_CHECK(hipHostFree(A_h)); + HIP_CHECK(hipFree(result)); + free(B_h); + } + } else { + SUCCEED("Memory model feature is only supported for gfx90a, Hence" + "skipping the testcase for this GPU " << device); + } +} +/* +This testcase verifies the builtinAtomic fmax API on Coherent memory +with RTC and memory type as global +Input: A_h with INITIAL_VAL +Output: Return val would be 0 and the input value to API will not + get updated. A_h would be INITIAL_VAL, B_h is 0 +*/ +TEST_CASE("Unit_BuiltinAtomicsRTC_fmaxCoherentGlobalMem") { + hipDeviceProp_t prop; + int device; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + std::string gfxName(prop.gcnArchName); + if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) { + if (prop.canMapHostMemory != 1) { + SUCCEED("Does not support HostPinned Memory"); + } else { + hiprtcProgram prog; + hiprtcCreateProgram(&prog, // prog + fmaxGlobalMem, // buffer + "kernel.cu", // name + 0, nullptr, nullptr); + std::string sarg = std::string("--gpu-architecture=") + prop.gcnArchName; + const char* options[] = {sarg.c_str()}; + hiprtcResult compileResult{hiprtcCompileProgram(prog, 1, options)}; + + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + INFO(log); + } + + REQUIRE(compileResult == HIPRTC_SUCCESS); + size_t codeSize; + HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize)); + + std::vector code(codeSize); + HIPRTC_CHECK(hiprtcGetCode(prog, code.data())); + HIPRTC_CHECK(hiprtcDestroyProgram(&prog)); + + hipModule_t module; + hipFunction_t fmaxkernel; + HIP_CHECK(hipModuleLoadData(&module, code.data())); + HIP_CHECK(hipModuleGetFunction(&fmaxkernel, module, + "unsafeAtomicMax_GlobalMem")); + + double *A_h, *B_h; + double *A_d; + double *result; + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_h), sizeof(double), + hipHostMallocCoherent)); + A_h[0] = INITIAL_VAL; + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&A_d), + A_h, 0)); + B_h = reinterpret_cast(malloc(sizeof(double))); + HIP_CHECK(hipMalloc(reinterpret_cast(&result), sizeof(double))); + struct { + double* p; + double* res; + } args_f{A_d, result}; + auto size = sizeof(args_f); + void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f, + HIP_LAUNCH_PARAM_BUFFER_SIZE, + &size, HIP_LAUNCH_PARAM_END}; + hipModuleLaunchKernel(fmaxkernel, 1, 1, 1, 1, 1, 1, 0, + nullptr, nullptr, config_d); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipMemcpy(B_h, result, sizeof(double), hipMemcpyDeviceToHost)); + REQUIRE(*B_h == 0); + REQUIRE(A_h[0] == INITIAL_VAL); + HIP_CHECK(hipHostFree(A_h)); + HIP_CHECK(hipFree(result)); + free(B_h); + } + } else { + SUCCEED("Memory model feature is only supported for gfx90a, Hence" + "skipping the testcase for this GPU " << device); + } +} +/* +This testcase verifies the builtinAtomic fmax API with RTC +1. Non Coherent memory with memory type as global +2. Memory type as flat +Input: A_h with INITIAL_VAL +Output: Return val would be initial val of A_h and the input value of + API would be updated with the max value + A_h would be 10, B_h would be INITIAL_VAL +*/ +TEST_CASE("Unit_BuiltinAtomicsRTC_fmaxNonCoherentGlobalFlatMem") { + int mem_type = GENERATE(0, 1); + hipDeviceProp_t prop; + int device; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + std::string gfxName(prop.gcnArchName); + if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) { + if (prop.canMapHostMemory != 1) { + SUCCEED("Does not support HostPinned Memory"); + } else { + hiprtcProgram prog; + if (mem_type) { + hiprtcCreateProgram(&prog, // prog + fmaxGlobalMem, // buffer + "kernel.cu", // name + 0, nullptr, nullptr); + } else { + hiprtcCreateProgram(&prog, // prog + fmaxFlatMem, // buffer + "kernel.cu", // name + 0, nullptr, nullptr); + } + std::string sarg = std::string("--gpu-architecture=") + prop.gcnArchName; + const char* options[] = {sarg.c_str()}; + hiprtcResult compileResult{hiprtcCompileProgram(prog, 1, options)}; + + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + INFO(log); + } + REQUIRE(compileResult == HIPRTC_SUCCESS); + size_t codeSize; + HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize)); + + std::vector code(codeSize); + HIPRTC_CHECK(hiprtcGetCode(prog, code.data())); + HIPRTC_CHECK(hiprtcDestroyProgram(&prog)); + + hipModule_t module; + hipFunction_t fmaxkernel; + HIP_CHECK(hipModuleLoadData(&module, code.data())); + if (mem_type) { + HIP_CHECK(hipModuleGetFunction(&fmaxkernel, module, + "unsafeAtomicMax_GlobalMem")); + } else { + HIP_CHECK(hipModuleGetFunction(&fmaxkernel, module, + "unsafeAtomicMax_FlatMem")); + } + + double *A_h, *B_h; + double *A_d; + double *result; + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_h), sizeof(double), + hipHostMallocNonCoherent)); + A_h[0] = INITIAL_VAL; + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&A_d), + A_h, 0)); + B_h = reinterpret_cast(malloc(sizeof(double))); + HIP_CHECK(hipMalloc(reinterpret_cast(&result), sizeof(double))); + struct { + double* p; + double* res; + } args_f{A_d, result}; + auto size = sizeof(args_f); + void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f, + HIP_LAUNCH_PARAM_BUFFER_SIZE, + &size, HIP_LAUNCH_PARAM_END}; + hipModuleLaunchKernel(fmaxkernel, 1, 1, 1, 1, 1, 1, 0, + nullptr, nullptr, config_d); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipMemcpy(B_h, result, sizeof(double), hipMemcpyDeviceToHost)); + REQUIRE(*B_h == INITIAL_VAL); + REQUIRE(A_h[0] == 10); + HIP_CHECK(hipHostFree(A_h)); + HIP_CHECK(hipFree(result)); + free(B_h); + } + } else { + SUCCEED("Memory model feature is only supported for gfx90a, Hence" + "skipping the testcase for this GPU " << device); + } +} diff --git a/projects/hip-tests/catch/unit/deviceLib/BuiltIns_fmin.cc b/projects/hip-tests/catch/unit/deviceLib/BuiltIns_fmin.cc new file mode 100644 index 0000000000..8ff1814db3 --- /dev/null +++ b/projects/hip-tests/catch/unit/deviceLib/BuiltIns_fmin.cc @@ -0,0 +1,357 @@ +/* + Copyright (c) 2021 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. + */ +/* +This testfile verifies Built fmin API scenarios +1. Builtin fmin on Coherent Memory with memory type as global +2. Builtin fmin on Non-Coherent Memory with memory type as global +3. Builtin fmin with memory type as flat +4. Builtin fmin on Coherent Memory with RTC and memory type as global +5. Builtin fmin on Non-Coherent Memory with RTC and memory type as global +6. Builtin fmin with RTC and memory type as flat +*/ + +#include +#include +#include + +#define INITIAL_VAL 5 + +static constexpr auto fminFlatMem{ +R"( +extern "C" +__global__ void unsafeAtomicMin_FlatMem(double* addr, double* result) { + __shared__ double int_val; + int_val = 5; + double comp = 10; + if (__builtin_amdgcn_is_shared( + (const __attribute__((address_space(0))) void*)(&int_val))) + *result = __builtin_amdgcn_flat_atomic_fmin_f64(&int_val, comp); + else + *result = __builtin_amdgcn_global_atomic_fmin_f64(&int_val, comp); + *addr = int_val; +} +)"}; + +static constexpr auto fminGlobalMem{ +R"( +extern "C" +__global__ void unsafeAtomicMin_GlobalMem(double* addr, double* result) { + double comp = 10; + if (__builtin_amdgcn_is_shared( + (const __attribute__((address_space(0))) void*)(addr))) + *result = __builtin_amdgcn_flat_atomic_fmin_f64(addr, comp); + else + *result = __builtin_amdgcn_global_atomic_fmin_f64(addr, comp); +} +)"}; + +__global__ void unsafeAtomicMin_FlatMem(double* addr, double* result) { + __shared__ double int_val; + int_val = 5; + double comp = 10; + if (__builtin_amdgcn_is_shared( + (const __attribute__((address_space(0))) void*)(&int_val))) + *result = __builtin_amdgcn_flat_atomic_fmin_f64(&int_val, comp); + else + *result = __builtin_amdgcn_global_atomic_fmin_f64(&int_val, comp); + *addr = int_val; +} +__global__ void unsafeAtomicMin_GlobalMem(double* addr, double* result) { + double comp = 10; + if (__builtin_amdgcn_is_shared( + (const __attribute__((address_space(0))) void*)(addr))) + *result = __builtin_amdgcn_flat_atomic_fmin_f64(addr, comp); + else + *result = __builtin_amdgcn_global_atomic_fmin_f64(addr, comp); +} + +/* +This testcase verifies the builtinAtomic fmin API on Coherent memory +with memory type as global +Input: A_h with INITIAL_VAL +Output: Return val would be 0 and the input value to API will not + get updated. A_h would be INITIAL_VAL, B_h is 0 +*/ +TEST_CASE("Unit_BuiltinAtomics_fminCoherentGlobalMem") { + hipDeviceProp_t prop; + int device; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + std::string gfxName(prop.gcnArchName); + if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) { + if (prop.canMapHostMemory != 1) { + SUCCEED("Does not support HostPinned Memory"); + } else { + double *A_h, *B_h; + double *A_d; + double *result; + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_h), sizeof(double), + hipHostMallocCoherent)); + A_h[0] = INITIAL_VAL; + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&A_d), + A_h, 0)); + B_h = reinterpret_cast(malloc(sizeof(double))); + HIP_CHECK(hipMalloc(reinterpret_cast(&result), sizeof(double))); + hipLaunchKernelGGL(unsafeAtomicMin_GlobalMem, dim3(1), dim3(1), + 0, 0, static_cast(A_d), result); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipMemcpy(B_h, result, sizeof(double), hipMemcpyDeviceToHost)); + REQUIRE(*B_h == 0); + REQUIRE(A_h[0] == INITIAL_VAL); + HIP_CHECK(hipHostFree(A_h)); + HIP_CHECK(hipFree(result)); + free(B_h); + } + } else { + SUCCEED("Memory model feature is only supported for gfx90a, Hence" + "skipping the testcase for this GPU " << device); + } +} + +/* +This testcase verifies the builtinAtomic fmin API +1. Non Coherent memory with memory type as global +2. Memory type as flat +Input: A_h with INITIAL_VAL +Output: Return val would be initial val of A_h and the input value of + API would be updated with the min value + A_h would be INITIAL_VAL, B_h would be INITIAL_VAL +*/ +TEST_CASE("Unit_BuiltinAtomics_fminNonCoherentGlobalFlatMem") { + auto mem_type = GENERATE(0, 1); + hipDeviceProp_t prop; + int device; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + std::string gfxName(prop.gcnArchName); + if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) { + if (prop.canMapHostMemory != 1) { + SUCCEED("Does not support HostPinned Memory"); + } else { + double *A_h, *B_h; + double *A_d; + double *result; + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_h), sizeof(double), + hipHostMallocNonCoherent)); + A_h[0] = INITIAL_VAL; + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&A_d), + A_h, 0)); + B_h = reinterpret_cast(malloc(sizeof(double))); + HIP_CHECK(hipMalloc(reinterpret_cast(&result), sizeof(double))); + if (mem_type) { + hipLaunchKernelGGL(unsafeAtomicMin_GlobalMem, dim3(1), dim3(1), + 0, 0, static_cast(A_d), result); + } else { + hipLaunchKernelGGL(unsafeAtomicMin_FlatMem, dim3(1), dim3(1), + 0, 0, static_cast(A_d), result); + } + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipMemcpy(B_h, result, sizeof(double), hipMemcpyDeviceToHost)); + REQUIRE(*B_h == INITIAL_VAL); + REQUIRE(A_h[0] == INITIAL_VAL); + HIP_CHECK(hipHostFree(A_h)); + HIP_CHECK(hipFree(result)); + free(B_h); + } + } else { + SUCCEED("Memory model feature is only supported for gfx90a, Hence" + "skipping the testcase for this GPU " << device); + } +} +/* +This testcase verifies the builtinAtomic fmin API on Coherent memory +with RTC and memory type as global +Input: A_h with INITIAL_VAL +Output: Return val would be 0 and the input value to API will not + get updated. A_h would be INITIAL_VAL, B_h is 0 +*/ + +TEST_CASE("Unit_BuiltinAtomicsRTC__fminCoherentGlobalMem") { + hipDeviceProp_t prop; + int device; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + std::string gfxName(prop.gcnArchName); + if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) { + if (prop.canMapHostMemory != 1) { + SUCCEED("Does not support HostPinned Memory"); + } else { + hiprtcProgram prog; + hiprtcCreateProgram(&prog, // prog + fminGlobalMem, // buffer + "kernel.cu", // name + 0, nullptr, nullptr); + std::string sarg = std::string("--gpu-architecture=") + prop.gcnArchName; + const char* options[] = {sarg.c_str()}; + hiprtcResult compileResult{hiprtcCompileProgram(prog, 1, options)}; + + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + INFO(log); + } + + REQUIRE(compileResult == HIPRTC_SUCCESS); + size_t codeSize; + HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize)); + + std::vector code(codeSize); + HIPRTC_CHECK(hiprtcGetCode(prog, code.data())); + HIPRTC_CHECK(hiprtcDestroyProgram(&prog)); + + hipModule_t module; + hipFunction_t fmaxkernel; + HIP_CHECK(hipModuleLoadData(&module, code.data())); + HIP_CHECK(hipModuleGetFunction(&fmaxkernel, module, + "unsafeAtomicMin_GlobalMem")); + + double *A_h, *B_h; + double *A_d; + double *result; + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_h), sizeof(double), + hipHostMallocCoherent)); + A_h[0] = INITIAL_VAL; + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&A_d), + A_h, 0)); + B_h = reinterpret_cast(malloc(sizeof(double))); + HIP_CHECK(hipMalloc(reinterpret_cast(&result), sizeof(double))); + struct { + double* p; + double* res; + } args_f{A_d, result}; + auto size = sizeof(args_f); + void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f, + HIP_LAUNCH_PARAM_BUFFER_SIZE, + &size, HIP_LAUNCH_PARAM_END}; + hipModuleLaunchKernel(fmaxkernel, 1, 1, 1, 1, 1, 1, 0, + nullptr, nullptr, config_d); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipMemcpy(B_h, result, sizeof(double), hipMemcpyDeviceToHost)); + REQUIRE(*B_h == 0); + REQUIRE(A_h[0] == INITIAL_VAL); + HIP_CHECK(hipHostFree(A_h)); + HIP_CHECK(hipFree(result)); + free(B_h); + } + } else { + SUCCEED("Memory model feature is only supported for gfx90a, Hence" + "skipping the testcase for this GPU " << device); + } +} + +/* +This testcase verifies the builtinAtomic fmin API with RTC +1. Non Coherent memory with memory type as global +2. Memory type as flat +Input: A_h with INITIAL_VAL +Output: Return val would be initial val of A_h and the input value of + API would be updated with the max value + A_h would be 10, B_h would be INITIAL_VAL +*/ +TEST_CASE("Unit_BuiltinAtomicsRTC_fminNonCoherentGlobalFlatMem") { + int mem_type = GENERATE(0, 1); + hipDeviceProp_t prop; + int device; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + std::string gfxName(prop.gcnArchName); + if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) { + if (prop.canMapHostMemory != 1) { + SUCCEED("Does not support HostPinned Memory"); + } else { + hiprtcProgram prog; + if (mem_type) { + hiprtcCreateProgram(&prog, // prog + fminGlobalMem, // buffer + "kernel.cu", // name + 0, nullptr, nullptr); + } else { + hiprtcCreateProgram(&prog, // prog + fminFlatMem, // buffer + "kernel.cu", // name + 0, nullptr, nullptr); + } + std::string sarg = std::string("--gpu-architecture=") + prop.gcnArchName; + const char* options[] = {sarg.c_str()}; + hiprtcResult compileResult{hiprtcCompileProgram(prog, 1, options)}; + + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + INFO(log); + } + + REQUIRE(compileResult == HIPRTC_SUCCESS); + size_t codeSize; + HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize)); + + std::vector code(codeSize); + HIPRTC_CHECK(hiprtcGetCode(prog, code.data())); + HIPRTC_CHECK(hiprtcDestroyProgram(&prog)); + + hipModule_t module; + hipFunction_t fmaxkernel; + HIP_CHECK(hipModuleLoadData(&module, code.data())); + if (mem_type) { + HIP_CHECK(hipModuleGetFunction(&fmaxkernel, module, + "unsafeAtomicMin_GlobalMem")); + } else { + HIP_CHECK(hipModuleGetFunction(&fmaxkernel, module, + "unsafeAtomicMin_FlatMem")); + } + + double *A_h, *B_h; + double *A_d; + double *result; + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_h), sizeof(double), + hipHostMallocNonCoherent)); + A_h[0] = INITIAL_VAL; + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&A_d), + A_h, 0)); + B_h = reinterpret_cast(malloc(sizeof(double))); + HIP_CHECK(hipMalloc(reinterpret_cast(&result), sizeof(double))); + struct { + double* p; + double* res; + } args_f{A_d, result}; + auto size = sizeof(args_f); + void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f, + HIP_LAUNCH_PARAM_BUFFER_SIZE, + &size, HIP_LAUNCH_PARAM_END}; + hipModuleLaunchKernel(fmaxkernel, 1, 1, 1, 1, 1, 1, 0, + nullptr, nullptr, config_d); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipMemcpy(B_h, result, sizeof(double), hipMemcpyDeviceToHost)); + REQUIRE(*B_h == INITIAL_VAL); + REQUIRE(A_h[0] == INITIAL_VAL); + HIP_CHECK(hipHostFree(A_h)); + HIP_CHECK(hipFree(result)); + free(B_h); + } + } else { + SUCCEED("Memory model feature is only supported for gfx90a, Hence" + "skipping the testcase for this GPU " << device); + } +} diff --git a/projects/hip-tests/catch/unit/deviceLib/CMakeLists.txt b/projects/hip-tests/catch/unit/deviceLib/CMakeLists.txt index 644569e49e..df74a6dbb5 100644 --- a/projects/hip-tests/catch/unit/deviceLib/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/deviceLib/CMakeLists.txt @@ -37,6 +37,16 @@ set(AMD_ARCH_SPEC_TEST_SRC AtomicAdd_NonCoherent_withoutflag.cc AtomicAdd_NonCoherent_withnoUnsafeflag.cc AtomicAdd_NonCoherent_withunsafeflag.cc + BuiltIns_fmax.cc + BuiltIns_fmin.cc + BuiltIns_fadd.cc + unsafeAtomicAdd_RTC.cc + unsafeAtomicAdd_Coherent_withunsafeflag.cc + unsafeAtomicAdd_Coherent_withoutflag.cc + unsafeAtomicAdd_Coherent_withnounsafeflag.cc + unsafeAtomicAdd_NonCoherent_withoutflag.cc + unsafeAtomicAdd_NonCoherent_withnounsafeflag.cc + unsafeAtomicAdd_NonCoherent_withunsafeflag.cc ) if(HIP_PLATFORM MATCHES "amd") @@ -49,8 +59,14 @@ if(${ARCH_CHECK} GREATER_EQUAL 0) set_source_files_properties(AtomicAdd_NonCoherent_withunsafeflag.cc PROPERTIES COMPILE_OPTIONS "-munsafe-fp-atomics") set_source_files_properties(AtomicAdd_Coherent_withnoUnsafeflag.cc PROPERTIES COMPILE_OPTIONS "-mno-unsafe-fp-atomics") set_source_files_properties(AtomicAdd_NonCoherent_withnoUnsafeflag.cc PROPERTIES COMPILE_OPTIONS "-mno-unsafe-fp-atomics") + set_source_files_properties(unsafeAtomicAdd_Coherent_withunsafeflag.cc PROPERTIES COMPILE_OPTIONS "-munsafe-fp-atomics") + 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") file(GLOB AtomicAdd_files *AtomicAdd_*_*.cc) set_property(SOURCE ${AtomicAdd_files} PROPERTY COMPILE_FLAGS --save-temps) + file(GLOB unsafeAtomicAdd_files *unsafeAtomicAdd_*_*.cc) + set_property(SOURCE ${unsafeAtomicAdd_files} PROPERTY COMPILE_FLAGS --save-temps) endif() hip_add_exe_to_target(NAME UnitDeviceTests TEST_SRC ${TEST_SRC} diff --git a/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withnounsafeflag.cc b/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withnounsafeflag.cc new file mode 100644 index 0000000000..384347de77 --- /dev/null +++ b/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withnounsafeflag.cc @@ -0,0 +1,99 @@ +/* + 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. + */ +/* +AtomicAdd on FineGrainMemory +1. The following test scenario verifies +unsafeatomicAdd on fineGrain memory with -mno-unsafe-fp-atomics flag +This testcase works only on gfx90a. +*/ + +#include +#include + +#define INC_VAL 10 +#define INITIAL_VAL 5 +template +static __global__ void AtomicCheck(T* Ad, T* result) { + T inc_val = 10; + *result = unsafeAtomicAdd(Ad, inc_val); +} + + +/*unsafeatomicAdd API for the fine grained memory variable + with -mno-unsafe-fp-atomics flag +Input: Ad{5}, INC_VAL{10} +Output: unsafeatomicAdd API would return 0 and the 0/P is 5 + Generate the assembly file and check whether + atomic add instruction is generated + or not */ + +TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_CoherentwithnoUnsafeflag", "", + float, double) { + hipDeviceProp_t prop; + int device; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + std::string gfxName(prop.gcnArchName); + if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) { + if (prop.canMapHostMemory != 1) { + SUCCEED("Does not support HostPinned Memory"); + } else { + TestType *A_h{nullptr}, *result{nullptr}; + TestType *A_d{nullptr}, *result_d{nullptr}; + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_h), sizeof(TestType), + hipHostMallocCoherent)); + A_h[0] = INITIAL_VAL; + HIP_CHECK(hipHostMalloc(reinterpret_cast(&result), + sizeof(TestType), + hipHostMallocCoherent)); + result[0] = INITIAL_VAL; + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&A_d), + A_h, 0)); + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&result_d), + result, 0)); + hipLaunchKernelGGL(AtomicCheck, dim3(1), dim3(1), + 0, 0, A_d, + result_d); + HIP_CHECK(hipDeviceSynchronize()); + bool testResult; + + if ((std::is_same::value)) { + testResult = HipTest::assemblyFile_Verification( + "unsafeAtomicAdd_Coherent_withnounsafeflag-hip-amdgcn(.*)\\.s", + "global_atomic_add_f32"); + REQUIRE(testResult == true); + } else { + testResult = HipTest::assemblyFile_Verification( + "unsafeAtomicAdd_Coherent_withnounsafeflag-hip-amdgcn(.*)\\.s", + "global_atomic_add_f64"); + REQUIRE(testResult == true); + } + + REQUIRE(A_h[0] == INITIAL_VAL); + REQUIRE(result[0] == 0); + HIP_CHECK(hipHostFree(A_h)); + HIP_CHECK(hipHostFree(result)); + } + } else { + SUCCEED("Memory model feature is only supported for gfx90a, Hence" + "skipping the testcase for this GPU " << device); + } +} diff --git a/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withoutflag.cc b/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withoutflag.cc new file mode 100644 index 0000000000..993705b764 --- /dev/null +++ b/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withoutflag.cc @@ -0,0 +1,99 @@ +/* + 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. + */ +/* +AtomicAdd on FineGrainMemory +1. The following test scenario verifies +unsafeatomicAdd on fineGrain memory without atomics flag +This testcase works only on gfx90a. +*/ + +#include +#include + +#define INC_VAL 10 +#define INITIAL_VAL 5 +template +static __global__ void AtomicCheck(T* Ad, T* result) { + T inc_val = 10; + *result = unsafeAtomicAdd(Ad, inc_val); +} + + +/*unsafeatomicAdd API for the fine grained memory variable + without atomics flag +Input: Ad{5}, INC_VAL{10} +Output: unsafeatomicAdd API would return 0 and the 0/P is 5 + Generate the assembly file and check whether + atomic add instruction is generated + or not */ + +TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_Coherentwithoutflag", "", + float, double) { + hipDeviceProp_t prop; + int device; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + std::string gfxName(prop.gcnArchName); + if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) { + if (prop.canMapHostMemory != 1) { + SUCCEED("Does not support HostPinned Memory"); + } else { + TestType *A_h{nullptr}, *result{nullptr}; + TestType *A_d{nullptr}, *result_d{nullptr}; + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_h), sizeof(TestType), + hipHostMallocCoherent)); + A_h[0] = INITIAL_VAL; + HIP_CHECK(hipHostMalloc(reinterpret_cast(&result), + sizeof(TestType), + hipHostMallocCoherent)); + result[0] = INITIAL_VAL; + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&A_d), + A_h, 0)); + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&result_d), + result, 0)); + hipLaunchKernelGGL(AtomicCheck, dim3(1), dim3(1), + 0, 0, A_d, + result_d); + HIP_CHECK(hipDeviceSynchronize()); + bool testResult; + + if ((std::is_same::value)) { + testResult = HipTest::assemblyFile_Verification( + "unsafeAtomicAdd_Coherent_withoutflag-hip-amdgcn(.*)\\.s", + "global_atomic_add_f32"); + REQUIRE(testResult == true); + } else { + testResult = HipTest::assemblyFile_Verification( + "unsafeAtomicAdd_Coherent_withoutflag-hip-amdgcn(.*)\\.s", + "global_atomic_add_f64"); + REQUIRE(testResult == true); + } + + REQUIRE(A_h[0] == INITIAL_VAL); + REQUIRE(result[0] == 0); + HIP_CHECK(hipHostFree(A_h)); + HIP_CHECK(hipHostFree(result)); + } + } else { + SUCCEED("Memory model feature is only supported for gfx90a, Hence" + "skipping the testcase for this GPU " << device); + } +} diff --git a/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withunsafeflag.cc b/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withunsafeflag.cc new file mode 100644 index 0000000000..868257384a --- /dev/null +++ b/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withunsafeflag.cc @@ -0,0 +1,100 @@ +/* + 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. + */ +/* +AtomicAdd on FineGrainMemory +1. The following test scenario verifies +unsafeatomicAdd on fineGrain memory with -munsafe-fp-atomics flag +This testcase works only on gfx90a. +*/ + +#include +#include + +#define INC_VAL 10 +#define INITIAL_VAL 5 + +template +static __global__ void AtomicCheck(T* Ad, T* result) { + T inc_val = 10; + *result = unsafeAtomicAdd(Ad, inc_val); +} + + +/*unsafeatomicAdd API for the fine grained memory variable + with -m-unsafe-atomics flag +Input: Ad{5}, INC_VAL{10} +Output: atomicAdd API would return 0 and the 0/P is 5 + Generate the assembly file and check whether + atomic add instruction is generated + or not */ + +TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_CoherentwithUnsafeflag", "", + float, double) { + hipDeviceProp_t prop; + int device; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + std::string gfxName(prop.gcnArchName); + if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) { + if (prop.canMapHostMemory != 1) { + SUCCEED("Does not support HostPinned Memory"); + } else { + TestType *A_h{nullptr}, *result{nullptr}; + TestType *A_d{nullptr}, *result_d{nullptr}; + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_h), sizeof(TestType), + hipHostMallocCoherent)); + A_h[0] = INITIAL_VAL; + HIP_CHECK(hipHostMalloc(reinterpret_cast(&result), + sizeof(TestType), + hipHostMallocCoherent)); + result[0] = INITIAL_VAL; + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&A_d), + A_h, 0)); + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&result_d), + result, 0)); + hipLaunchKernelGGL(AtomicCheck, dim3(1), dim3(1), + 0, 0, A_d, + result_d); + HIP_CHECK(hipDeviceSynchronize()); + bool testResult; + + if ((std::is_same::value)) { + testResult = HipTest::assemblyFile_Verification( + "unsafeAtomicAdd_Coherent_withunsafeflag-hip-amdgcn(.*)\\.s", + "global_atomic_add_f32"); + REQUIRE(testResult == true); + } else { + testResult = HipTest::assemblyFile_Verification( + "unsafeAtomicAdd_Coherent_withunsafeflag-hip-amdgcn(.*)\\.s", + "global_atomic_add_f64"); + REQUIRE(testResult == true); + } + + REQUIRE(A_h[0] == INITIAL_VAL); + REQUIRE(result[0] == 0); + HIP_CHECK(hipHostFree(A_h)); + HIP_CHECK(hipHostFree(result)); + } + } else { + SUCCEED("Memory model feature is only supported for gfx90a, Hence" + "skipping the testcase for this GPU " << device); + } +} diff --git a/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_NonCoherent_withnounsafeflag.cc b/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_NonCoherent_withnounsafeflag.cc new file mode 100644 index 0000000000..aff54abd14 --- /dev/null +++ b/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_NonCoherent_withnounsafeflag.cc @@ -0,0 +1,97 @@ +/* + 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. + */ +/* +AtomicAdd on CoarseGrainMemory +1. The following test scenario verifies +unsafeAtomicAdd on CoarseGrain memory with -mno-unsafe-fp-atomics flag +This testcase works only on gfx90a. +*/ + +#include +#include + +#define INC_VAL 10 +#define INITIAL_VAL 5 +template +static __global__ void AtomicCheck(T* Ad, T* result) { + T inc_val = 10; + *result = unsafeAtomicAdd(Ad, inc_val); +} + +/*unsafeAtomicAdd API for the coarse grained memory variable + with -mno-unsafe-fp-atomics flag +Input: Ad{5}, INC_VAL{10} +Output: unsafeAtomicAdd API would work and the 0/P is INITIAL_VAL + INC_VAL + Generate the assembly file and check whether + global_atomic_add instruction is generated + or not */ + +TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_NonCoherentnounsafeatomicsflag", "", + float, double) { + hipDeviceProp_t prop; + int device; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + std::string gfxName(prop.gcnArchName); + if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) { + if (prop.canMapHostMemory != 1) { + SUCCEED("Does not support HostPinned Memory"); + } else { + TestType *A_h{nullptr}, *result{nullptr}; + TestType *A_d{nullptr}, *result_d{nullptr}; + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_h), sizeof(TestType), + hipHostMallocNonCoherent)); + A_h[0] = INITIAL_VAL; + HIP_CHECK(hipHostMalloc(reinterpret_cast(&result), + sizeof(TestType), + hipHostMallocNonCoherent)); + result[0] = INITIAL_VAL; + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&A_d), + A_h, 0)); + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&result_d), + result, 0)); + hipLaunchKernelGGL(AtomicCheck, + dim3(1), dim3(1), + 0, 0, A_d, + result_d); + HIP_CHECK(hipDeviceSynchronize()); + bool testResult; + REQUIRE(A_h[0] == INITIAL_VAL + INC_VAL); + REQUIRE(result[0] == INITIAL_VAL); + if ((std::is_same::value)) { + testResult = HipTest::assemblyFile_Verification( + "unsafeAtomicAdd_NonCoherent_withnounsafeflag-hip-amdgcn(.*)\\.s", + "global_atomic_add_f32"); + REQUIRE(testResult == true); + } else { + testResult = HipTest::assemblyFile_Verification( + "unsafeAtomicAdd_NonCoherent_withnounsafeflag-hip-amdgcn(.*)\\.s", + "global_atomic_add_f64"); + REQUIRE(testResult == true); + } + HIP_CHECK(hipHostFree(A_h)); + HIP_CHECK(hipHostFree(result)); + } + } else { + SUCCEED("Memory model feature is only supported for gfx90a, Hence" + "skipping the testcase for this GPU " << device); + } +} diff --git a/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_NonCoherent_withoutflag.cc b/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_NonCoherent_withoutflag.cc new file mode 100644 index 0000000000..d0b611d89c --- /dev/null +++ b/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_NonCoherent_withoutflag.cc @@ -0,0 +1,97 @@ +/* + 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. + */ +/* +unsafeAtomicAdd on CoarseGrainMemory +1. The following test scenario verifies +unsafeAtomicAdd on CoarseGrain memory without any unsafeatomics flag +This testcase works only on gfx90a. +*/ + +#include +#include + +#define INC_VAL 10 +#define INITIAL_VAL 5 +template +static __global__ void AtomicCheck(T* Ad, T* result) { + T inc_val = 10; + *result = unsafeAtomicAdd(Ad, inc_val); +} + +/*unsafeAtomicAdd API for the coarse grained memory variable + without any flag +Input: Ad{5}, INC_VAL{10} +Output: unsafeAtomicAdd API would work and the 0/P is INITIAL_VAL + INC_VAL + Generate the assembly file and check whether + global_atomic_cmpswap instruction is generated + or not */ + +TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_NonCoherentwithoutflag", "", + float, double) { + hipDeviceProp_t prop; + int device; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + std::string gfxName(prop.gcnArchName); + if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) { + if (prop.canMapHostMemory != 1) { + SUCCEED("Does not support HostPinned Memory"); + } else { + TestType *A_h{nullptr}, *result{nullptr}; + TestType *A_d{nullptr}, *result_d{nullptr}; + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_h), sizeof(TestType), + hipHostMallocNonCoherent)); + A_h[0] = INITIAL_VAL; + HIP_CHECK(hipHostMalloc(reinterpret_cast(&result), + sizeof(TestType), + hipHostMallocNonCoherent)); + result[0] = INITIAL_VAL; + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&A_d), + A_h, 0)); + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&result_d), + result, 0)); + hipLaunchKernelGGL(AtomicCheck, + dim3(1), dim3(1), + 0, 0, A_d, + result_d); + HIP_CHECK(hipDeviceSynchronize()); + bool testResult; + REQUIRE(A_h[0] == INITIAL_VAL + INC_VAL); + REQUIRE(result[0] == INITIAL_VAL); + if ((std::is_same::value)) { + testResult = HipTest::assemblyFile_Verification( + "unsafeAtomicAdd_NonCoherent_withoutflag-hip-amdgcn(.*)\\.s", + "global_atomic_add_f32"); + REQUIRE(testResult == true); + } else { + testResult = HipTest::assemblyFile_Verification( + "unsafeAtomicAdd_NonCoherent_withoutflag-hip-amdgcn(.*)\\.s", + "global_atomic_add_f64"); + REQUIRE(testResult == true); + } + HIP_CHECK(hipHostFree(A_h)); + HIP_CHECK(hipHostFree(result)); + } + } else { + SUCCEED("Memory model feature is only supported for gfx90a, Hence" + "skipping the testcase for this GPU " << device); + } +} diff --git a/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_NonCoherent_withunsafeflag.cc b/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_NonCoherent_withunsafeflag.cc new file mode 100644 index 0000000000..ba4af3df96 --- /dev/null +++ b/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_NonCoherent_withunsafeflag.cc @@ -0,0 +1,97 @@ +/* + 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. + */ +/* +unsafeAtomicAdd on CoarseGrainMemory +1. The following test scenario verifies +unsafeAtomicAdd on CoarseGrain memory with unsafeatomics flag +This testcase works only on gfx90a. +*/ + +#include +#include + +#define INC_VAL 10 +#define INITIAL_VAL 5 +template +static __global__ void AtomicCheck(T* Ad, T* result) { + T inc_val = 10; + *result = unsafeAtomicAdd(Ad, inc_val); +} + +/*unsafeAtomicAdd API for the coarse grained memory variable + with -munsafe-fp-atomics flag +Input: Ad{5}, INC_VAL{10} +Output: unsafeAtomicAdd API would work and the 0/P is INITIAL_VAL + INC_VAL + Generate the assembly file and check whether + global_atomic_add instruction is generated + or not */ + +TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_NonCoherentwithunsafeatomicsflag", "", + float, double) { + hipDeviceProp_t prop; + int device; + HIP_CHECK(hipGetDevice(&device)); + HIP_CHECK(hipGetDeviceProperties(&prop, device)); + std::string gfxName(prop.gcnArchName); + if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) { + if (prop.canMapHostMemory != 1) { + SUCCEED("Does not support HostPinned Memory"); + } else { + TestType *A_h{nullptr}, *result{nullptr}; + TestType *A_d{nullptr}, *result_d{nullptr}; + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_h), sizeof(TestType), + hipHostMallocNonCoherent)); + A_h[0] = INITIAL_VAL; + HIP_CHECK(hipHostMalloc(reinterpret_cast(&result), + sizeof(TestType), + hipHostMallocNonCoherent)); + result[0] = INITIAL_VAL; + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&A_d), + A_h, 0)); + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&result_d), + result, 0)); + hipLaunchKernelGGL(AtomicCheck, + dim3(1), dim3(1), + 0, 0, A_d, + result_d); + HIP_CHECK(hipDeviceSynchronize()); + bool testResult; + REQUIRE(A_h[0] == INITIAL_VAL + INC_VAL); + REQUIRE(result[0] == INITIAL_VAL); + if ((std::is_same::value)) { + testResult = HipTest::assemblyFile_Verification( + "unsafeAtomicAdd_NonCoherent_withunsafeflag-hip-amdgcn(.*)\\.s", + "global_atomic_add_f32"); + REQUIRE(testResult == true); + } else { + testResult = HipTest::assemblyFile_Verification( + "unsafeAtomicAdd_NonCoherent_withunsafeflag-hip-amdgcn(.*)\\.s", + "global_atomic_add_f64"); + REQUIRE(testResult == true); + } + HIP_CHECK(hipHostFree(A_h)); + HIP_CHECK(hipHostFree(result)); + } + } else { + SUCCEED("Memory model feature is only supported for gfx90a, Hence" + "skipping the testcase for this GPU " << device); + } +} diff --git a/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_RTC.cc b/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_RTC.cc new file mode 100644 index 0000000000..e0003e5528 --- /dev/null +++ b/projects/hip-tests/catch/unit/deviceLib/unsafeAtomicAdd_RTC.cc @@ -0,0 +1,583 @@ +/* + Copyright (c) 2021 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. + */ + +/* +unsafeAtomicAdd Scenarios with hipRTC: +1. FineGrainMemory with -m-nounsafe-fp-atomics flag +2. FineGrainMemory without compilation flag +3. FineGrainMemory without -munsafe-fp-atomics flag +4. CoarseGrainMemory with -m-nounsafe-fp-atomics flag +5. CoarseGrainMemory without compilation flag +6. CoarseGrainMemory without -munsafe-fp-atomics flag +*/ + +#include +#include +#include +#define INCREMENT_VAL 10 +#define INITIAL_VAL 5 + +static constexpr auto fkernel{ +R"( +extern "C" +__global__ void AtomicCheck(float* Ad, float *result) { +*result = unsafeAtomicAdd(Ad, 10); +} +)"}; + +static constexpr auto dkernel{ +R"( +extern "C" +__global__ void AtomicCheck(double* Ad, double *result) { +*result = unsafeAtomicAdd(Ad, 10); +} +)"}; + +/* + Test unsafeAtomicAdd API for the fine grained memory variable + where kernel is compiled using hipRTC and with + compilation flag -mno-unsafe-fp-atomics. + Input: Ad{5}, INCREMENT_VAL{10} + Output: unsafeAtomicAdd API will not work and returns 0 so + the initial value will be intact. expected O/P is 5 +*/ +TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_CoherentRTCnounsafeatomicflag", "", + float, double) { + 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; + if (std::is_same::value) { + hiprtcCreateProgram(&prog, // prog + fkernel, // buffer + "kernel.cu", // name + 0, nullptr, nullptr); + } else { + hiprtcCreateProgram(&prog, // prog + dkernel, // buffer + "kernel.cu", // name + 0, nullptr, nullptr); + } + std::string sarg = std::string("--gpu-architecture=") + props.gcnArchName; + const char* options[] = {sarg.c_str(), "-mno-unsafe-fp-atomics"}; + hiprtcResult compileResult{hiprtcCompileProgram(prog, 2, options)}; + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + INFO(log); + } + + REQUIRE(compileResult == HIPRTC_SUCCESS); + size_t codeSize; + HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize)); + + std::vector code(codeSize); + HIPRTC_CHECK(hiprtcGetCode(prog, code.data())); + HIPRTC_CHECK(hiprtcDestroyProgram(&prog)); + + hipModule_t module; + hipFunction_t f_kernel; + HIP_CHECK(hipModuleLoadData(&module, code.data())); + HIP_CHECK(hipModuleGetFunction(&f_kernel, module, "AtomicCheck")); + if (props.canMapHostMemory != 1) { + SUCCEED("Does not support HostPinned Memory"); + } else { + TestType *A_h, *result; + TestType *A_d, *result_d; + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_h), sizeof(TestType), + hipHostMallocCoherent)); + HIP_CHECK(hipHostMalloc(reinterpret_cast(&result), + sizeof(TestType), + hipHostMallocCoherent)); + A_h[0] = INITIAL_VAL; + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&A_d), + A_h, 0)); + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&result_d), + result, 0)); + struct { + TestType* p; + TestType* result; + } args_f{A_d, result_d}; + auto size = sizeof(args_f); + void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f, + HIP_LAUNCH_PARAM_BUFFER_SIZE, + &size, HIP_LAUNCH_PARAM_END}; + hipModuleLaunchKernel(f_kernel, 1, 1, 1, 1, 1, 1, 0, + nullptr, nullptr, config_d); + HIP_CHECK(hipDeviceSynchronize()); + REQUIRE(A_h[0] == INITIAL_VAL); + REQUIRE(*result == 0); + HIP_CHECK(hipHostFree(A_h)); + HIP_CHECK(hipHostFree(result)); + } + HIP_CHECK(hipModuleUnload(module)); + } else { + SUCCEED("Memory model feature is only supported for gfx90a, Hence" + "skipping the testcase for this GPU " << device); + } +} + + +/* + Test unsafeAtomicAdd API for the fine grained memory variable + where kernel is compiled using hipRTC and with + compilation flag -munsafe-fp-atomics. + Input: Ad{5}, INCREMENT_VAL{10} + Output: unsafeAtomicAdd API will not work and r`eturns 0 so + the initial value will be intact. expected O/P is 5 +*/ +TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_CoherentRTCunsafeatomicflag", "", + float, double) { + 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; + if (std::is_same::value) { + hiprtcCreateProgram(&prog, // prog + fkernel, // buffer + "kernel.cu", // name + 0, nullptr, nullptr); + } else { + hiprtcCreateProgram(&prog, // prog + dkernel, // buffer + "kernel.cu", // name + 0, nullptr, nullptr); + } + std::string sarg = std::string("--gpu-architecture=") + props.gcnArchName; + const char* options[] = {sarg.c_str(), "-munsafe-fp-atomics"}; + hiprtcResult compileResult{hiprtcCompileProgram(prog, 2, options)}; + + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + INFO(log); + } + + REQUIRE(compileResult == HIPRTC_SUCCESS); + size_t codeSize; + HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize)); + + std::vector code(codeSize); + HIPRTC_CHECK(hiprtcGetCode(prog, code.data())); + HIPRTC_CHECK(hiprtcDestroyProgram(&prog)); + + hipModule_t module; + hipFunction_t f_kernel; + HIP_CHECK(hipModuleLoadData(&module, code.data())); + HIP_CHECK(hipModuleGetFunction(&f_kernel, module, "AtomicCheck")); + + if (props.canMapHostMemory != 1) { + SUCCEED("Does not support HostPinned Memory"); + } else { + TestType *A_h, *result; + TestType *A_d, *result_d; + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_h), sizeof(TestType), + hipHostMallocCoherent)); + HIP_CHECK(hipHostMalloc(reinterpret_cast(&result), + sizeof(TestType), + hipHostMallocCoherent)); + A_h[0] = INITIAL_VAL; + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&A_d), + A_h, 0)); + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&result_d), + result, 0)); + struct { + TestType* p; + TestType* result; + } args_f{A_d, result_d}; + auto size = sizeof(args_f); + void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f, + HIP_LAUNCH_PARAM_BUFFER_SIZE, + &size, HIP_LAUNCH_PARAM_END}; + hipModuleLaunchKernel(f_kernel, 1, 1, 1, 1, 1, 1, 0, + nullptr, nullptr, config_d); + HIP_CHECK(hipDeviceSynchronize()); + REQUIRE(A_h[0] == INITIAL_VAL); + REQUIRE(*result == 0); + HIP_CHECK(hipHostFree(A_h)); + HIP_CHECK(hipHostFree(result)); + } + HIP_CHECK(hipModuleUnload(module)); + } else { + SUCCEED("Memory model feature is only supported for gfx90a, Hence" + "skipping the testcase for this GPU " << device); + } +} + +/* Test unsafeAtomicAdd API for the fine grained memory variable + where kernel is compiled using hipRTC and without compilation flag + Input: Ad{5}, INCREMENT_VAL{10} + Output: unsafeAtomicAdd API will not work and returns 0 so + the initial value will be intact. expected O/P is 5*/ + +TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_CoherentRTCwithoutflag", "", + float, double) { + 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; + if (std::is_same::value) { + hiprtcCreateProgram(&prog, // prog + fkernel, // buffer + "kernel.cu", // name + 0, nullptr, nullptr); + } else { + hiprtcCreateProgram(&prog, // prog + dkernel, // 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) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + INFO(log); + } + + REQUIRE(compileResult == HIPRTC_SUCCESS); + size_t codeSize; + HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize)); + + std::vector code(codeSize); + HIPRTC_CHECK(hiprtcGetCode(prog, code.data())); + HIPRTC_CHECK(hiprtcDestroyProgram(&prog)); + + hipModule_t module; + hipFunction_t f_kernel; + HIP_CHECK(hipModuleLoadData(&module, code.data())); + HIP_CHECK(hipModuleGetFunction(&f_kernel, module, "AtomicCheck")); + + if (props.canMapHostMemory != 1) { + SUCCEED("Does not support HostPinned Memory"); + } else { + TestType *A_h, *result; + TestType *A_d, *result_d; + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_h), sizeof(float), + hipHostMallocCoherent)); + HIP_CHECK(hipHostMalloc(reinterpret_cast(&result), sizeof(float), + hipHostMallocCoherent)); + A_h[0] = INITIAL_VAL; + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&A_d), + A_h, 0)); + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&result_d), + result, 0)); + struct { + TestType* p; + TestType* result; + } args_f{A_d, result_d}; + auto size = sizeof(args_f); + void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f, + HIP_LAUNCH_PARAM_BUFFER_SIZE, + &size, HIP_LAUNCH_PARAM_END}; + hipModuleLaunchKernel(f_kernel, 1, 1, 1, 1, 1, + 1, 0, nullptr, nullptr, config_d); + HIP_CHECK(hipDeviceSynchronize()); + REQUIRE(A_h[0] == INITIAL_VAL); + REQUIRE(*result == 0); + HIP_CHECK(hipHostFree(A_h)); + HIP_CHECK(hipHostFree(result)); + } + HIP_CHECK(hipModuleUnload(module)); + } else { + SUCCEED("Memory model feature is only supported for gfx90a, Hence" + "skipping the testcase for this GPU " << device); + } +} + +/* + Test unsafeAtomicAdd API for the coarse grained memory variable where kernel + is compiled using hipRTC and with compilation flag -mno-unsafe-fp-atomics + Input: Ad{5}, INCREMENT_VAL{10} + Output: Expected O/P is 15 */ +TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_NonCoherentRTCnounsafeatomicflag", "", + float, double) { + 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; + if (std::is_same::value) { + hiprtcCreateProgram(&prog, // prog + fkernel, // buffer + "kernel.cu", // name + 0, nullptr, nullptr); + } else { + hiprtcCreateProgram(&prog, // prog + dkernel, // buffer + "kernel.cu", // name + 0, nullptr, nullptr); + } + std::string sarg = std::string("--gpu-architecture=") + props.gcnArchName; + const char* options[] = {sarg.c_str(), "-mno-unsafe-fp-atomics"}; + hiprtcResult compileResult{hiprtcCompileProgram(prog, 2, options)}; + + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + INFO(log); + } + + REQUIRE(compileResult == HIPRTC_SUCCESS); + size_t codeSize; + HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize)); + + std::vector code(codeSize); + HIPRTC_CHECK(hiprtcGetCode(prog, code.data())); + HIPRTC_CHECK(hiprtcDestroyProgram(&prog)); + + hipModule_t module; + hipFunction_t f_kernel; + HIP_CHECK(hipModuleLoadData(&module, code.data())); + HIP_CHECK(hipModuleGetFunction(&f_kernel, module, "AtomicCheck")); + if (props.canMapHostMemory != 1) { + SUCCEED("Does not support HostPinned Memory"); + } else { + TestType *A_h, *result; + TestType *A_d, *result_d; + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_h), sizeof(TestType), + hipHostMallocNonCoherent)); + HIP_CHECK(hipHostMalloc(reinterpret_cast(&result), + sizeof(TestType))); + A_h[0] = INITIAL_VAL; + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&A_d), + A_h, 0)); + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&result_d), + result, 0)); + struct { + TestType* p; + TestType* result; + } args_f{A_d, result_d}; + auto size = sizeof(args_f); + void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f, + HIP_LAUNCH_PARAM_BUFFER_SIZE, + &size, HIP_LAUNCH_PARAM_END}; + hipModuleLaunchKernel(f_kernel, 1, 1, 1, 1, 1, 1, 0, + nullptr, nullptr, config_d); + HIP_CHECK(hipDeviceSynchronize()); + REQUIRE(A_h[0] == INITIAL_VAL + INCREMENT_VAL); + REQUIRE(*result == INITIAL_VAL); + HIP_CHECK(hipHostFree(A_h)); + HIP_CHECK(hipHostFree(result)); + } + HIP_CHECK(hipModuleUnload(module)); + } else { + SUCCEED("Memory model feature is only supported for gfx90a, Hence" + "skipping the testcase for this GPU " << device); + } +} + +/* + Test unsafeAtomicAdd API for the coarse grained memory variable where kernel + is compiled using hipRTC and with compilation flag -munsafe-fp-atomics + Input: Ad{5}, INCREMENT_VAL{10} + Output: Expected O/P is 15 */ + +TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_NonCoherentRTCunsafeatomicflag", "", + float, double) { + 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; + if (std::is_same::value) { + hiprtcCreateProgram(&prog, // prog + fkernel, // buffer + "kernel.cu", // name + 0, nullptr, nullptr); + } else { + hiprtcCreateProgram(&prog, // prog + dkernel, // buffer + "kernel.cu", // name + 0, nullptr, nullptr); + } + std::string sarg = std::string("--gpu-architecture=") + props.gcnArchName; + const char* options[] = {sarg.c_str(), "-munsafe-fp-atomics"}; + hiprtcResult compileResult{hiprtcCompileProgram(prog, 2, options)}; + + size_t logSize; + HIPRTC_CHECK(hiprtcGetProgramLogSize(prog, &logSize)); + if (logSize) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + INFO(log); + } + + REQUIRE(compileResult == HIPRTC_SUCCESS); + size_t codeSize; + HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize)); + + std::vector code(codeSize); + HIPRTC_CHECK(hiprtcGetCode(prog, code.data())); + HIPRTC_CHECK(hiprtcDestroyProgram(&prog)); + + hipModule_t module; + hipFunction_t f_kernel; + HIP_CHECK(hipModuleLoadData(&module, code.data())); + HIP_CHECK(hipModuleGetFunction(&f_kernel, module, "AtomicCheck")); + + if (props.canMapHostMemory != 1) { + SUCCEED("Does not support HostPinned Memory"); + } else { + TestType *A_h, *result; + TestType *A_d, *result_d; + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_h), sizeof(TestType), + hipHostMallocNonCoherent)); + HIP_CHECK(hipHostMalloc(reinterpret_cast(&result), + sizeof(TestType))); + A_h[0] = INITIAL_VAL; + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&A_d), + A_h, 0)); + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&result_d), + result, 0)); + struct { + TestType* p; + TestType* result; + } args_f{A_d, result_d}; + auto size = sizeof(args_f); + void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f, + HIP_LAUNCH_PARAM_BUFFER_SIZE, + &size, HIP_LAUNCH_PARAM_END}; + hipModuleLaunchKernel(f_kernel, 1, 1, 1, 1, 1, 1, 0, + nullptr, nullptr, config_d); + HIP_CHECK(hipDeviceSynchronize()); + REQUIRE(A_h[0] == INITIAL_VAL + INCREMENT_VAL); + REQUIRE(*result == INITIAL_VAL); + HIP_CHECK(hipHostFree(A_h)); + HIP_CHECK(hipHostFree(result)); + } + HIP_CHECK(hipModuleUnload(module)); + } else { + SUCCEED("Memory model feature is only supported for gfx90a, Hence" + "skipping the testcase for this GPU " << device); + } +} + +/* + Test unsafeAtomicAdd API for the coarse grained memory variable + where kernel is compiled using hipRTC and without compilation flag + Input: Ad{5}, INCREMENT_VAL{10} + Output: O/P is 15 */ + +TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_NonCoherentRTC", "", + float, double) { + 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; + if (std::is_same::value) { + hiprtcCreateProgram(&prog, // prog + fkernel, // buffer + "kernel.cu", // name + 0, nullptr, nullptr); + } else { + hiprtcCreateProgram(&prog, // prog + dkernel, // 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) { + std::string log(logSize, '\0'); + HIPRTC_CHECK(hiprtcGetProgramLog(prog, &log[0])); + INFO(log); + } + + REQUIRE(compileResult == HIPRTC_SUCCESS); + size_t codeSize; + HIPRTC_CHECK(hiprtcGetCodeSize(prog, &codeSize)); + std::vector code(codeSize); + HIPRTC_CHECK(hiprtcGetCode(prog, code.data())); + HIPRTC_CHECK(hiprtcDestroyProgram(&prog)); + + hipModule_t module; + hipFunction_t f_kernel; + HIP_CHECK(hipModuleLoadData(&module, code.data())); + HIP_CHECK(hipModuleGetFunction(&f_kernel, module, "AtomicCheck")); + + if (props.canMapHostMemory != 1) { + SUCCEED("Does not support HostPinned Memory"); + } else { + TestType *A_h, *result; + TestType *A_d, *result_d; + HIP_CHECK(hipHostMalloc(reinterpret_cast(&A_h), sizeof(TestType), + hipHostMallocNonCoherent)); + HIP_CHECK(hipHostMalloc(reinterpret_cast(&result), + sizeof(TestType))); + A_h[0] = INITIAL_VAL; + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&A_d), + A_h, 0)); + HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast(&result_d), + result, 0)); + struct { + TestType* p; + TestType* result; + } args_f{A_d, result_d}; + auto size = sizeof(args_f); + void* config_d[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args_f, + HIP_LAUNCH_PARAM_BUFFER_SIZE, + &size, HIP_LAUNCH_PARAM_END}; + hipModuleLaunchKernel(f_kernel, 1, 1, 1, 1, 1, 1, 0, + nullptr, nullptr, config_d); + HIP_CHECK(hipDeviceSynchronize()); + REQUIRE(A_h[0] == INITIAL_VAL + INCREMENT_VAL); + REQUIRE(*result == INITIAL_VAL); + HIP_CHECK(hipHostFree(A_h)); + HIP_CHECK(hipHostFree(result)); + } + HIP_CHECK(hipModuleUnload(module)); + } else { + SUCCEED("Memory model feature is only supported for gfx90a, Hence" + "skipping the testcase for this GPU " << device); + } +}