AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-393910 - Port gfx940 changes to mainline. (#321)

Browse Source 
Change-Id: I77e70bcb15fbd38b1c79c1f68b7e483b2b3e0daf
This commit is contained in:
ROCm CI Service Account
2023-06-29 08:57:51 +05:30
committed by GitHub
parent a7ab47589b
commit d33ccb3569
23 changed files with 190 additions and 74 deletions
Download Patch File Download Diff File Expand all files Collapse all files
catch/unit/deviceLib/AtomicAdd_Coherent_withnoUnsafeflag.cc
+7 -5
View File
@@ -22,11 +22,12 @@
AtomicAdd on FineGrainMemory
1. The following test scenario verifies
atomicAdd on fineGrain memory with -mno-unsafe-atomics flag
This testcase works only on gfx90a.
This testcase works only on gfx90a, gfx940, gfx941, gfx942.
*/
#include<hip_test_checkers.hh>
#include<hip_test_common.hh>
#include <hip_test_checkers.hh>
#include <hip_test_common.hh>
#include <hip_test_features.hh>
#define INC_VAL 10
@@ -53,7 +54,8 @@ TEMPLATE_TEST_CASE("Unit_AtomicAdd_Coherentwithnounsafeflag", "",
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (CheckIfFeatSupported(CTFeatures::CT_FEATURE_FINEGRAIN_HWSUPPORT, gfxName)) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
@@ -86,7 +88,7 @@ TEMPLATE_TEST_CASE("Unit_AtomicAdd_Coherentwithnounsafeflag", "",
HIP_CHECK(hipHostFree(result));
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
SUCCEED("Memory model feature is only supported for gfx90a, gfx940, gfx941, gfx942, Hence"
"skipping the testcase for this GPU " << device);
}
}
catch/unit/deviceLib/AtomicAdd_Coherent_withoutflag.cc
+4 -4
View File
@@ -22,12 +22,12 @@
AtomicAdd on FineGrainMemory
1. The following test scenario verifies
atomicAdd on fineGrain memory without any unsafeatomics flag
This testcase works only on gfx90a.
This testcase works only on gfx90a, gfx940, gfx941, gfx942.
*/
#include<hip_test_checkers.hh>
#include<hip_test_common.hh>
#include <hip_test_features.hh>
#define INC_VAL 10
#define INITIAL_VAL 5
@@ -52,7 +52,7 @@ TEMPLATE_TEST_CASE("Unit_AtomicAdd_Coherentwithoutflag", "",
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (CheckIfFeatSupported(CTFeatures::CT_FEATURE_FINEGRAIN_HWSUPPORT, gfxName)) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
@@ -85,7 +85,7 @@ TEMPLATE_TEST_CASE("Unit_AtomicAdd_Coherentwithoutflag", "",
HIP_CHECK(hipHostFree(result));
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
SUCCEED("Memory model feature is only supported for gfx90a, gfx940, gfx941, gfx942, Hence"
"skipping the testcase for this GPU " << device);
}
}
catch/unit/deviceLib/AtomicAdd_Coherent_withunsafeflag.cc
+4 -3
View File
@@ -22,11 +22,12 @@
AtomicAdd on FineGrainMemory
1. The following test scenario verifies
atomicAdd on fineGrain memory with -munsafe-fp-atomics flag
This testcase works only on gfx90a.
This testcase works only on gfx90a, gfx940, gfx941, gfx942.
*/
#include<hip_test_checkers.hh>
#include<hip_test_common.hh>
#include<hip_test_features.hh>
#define INC_VAL 10
#define INITIAL_VAL 5
@@ -52,7 +53,7 @@ TEMPLATE_TEST_CASE("Unit_AtomicAdd_CoherentwithUnsafeflag", "",
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (CheckIfFeatSupported(CTFeatures::CT_FEATURE_FINEGRAIN_HWSUPPORT, gfxName)) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
@@ -94,7 +95,7 @@ TEMPLATE_TEST_CASE("Unit_AtomicAdd_CoherentwithUnsafeflag", "",
HIP_CHECK(hipHostFree(result));
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
SUCCEED("Memory model feature is only supported for gfx90a, gfx940, gfx941, gfx942, Hence"
"skipping the testcase for this GPU " << device);
}
}
catch/unit/deviceLib/AtomicAdd_NonCoherent_withnoUnsafeflag.cc
+4 -4
View File
@@ -22,12 +22,12 @@
AtomicAdd on CoarseGrainMemory
1. The following test scenario verifies
atomicAdd on CoarseGrain memory with -mno-unsafe-atomics flag
This testcase works only on gfx90a.
This testcase works only on gfx90a, gfx940, gfx941, gfx942.
*/
#include<hip_test_checkers.hh>
#include<hip_test_common.hh>
#include<hip_test_features.hh>
#define INC_VAL 10
#define INITIAL_VAL 5
@@ -52,7 +52,7 @@ TEMPLATE_TEST_CASE("Unit_AtomicAdd_NonCoherentwithnounsafeflag", "",
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (CheckIfFeatSupported(CTFeatures::CT_FEATURE_FINEGRAIN_HWSUPPORT, gfxName)) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
@@ -86,7 +86,7 @@ TEMPLATE_TEST_CASE("Unit_AtomicAdd_NonCoherentwithnounsafeflag", "",
HIP_CHECK(hipHostFree(result));
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
SUCCEED("Memory model feature is only supported for gfx90a, gfx940, gfx941, gfx942, Hence"
"skipping the testcase for this GPU " << device);
}
}
catch/unit/deviceLib/AtomicAdd_NonCoherent_withoutflag.cc
+4 -3
View File
@@ -22,11 +22,12 @@
AtomicAdd on CoarseGrainMemory
1. The following test scenario verifies
atomicAdd on CoarseGrain memory without any unsafeatomics flag
This testcase works only on gfx90a.
This testcase works only on gfx90a, gfx940, gfx941, gfx942.
*/
#include<hip_test_checkers.hh>
#include<hip_test_common.hh>
#include<hip_test_features.hh>
#define INC_VAL 10
#define INITIAL_VAL 5
@@ -52,7 +53,7 @@ TEMPLATE_TEST_CASE("Unit_AtomicAdd_NonCoherentwithoutflag", "",
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (CheckIfFeatSupported(CTFeatures::CT_FEATURE_FINEGRAIN_HWSUPPORT, gfxName)) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
@@ -86,7 +87,7 @@ TEMPLATE_TEST_CASE("Unit_AtomicAdd_NonCoherentwithoutflag", "",
HIP_CHECK(hipHostFree(result));
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
SUCCEED("Memory model feature is only supported for gfx90a, gfx940, gfx941, gfx942, Hence"
"skipping the testcase for this GPU " << device);
}
}
catch/unit/deviceLib/AtomicAdd_NonCoherent_withunsafeflag.cc
+4 -4
View File
@@ -22,12 +22,12 @@
AtomicAdd on CoarseGrainMemory
1. The following test scenario verifies
atomicAdd on CoarseGrain memory with -munsafe-fp-atomics flag
This testcase works only on gfx90a.
This testcase works only on gfx90a, gfx940, gfx941, gfx942.
*/
#include<hip_test_checkers.hh>
#include<hip_test_common.hh>
#include<hip_test_features.hh>
#define INC_VAL 10
#define INITIAL_VAL 5
@@ -52,7 +52,7 @@ TEMPLATE_TEST_CASE("Unit_AtomicAdd_NonCoherentwithUnsafeflag", "",
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (CheckIfFeatSupported(CTFeatures::CT_FEATURE_FINEGRAIN_HWSUPPORT, gfxName)) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
@@ -93,7 +93,7 @@ TEMPLATE_TEST_CASE("Unit_AtomicAdd_NonCoherentwithUnsafeflag", "",
HIP_CHECK(hipHostFree(result));
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
SUCCEED("Memory model feature is only supported for gfx90a, gfx940, gfx941, gfx942, Hence"
"skipping the testcase for this GPU " << device);
}
}
catch/unit/deviceLib/BuiltIns_fadd.cc
+9 -8
View File
@@ -28,6 +28,7 @@ This testfile verifies __builtin_amdgcn_global_atomic_fadd_f64 API scenarios
#include<hip_test_checkers.hh>
#include<hip_test_common.hh>
#include<hip_test_features.hh>
#include <hip/hiprtc.h>
#define INC_VAL 10
@@ -57,7 +58,7 @@ TEST_CASE("Unit_BuiltInAtomicAdd_CoherentGlobalMem") {
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (CheckIfFeatSupported(CTFeatures::CT_FEATURE_FINEGRAIN_HWSUPPORT, gfxName)) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does support HostPinned Memory");
} else {
@@ -85,7 +86,7 @@ TEST_CASE("Unit_BuiltInAtomicAdd_CoherentGlobalMem") {
HIP_CHECK(hipFree(result));
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a Hence"
SUCCEED("Memory model feature is only supported for gfx90a, gfx940, gfx941, gfx942 Hence"
"skipping the testcase for this GPU " << device);
}
}
@@ -103,7 +104,7 @@ TEST_CASE("Unit_BuiltInAtomicAdd_NonCoherentGlobalMem") {
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (CheckIfFeatSupported(CTFeatures::CT_FEATURE_FINEGRAIN_HWSUPPORT, gfxName)) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
@@ -129,7 +130,7 @@ TEST_CASE("Unit_BuiltInAtomicAdd_NonCoherentGlobalMem") {
free(B_h);
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a"
SUCCEED("Memory model feature is only supported for gfx90a, gfx940, gfx941, gfx942"
"Hence skipping the testcase for GPU-0");
}
}
@@ -146,7 +147,7 @@ TEST_CASE("Unit_BuiltInAtomicAdd_CoherentGlobalMemWithRtc") {
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (CheckIfFeatSupported(CTFeatures::CT_FEATURE_FINEGRAIN_HWSUPPORT, gfxName)) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
@@ -208,7 +209,7 @@ TEST_CASE("Unit_BuiltInAtomicAdd_CoherentGlobalMemWithRtc") {
free(B_h);
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
SUCCEED("Memory model feature is only supported for gfx90a, gfx940, gfx941, gfx942, Hence"
"skipping the testcase for this GPU " << device);
}
}
@@ -226,7 +227,7 @@ TEST_CASE("Unit_BuiltInAtomicAdd_NonCoherentGlobalMemWithRtc") {
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (CheckIfFeatSupported(CTFeatures::CT_FEATURE_FINEGRAIN_HWSUPPORT, gfxName)) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does support HostPinned Memory");
} else {
@@ -288,7 +289,7 @@ TEST_CASE("Unit_BuiltInAtomicAdd_NonCoherentGlobalMemWithRtc") {
free(B_h);
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
SUCCEED("Memory model feature is only supported for gfx90a, gfx940, gfx941, gfx942, Hence"
"skipping the testcase for this GPU " << device);
}
}
catch/unit/deviceLib/CMakeLists.txt
+17 -2
View File
@@ -86,11 +86,26 @@ add_custom_target(kerDevAllocSingleKer.code
-I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/
-I${CMAKE_CURRENT_SOURCE_DIR}/../../include --rocm-path=${ROCM_PATH})
# Accepted archs to compile this cmake file
set(ACCEPTED_OFFLOAD_ARCHS gfx90a gfx940 gfx941 gfx942)
function(CheckAcceptedArchs OFFLOAD_ARCH_STR_LOCAL)
set(ARCH_CHECK -1 PARENT_SCOPE)
string(REGEX MATCHALL "--offload-arch=gfx[0-9a-z]+" OFFLOAD_ARCH_LIST ${OFFLOAD_ARCH_STR_LOCAL})
foreach(OFFLOAD_ARCH IN LISTS OFFLOAD_ARCH_LIST)
string(REGEX MATCHALL "--offload-arch=(gfx[0-9a-z]+)" matches ${OFFLOAD_ARCH})
if (CMAKE_MATCH_COUNT EQUAL 1)
if (CMAKE_MATCH_1 IN_LIST ACCEPTED_OFFLOAD_ARCHS)
set(ARCH_CHECK 1 PARENT_SCOPE)
endif() # CMAKE_MATCH_1
endif() # CMAKE_MATCH_COUNT
endforeach() # OFFLOAD_ARCH_LIST
endfunction() # CheckAcceptedArchs
if(HIP_PLATFORM MATCHES "amd")
if (DEFINED OFFLOAD_ARCH_STR)
string(FIND ${OFFLOAD_ARCH_STR} "gfx90a" ARCH_CHECK)
CheckAcceptedArchs(${OFFLOAD_ARCH_STR})
elseif(DEFINED $ENV{HCC_AMDGPU_TARGET})
string(FIND $ENV{HCC_AMDGPU_TARGET} "gfx90a" ARCH_CHECK)
CheckAcceptedArchs($ENV{HCC_AMDGPU_TARGET})
else()
set(ARCH_CHECK -1)
endif()
catch/unit/deviceLib/unsafeAtomicAdd.cc
+2 -1
View File
@@ -21,6 +21,7 @@ THE SOFTWARE.
*/
#include <hip_test_common.hh>
#include <hip_test_features.hh>
#include <hip/hiprtc.h>
#include <hip/hip_runtime.h>
@@ -52,7 +53,7 @@ TEST_CASE("Unit_unsafeAtomicAdd") {
HIP_CHECK(hipGetDeviceProperties(&props, device));
std::string gfxName(props.gcnArchName);
if (gfxName == "gfx90a" || gfxName.find("gfx90a:") == 0) {
if (CheckIfFeatSupported(CTFeatures::CT_FEATURE_FINEGRAIN_HWSUPPORT, gfxName)) {
hiprtcProgram prog;
hiprtcCreateProgram(&prog, // prog
kernel, // buffer
catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withnounsafeflag.cc
+4 -3
View File
@@ -22,11 +22,12 @@
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.
This testcase works only on gfx90a, gfx940, gfx941, gfx942.
*/
#include<hip_test_common.hh>
#include<hip_test_checkers.hh>
#include<hip_test_features.hh>
#define INC_VAL 10
#define INITIAL_VAL 5
@@ -52,7 +53,7 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_CoherentwithnoUnsafeflag", "",
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (CheckIfFeatSupported(CTFeatures::CT_FEATURE_FINEGRAIN_HWSUPPORT, gfxName)) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
@@ -94,7 +95,7 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_CoherentwithnoUnsafeflag", "",
HIP_CHECK(hipHostFree(result));
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
SUCCEED("Memory model feature is only supported for gfx90a, gfx940, gfx941, gfx942, Hence"
"skipping the testcase for this GPU " << device);
}
}
catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withoutflag.cc
+4 -3
View File
@@ -22,11 +22,12 @@
AtomicAdd on FineGrainMemory
1. The following test scenario verifies
unsafeatomicAdd on fineGrain memory without atomics flag
This testcase works only on gfx90a.
This testcase works only on gfx90a, gfx940, gfx941, gfx942.
*/
#include<hip_test_checkers.hh>
#include<hip_test_common.hh>
#include<hip_test_features.hh>
#define INC_VAL 10
#define INITIAL_VAL 5
@@ -52,7 +53,7 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_Coherentwithoutflag", "",
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (CheckIfFeatSupported(CTFeatures::CT_FEATURE_FINEGRAIN_HWSUPPORT, gfxName)) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
@@ -94,7 +95,7 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_Coherentwithoutflag", "",
HIP_CHECK(hipHostFree(result));
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
SUCCEED("Memory model feature is only supported for gfx90a, gfx940, gfx941, gfx942, Hence"
"skipping the testcase for this GPU " << device);
}
}
catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withunsafeflag.cc
+3 -2
View File
@@ -27,6 +27,7 @@ This testcase works only on gfx90a.
#include<hip_test_checkers.hh>
#include<hip_test_common.hh>
#include<hip_test_features.hh>
#define INC_VAL 10
#define INITIAL_VAL 5
@@ -53,7 +54,7 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_CoherentwithUnsafeflag", "",
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (CheckIfFeatSupported(CTFeatures::CT_FEATURE_FINEGRAIN_HWSUPPORT, gfxName)) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
@@ -95,7 +96,7 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_CoherentwithUnsafeflag", "",
HIP_CHECK(hipHostFree(result));
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
SUCCEED("Memory model feature is only supported for gfx90a, gfx940, gfx941, gfx942, Hence"
"skipping the testcase for this GPU " << device);
}
}
catch/unit/deviceLib/unsafeAtomicAdd_NonCoherent_withnounsafeflag.cc
+4 -3
View File
@@ -22,11 +22,12 @@
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.
This testcase works only on gfx90a, gfx940, gfx941, gfx942.
*/
#include<hip_test_checkers.hh>
#include<hip_test_common.hh>
#include<hip_test_features.hh>
#define INC_VAL 10
#define INITIAL_VAL 5
@@ -51,7 +52,7 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_NonCoherentnounsafeatomicsflag", "",
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (CheckIfFeatSupported(CTFeatures::CT_FEATURE_FINEGRAIN_HWSUPPORT, gfxName)) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
@@ -92,7 +93,7 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_NonCoherentnounsafeatomicsflag", "",
HIP_CHECK(hipHostFree(result));
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
SUCCEED("Memory model feature is only supported for gfx90a, gfx940, gfx941, gfx942, Hence"
"skipping the testcase for this GPU " << device);
}
}
catch/unit/deviceLib/unsafeAtomicAdd_NonCoherent_withoutflag.cc
+4 -3
View File
@@ -22,11 +22,12 @@
unsafeAtomicAdd on CoarseGrainMemory
1. The following test scenario verifies
unsafeAtomicAdd on CoarseGrain memory without any unsafeatomics flag
This testcase works only on gfx90a.
This testcase works only on gfx90a, gfx940, gfx941, gfx942.
*/
#include<hip_test_checkers.hh>
#include<hip_test_common.hh>
#include<hip_test_features.hh>
#define INC_VAL 10
#define INITIAL_VAL 5
@@ -51,7 +52,7 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_NonCoherentwithoutflag", "",
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (CheckIfFeatSupported(CTFeatures::CT_FEATURE_FINEGRAIN_HWSUPPORT, gfxName)) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
@@ -92,7 +93,7 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_NonCoherentwithoutflag", "",
HIP_CHECK(hipHostFree(result));
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
SUCCEED("Memory model feature is only supported for gfx90a, gfx940, gfx941, gfx942, Hence"
"skipping the testcase for this GPU " << device);
}
}
catch/unit/deviceLib/unsafeAtomicAdd_NonCoherent_withunsafeflag.cc
+4 -3
View File
@@ -22,11 +22,12 @@
unsafeAtomicAdd on CoarseGrainMemory
1. The following test scenario verifies
unsafeAtomicAdd on CoarseGrain memory with unsafeatomics flag
This testcase works only on gfx90a.
This testcase works only on gfx90a, gfx940, gfx941, gfx942.
*/
#include<hip_test_checkers.hh>
#include<hip_test_common.hh>
#include<hip_test_features.hh>
#define INC_VAL 10
#define INITIAL_VAL 5
@@ -51,7 +52,7 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_NonCoherentwithunsafeatomicsflag", "",
HIP_CHECK(hipGetDevice(&device));
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (CheckIfFeatSupported(CTFeatures::CT_FEATURE_FINEGRAIN_HWSUPPORT, gfxName)) {
if (prop.canMapHostMemory != 1) {
SUCCEED("Does not support HostPinned Memory");
} else {
@@ -92,7 +93,7 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_NonCoherentwithunsafeatomicsflag", "",
HIP_CHECK(hipHostFree(result));
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
SUCCEED("Memory model feature is only supported for gfx90a, gfx940, gfx941, gfx942, Hence"
"skipping the testcase for this GPU " << device);
}
}
catch/unit/deviceLib/unsafeAtomicAdd_RTC.cc
+13 -12
View File
@@ -31,6 +31,7 @@ unsafeAtomicAdd Scenarios with hipRTC:
#include<hip_test_checkers.hh>
#include<hip_test_common.hh>
#include<hip_test_features.hh>
#include <hip/hiprtc.h>
#define INCREMENT_VAL 10
#define INITIAL_VAL 5
@@ -66,7 +67,7 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_CoherentRTCnounsafeatomicflag", "",
HIP_CHECK(hipGetDeviceProperties(&props, device));
std::string gfxName(props.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (CheckIfFeatSupported(CTFeatures::CT_FEATURE_FINEGRAIN_HWSUPPORT, gfxName)) {
hiprtcProgram prog;
if (std::is_same<TestType, float>::value) {
hiprtcCreateProgram(&prog, // prog
@@ -135,7 +136,7 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_CoherentRTCnounsafeatomicflag", "",
}
HIP_CHECK(hipModuleUnload(module));
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
SUCCEED("Memory model feature is only supported for gfx90a, gfx940, gfx941, gfx942, Hence"
"skipping the testcase for this GPU " << device);
}
}
@@ -156,7 +157,7 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_CoherentRTCunsafeatomicflag", "",
HIP_CHECK(hipGetDeviceProperties(&props, device));
std::string gfxName(props.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (CheckIfFeatSupported(CTFeatures::CT_FEATURE_FINEGRAIN_HWSUPPORT, gfxName)) {
hiprtcProgram prog;
if (std::is_same<TestType, float>::value) {
hiprtcCreateProgram(&prog, // prog
@@ -227,7 +228,7 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_CoherentRTCunsafeatomicflag", "",
}
HIP_CHECK(hipModuleUnload(module));
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
SUCCEED("Memory model feature is only supported for gfx90a, gfx940, gfx941, gfx942, Hence"
"skipping the testcase for this GPU " << device);
}
}
@@ -245,7 +246,7 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_CoherentRTCwithoutflag", "",
HIP_CHECK(hipGetDeviceProperties(&props, device));
std::string gfxName(props.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if(CheckIfFeatSupported(CTFeatures::CT_FEATURE_FINEGRAIN_HWSUPPORT, gfxName)) {
hiprtcProgram prog;
if (std::is_same<TestType, float>::value) {
hiprtcCreateProgram(&prog, // prog
@@ -315,7 +316,7 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_CoherentRTCwithoutflag", "",
}
HIP_CHECK(hipModuleUnload(module));
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
SUCCEED("Memory model feature is only supported for gfx90a, gfx940, gfx941, gfx942, Hence"
"skipping the testcase for this GPU " << device);
}
}
@@ -332,7 +333,7 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_NonCoherentRTCnounsafeatomicflag", "",
HIP_CHECK(hipGetDeviceProperties(&props, device));
std::string gfxName(props.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (CheckIfFeatSupported(CTFeatures::CT_FEATURE_FINEGRAIN_HWSUPPORT, gfxName)) {
hiprtcProgram prog;
if (std::is_same<TestType, float>::value) {
hiprtcCreateProgram(&prog, // prog
@@ -401,7 +402,7 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_NonCoherentRTCnounsafeatomicflag", "",
}
HIP_CHECK(hipModuleUnload(module));
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
SUCCEED("Memory model feature is only supported for gfx90a, gfx940, gfx941, gfx942, Hence"
"skipping the testcase for this GPU " << device);
}
}
@@ -419,7 +420,7 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_NonCoherentRTCunsafeatomicflag", "",
HIP_CHECK(hipGetDeviceProperties(&props, device));
std::string gfxName(props.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if(CheckIfFeatSupported(CTFeatures::CT_FEATURE_FINEGRAIN_HWSUPPORT, gfxName)) {
hiprtcProgram prog;
if (std::is_same<TestType, float>::value) {
hiprtcCreateProgram(&prog, // prog
@@ -489,7 +490,7 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_NonCoherentRTCunsafeatomicflag", "",
}
HIP_CHECK(hipModuleUnload(module));
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
SUCCEED("Memory model feature is only supported for gfx90a, gfx940, gfx941, gfx942, Hence"
"skipping the testcase for this GPU " << device);
}
}
@@ -507,7 +508,7 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_NonCoherentRTC", "",
HIP_CHECK(hipGetDeviceProperties(&props, device));
std::string gfxName(props.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (CheckIfFeatSupported(CTFeatures::CT_FEATURE_FINEGRAIN_HWSUPPORT, gfxName)) {
hiprtcProgram prog;
if (std::is_same<TestType, float>::value) {
hiprtcCreateProgram(&prog, // prog
@@ -577,7 +578,7 @@ TEMPLATE_TEST_CASE("Unit_unsafeAtomicAdd_NonCoherentRTC", "",
}
HIP_CHECK(hipModuleUnload(module));
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
SUCCEED("Memory model feature is only supported for gfx90a, gfx940, gfx941, gfx942, Hence"
"skipping the testcase for this GPU " << device);
}
}
catch/unit/memory/hipMemAdvise_old.cc
+5 -4
View File
@@ -67,6 +67,7 @@ THE SOFTWARE.
*/
#include <hip_test_common.hh>
#include <hip_test_features.hh>
#if __linux__
#include <unistd.h>
#include <sys/mman.h>
@@ -661,7 +662,7 @@ TEST_CASE("Unit_hipMemAdvise_TstAlignedAllocMem") {
WARN("Unable to turn on HSA_XNACK, hence terminating the Test case!");
REQUIRE(false);
}
// The following code block checks for gfx90a so as to skip if the device is not MI200
// The following code block checks for gfx90a,940,941,942 so as to skip if the device is not
hipDeviceProp_t prop;
int device;
@@ -669,7 +670,7 @@ TEST_CASE("Unit_hipMemAdvise_TstAlignedAllocMem") {
HIP_CHECK(hipGetDeviceProperties(&prop, device));
std::string gfxName(prop.gcnArchName);
if ((gfxName == "gfx90a" || gfxName.find("gfx90a:")) == 0) {
if (CheckIfFeatSupported(CTFeatures::CT_FEATURE_HMM, prop.gcnArchName)) {
int stat = 0;
if (fork() == 0) {
// The below part should be inside fork
@@ -726,9 +727,9 @@ TEST_CASE("Unit_hipMemAdvise_TstAlignedAllocMem") {
}
}
} else {
SUCCEED("Memory model feature is only supported for gfx90a, Hence"
SUCCEED("Memory model feature is only supported for gfx90a, gfx940, gx941, gfx942, Hence"
"skipping the testcase for this GPU " << device);
WARN("Memory model feature is only supported for gfx90a, Hence"
WARN("Memory model feature is only supported for gfx90a, gfx940, gx941, gfx942, Hence"
"skipping the testcase for this GPU " << device);
}
catch/unit/texture/hipTextureObj3DCheckModes.cc
+5 -4
View File
@@ -18,10 +18,11 @@ THE SOFTWARE.
*/
#include <hip_test_common.hh>
#include <hip_test_features.hh>
#include <hip_test_checkers.hh>
#include <hip_texture_helper.hh>
bool isGfx90a = false;
bool LinearFilter3D = false;
template<bool normalizedCoords>
__global__ void tex3DKernel(float *outputData, hipTextureObject_t textureObject,
@@ -95,7 +96,7 @@ static void runTest(const int width, const int height, const int depth, const fl
if (res != hipSuccess) {
HIP_CHECK(hipFreeArray(arr));
free(hData);
if (res == hipErrorNotSupported && isGfx90a) {
if (res == hipErrorNotSupported && LinearFilter3D) {
printf("gfx90a doesn't support 3D linear filter! Skipped!\n");
} else {
result = false;
@@ -153,8 +154,8 @@ TEST_CASE("Unit_hipTextureObj3DCheckModes") {
int device = 0;
hipDeviceProp_t props;
HIPCHECK(hipGetDeviceProperties(&props, device));
if (!strncmp(props.gcnArchName, "gfx90a", strlen("gfx90a"))) {
isGfx90a = true;
if (CheckIfFeatSupported(CTFeatures::CT_FEATURE_TEXTURES_NOT_SUPPORTED, props.gcnArchName)) {
LinearFilter3D = true;
}
SECTION("hipAddressModeClamp, hipFilterModePoint, regularCoords") {