From f58393108f0c703068d944840571c1f5e37c3363 Mon Sep 17 00:00:00 2001 From: Sourabh U Betigeri Date: Wed, 12 Nov 2025 13:37:00 -0800 Subject: [PATCH] SWDEV-564408 - Reduces hip-tests runtime Pt2 (#1724) --- .../catch/unit/atomics/atomicCAS_system.cc | 10 ---------- .../unit/cooperativeGrps/thread_block_tile.cc | 18 +++++++++--------- .../multiThread/hipMemsetAsyncMultiThread.cc | 2 +- 3 files changed, 10 insertions(+), 20 deletions(-) diff --git a/projects/hip-tests/catch/unit/atomics/atomicCAS_system.cc b/projects/hip-tests/catch/unit/atomics/atomicCAS_system.cc index 326ba41557..9bb84f672d 100644 --- a/projects/hip-tests/catch/unit/atomics/atomicCAS_system.cc +++ b/projects/hip-tests/catch/unit/atomics/atomicCAS_system.cc @@ -72,11 +72,6 @@ TEMPLATE_TEST_CASE("Unit_atomicCAS_system_Positive_Peer_GPUs", "", int, unsigned 2, 2, 1, sizeof(TestType)); } - DYNAMIC_SECTION("Adjacent addresses " << current) { - MultipleDeviceMultipleKernelAndHostTest( - 2, 2, warp_size, sizeof(TestType)); - } - DYNAMIC_SECTION("Scattered addresses " << current) { MultipleDeviceMultipleKernelAndHostTest( 2, 2, warp_size, cache_line_size); @@ -172,11 +167,6 @@ TEMPLATE_TEST_CASE("Unit_atomicCAS_system_Positive_Host_And_Peer_GPUs", "", int, 2, 2, 1, sizeof(TestType), 4); } - DYNAMIC_SECTION("Adjacent addresses " << current) { - MultipleDeviceMultipleKernelAndHostTest( - 2, 2, warp_size, sizeof(TestType), 4); - } - DYNAMIC_SECTION("Scattered addresses " << current) { MultipleDeviceMultipleKernelAndHostTest( 2, 2, warp_size, cache_line_size, 4); diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/thread_block_tile.cc b/projects/hip-tests/catch/unit/cooperativeGrps/thread_block_tile.cc index 9d75c5b867..0f9396e07c 100644 --- a/projects/hip-tests/catch/unit/cooperativeGrps/thread_block_tile.cc +++ b/projects/hip-tests/catch/unit/cooperativeGrps/thread_block_tile.cc @@ -195,7 +195,7 @@ template void BlockTileShflUpTest() { */ TEMPLATE_TEST_CASE("Unit_Thread_Block_Tile_Shfl_Up_Positive_Basic", "", int, unsigned int, long, unsigned long, long long, unsigned long long, float, double) { - BlockTileShflUpTest(); + BlockTileShflUpTest(); } @@ -252,7 +252,7 @@ template void BlockTileShflDownTest() { /** * Test Description * ------------------------ - * - Validates the shuffle down behavior of thread block tiles of all valid sizes{2, 4, 8, 16, + * - Validates the shuffle down behavior of thread block tiles of all valid sizes{2, 16, * 32, 64(if AMD)} for delta values of [0, tile size). The test is run for all overloads of * shfl_down. * Test source @@ -264,7 +264,7 @@ template void BlockTileShflDownTest() { */ TEMPLATE_TEST_CASE("Unit_Thread_Block_Tile_Shfl_Down_Positive_Basic", "", int, unsigned int, long, unsigned long, long long, unsigned long long, float, double) { - BlockTileShflDownTest(); + BlockTileShflDownTest(); } @@ -317,7 +317,7 @@ template void BlockTileShflXORTest() { /** * Test Description * ------------------------ - * - Validates the shuffle xor behavior of thread block tiles of all valid sizes{2, 4, 8, 16, 32, + * - Validates the shuffle xor behavior of thread block tiles of all valid sizes{2, 16, 32, * 64(if AMD)} for mask values of [0, tile size). The test is run for all overloads of shfl_xor. * Test source * ------------------------ @@ -328,7 +328,7 @@ template void BlockTileShflXORTest() { */ TEMPLATE_TEST_CASE("Unit_Thread_Block_Tile_Shfl_XOR_Positive_Basic", "", int, unsigned int, long, unsigned long, long long, unsigned long long, float, double) { - BlockTileShflXORTest(); + BlockTileShflXORTest(); } template @@ -397,7 +397,7 @@ template void BlockTileShflTest() { /** * Test Description * ------------------------ - * - Validates the shuffle behavior of thread block tiles of all valid sizes{2, 4, 8, 16, 32, + * - Validates the shuffle behavior of thread block tiles of all valid sizes{2, 16, 32, * 64(if AMD)} for generated shuffle target lanes. The test is run for all overloads of shfl. Test * source * ------------------------ @@ -408,7 +408,7 @@ template void BlockTileShflTest() { */ TEMPLATE_TEST_CASE("Unit_Thread_Block_Tile_Shfl_Positive_Basic", "", int, unsigned int, long, unsigned long, long long, unsigned long long, float, double) { - BlockTileShflTest(); + BlockTileShflTest(); } @@ -521,8 +521,8 @@ template void BlockTileSy * - HIP_VERSION >= 5.2 */ TEMPLATE_TEST_CASE("Unit_Thread_Block_Tile_Sync_Positive_Basic", "", uint8_t, uint16_t, uint32_t) { - SECTION("Global memory") { BlockTileSyncTest(); } - SECTION("Shared memory") { BlockTileSyncTest(); } + SECTION("Global memory") { BlockTileSyncTest(); } + SECTION("Shared memory") { BlockTileSyncTest(); } } /** diff --git a/projects/hip-tests/catch/unit/multiThread/hipMemsetAsyncMultiThread.cc b/projects/hip-tests/catch/unit/multiThread/hipMemsetAsyncMultiThread.cc index f6a1887583..37a6f966f5 100644 --- a/projects/hip-tests/catch/unit/multiThread/hipMemsetAsyncMultiThread.cc +++ b/projects/hip-tests/catch/unit/multiThread/hipMemsetAsyncMultiThread.cc @@ -62,7 +62,7 @@ template void threadCall(Func f, hipStream_t stream) T* ptr{nullptr}; constexpr size_t size = 1024; - constexpr size_t iter = 1024; + constexpr size_t iter = 512; HIP_CHECK_THREAD(hipMalloc(&ptr, sizeof(T) * size)); hipEvent_t event{}; HIP_CHECK_THREAD(hipEventCreate(&event));