From 3a20cff19b80cc0da20dccbcc2785a7f00055b2c Mon Sep 17 00:00:00 2001 From: Branislav Brzak Date: Wed, 16 Oct 2024 12:28:16 +0000 Subject: [PATCH] SWDEV-441986 - Create tests for divergent execution of warp_sync builtins Change-Id: If9ce0bdb0974732b692c779c7cae5527cec2fc83 --- catch/unit/warp/hipMatchSyncAllTests.cc | 51 ++++++++++++++++++++++ catch/unit/warp/hipMatchSyncAnyTests.cc | 50 +++++++++++++++++++++ catch/unit/warp/hipShflSyncDownTests.cc | 58 +++++++++++++++++++++++++ catch/unit/warp/hipShflSyncTests.cc | 46 ++++++++++++++++++++ catch/unit/warp/hipShflSyncUpTests.cc | 46 ++++++++++++++++++++ catch/unit/warp/hipShflSyncXorTests.cc | 46 ++++++++++++++++++++ catch/unit/warp/warp_common.hh | 9 ++++ 7 files changed, 306 insertions(+) diff --git a/catch/unit/warp/hipMatchSyncAllTests.cc b/catch/unit/warp/hipMatchSyncAllTests.cc index 3d8e925649..570f44fde7 100644 --- a/catch/unit/warp/hipMatchSyncAllTests.cc +++ b/catch/unit/warp/hipMatchSyncAllTests.cc @@ -215,6 +215,54 @@ static void runTestMatchAll_3() { } } +__global__ void matchAll_4(int *Input, int *Output) { + int tid = threadIdx.x; + unsigned long long masks[2] = { Every5thBut9th, Every9thBit }; + + Output[tid] = -1; + if (tid % 5 == 0 || tid % 9 == 0) { + Output[tid] = !!__match_all_sync(masks[tid % 9 == 0], Input[tid], &Input[tid]); + } +} + +static void runTestMatchAll_4() { + size_t warpSize = getWarpSize(); + + auto Input = std::vector(warpSize); + + for (size_t i = 0; i < Input.size(); i++) { + if (i % 9 == 0 || i % 5 == 0) { + Input[i] = 0x55; + } else { + Input[i] = i; + } + } + + auto Output = std::vector(warpSize); + auto Expected = std::vector(warpSize); + + for (size_t i = 0; i < Expected.size(); i++) { + if (i % 9 == 0 || i % 5 == 0) { + Expected[i] = 1; + } else { + Expected[i] = -1; + } + } + + int* d_Input; + int* d_Output; + HIP_CHECK(hipMalloc(&d_Input, Input.size() * sizeof(Input[0]))); + HIP_CHECK(hipMalloc(&d_Output, Output.size() * sizeof(Output[0]))); + + HIP_CHECK(hipMemcpy(d_Input, Input.data(), Input.size() * sizeof(Input[0]), hipMemcpyDefault)); + hipLaunchKernelGGL(matchAll_4, 1, warpSize, 0, 0, d_Input, d_Output); + + HIP_CHECK(hipMemcpy(Output.data(), d_Output, Output.size() * sizeof(Output[0]), hipMemcpyDefault)); + for (size_t i = 0; i < Output.size(); i++) { + REQUIRE(Output[i] == Expected[i]); + } +} + /** * @addtogroup __match_sync * @{ @@ -280,4 +328,7 @@ TEST_CASE("Unit_hipMatchSync_All") { runTestMatchAll_2(); runTestMatchAll_3(); } + SECTION("run divergent execution tests") { + runTestMatchAll_4(); + } } diff --git a/catch/unit/warp/hipMatchSyncAnyTests.cc b/catch/unit/warp/hipMatchSyncAnyTests.cc index 12e7d5a600..2e492716b3 100644 --- a/catch/unit/warp/hipMatchSyncAnyTests.cc +++ b/catch/unit/warp/hipMatchSyncAnyTests.cc @@ -138,6 +138,53 @@ static void runTestMatchAny_2() { } } +__global__ void matchAny_3(int *Input, int *Output) { + auto tid = threadIdx.x; + unsigned long long masks[2] = { Every5thBut9th, Every9thBit }; + + Output[tid] = -1; + if (tid % 5 == 0 || tid % 9 == 0) + Output[tid] = !!__match_any_sync(masks[tid % 9 == 0], Input[tid]); +} + +static void runTestMatchAny_3() { + size_t warpSize = getWarpSize(); + + auto Input = std::vector(warpSize); + + for (size_t i = 0; i < Input.size(); i++) { + if (i % 9 == 0 || i % 5 == 0) { + Input[i] = 0x55; + } else { + Input[i] = i; + } + } + + auto Output = std::vector(warpSize); + auto Expected = std::vector(warpSize); + + for (size_t i = 0; i < Expected.size(); i++) { + if (i % 9 == 0 || i % 5 == 0) { + Expected[i] = 1; + } else { + Expected[i] = -1; + } + } + + int* d_Input; + int* d_Output; + HIP_CHECK(hipMalloc(&d_Input, Input.size() * sizeof(Input[0]))); + HIP_CHECK(hipMalloc(&d_Output, Output.size() * sizeof(Output[0]))); + + HIP_CHECK(hipMemcpy(d_Input, Input.data(), Input.size() * sizeof(Input[0]), hipMemcpyDefault)); + hipLaunchKernelGGL(matchAny_3, 1, warpSize, 0, 0, d_Input, d_Output); + + HIP_CHECK(hipMemcpy(Output.data(), d_Output, Output.size() * sizeof(Output[0]), hipMemcpyDefault)); + for (size_t i = 0; i < Output.size(); i++) { + REQUIRE(Output[i] == Expected[i]); + } +} + /** * @addtogroup __match_sync * @{ @@ -195,4 +242,7 @@ TEST_CASE("Unit_hipMatchSync_Any") { runTestMatchAny_1(); runTestMatchAny_2(); } + SECTION("run divergent execution tests") { + runTestMatchAny_3(); + } } diff --git a/catch/unit/warp/hipShflSyncDownTests.cc b/catch/unit/warp/hipShflSyncDownTests.cc index 12cfba61c9..51244abd2a 100644 --- a/catch/unit/warp/hipShflSyncDownTests.cc +++ b/catch/unit/warp/hipShflSyncDownTests.cc @@ -165,6 +165,61 @@ static void runTestShflDown_3() { } } +__global__ void shflDown_4(int *Input, int *Output) { + auto tid = threadIdx.x; + unsigned long long masks[2] = { Every5thBut9th, Every9thBit }; + + Output[tid] = -1; + if (tid % 5 == 0 || tid % 9 == 0) + Output[tid] = __shfl_down_sync(masks[tid % 9 == 0], Input[tid], tid); +} + +static void runTestShflDown_4() { + size_t warpSize = getWarpSize(); + + auto Input = std::vector(warpSize); + + for (size_t i = 0; i < Input.size(); i++) { + Input[i] = 0x55 * i; + } + + auto Output = std::vector(warpSize); + auto Expected = std::vector(warpSize); + + for (size_t i = 0; i < Expected.size() / 2; i++) { + if (i % 5 == 0) { + Expected[i] = 0x55 * (alignUp(i, 5) * 2); + } else if (i % 9 == 0) { + Expected[i] = 0x55 * (alignUp(i, 9) * 2); + } else { + Expected[i] = -1; + } + } + + for (size_t i = Expected.size() / 2; i < Expected.size(); i++) { + if (i % 5 == 0) { + Expected[i] = 0x55 * alignUp(i, 5); + } else if (i % 9 == 0) { + Expected[i] = 0x55 * alignUp(i, 9); + } else { + Expected[i] = -1; + } + } + + int* d_Input; + int* d_Output; + HIP_CHECK(hipMalloc(&d_Input, Input.size() * sizeof(Input[0]))); + HIP_CHECK(hipMalloc(&d_Output, Output.size() * sizeof(Output[0]))); + + HIP_CHECK(hipMemcpy(d_Input, Input.data(), Input.size() * sizeof(Input[0]), hipMemcpyDefault)); + hipLaunchKernelGGL(shflDown_4, 1, warpSize, 0, 0, d_Input, d_Output); + + HIP_CHECK(hipMemcpy(Output.data(), d_Output, Output.size() * sizeof(Output[0]), hipMemcpyDefault)); + for (size_t i = 0; i < Output.size(); i++) { + REQUIRE(Output[i] == Expected[i]); + } +} + /** * @addtogroup __shfl_sync * @{ @@ -250,4 +305,7 @@ TEST_CASE("Unit_hipShflSync_Down") { runTestShflDown_2<__half2>(); runTestShflDown_3<__half2>(); } + SECTION("run divergent execution tests") { + runTestShflDown_4(); + } } diff --git a/catch/unit/warp/hipShflSyncTests.cc b/catch/unit/warp/hipShflSyncTests.cc index 6a37c30727..7c552a1092 100644 --- a/catch/unit/warp/hipShflSyncTests.cc +++ b/catch/unit/warp/hipShflSyncTests.cc @@ -110,6 +110,49 @@ static void runTestShfl_2() { } } +__global__ void shfl_3(int *Input, int *Output) { + auto tid = threadIdx.x; + unsigned long long masks[2] = { Every5thBut9th, Every9thBit }; + + Output[tid] = -1; + if (tid % 5 == 0 || tid % 9 == 0) + Output[tid] = __shfl_sync(masks[tid % 9 == 0], Input[tid], tid); +} + +static void runTestShfl_3() { + size_t warpSize = getWarpSize(); + + auto Input = std::vector(warpSize); + + for (size_t i = 0; i < Input.size(); i++) { + Input[i] = i; + } + + auto Output = std::vector(warpSize); + auto Expected = std::vector(warpSize); + + for (size_t i = 0; i < Expected.size(); i++) { + if (i % 9 == 0 || i % 5 == 0) { + Expected[i] = i; + } else { + Expected[i] = -1; + } + } + + int* d_Input; + int* d_Output; + HIP_CHECK(hipMalloc(&d_Input, Input.size() * sizeof(Input[0]))); + HIP_CHECK(hipMalloc(&d_Output, Output.size() * sizeof(Output[0]))); + + HIP_CHECK(hipMemcpy(d_Input, Input.data(), Input.size() * sizeof(Input[0]), hipMemcpyDefault)); + hipLaunchKernelGGL(shfl_3, 1, warpSize, 0, 0, d_Input, d_Output); + + HIP_CHECK(hipMemcpy(Output.data(), d_Output, Output.size() * sizeof(Output[0]), hipMemcpyDefault)); + for (size_t i = 0; i < Output.size(); i++) { + REQUIRE(Output[i] == Expected[i]); + } +} + /** * @addtogroup __shfl_sync * @{ @@ -175,4 +218,7 @@ TEST_CASE("Unit_hipShflSync") { runTestShfl_1(); runTestShfl_2(); } + SECTION("divergent execution test") { + runTestShfl_3(); + } } diff --git a/catch/unit/warp/hipShflSyncUpTests.cc b/catch/unit/warp/hipShflSyncUpTests.cc index ed4c4a640d..14a140274f 100644 --- a/catch/unit/warp/hipShflSyncUpTests.cc +++ b/catch/unit/warp/hipShflSyncUpTests.cc @@ -157,6 +157,49 @@ static void runTestShflUp_3() { } } +__global__ void shflUp_4(int *Input, int *Output) { + auto tid = threadIdx.x; + unsigned long long masks[2] = { Every5thBut9th, Every9thBit }; + + Output[tid] = -1; + if (tid % 5 == 0 || tid % 9 == 0) + Output[tid] = __shfl_up_sync(masks[tid % 9 == 0], Input[tid], tid); +} + +static void runTestShflUp_4() { + size_t warpSize = getWarpSize(); + + auto Input = std::vector(warpSize); + + for (size_t i = 0; i < Input.size(); i++) { + Input[i] = 0x55 * (i + 1); + } + + auto Output = std::vector(warpSize); + auto Expected = std::vector(warpSize); + + for (size_t i = 0; i < Expected.size(); i++) { + if (i % 9 == 0 || i % 5 == 0) { + Expected[i] = 0x55; + } else { + Expected[i] = -1; + } + } + + int* d_Input; + int* d_Output; + HIP_CHECK(hipMalloc(&d_Input, Input.size() * sizeof(Input[0]))); + HIP_CHECK(hipMalloc(&d_Output, Output.size() * sizeof(Output[0]))); + + HIP_CHECK(hipMemcpy(d_Input, Input.data(), Input.size() * sizeof(Input[0]), hipMemcpyDefault)); + hipLaunchKernelGGL(shflUp_4, 1, warpSize, 0, 0, d_Input, d_Output); + + HIP_CHECK(hipMemcpy(Output.data(), d_Output, Output.size() * sizeof(Output[0]), hipMemcpyDefault)); + for (size_t i = 0; i < Output.size(); i++) { + REQUIRE(Output[i] == Expected[i]); + } +} + /** * @addtogroup __shfl_sync * @{ @@ -242,4 +285,7 @@ TEST_CASE("Unit_hipShflSync_Up") { runTestShflUp_2<__half2>(); runTestShflUp_3<__half2>(); } + SECTION("run divergent execution test") { + runTestShflUp_4(); + } } diff --git a/catch/unit/warp/hipShflSyncXorTests.cc b/catch/unit/warp/hipShflSyncXorTests.cc index 9a48c612e1..e1f7326b76 100644 --- a/catch/unit/warp/hipShflSyncXorTests.cc +++ b/catch/unit/warp/hipShflSyncXorTests.cc @@ -148,6 +148,49 @@ static void runTestShflXor_3() { } } +__global__ void shflXor_4(int *Input, int *Output) { + auto tid = threadIdx.x; + unsigned long long masks[2] = { Every5thBut9th, Every9thBit }; + + Output[tid] = -1; + if (tid % 5 == 0 || tid % 9 == 0) + Output[tid] = __shfl_xor_sync(masks[tid % 9 == 0], Input[tid], tid); +} + +static void runTestShflXor_4() { + size_t warpSize = getWarpSize(); + + auto Input = std::vector(warpSize); + + for (size_t i = 0; i < Input.size(); i++) { + Input[i] = 0x55 * (i + 1); + } + + auto Output = std::vector(warpSize); + auto Expected = std::vector(warpSize); + + for (size_t i = 0; i < Expected.size(); i++) { + if (i % 9 == 0 || i % 5 == 0) { + Expected[i] = 0x55; + } else { + Expected[i] = -1; + } + } + + int* d_Input; + int* d_Output; + HIP_CHECK(hipMalloc(&d_Input, Input.size() * sizeof(Input[0]))); + HIP_CHECK(hipMalloc(&d_Output, Output.size() * sizeof(Output[0]))); + + HIP_CHECK(hipMemcpy(d_Input, Input.data(), Input.size() * sizeof(Input[0]), hipMemcpyDefault)); + hipLaunchKernelGGL(shflXor_4, 1, warpSize, 0, 0, d_Input, d_Output); + + HIP_CHECK(hipMemcpy(Output.data(), d_Output, Output.size() * sizeof(Output[0]), hipMemcpyDefault)); + for (size_t i = 0; i < Output.size(); i++) { + REQUIRE(Output[i] == Expected[i]); + } +} + /** * @addtogroup __shfl_sync * @{ @@ -233,4 +276,7 @@ TEST_CASE("Unit_hipShflSync_Xor") { runTestShflXor_2<__half2>(); runTestShflXor_3<__half2>(); } + SECTION("run divergent exec tests") { + runTestShflXor_4(); + } } diff --git a/catch/unit/warp/warp_common.hh b/catch/unit/warp/warp_common.hh index 9fe207c030..ae2b456bc1 100644 --- a/catch/unit/warp/warp_common.hh +++ b/catch/unit/warp/warp_common.hh @@ -167,3 +167,12 @@ inline bool compareMaskEqual(unsigned long long *Actual, unsigned long long *Exp return (unsigned)Actual[i] == (unsigned)Expected[i]; return Actual[i] == Expected[i]; } + +template +inline T alignUp(T num, size_t n) { + if (num % n == 0) { + return num; + } + + return ((num + n - 1) / n) * n; +}