AI/rocm-systems
2
0
Rozštěpit 0
Zdrojový kód Úkoly Pull requesty Akce Balíčky Projekty Vydání Wiki Aktivita

SWDEV-441986 - Create tests for divergent execution of warp_sync builtins

Procházet zdrojové kódy 
Change-Id: If9ce0bdb0974732b692c779c7cae5527cec2fc83
Tento commit je obsažen v:
Branislav Brzak
2024-10-16 12:28:16 +00:00
odevzdal Branislav Brzak
rodič 533db3b042
revize 3a20cff19b
7 změnil soubory, kde provedl 306 přidání a 0 odebrání
Stáhněte soubor záplaty Stáhněte rozdílový soubor Rozbalit všechny soubory Sbalit všechny soubory
catch/unit/warp/hipMatchSyncAllTests.cc
+51
Zobrazit soubor
@@ -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<int>(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<int>(warpSize);
auto Expected = std::vector<int>(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<double>();
runTestMatchAll_3<double>();
}
SECTION("run divergent execution tests") {
runTestMatchAll_4();
}
}
catch/unit/warp/hipMatchSyncAnyTests.cc
+50
Zobrazit soubor
@@ -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<int>(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<int>(warpSize);
auto Expected = std::vector<int>(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<double>();
runTestMatchAny_2<double>();
}
SECTION("run divergent execution tests") {
runTestMatchAny_3();
}
}
catch/unit/warp/hipShflSyncDownTests.cc
+58
Zobrazit soubor
@@ -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<int>(warpSize);
for (size_t i = 0; i < Input.size(); i++) {
Input[i] = 0x55 * i;
}
auto Output = std::vector<int>(warpSize);
auto Expected = std::vector<int>(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();
}
}
catch/unit/warp/hipShflSyncTests.cc
+46
Zobrazit soubor
@@ -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<int>(warpSize);
for (size_t i = 0; i < Input.size(); i++) {
Input[i] = i;
}
auto Output = std::vector<int>(warpSize);
auto Expected = std::vector<int>(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<double>();
runTestShfl_2<double>();
}
SECTION("divergent execution test") {
runTestShfl_3();
}
}
catch/unit/warp/hipShflSyncUpTests.cc
+46
Zobrazit soubor
@@ -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<int>(warpSize);
for (size_t i = 0; i < Input.size(); i++) {
Input[i] = 0x55 * (i + 1);
}
auto Output = std::vector<int>(warpSize);
auto Expected = std::vector<int>(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();
}
}
catch/unit/warp/hipShflSyncXorTests.cc
+46
Zobrazit soubor
@@ -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<int>(warpSize);
for (size_t i = 0; i < Input.size(); i++) {
Input[i] = 0x55 * (i + 1);
}
auto Output = std::vector<int>(warpSize);
auto Expected = std::vector<int>(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();
}
}
catch/unit/warp/warp_common.hh
+9
Zobrazit soubor
@@ -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 <typename T>
inline T alignUp(T num, size_t n) {
if (num % n == 0) {
return num;
}
return ((num + n - 1) / n) * n;
}