From df2bcdb90f46cd50f96957171e3f3eddbf297f1e Mon Sep 17 00:00:00 2001 From: taosang2 Date: Wed, 20 Dec 2023 14:32:58 -0500 Subject: [PATCH] SWDEV-425830 - Add pcie atomic query in atomic tests Add code to query pcie-atomics in atomic tests. If pcie-atomics is not supported, skip the tests. Enable the tests for Nvidia GPUs as well. Change-Id: I7bfc8600511d78d8c73ab526f9846cd268651278 [ROCm/hip-tests commit: fafcad91629a706d6daa849979d34098ed9c1733] --- .../catch/unit/memory/CMakeLists.txt | 30 ++++++++++--------- .../catch/unit/memory/hipMemCoherencyTst.cc | 13 ++++---- .../unit/memory/hipSVMTestByteGranularity.cpp | 7 +++++ .../hipSVMTestFineGrainMemoryConsistency.cpp | 7 +++++ .../memory/hipSVMTestFineGrainSyncBuffers.cpp | 7 +++++ .../hipSVMTestSharedAddressSpaceFineGrain.cpp | 4 ++- .../hipDestroyExternalSemaphore.cc | 2 +- 7 files changed, 49 insertions(+), 21 deletions(-) diff --git a/projects/hip-tests/catch/unit/memory/CMakeLists.txt b/projects/hip-tests/catch/unit/memory/CMakeLists.txt index 2283d5e558..e934cbee5f 100644 --- a/projects/hip-tests/catch/unit/memory/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/memory/CMakeLists.txt @@ -72,14 +72,14 @@ set(TEST_SRC hipMemAdviseMmap.cc hipMallocManaged.cc hipMemRangeGetAttribute.cc - hipMemRangeGetAttribute_old.cc) + hipMemRangeGetAttribute_old.cc + hipMemCoherencyTst.cc) if(HIP_PLATFORM MATCHES "amd") set(TEST_SRC ${TEST_SRC} hipMemPtrGetInfo.cc hipPointerGetAttributes.cc - hipMemCoherencyTst.cc hipExtMallocWithFlags.cc hipMallocMngdMultiThread.cc hipMemVmm.cc @@ -101,6 +101,7 @@ hip_add_exe_to_target(NAME MemoryTest1 if(HIP_PLATFORM MATCHES "amd") set_source_files_properties(hipHostRegister.cc PROPERTIES COMPILE_FLAGS -std=c++17) add_executable(hipHostRegisterPerf EXCLUDE_FROM_ALL hipHostRegister_exe.cc) + add_dependencies(build_tests hipHostRegisterPerf) if(UNIX) add_executable(hipMemAdviseTstAlignedAllocMem EXCLUDE_FROM_ALL hipMemAdvise_AlignedAllocMem_Exe.cc) add_dependencies(MemoryTest1 hipMemAdviseTstAlignedAllocMem) @@ -207,17 +208,18 @@ hip_add_exe_to_target(NAME MemoryTest2 TEST_SRC ${TEST_SRC} TEST_TARGET_NAME build_tests COMMON_SHARED_SRC ${COMMON_SHARED_SRC}) -if(HIP_PLATFORM MATCHES "amd") - set(TEST_SRC - hipSVMTestByteGranularity.cpp - hipSVMTestFineGrainMemoryConsistency.cpp - hipSVMTestFineGrainSyncBuffers.cpp - hipSVMTestSharedAddressSpaceFineGrain.cpp - ) +set(TEST_SRC + hipSVMTestByteGranularity.cpp + hipSVMTestFineGrainMemoryConsistency.cpp + hipSVMTestFineGrainSyncBuffers.cpp + hipSVMTestSharedAddressSpaceFineGrain.cpp +) - hip_add_exe_to_target(NAME SVMAtomicTest - TEST_SRC ${TEST_SRC} - TEST_TARGET_NAME build_tests COMMON_SHARED_SRC ${COMMON_SHARED_SRC}) +hip_add_exe_to_target(NAME SVMAtomicTest + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME build_tests COMMON_SHARED_SRC ${COMMON_SHARED_SRC}) - add_dependencies(build_tests hipHostRegisterPerf) -endif() \ No newline at end of file +if(HIP_PLATFORM MATCHES "nvidia") + set_target_properties(SVMAtomicTest PROPERTIES COMPILE_FLAGS -arch=sm_70) + set_target_properties(MemoryTest1 PROPERTIES COMPILE_FLAGS -arch=sm_70) +endif() diff --git a/projects/hip-tests/catch/unit/memory/hipMemCoherencyTst.cc b/projects/hip-tests/catch/unit/memory/hipMemCoherencyTst.cc index 6e2c2b588d..6879456ee2 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemCoherencyTst.cc +++ b/projects/hip-tests/catch/unit/memory/hipMemCoherencyTst.cc @@ -94,10 +94,15 @@ static void TstCoherency(int* ptr, bool hmmMem) { /* Test case description: The following test validates if fine grain behavior is observed or not with memory allocated using hipHostMalloc()*/ -// The following tests are disabled for Nvidia as they are not consistently -// passing -#if HT_AMD TEST_CASE("Unit_hipHostMalloc_CoherentTst") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + fprintf(stderr, "Device doesn't support pcie atomic, Skipped\n"); + REQUIRE(true); + return; + } + int *Ptr = nullptr, SIZE = sizeof(int); bool HmmMem = false; YES_COHERENT = false; @@ -117,8 +122,6 @@ TEST_CASE("Unit_hipHostMalloc_CoherentTst") { HIP_CHECK(hipHostFree(Ptr)); REQUIRE(YES_COHERENT); } -#endif - /* Test case description: The following test validates if fine grain behavior is observed or not with memory allocated using hipMallocManaged()*/ diff --git a/projects/hip-tests/catch/unit/memory/hipSVMTestByteGranularity.cpp b/projects/hip-tests/catch/unit/memory/hipSVMTestByteGranularity.cpp index 8b682bb10d..6fee49758d 100644 --- a/projects/hip-tests/catch/unit/memory/hipSVMTestByteGranularity.cpp +++ b/projects/hip-tests/catch/unit/memory/hipSVMTestByteGranularity.cpp @@ -74,6 +74,13 @@ __global__ void sum_neighbor_locations(char* a, unsigned int num_devices, * - HIP_VERSION >= 5.7 */ TEST_CASE("test_svm_byte_granularity") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + fprintf(stderr, "Device doesn't support pcie atomic, Skipped\n"); + REQUIRE(true); + return; + } const int num_elements = 2048; int num_devices = 0; HIP_CHECK(hipGetDeviceCount(&num_devices)); diff --git a/projects/hip-tests/catch/unit/memory/hipSVMTestFineGrainMemoryConsistency.cpp b/projects/hip-tests/catch/unit/memory/hipSVMTestFineGrainMemoryConsistency.cpp index 29d4cbc19a..a4eb2d60d6 100644 --- a/projects/hip-tests/catch/unit/memory/hipSVMTestFineGrainMemoryConsistency.cpp +++ b/projects/hip-tests/catch/unit/memory/hipSVMTestFineGrainMemoryConsistency.cpp @@ -234,6 +234,13 @@ void launch_kernels_and_verify(std::vector &streams, unsigned int n * - HIP_VERSION >= 5.7 */ TEST_CASE("test_svm_fine_grain_memory_consistency") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + fprintf(stderr, "Device doesn't support pcie atomic, Skipped\n"); + REQUIRE(true); + return; + } const int num_elements = 2167; int num_devices = 0; HIP_CHECK(hipGetDeviceCount(&num_devices)); diff --git a/projects/hip-tests/catch/unit/memory/hipSVMTestFineGrainSyncBuffers.cpp b/projects/hip-tests/catch/unit/memory/hipSVMTestFineGrainSyncBuffers.cpp index 1e47818b0c..34f7ff37a6 100644 --- a/projects/hip-tests/catch/unit/memory/hipSVMTestFineGrainSyncBuffers.cpp +++ b/projects/hip-tests/catch/unit/memory/hipSVMTestFineGrainSyncBuffers.cpp @@ -76,6 +76,13 @@ void spawnAnalysisTask(int location) * - HIP_VERSION >= 5.7 */ TEST_CASE("test_svm_fine_grain_sync_buffers") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + fprintf(stderr, "Device doesn't support pcie atomic, Skipped\n"); + REQUIRE(true); + return; + } size_t num_pixels = 1024 * 1024 * 2; hipStream_t stream; HIP_CHECK(hipSetDevice(0)); diff --git a/projects/hip-tests/catch/unit/memory/hipSVMTestSharedAddressSpaceFineGrain.cpp b/projects/hip-tests/catch/unit/memory/hipSVMTestSharedAddressSpaceFineGrain.cpp index 0cb35e912e..e3144dd596 100644 --- a/projects/hip-tests/catch/unit/memory/hipSVMTestSharedAddressSpaceFineGrain.cpp +++ b/projects/hip-tests/catch/unit/memory/hipSVMTestSharedAddressSpaceFineGrain.cpp @@ -129,7 +129,8 @@ void verify_linked_lists_on_device(hipStream_t stream, Node* pNodes, int correct_count = *pNumCorrect; if(correct_count != ListLength * numLists) { - fprintf(stderr,"Failed\n"); + fprintf(stderr, "Failed: correct_count = %d, ListLength=%u, numLists = %u\n", correct_count, + ListLength, numLists); REQUIRE(false); } } @@ -239,6 +240,7 @@ TEST_CASE("test_svm_shared_address_space_fine_grain_buffers") { * Test requirements * ------------------------ * - Host specific (WINDOWS and LINUX) +* - Unified address supported on devices * - System fine grain access supported on devices * - HIP_VERSION >= 5.7 */ diff --git a/projects/hip-tests/catch/unit/vulkan_interop/hipDestroyExternalSemaphore.cc b/projects/hip-tests/catch/unit/vulkan_interop/hipDestroyExternalSemaphore.cc index e4b7dbb082..cf870d8918 100644 --- a/projects/hip-tests/catch/unit/vulkan_interop/hipDestroyExternalSemaphore.cc +++ b/projects/hip-tests/catch/unit/vulkan_interop/hipDestroyExternalSemaphore.cc @@ -36,4 +36,4 @@ TEST_CASE("Unit_hipDestroyExternalSemaphore_Vulkan_Negative_Parameters") { HIP_CHECK_ERROR(hipDestroyExternalSemaphore(ext_semaphore), hipErrorInvalidValue); } #endif -} \ No newline at end of file +}