SWDEV-293806 - Added tests for AtomicAdd with and without unsafeAtomics flag (#2603)
Added tests for AtomicAdd API for coherent/non-coherent memory with and without unsafeAtomics flags Change-Id: I1937d7e936ac062b8d93feb59a50a6dd8ae9feca
Este cometimento está contido em:
ROCm CI Service Account
cometido por
GitHub
GitHub
ascendente
dd31050d4c
cometimento
d9b40f34a4
@@ -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 <iostream>
|
||||
#include<fstream>
|
||||
#include<regex>
|
||||
#include <type_traits>
|
||||
|
||||
#define guarantee(cond, str) \
|
||||
{ \
|
||||
if (!(cond)) { \
|
||||
@@ -235,5 +239,55 @@ unsigned setNumBlocks(T blocksPerCU, T threadsPerBlock,
|
||||
}
|
||||
return blocks;
|
||||
}
|
||||
|
||||
template<typename T>
|
||||
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<T, float>::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
|
||||
|
||||
@@ -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<hip_test_checkers.hh>
|
||||
#include<hip_test_common.hh>
|
||||
|
||||
|
||||
#define INC_VAL 10
|
||||
#define INITIAL_VAL 5
|
||||
|
||||
template<typename T>
|
||||
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<void**>(&A_h), sizeof(TestType),
|
||||
hipHostMallocCoherent));
|
||||
A_h[0] = INITIAL_VAL;
|
||||
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&result),
|
||||
sizeof(TestType),
|
||||
hipHostMallocCoherent));
|
||||
result[0] = INITIAL_VAL;
|
||||
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d),
|
||||
A_h, 0));
|
||||
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&result_d),
|
||||
result, 0));
|
||||
hipLaunchKernelGGL(AtomicCheck<TestType>, dim3(1), dim3(1),
|
||||
0, 0, A_d,
|
||||
result_d);
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
bool testResult;
|
||||
testResult = HipTest::assemblyFile_Verification<TestType>(
|
||||
"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);
|
||||
}
|
||||
}
|
||||
@@ -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<hip_test_checkers.hh>
|
||||
#include<hip_test_common.hh>
|
||||
|
||||
|
||||
#define INC_VAL 10
|
||||
#define INITIAL_VAL 5
|
||||
template<typename T>
|
||||
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<void**>(&A_h), sizeof(TestType),
|
||||
hipHostMallocCoherent));
|
||||
A_h[0] = INITIAL_VAL;
|
||||
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&result),
|
||||
sizeof(TestType),
|
||||
hipHostMallocCoherent));
|
||||
result[0] = INITIAL_VAL;
|
||||
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d),
|
||||
A_h, 0));
|
||||
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&result_d),
|
||||
result, 0));
|
||||
hipLaunchKernelGGL(AtomicCheck<TestType>, dim3(1), dim3(1),
|
||||
0, 0, A_d,
|
||||
result_d);
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
bool testResult;
|
||||
testResult = HipTest::assemblyFile_Verification<TestType>(
|
||||
"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);
|
||||
}
|
||||
}
|
||||
@@ -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<hip_test_checkers.hh>
|
||||
#include<hip_test_common.hh>
|
||||
|
||||
#define INC_VAL 10
|
||||
#define INITIAL_VAL 5
|
||||
template<typename T>
|
||||
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<void**>(&A_h), sizeof(TestType),
|
||||
hipHostMallocCoherent));
|
||||
A_h[0] = INITIAL_VAL;
|
||||
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&result),
|
||||
sizeof(TestType),
|
||||
hipHostMallocCoherent));
|
||||
result[0] = INITIAL_VAL;
|
||||
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d),
|
||||
A_h, 0));
|
||||
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&result_d),
|
||||
result, 0));
|
||||
hipLaunchKernelGGL(AtomicCheck<TestType>, dim3(1), dim3(1),
|
||||
0, 0, A_d,
|
||||
result_d);
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
bool testResult;
|
||||
|
||||
if ((std::is_same<TestType, float>::value)) {
|
||||
testResult = HipTest::assemblyFile_Verification<TestType>(
|
||||
"AtomicAdd_Coherent_withunsafeflag-hip-amdgcn(.*)\\.s",
|
||||
"global_atomic_add_f32");
|
||||
REQUIRE(testResult == true);
|
||||
} else {
|
||||
testResult = HipTest::assemblyFile_Verification<TestType>(
|
||||
"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);
|
||||
}
|
||||
}
|
||||
@@ -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<hip_test_checkers.hh>
|
||||
#include<hip_test_common.hh>
|
||||
|
||||
|
||||
#define INC_VAL 10
|
||||
#define INITIAL_VAL 5
|
||||
template<typename T>
|
||||
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<void**>(&A_h), sizeof(TestType),
|
||||
hipHostMallocNonCoherent));
|
||||
A_h[0] = INITIAL_VAL;
|
||||
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&result),
|
||||
sizeof(TestType),
|
||||
hipHostMallocNonCoherent));
|
||||
result[0] = INITIAL_VAL;
|
||||
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d),
|
||||
A_h, 0));
|
||||
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&result_d),
|
||||
result, 0));
|
||||
hipLaunchKernelGGL(AtomicCheck<TestType>,
|
||||
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<TestType>(
|
||||
"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);
|
||||
}
|
||||
}
|
||||
@@ -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<hip_test_checkers.hh>
|
||||
#include<hip_test_common.hh>
|
||||
|
||||
#define INC_VAL 10
|
||||
#define INITIAL_VAL 5
|
||||
|
||||
template<typename T>
|
||||
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<void**>(&A_h), sizeof(TestType),
|
||||
hipHostMallocNonCoherent));
|
||||
A_h[0] = INITIAL_VAL;
|
||||
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&result),
|
||||
sizeof(TestType),
|
||||
hipHostMallocNonCoherent));
|
||||
result[0] = INITIAL_VAL;
|
||||
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d),
|
||||
A_h, 0));
|
||||
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&result_d),
|
||||
result, 0));
|
||||
hipLaunchKernelGGL(AtomicCheck<TestType>,
|
||||
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<TestType>(
|
||||
"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);
|
||||
}
|
||||
}
|
||||
@@ -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<hip_test_checkers.hh>
|
||||
#include<hip_test_common.hh>
|
||||
|
||||
|
||||
#define INC_VAL 10
|
||||
#define INITIAL_VAL 5
|
||||
template<typename T>
|
||||
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<void**>(&A_h), sizeof(TestType),
|
||||
hipHostMallocNonCoherent));
|
||||
A_h[0] = INITIAL_VAL;
|
||||
HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&result),
|
||||
sizeof(TestType),
|
||||
hipHostMallocNonCoherent));
|
||||
result[0] = INITIAL_VAL;
|
||||
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&A_d),
|
||||
A_h, 0));
|
||||
HIP_CHECK(hipHostGetDevicePointer(reinterpret_cast<void**>(&result_d),
|
||||
result, 0));
|
||||
hipLaunchKernelGGL(AtomicCheck<TestType>, 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<TestType, float>::value)) {
|
||||
testResult = HipTest::assemblyFile_Verification<TestType>(
|
||||
"AtomicAdd_NonCoherent_withunsafeflag-hip-amdgcn(.*)\\.s",
|
||||
"global_atomic_add_f32");
|
||||
REQUIRE(testResult == true);
|
||||
} else {
|
||||
testResult = HipTest::assemblyFile_Verification<TestType>(
|
||||
"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);
|
||||
}
|
||||
}
|
||||
@@ -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)
|
||||
|
||||
Criar uma nova questão referindo esta
Bloquear um utilizador