diff --git a/tests/catch/include/hip_test_checkers.hh b/tests/catch/include/hip_test_checkers.hh index 369b8b810a..77ac0f4d74 100644 --- a/tests/catch/include/hip_test_checkers.hh +++ b/tests/catch/include/hip_test_checkers.hh @@ -1,5 +1,5 @@ /* -Copyright (c) 2021 Advanced Micro Devices, Inc. All rights reserved. +Copyright (c) 2021 - 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 @@ -23,6 +23,10 @@ THE SOFTWARE. #pragma once #include "hip_test_common.hh" #include +#include +#include +#include + #define guarantee(cond, str) \ { \ if (!(cond)) { \ @@ -235,5 +239,55 @@ unsigned setNumBlocks(T blocksPerCU, T threadsPerBlock, } return blocks; } - +template +static bool assemblyFile_Verification(std::string assemfilename, std::string inst) { + std::string filePath = "./catch/unit/deviceLib/"; + bool result = false; + std::string filename; + filename = filePath + assemfilename; + std::ifstream file(filename.c_str(), std::ios::out); + if (file) { + std::string line; + int line_pos = 0, start_pos = 0; + int last_pos = 0; + int start_match = 0; + while (getline(file, line)) { + line_pos++; + if ((std::is_same::value)) { + if (!start_pos && + std::regex_search(line, + std::regex("Begin function (.*)AtomicCheck"))) { + start_pos = line_pos; + } + if (!last_pos && + std::regex_search(line, + std::regex(".Lfunc_end0-(.*)AtomicCheck"))) { + last_pos = line_pos; + break; + } + } else { + if ((start_match != 2) && std::regex_search(line, + std::regex("Begin function (.*)AtomicCheck"))) { + start_match++; + if (start_match == 2) + start_pos = line_pos; + } + if (!last_pos && std::regex_search(line, + std::regex("func_end1-(.*)AtomicCheck"))) { + last_pos = line_pos; + break; + } + } + if (start_pos) { + result = std::regex_search(line, std::regex(inst)); + if (result) + break; + } + } + } else { + result = true; + SUCCEED("Assembly file does not exist"); + } + return result; +} } // namespace HipTest diff --git a/tests/catch/unit/deviceLib/AtomicAdd_Coherent_withnoUnsafeflag.cc b/tests/catch/unit/deviceLib/AtomicAdd_Coherent_withnoUnsafeflag.cc new file mode 100644 index 0000000000..3622300514 --- /dev/null +++ b/tests/catch/unit/deviceLib/AtomicAdd_Coherent_withnoUnsafeflag.cc @@ -0,0 +1,91 @@ +/* + 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 +atomicAdd on fineGrain memory with -mno-unsafe-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 = atomicAdd(Ad, inc_val); +} + +/*atomicAdd API for the fine grained memory variable + with -mno-unsafe-atomics flag +Input: Ad{5}, INC_VAL{10} +Output: atomicAdd 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_AtomicAdd_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; + testResult = HipTest::assemblyFile_Verification( + "AtomicAdd_Coherent_withnoUnsafeflag-hip-amdgcn(.*)\\.s", + "global_atomic_cmpswap"); + REQUIRE(testResult == true); + REQUIRE(A_h[0] == INITIAL_VAL + INC_VAL); + REQUIRE(result[0] == INITIAL_VAL); + 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/tests/catch/unit/deviceLib/AtomicAdd_Coherent_withoutflag.cc b/tests/catch/unit/deviceLib/AtomicAdd_Coherent_withoutflag.cc new file mode 100644 index 0000000000..4d12ea1f2c --- /dev/null +++ b/tests/catch/unit/deviceLib/AtomicAdd_Coherent_withoutflag.cc @@ -0,0 +1,90 @@ +/* + 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 +atomicAdd on fineGrain 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 = atomicAdd(Ad, inc_val); +} + +/*atomicAdd API for the fine grained memory variable + without any flag +Input: Ad{5}, INC_VAL{10} +Output: atomicAdd 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_AtomicAdd_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; + testResult = HipTest::assemblyFile_Verification( + "AtomicAdd_Coherent_withoutflag-hip-amdgcn(.*)\\.s", + "global_atomic_cmpswap"); + REQUIRE(result[0] == INITIAL_VAL); + REQUIRE(A_h[0] == INITIAL_VAL + INC_VAL); + 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/tests/catch/unit/deviceLib/AtomicAdd_Coherent_withunsafeflag.cc b/tests/catch/unit/deviceLib/AtomicAdd_Coherent_withunsafeflag.cc new file mode 100644 index 0000000000..27120d930e --- /dev/null +++ b/tests/catch/unit/deviceLib/AtomicAdd_Coherent_withunsafeflag.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 +atomicAdd 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 = atomicAdd(Ad, inc_val); +} + + +/*atomicAdd 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 + global_atomic_cmpswap instruction is generated + or not */ + +TEMPLATE_TEST_CASE("Unit_AtomicAdd_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( + "AtomicAdd_Coherent_withunsafeflag-hip-amdgcn(.*)\\.s", + "global_atomic_add_f32"); + REQUIRE(testResult == true); + } else { + testResult = HipTest::assemblyFile_Verification( + "AtomicAdd_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/tests/catch/unit/deviceLib/AtomicAdd_NonCoherent_withnoUnsafeflag.cc b/tests/catch/unit/deviceLib/AtomicAdd_NonCoherent_withnoUnsafeflag.cc new file mode 100644 index 0000000000..01f6e6ec07 --- /dev/null +++ b/tests/catch/unit/deviceLib/AtomicAdd_NonCoherent_withnoUnsafeflag.cc @@ -0,0 +1,91 @@ +/* + 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 +atomicAdd on CoarseGrain memory with -mno-unsafe-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 = atomicAdd(Ad, inc_val); +} + +/*atomicAdd API for the coarse grained memory variable + with -mno-unsafe-atomics flag +Input: Ad{5}, INC_VAL{10} +Output: atomicAdd 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_AtomicAdd_NonCoherentwithnounsafeflag", "", + 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); + testResult = HipTest::assemblyFile_Verification( + "AtomicAdd_NonCoherent_withnounsafeflag-hip-amdgcn(.*)\\.s", + "global_atomic_cmpswap"); + 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/tests/catch/unit/deviceLib/AtomicAdd_NonCoherent_withoutflag.cc b/tests/catch/unit/deviceLib/AtomicAdd_NonCoherent_withoutflag.cc new file mode 100644 index 0000000000..a5d4c58e94 --- /dev/null +++ b/tests/catch/unit/deviceLib/AtomicAdd_NonCoherent_withoutflag.cc @@ -0,0 +1,91 @@ +/* + 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 +atomicAdd 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 = atomicAdd(Ad, inc_val); +} + +/*atomicAdd API for the coarse grained memory variable + without any flag +Input: Ad{5}, INC_VAL{10} +Output: atomicAdd 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_AtomicAdd_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); + testResult = HipTest::assemblyFile_Verification( + "AtomicAdd_NonCoherent_withoutflag-hip-amdgcn(.*)\\.s", + "global_atomic_cmpswap"); + 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/tests/catch/unit/deviceLib/AtomicAdd_NonCoherent_withunsafeflag.cc b/tests/catch/unit/deviceLib/AtomicAdd_NonCoherent_withunsafeflag.cc new file mode 100644 index 0000000000..59e4ccfeb4 --- /dev/null +++ b/tests/catch/unit/deviceLib/AtomicAdd_NonCoherent_withunsafeflag.cc @@ -0,0 +1,98 @@ +/* + 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 +atomicAdd on CoarseGrain 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 = atomicAdd(Ad, inc_val); +} + +/*atomicAdd API for the fine grained memory variable + with -m-unsafe-atomics flag +Input: Ad{5}, INC_VAL{10} +Output: atomicAdd API would work and the 0/P is 15 + Generate the assembly file and check whether + global_atomic_add_float/double instruction is generated + or not */ + +TEMPLATE_TEST_CASE("Unit_AtomicAdd_NonCoherentwithUnsafeflag", "", + 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( + "AtomicAdd_NonCoherent_withunsafeflag-hip-amdgcn(.*)\\.s", + "global_atomic_add_f32"); + REQUIRE(testResult == true); + } else { + testResult = HipTest::assemblyFile_Verification( + "AtomicAdd_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/tests/catch/unit/deviceLib/CMakeLists.txt b/tests/catch/unit/deviceLib/CMakeLists.txt index 2f6819f43f..644569e49e 100644 --- a/tests/catch/unit/deviceLib/CMakeLists.txt +++ b/tests/catch/unit/deviceLib/CMakeLists.txt @@ -10,7 +10,6 @@ set(TEST_SRC popc.cc ldg.cc threadfence_system.cc - hipTestDeviceSymbol.cc ) # skipped for windows compiler issue - Illegal instruction detected @@ -31,10 +30,28 @@ set(AMD_TEST_SRC bitInsert.cc floatTM.cc ) +set(AMD_ARCH_SPEC_TEST_SRC + AtomicAdd_Coherent_withunsafeflag.cc + AtomicAdd_Coherent_withoutflag.cc + AtomicAdd_Coherent_withnoUnsafeflag.cc + AtomicAdd_NonCoherent_withoutflag.cc + AtomicAdd_NonCoherent_withnoUnsafeflag.cc + AtomicAdd_NonCoherent_withunsafeflag.cc +) if(HIP_PLATFORM MATCHES "amd") + string(FIND ${OFFLOAD_ARCH_STR} "gfx90a" ARCH_CHECK) set(TEST_SRC ${TEST_SRC} ${AMD_TEST_SRC}) set_source_files_properties(floatTM.cc PROPERTIES COMPILE_FLAGS -std=c++17) +if(${ARCH_CHECK} GREATER_EQUAL 0) + set(TEST_SRC ${TEST_SRC} ${AMD_ARCH_SPEC_TEST_SRC}) + set_source_files_properties(AtomicAdd_Coherent_withunsafeflag.cc PROPERTIES COMPILE_OPTIONS "-munsafe-fp-atomics") + 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") + file(GLOB AtomicAdd_files *AtomicAdd_*_*.cc) + set_property(SOURCE ${AtomicAdd_files} PROPERTY COMPILE_FLAGS --save-temps) +endif() hip_add_exe_to_target(NAME UnitDeviceTests TEST_SRC ${TEST_SRC} TEST_TARGET_NAME build_tests)