diff --git a/catch/hipTestMain/CMakeLists.txt b/catch/hipTestMain/CMakeLists.txt index 13407032a8..95b7a09547 100644 --- a/catch/hipTestMain/CMakeLists.txt +++ b/catch/hipTestMain/CMakeLists.txt @@ -22,7 +22,7 @@ if(CMAKE_BUILD_TYPE MATCHES "^Debug$") add_definitions(-DHT_LOG_ENABLE) endif() -add_library(Main_Object EXCLUDE_FROM_ALL OBJECT main.cc hip_test_context.cc) +add_library(Main_Object EXCLUDE_FROM_ALL OBJECT main.cc hip_test_context.cc hip_test_features.cc) if(HIP_PLATFORM MATCHES "amd") set_property(TARGET Main_Object PROPERTY CXX_STANDARD 17) else() diff --git a/catch/hipTestMain/hip_test_context.cc b/catch/hipTestMain/hip_test_context.cc index eb2893a00c..f30a747d59 100644 --- a/catch/hipTestMain/hip_test_context.cc +++ b/catch/hipTestMain/hip_test_context.cc @@ -6,6 +6,7 @@ #include #include "hip_test_context.hh" #include "hip_test_filesystem.hh" +#include "hip_test_features.hh" void TestContext::detectOS() { #if (HT_WIN == 1) diff --git a/catch/hipTestMain/hip_test_features.cc b/catch/hipTestMain/hip_test_features.cc new file mode 100644 index 0000000000..b53811be03 --- /dev/null +++ b/catch/hipTestMain/hip_test_features.cc @@ -0,0 +1,45 @@ +#include "hip_test_features.hh" + +#include +#include + +#include "hip_test_context.hh" + +std::vector> GCNArchFeatMap = { + {"gfx90a", "gfx940", "gfx941", "gfx942"}, // CT_FEATURE_FINEGRAIN_HWSUPPORT + {"gfx90a", "gfx940", "gfx941", "gfx942"}, // CT_FEATURE_HMM + {"gfx90a", "gfx940", "gfx941", "gfx942"}, // CT_FEATURE_TEXTURES_NOT_SUPPORTED +}; + +#if HT_AMD +std::string TrimAndGetGFXName(const std::string& full_gfx_name) { + std::string gfx_name(""); + + // Split the first part of the delimiter + std::string delimiter = ":"; + auto pos = full_gfx_name.find(delimiter); + if (pos == std::string::npos) { + gfx_name = full_gfx_name; + } else { + gfx_name = full_gfx_name.substr(0, pos); + } + + assert(gfx_name.substr(0,3) == "gfx"); + return gfx_name; +} +#endif + +// Check if the GCN Maps +bool CheckIfFeatSupported(enum CTFeatures test_feat, std::string gcn_arch) { +#if HT_NVIDIA + return true; // returning true since feature check does not exist for NV. +#elif HT_AMD + assert(test_feat >= 0 && test_feat < CTFeatures::CT_FEATURE_LAST); + gcn_arch = TrimAndGetGFXName(gcn_arch); + assert(gcn_arch != ""); + return (GCNArchFeatMap[test_feat].find(gcn_arch) != GCNArchFeatMap[test_feat].cend()); +#else + std::cout<<"Platform has to be either AMD or NVIDIA, asserting..."< +#include +#include +#include + +// Catch Test Features +typedef enum CTFeatures { + CT_FEATURE_FINEGRAIN_HWSUPPORT = 0x0, // FINEGRAIN Supported Hardware. + CT_FEATURE_HMM = 0x1, // HMM Enabled + CT_FEATURE_TEXTURES_NOT_SUPPORTED = 0x2, // Textures not supported + CT_FEATURE_LAST = 0x3 +} CTFeatures; + +bool CheckIfFeatSupported(enum CTFeatures test_feat, std::string gcn_arch); diff --git a/catch/multiproc/hipMemCoherencyTstMProc.cc b/catch/multiproc/hipMemCoherencyTstMProc.cc index 8579aabc27..576f8d8bc2 100644 --- a/catch/multiproc/hipMemCoherencyTstMProc.cc +++ b/catch/multiproc/hipMemCoherencyTstMProc.cc @@ -33,6 +33,7 @@ */ #include +#include #include #include #include @@ -158,11 +159,12 @@ TEST_CASE("Unit_malloc_CoherentTst") { hipDeviceProp_t prop; HIPCHECK(hipGetDeviceProperties(&prop, 0)); char *p = NULL; - p = strstr(prop.gcnArchName, "gfx90a"); - if (p) { + + if (CheckIfFeatSupported(CTFeatures::CT_FEATURE_FINEGRAIN_HWSUPPORT, prop.gcnArchName)) { WARN("gfx90a gpu found on this system!!"); GpuId[0] = 1; } + // Write concatenated string and close writing end write(fd1[1], GpuId, 2 * sizeof(int)); close(fd1[1]); diff --git a/catch/unit/deviceLib/AtomicAdd_Coherent_withnoUnsafeflag.cc b/catch/unit/deviceLib/AtomicAdd_Coherent_withnoUnsafeflag.cc index 58e0d1d6a2..2c6effaa89 100644 --- a/catch/unit/deviceLib/AtomicAdd_Coherent_withnoUnsafeflag.cc +++ b/catch/unit/deviceLib/AtomicAdd_Coherent_withnoUnsafeflag.cc @@ -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 -#include +#include +#include +#include #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); } } diff --git a/catch/unit/deviceLib/AtomicAdd_Coherent_withoutflag.cc b/catch/unit/deviceLib/AtomicAdd_Coherent_withoutflag.cc index 300b84fed0..e068b94cc2 100644 --- a/catch/unit/deviceLib/AtomicAdd_Coherent_withoutflag.cc +++ b/catch/unit/deviceLib/AtomicAdd_Coherent_withoutflag.cc @@ -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 #include - +#include #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); } } diff --git a/catch/unit/deviceLib/AtomicAdd_Coherent_withunsafeflag.cc b/catch/unit/deviceLib/AtomicAdd_Coherent_withunsafeflag.cc index 5472b6225f..b9fffb5de5 100644 --- a/catch/unit/deviceLib/AtomicAdd_Coherent_withunsafeflag.cc +++ b/catch/unit/deviceLib/AtomicAdd_Coherent_withunsafeflag.cc @@ -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 #include +#include #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); } } diff --git a/catch/unit/deviceLib/AtomicAdd_NonCoherent_withnoUnsafeflag.cc b/catch/unit/deviceLib/AtomicAdd_NonCoherent_withnoUnsafeflag.cc index 98df491998..00ded36a21 100644 --- a/catch/unit/deviceLib/AtomicAdd_NonCoherent_withnoUnsafeflag.cc +++ b/catch/unit/deviceLib/AtomicAdd_NonCoherent_withnoUnsafeflag.cc @@ -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 #include - +#include #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); } } diff --git a/catch/unit/deviceLib/AtomicAdd_NonCoherent_withoutflag.cc b/catch/unit/deviceLib/AtomicAdd_NonCoherent_withoutflag.cc index 38ba5a5690..6e88a26afc 100644 --- a/catch/unit/deviceLib/AtomicAdd_NonCoherent_withoutflag.cc +++ b/catch/unit/deviceLib/AtomicAdd_NonCoherent_withoutflag.cc @@ -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 #include +#include #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); } } diff --git a/catch/unit/deviceLib/AtomicAdd_NonCoherent_withunsafeflag.cc b/catch/unit/deviceLib/AtomicAdd_NonCoherent_withunsafeflag.cc index 6bfff7262c..5cfb70c816 100644 --- a/catch/unit/deviceLib/AtomicAdd_NonCoherent_withunsafeflag.cc +++ b/catch/unit/deviceLib/AtomicAdd_NonCoherent_withunsafeflag.cc @@ -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 #include - +#include #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); } } diff --git a/catch/unit/deviceLib/BuiltIns_fadd.cc b/catch/unit/deviceLib/BuiltIns_fadd.cc index 0819f90b2c..769b0dc97a 100644 --- a/catch/unit/deviceLib/BuiltIns_fadd.cc +++ b/catch/unit/deviceLib/BuiltIns_fadd.cc @@ -28,6 +28,7 @@ This testfile verifies __builtin_amdgcn_global_atomic_fadd_f64 API scenarios #include #include +#include #include #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); } } diff --git a/catch/unit/deviceLib/CMakeLists.txt b/catch/unit/deviceLib/CMakeLists.txt index 60fafe7d99..b2718f39bb 100644 --- a/catch/unit/deviceLib/CMakeLists.txt +++ b/catch/unit/deviceLib/CMakeLists.txt @@ -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() diff --git a/catch/unit/deviceLib/unsafeAtomicAdd.cc b/catch/unit/deviceLib/unsafeAtomicAdd.cc index a4e005404b..9d6ce388db 100644 --- a/catch/unit/deviceLib/unsafeAtomicAdd.cc +++ b/catch/unit/deviceLib/unsafeAtomicAdd.cc @@ -21,6 +21,7 @@ THE SOFTWARE. */ #include +#include #include #include @@ -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 diff --git a/catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withnounsafeflag.cc b/catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withnounsafeflag.cc index 596fb1ef04..91b2fc433f 100644 --- a/catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withnounsafeflag.cc +++ b/catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withnounsafeflag.cc @@ -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 #include +#include #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); } } diff --git a/catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withoutflag.cc b/catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withoutflag.cc index 685f91c6ca..9999c39344 100644 --- a/catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withoutflag.cc +++ b/catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withoutflag.cc @@ -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 #include +#include #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); } } diff --git a/catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withunsafeflag.cc b/catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withunsafeflag.cc index 4263e412e1..cbe614995f 100644 --- a/catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withunsafeflag.cc +++ b/catch/unit/deviceLib/unsafeAtomicAdd_Coherent_withunsafeflag.cc @@ -27,6 +27,7 @@ This testcase works only on gfx90a. #include #include +#include #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); } } diff --git a/catch/unit/deviceLib/unsafeAtomicAdd_NonCoherent_withnounsafeflag.cc b/catch/unit/deviceLib/unsafeAtomicAdd_NonCoherent_withnounsafeflag.cc index b0f1d86421..c12b30d9f3 100644 --- a/catch/unit/deviceLib/unsafeAtomicAdd_NonCoherent_withnounsafeflag.cc +++ b/catch/unit/deviceLib/unsafeAtomicAdd_NonCoherent_withnounsafeflag.cc @@ -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 #include +#include #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); } } diff --git a/catch/unit/deviceLib/unsafeAtomicAdd_NonCoherent_withoutflag.cc b/catch/unit/deviceLib/unsafeAtomicAdd_NonCoherent_withoutflag.cc index fc298a8592..21e071493b 100644 --- a/catch/unit/deviceLib/unsafeAtomicAdd_NonCoherent_withoutflag.cc +++ b/catch/unit/deviceLib/unsafeAtomicAdd_NonCoherent_withoutflag.cc @@ -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 #include +#include #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); } } diff --git a/catch/unit/deviceLib/unsafeAtomicAdd_NonCoherent_withunsafeflag.cc b/catch/unit/deviceLib/unsafeAtomicAdd_NonCoherent_withunsafeflag.cc index bff5e00483..f8d6cf0b5e 100644 --- a/catch/unit/deviceLib/unsafeAtomicAdd_NonCoherent_withunsafeflag.cc +++ b/catch/unit/deviceLib/unsafeAtomicAdd_NonCoherent_withunsafeflag.cc @@ -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 #include +#include #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); } } diff --git a/catch/unit/deviceLib/unsafeAtomicAdd_RTC.cc b/catch/unit/deviceLib/unsafeAtomicAdd_RTC.cc index e0003e5528..b91d15cf62 100644 --- a/catch/unit/deviceLib/unsafeAtomicAdd_RTC.cc +++ b/catch/unit/deviceLib/unsafeAtomicAdd_RTC.cc @@ -31,6 +31,7 @@ unsafeAtomicAdd Scenarios with hipRTC: #include #include +#include #include #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::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::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::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::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::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::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); } } diff --git a/catch/unit/memory/hipMemAdvise_old.cc b/catch/unit/memory/hipMemAdvise_old.cc index b951998053..8b7350a7a6 100644 --- a/catch/unit/memory/hipMemAdvise_old.cc +++ b/catch/unit/memory/hipMemAdvise_old.cc @@ -67,6 +67,7 @@ THE SOFTWARE. */ #include +#include #if __linux__ #include #include @@ -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); } diff --git a/catch/unit/texture/hipTextureObj3DCheckModes.cc b/catch/unit/texture/hipTextureObj3DCheckModes.cc index 04d8433566..538237b4ad 100644 --- a/catch/unit/texture/hipTextureObj3DCheckModes.cc +++ b/catch/unit/texture/hipTextureObj3DCheckModes.cc @@ -18,10 +18,11 @@ THE SOFTWARE. */ #include +#include #include #include -bool isGfx90a = false; +bool LinearFilter3D = false; template __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") {