AI/rocm-systems
2
0
Garpu 0
Kode Masalah Tarik Permintaan Actions Packages Projects Rilis Wiki Kegiatan

SWDEV-549708 - WarpTest memory leak fix (#838)

Telusuri Sumber 
Signed-off-by: Sebastian Luzynski <Sebastian.Luzynski@amd.com>
This commit is contained in:
sluzynsk-amd
2025-09-11 05:22:38 -04:00
committed by GitHub
orang tua dd44ae3295
melakukan a141df3518
7 mengubah file dengan 99 tambahan dan 0 penghapusan
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip-tests/catch/unit/warp/hipMatchSyncAllTests.cc
+15
Melihat File
@@ -67,6 +67,10 @@ template <typename T> static void runTestMatchAll_1() {
for (int i = 0; i != warpSize; ++i) {
REQUIRE(Predicate[i] == ExpPredicate[i]);
}
HIP_CHECK(hipFree(d_Input));
HIP_CHECK(hipFree(d_Output));
HIP_CHECK(hipFree(d_Predicate));
}
template <typename T>
@@ -123,6 +127,10 @@ template <typename T> static void runTestMatchAll_2() {
for (int i = 0; i != warpSize; ++i) {
REQUIRE(Predicate[i] == ExpPredicate[i]);
}
HIP_CHECK(hipFree(d_Input));
HIP_CHECK(hipFree(d_Output));
HIP_CHECK(hipFree(d_Predicate));
}
template <typename T>
@@ -242,6 +250,10 @@ template <typename T> static void runTestMatchAll_3() {
for (int i = 0; i != warpSize; ++i) {
REQUIRE(Predicate[i] == ExpPredicate[i]);
}
HIP_CHECK(hipFree(d_Input));
HIP_CHECK(hipFree(d_Output));
HIP_CHECK(hipFree(d_Predicate));
}
__global__ void matchAll_4(int* Input, int* Output) {
@@ -291,6 +303,9 @@ static void runTestMatchAll_4() {
for (size_t i = 0; i < Output.size(); i++) {
REQUIRE(Output[i] == Expected[i]);
}
HIP_CHECK(hipFree(d_Input));
HIP_CHECK(hipFree(d_Output));
}
/**
projects/hip-tests/catch/unit/warp/hipMatchSyncAnyTests.cc
+9
Melihat File
@@ -67,6 +67,9 @@ template <typename T> static void runTestMatchAny_1() {
for (int i = 0; i != warpSize; ++i) {
REQUIRE(compareMaskEqual(Output, Expected, i, warpSize));
}
HIP_CHECK(hipFree(d_Input));
HIP_CHECK(hipFree(d_Output));
}
template <typename T> __global__ void matchAny_2(T* Input, unsigned long long* Output) {
@@ -124,6 +127,9 @@ template <typename T> static void runTestMatchAny_2() {
for (int i = 0; i != warpSize; ++i) {
REQUIRE(compareMaskEqual(Output, Expected, i, warpSize));
}
HIP_CHECK(hipFree(d_Input));
HIP_CHECK(hipFree(d_Output));
}
__global__ void matchAny_3(int* Input, int* Output) {
@@ -172,6 +178,9 @@ static void runTestMatchAny_3() {
for (size_t i = 0; i < Output.size(); i++) {
REQUIRE(Output[i] == Expected[i]);
}
HIP_CHECK(hipFree(d_Input));
HIP_CHECK(hipFree(d_Output));
}
/**
projects/hip-tests/catch/unit/warp/hipShflSyncDownTests.cc
+12
Melihat File
@@ -56,6 +56,9 @@ template <typename T> static void runTestShflDown_1() {
for (int i = 0; i != warpSize; ++i) {
REQUIRE(compareEqual(Output[i], Expected[i]));
}
HIP_CHECK(hipFree(d_Input));
HIP_CHECK(hipFree(d_Output));
}
// Use the mask argument to divide the warp into groups of 12 threads, and then
@@ -96,6 +99,9 @@ template <typename T> static void runTestShflDown_2() {
for (int i = 0; i != warpSize; ++i) {
REQUIRE(compareEqual(Output[i], Expected[i]));
}
HIP_CHECK(hipFree(d_Input));
HIP_CHECK(hipFree(d_Output));
}
template <typename T> __global__ void shflDown_3(T* Input, T* Output) {
@@ -137,6 +143,9 @@ template <typename T> static void runTestShflDown_3() {
for (int i = 0; i != warpSize; ++i) {
REQUIRE(compareEqual(Output[i], Expected[i]));
}
HIP_CHECK(hipFree(d_Input));
HIP_CHECK(hipFree(d_Output));
}
__global__ void shflDown_4(int* Input, int* Output) {
@@ -193,6 +202,9 @@ static void runTestShflDown_4() {
for (size_t i = 0; i < Output.size(); i++) {
REQUIRE(Output[i] == Expected[i]);
}
HIP_CHECK(hipFree(d_Input));
HIP_CHECK(hipFree(d_Output));
}
/**
projects/hip-tests/catch/unit/warp/hipShflSyncTests.cc
+9
Melihat File
@@ -56,6 +56,9 @@ template <typename T> static void runTestShfl_1() {
for (int i = 0; i != warpSize; ++i) {
REQUIRE(compareEqual(Output[i], Expected[i]));
}
HIP_CHECK(hipFree(d_Input));
HIP_CHECK(hipFree(d_Output));
}
template <typename T> __global__ void shfl_2(T* Input, T* Output) {
@@ -95,6 +98,9 @@ template <typename T> static void runTestShfl_2() {
for (int i = 0; i != warpSize; ++i) {
REQUIRE(compareEqual(Output[i], Expected[i]));
}
HIP_CHECK(hipFree(d_Input));
HIP_CHECK(hipFree(d_Output));
}
__global__ void shfl_3(int* Input, int* Output) {
@@ -138,6 +144,9 @@ static void runTestShfl_3() {
for (size_t i = 0; i < Output.size(); i++) {
REQUIRE(Output[i] == Expected[i]);
}
HIP_CHECK(hipFree(d_Input));
HIP_CHECK(hipFree(d_Output));
}
/**
projects/hip-tests/catch/unit/warp/hipShflSyncUpTests.cc
+12
Melihat File
@@ -53,6 +53,9 @@ template <typename T> static void runTestShflUp_1() {
for (int i = 0; i != warpSize; ++i) {
REQUIRE(compareEqual(Output[i], Expected[i]));
}
HIP_CHECK(hipFree(d_Input));
HIP_CHECK(hipFree(d_Output));
}
template <typename T> __global__ void shflUp_2(T* Input, T* Output) {
@@ -89,6 +92,9 @@ template <typename T> static void runTestShflUp_2() {
for (int i = 0; i != warpSize; ++i) {
REQUIRE(compareEqual(Output[i], Expected[i]));
}
HIP_CHECK(hipFree(d_Input));
HIP_CHECK(hipFree(d_Output));
}
template <typename T> __global__ void shflUp_3(T* Input, T* Output) {
@@ -131,6 +137,9 @@ template <typename T> static void runTestShflUp_3() {
for (int i = 0; i != warpSize; ++i) {
REQUIRE(compareEqual(Output[i], Expected[i]));
}
HIP_CHECK(hipFree(d_Input));
HIP_CHECK(hipFree(d_Output));
}
__global__ void shflUp_4(int* Input, int* Output) {
@@ -175,6 +184,9 @@ static void runTestShflUp_4() {
for (size_t i = 0; i < Output.size(); i++) {
REQUIRE(Output[i] == Expected[i]);
}
HIP_CHECK(hipFree(d_Input));
HIP_CHECK(hipFree(d_Output));
}
/**
projects/hip-tests/catch/unit/warp/hipShflSyncXorTests.cc
+3
Melihat File
@@ -261,6 +261,9 @@ static void runTestShflXor_4() {
for (size_t i = 0; i < Output.size(); i++) {
REQUIRE(Output[i] == Expected[i]);
}
HIP_CHECK(hipFree(d_Input));
HIP_CHECK(hipFree(d_Output));
}
/**
projects/hip-tests/catch/unit/warp/hipVoteSyncTests.cc
+39
Melihat File
@@ -49,6 +49,9 @@ static void runTestAny_1() {
for (int i = 0; i != warpSize; ++i) {
REQUIRE(Output[i] == Expected[i]);
}
HIP_CHECK(hipFree(d_Input));
HIP_CHECK(hipFree(d_Output));
}
__global__ void any_2(int* Input, int* Output) {
@@ -83,6 +86,9 @@ static void runTestAny_2_w64() {
for (int i = 0; i != warpSize; ++i) {
REQUIRE(Output[i] == Expected[i]);
}
HIP_CHECK(hipFree(d_Input));
HIP_CHECK(hipFree(d_Output));
}
static void runTestAny_2_w32() {
@@ -112,6 +118,9 @@ static void runTestAny_2_w32() {
for (int i = 0; i != warpSize; ++i) {
REQUIRE(Output[i] == Expected[i]);
}
HIP_CHECK(hipFree(d_Input));
HIP_CHECK(hipFree(d_Output));
}
__global__ void any_3(int* Input, int* Output) {
@@ -144,6 +153,9 @@ static void runTestAny_3() {
for (int i = 0; i != warpSize; ++i) {
REQUIRE(Output[i] == Expected[i]);
}
HIP_CHECK(hipFree(d_Input));
HIP_CHECK(hipFree(d_Output));
}
__global__ void any_4(int* Input, int* Output) {
@@ -192,6 +204,9 @@ static void runTestAny_4() {
for (int i = 0; i != warpSize; ++i) {
REQUIRE(Output[i] == Expected[i]);
}
HIP_CHECK(hipFree(d_Input));
HIP_CHECK(hipFree(d_Output));
}
__global__ void all_1(int* Input, int* Output) {
@@ -225,6 +240,9 @@ static void runTestAll_1_w64() {
for (int i = 0; i != warpSize; ++i) {
REQUIRE(Output[i] == Expected[i]);
}
HIP_CHECK(hipFree(d_Input));
HIP_CHECK(hipFree(d_Output));
}
static void runTestAll_1_w32() {
@@ -252,6 +270,9 @@ static void runTestAll_1_w32() {
for (int i = 0; i != warpSize; ++i) {
REQUIRE(Output[i] == Expected[i]);
}
HIP_CHECK(hipFree(d_Input));
HIP_CHECK(hipFree(d_Output));
}
__global__ void all_2(int* Input, int* Output) {
@@ -284,6 +305,9 @@ static void runTestAll_2() {
for (int i = 0; i != warpSize; ++i) {
REQUIRE(Output[i] == Expected[i]);
}
HIP_CHECK(hipFree(d_Input));
HIP_CHECK(hipFree(d_Output));
}
__global__ void all_3(int* Input, int* Output) {
@@ -316,6 +340,9 @@ static void runTestAll_3() {
for (int i = 0; i != warpSize; ++i) {
REQUIRE(Output[i] == Expected[i]);
}
HIP_CHECK(hipFree(d_Input));
HIP_CHECK(hipFree(d_Output));
}
__global__ void all_4(int* Input, int* Output) {
@@ -363,6 +390,9 @@ static void runTestAll_4() {
for (int i = 0; i != warpSize; ++i) {
REQUIRE(Output[i] == Expected[i]);
}
HIP_CHECK(hipFree(d_Input));
HIP_CHECK(hipFree(d_Output));
}
__global__ void ballot_1(int* Input, unsigned long long* Output) {
@@ -408,6 +438,9 @@ static void runTestBallot_1() {
for (int i = 0; i != warpSize; ++i) {
REQUIRE(compareMaskEqual(Output, Expected, i, warpSize));
}
HIP_CHECK(hipFree(d_Input));
HIP_CHECK(hipFree(d_Output));
}
__global__ void ballot_2(int* Input, unsigned long long* Output) {
@@ -454,6 +487,9 @@ static void runTestBallot_2() {
for (int i = 0; i != warpSize; ++i) {
REQUIRE(compareMaskEqual(Output, Expected, i, warpSize));
}
HIP_CHECK(hipFree(d_Input));
HIP_CHECK(hipFree(d_Output));
}
__global__ void ballot_3(int* Input, unsigned long long* Output) {
@@ -500,6 +536,9 @@ static void runTestBallot_3() {
for (int i = 0; i != warpSize; ++i) {
REQUIRE(compareMaskEqual(Output, Expected, i, warpSize));
}
HIP_CHECK(hipFree(d_Input));
HIP_CHECK(hipFree(d_Output));
}
/**