From 9de3979083e34d20f33c5b47a26f58e27502deb0 Mon Sep 17 00:00:00 2001 From: Rakesh Roy Date: Thu, 4 Apr 2024 12:00:33 +0530 Subject: [PATCH] SWDEV-1 - Skip tests hanging/failing in Windows PSDB - Device side memory allocation & printf testcases are hanging because pcie atomics is disabled in recent builds - Few Coalesced Group testcases also hanging - Unit_Thread_Block_Getters_Positive_Basic testcase is failing randomly Change-Id: Ie53515db05c41e879248bb02655a74f0c296aa5f --- catch/hipTestMain/config/config_amd_windows | 7 +- catch/unit/deviceLib/deviceAllocation.cc | 146 +++++++++++++++++++ catch/unit/printf/printfFlags.cc | 6 + catch/unit/printf/printfFlagsNonHost.cc | 6 + catch/unit/printf/printfHost.cc | 6 + catch/unit/printf/printfLength.cc | 6 + catch/unit/printf/printfNonHost.cc | 26 +++- catch/unit/printf/printfSpecifiers.cc | 12 ++ catch/unit/printf/printfSpecifiersNonHost.cc | 6 + catch/unit/rtc/hipRTCDeviceMalloc.cc | 7 + 10 files changed, 224 insertions(+), 4 deletions(-) diff --git a/catch/hipTestMain/config/config_amd_windows b/catch/hipTestMain/config/config_amd_windows index bdcd27c18e..fe1617e4b7 100644 --- a/catch/hipTestMain/config/config_amd_windows +++ b/catch/hipTestMain/config/config_amd_windows @@ -670,6 +670,7 @@ "Unit_hipMemPoolGetAccess_Negative_Parameters", "Unit_hipMemPoolSetAttribute_Negative_Parameters", "Unit_hipMemPoolGetAttribute_Negative_Parameters", + "Unit_Thread_Block_Getters_Positive_Basic", "Note: Test disabled due to defect - EXSWHTEC-151", "Unit_hipModuleLoad_Negative_Load_From_A_File_That_Is_Not_A_Module", "Note: Test disabled due to defect - EXSWHTEC-152", @@ -691,9 +692,9 @@ "Unit_Thread_Block_Tile_Sync_Positive_Basic - uint8_t", "Unit_Thread_Block_Tile_Sync_Positive_Basic - uint16_t", "Unit_Thread_Block_Tile_Sync_Positive_Basic - uint32_t", - "Unit_hiprtc_devicemalloc", - "Unit_deviceAllocation_InOneThread_AccessInAllThreads", - "Unit_deviceAllocation_Malloc_ComplexDataType", + "Unit_coalesced_groups", + "Unit_coalesced_groups_shfl_down", + "Unit_coalesced_groups_shfl_up", "=== SWDEV-441604: Below tests take long time to run in stress test on 12/01/24 ===", "Unit_Thread_Block_Tile_Shfl_Up_Positive_Basic - int", "Unit_Thread_Block_Tile_Shfl_Up_Positive_Basic - unsigned int", diff --git a/catch/unit/deviceLib/deviceAllocation.cc b/catch/unit/deviceLib/deviceAllocation.cc index 63cba2171e..cd6d2eb442 100644 --- a/catch/unit/deviceLib/deviceAllocation.cc +++ b/catch/unit/deviceLib/deviceAllocation.cc @@ -935,6 +935,12 @@ static bool TestAllocInDeviceFunc(int test_type) { * types like char, short, int etc. */ TEST_CASE("Unit_deviceAllocation_Malloc_PerThread_PrimitiveDataType") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } constexpr size_t sizePerThread = 128; // malloc()/free() tests @@ -970,6 +976,12 @@ TEST_CASE("Unit_deviceAllocation_Malloc_PerThread_PrimitiveDataType") { * types like char, short, int etc. */ TEST_CASE("Unit_deviceAllocation_New_PerThread_PrimitiveDataType") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } constexpr size_t sizePerThread = 128; // new/delete tests @@ -1004,6 +1016,12 @@ TEST_CASE("Unit_deviceAllocation_New_PerThread_PrimitiveDataType") { * using malloc/free in every gpu thread and block for structure. */ TEST_CASE("Unit_deviceAllocation_Malloc_PerThread_StructDataType") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } constexpr size_t sizePerThread = 64; struct simpleStruct sampleStr{INT_MAX, DBL_MAX, FLT_MAX, SHRT_MAX, SCHAR_MAX, {1, 2, 3, 4, 5, 6, 7, 8}}; @@ -1016,6 +1034,12 @@ TEST_CASE("Unit_deviceAllocation_Malloc_PerThread_StructDataType") { * using new/delete in every gpu thread and block for structure. */ TEST_CASE("Unit_deviceAllocation_New_PerThread_StructDataType") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } constexpr size_t sizePerThread = 64; struct simpleStruct sampleStr{INT_MAX, DBL_MAX, FLT_MAX, SHRT_MAX, SCHAR_MAX, {1, 2, 3, 4, 5, 6, 7, 8}}; @@ -1029,6 +1053,13 @@ TEST_CASE("Unit_deviceAllocation_New_PerThread_StructDataType") { * char, short, int etc. */ TEST_CASE("Unit_deviceAllocation_InOneThread_AccessInAllThreads") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } + // malloc()/free() tests SECTION("Test char datatype allocation with malloc") { REQUIRE(true == TestMemoryAccessInAllThread(TEST_MALLOC_FREE, 0)); @@ -1078,6 +1109,12 @@ TEST_CASE("Unit_deviceAllocation_InOneThread_AccessInAllThreads") { * int etc. */ TEST_CASE("Unit_deviceAllocation_Malloc_AcrossKernels") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } // malloc()/free() tests SECTION("Test char datatype allocation with malloc") { REQUIRE(true == TestMemoryAcrossMulKernels(TEST_MALLOC_FREE)); @@ -1106,6 +1143,12 @@ TEST_CASE("Unit_deviceAllocation_Malloc_AcrossKernels") { * int etc. */ TEST_CASE("Unit_deviceAllocation_New_AcrossKernels") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } // new/delete tests SECTION("Test char datatype allocation with new") { REQUIRE(true == TestMemoryAcrossMulKernels(TEST_NEW_DELETE)); @@ -1136,6 +1179,13 @@ TEST_CASE("Unit_deviceAllocation_New_AcrossKernels") { * __device__ functions. */ TEST_CASE("Unit_deviceAllocation_Malloc_ComplexDataType") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } + // malloc()/free() tests REQUIRE(true == TestMemoryAccessInAllThread_CmplxStr(TEST_MALLOC_FREE)); } @@ -1148,6 +1198,12 @@ TEST_CASE("Unit_deviceAllocation_Malloc_ComplexDataType") { * __device__ functions. */ TEST_CASE("Unit_deviceAllocation_New_ComplexDataType") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } // new/delete tests REQUIRE(true == TestMemoryAccessInAllThread_CmplxStr(TEST_NEW_DELETE)); } @@ -1157,6 +1213,12 @@ TEST_CASE("Unit_deviceAllocation_New_ComplexDataType") { * across multiple kernels for Union data type. */ TEST_CASE("Unit_deviceAllocation_Malloc_UnionType") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } // malloc()/free() tests REQUIRE(true == TestMemoryAccessInAllThread_Union(TEST_MALLOC_FREE)); } @@ -1166,6 +1228,12 @@ TEST_CASE("Unit_deviceAllocation_Malloc_UnionType") { * across multiple kernels for Union data type. */ TEST_CASE("Unit_deviceAllocation_New_UnionType") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } // new/delete tests REQUIRE(true == TestMemoryAccessInAllThread_Union(TEST_NEW_DELETE)); } @@ -1176,6 +1244,12 @@ TEST_CASE("Unit_deviceAllocation_New_UnionType") { * Code Object kernel. */ TEST_CASE("Unit_deviceAllocation_Malloc_SingleCodeObj") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } constexpr size_t sizePerThread = 128; REQUIRE(true == TestAlloc_Load_SingleKer_AllocFree(TEST_MALLOC_FREE, @@ -1188,6 +1262,12 @@ TEST_CASE("Unit_deviceAllocation_Malloc_SingleCodeObj") { * Code Object kernel. */ TEST_CASE("Unit_deviceAllocation_New_SingleCodeObj") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } constexpr size_t sizePerThread = 128; REQUIRE(true == TestAlloc_Load_SingleKer_AllocFree(TEST_NEW_DELETE, @@ -1265,6 +1345,12 @@ TEST_CASE("Unit_deviceAllocation_New_PerThread_MultKerMultStrm") { * using malloc/free in graph. */ TEST_CASE("Unit_deviceAllocation_Malloc_PerThread_Graph") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } // malloc()/free() tests SECTION("Test char datatype allocation with malloc") { REQUIRE(true == @@ -1297,6 +1383,12 @@ TEST_CASE("Unit_deviceAllocation_Malloc_PerThread_Graph") { * using new/delete in graph. */ TEST_CASE("Unit_deviceAllocation_New_PerThread_Graph") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } // new/delete tests SECTION("Test char datatype allocation with new") { REQUIRE(true == @@ -1329,6 +1421,12 @@ TEST_CASE("Unit_deviceAllocation_New_PerThread_Graph") { * using pointers to device functions. */ TEST_CASE("Unit_deviceAllocation_Malloc_DeviceFunc") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } // malloc/free tests REQUIRE(true == TestAllocInDeviceFunc(TEST_MALLOC_FREE)); } @@ -1338,6 +1436,12 @@ TEST_CASE("Unit_deviceAllocation_Malloc_DeviceFunc") { * using pointers to device functions. */ TEST_CASE("Unit_deviceAllocation_New_DeviceFunc") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } // new/delete tests REQUIRE(true == TestAllocInDeviceFunc(TEST_NEW_DELETE)); } @@ -1346,6 +1450,12 @@ TEST_CASE("Unit_deviceAllocation_New_DeviceFunc") { * Scenario: This test validates device allocation using vitual functions */ TEST_CASE("Unit_deviceAllocation_VirtualFunction") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } int *outputVec_d{nullptr}, *outputVec_h{nullptr}; constexpr size_t sizeBufferPerThread = 8; size_t arraysize = (sizeBufferPerThread * BLOCKSIZE * GRIDSIZE); @@ -1376,6 +1486,12 @@ TEST_CASE("Unit_deviceAllocation_VirtualFunction") { * across multiple kernels launched using threads. */ TEST_CASE("Unit_deviceAllocation_Malloc_MulKernels_MulThreads") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } // malloc()/free() tests SECTION("Test char datatype allocation with malloc") { REQUIRE(true == TestDevMemAllocMulKerMulThrd(TEST_MALLOC_FREE)); @@ -1403,6 +1519,12 @@ TEST_CASE("Unit_deviceAllocation_Malloc_MulKernels_MulThreads") { * across multiple kernels launched using threads. */ TEST_CASE("Unit_deviceAllocation_New_MulKernels_MulThreads") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } // new/delete tests SECTION("Test char datatype allocation with new") { REQUIRE(true == TestDevMemAllocMulKerMulThrd(TEST_NEW_DELETE)); @@ -1433,6 +1555,12 @@ TEST_CASE("Unit_deviceAllocation_New_MulKernels_MulThreads") { * in a single kernel launched using threads. */ TEST_CASE("Unit_deviceAllocation_Malloc_SingKernels_MulThreads") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } // malloc()/free() tests std::vector tests; // Spawn the test threads @@ -1456,6 +1584,12 @@ TEST_CASE("Unit_deviceAllocation_Malloc_SingKernels_MulThreads") { * in a single kernel launched using threads. */ TEST_CASE("Unit_deviceAllocation_New_SingKernels_MulThreads") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } // new/delete tests std::vector tests; // Spawn the test threads @@ -1480,6 +1614,12 @@ TEST_CASE("Unit_deviceAllocation_New_SingKernels_MulThreads") { * code object kernels defined in different source files. */ TEST_CASE("Unit_deviceAllocation_Malloc_MulCodeObj") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } REQUIRE(true == TestAlloc_Load_MultKernels(TEST_MALLOC_FREE, INT_MAX)); } @@ -1489,6 +1629,12 @@ TEST_CASE("Unit_deviceAllocation_Malloc_MulCodeObj") { * code object kernels defined in different source files. */ TEST_CASE("Unit_deviceAllocation_New_MulCodeObj") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } REQUIRE(true == TestAlloc_Load_MultKernels(TEST_NEW_DELETE, INT_MAX)); } diff --git a/catch/unit/printf/printfFlags.cc b/catch/unit/printf/printfFlags.cc index 5d2cfc70e7..441c9cdb74 100644 --- a/catch/unit/printf/printfFlags.cc +++ b/catch/unit/printf/printfFlags.cc @@ -45,6 +45,12 @@ THE SOFTWARE. * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_Printf_flags_Sanity_Positive") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } std::string reference(R"here(00000042 -0000042 00000042 diff --git a/catch/unit/printf/printfFlagsNonHost.cc b/catch/unit/printf/printfFlagsNonHost.cc index 771087692b..0eb21ee72d 100644 --- a/catch/unit/printf/printfFlagsNonHost.cc +++ b/catch/unit/printf/printfFlagsNonHost.cc @@ -41,6 +41,12 @@ THE SOFTWARE. */ TEST_CASE("Unit_Buffered_Printf_Flags") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } std::string reference(R"here(00000042 -0000042 00000042 diff --git a/catch/unit/printf/printfHost.cc b/catch/unit/printf/printfHost.cc index 1b9ef078da..1890c60a1f 100644 --- a/catch/unit/printf/printfHost.cc +++ b/catch/unit/printf/printfHost.cc @@ -44,6 +44,12 @@ __global__ void run_printf(int *count) { * - HIP_VERSION >= 5.7 */ TEST_CASE("Unit_Host_Printf") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } int *count{nullptr}, *count_d{nullptr}; count = reinterpret_cast(malloc(sizeof(int))); HIP_CHECK(hipMalloc(&count_d, sizeof(int))); diff --git a/catch/unit/printf/printfLength.cc b/catch/unit/printf/printfLength.cc index af3dad60ef..f23d8c21b3 100644 --- a/catch/unit/printf/printfLength.cc +++ b/catch/unit/printf/printfLength.cc @@ -44,6 +44,12 @@ THE SOFTWARE. * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_Printf_length_Sanity_Positive") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } #if HT_NVIDIA std::string reference(R"here(-42 -42 -42 -42 diff --git a/catch/unit/printf/printfNonHost.cc b/catch/unit/printf/printfNonHost.cc index 903aeedb78..9773671aa0 100644 --- a/catch/unit/printf/printfNonHost.cc +++ b/catch/unit/printf/printfNonHost.cc @@ -19,7 +19,7 @@ THE SOFTWARE. */ #include - + #define ITER_COUNT 61681 #define KERNEL_ITERATIONS 15 @@ -65,6 +65,12 @@ __global__ void kernel_printf_thread(int *count) { */ TEST_CASE("Unit_NonHost_Printf_basic") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } int *count{nullptr}, *count_d{nullptr}; count = reinterpret_cast(malloc(sizeof(int))); @@ -93,6 +99,12 @@ TEST_CASE("Unit_NonHost_Printf_basic") { */ TEST_CASE("Unit_NonHost_Printf_loop") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } int *count{nullptr}, *count_d{nullptr}; count = reinterpret_cast(malloc(ITER_COUNT * sizeof(int))); @@ -131,6 +143,12 @@ TEST_CASE("Unit_NonHost_Printf_loop") { */ TEST_CASE("Unit_NonHost_Printf_multiple_Threads") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } int *count{nullptr}, *count_d{nullptr}; count = reinterpret_cast(malloc(ITER_COUNT_FOR_THREAD * sizeof(int))); @@ -172,6 +190,12 @@ TEST_CASE("Unit_NonHost_Printf_multiple_Threads") { */ TEST_CASE("Unit_NonHost_Printf_BufferAvailability") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } int *count{nullptr}, *count_d{nullptr}; count = reinterpret_cast(malloc((ITER_COUNT-1) * sizeof(int))); diff --git a/catch/unit/printf/printfSpecifiers.cc b/catch/unit/printf/printfSpecifiers.cc index 6eddfcec6e..c50ed751d1 100644 --- a/catch/unit/printf/printfSpecifiers.cc +++ b/catch/unit/printf/printfSpecifiers.cc @@ -47,6 +47,12 @@ THE SOFTWARE. * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_Printf_specifier_Sanity_Positive") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } #if HT_NVIDIA std::string reference(R"here(xyzzy % @@ -147,6 +153,12 @@ x * - HIP_VERSION >= 5.2 */ TEST_CASE("Unit_Printf_Negative_Parameters_RTC") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } hiprtcProgram program{}; const auto program_source = kPrintfParam; diff --git a/catch/unit/printf/printfSpecifiersNonHost.cc b/catch/unit/printf/printfSpecifiersNonHost.cc index cc372f1714..32d7040fc2 100644 --- a/catch/unit/printf/printfSpecifiersNonHost.cc +++ b/catch/unit/printf/printfSpecifiersNonHost.cc @@ -45,6 +45,12 @@ THE SOFTWARE. */ TEST_CASE("Unit_Buffered_Printf_Specifier") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } #ifdef __HIP_PLATFORM_NVIDIA__ std::string reference(R"here(xyzzy % diff --git a/catch/unit/rtc/hipRTCDeviceMalloc.cc b/catch/unit/rtc/hipRTCDeviceMalloc.cc index 9262f35aff..7b2ef2de17 100644 --- a/catch/unit/rtc/hipRTCDeviceMalloc.cc +++ b/catch/unit/rtc/hipRTCDeviceMalloc.cc @@ -37,6 +37,13 @@ void devicemalloc(float* x, float* y, float* out, float** px, float** py, size_t )"}; TEST_CASE("Unit_hiprtc_devicemalloc") { + int pcieAtomic = 0; + HIP_CHECK(hipDeviceGetAttribute(&pcieAtomic, hipDeviceAttributeHostNativeAtomicSupported, 0)); + if (!pcieAtomic) { + HipTest::HIP_SKIP_TEST("Device doesn't support pcie atomic, Skipped"); + return; + } + using namespace std; hiprtcProgram prog; hiprtcCreateProgram(&prog, // prog