AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-417765 - Modifies test to verify dimensions returned by group_dim()

Browse Source 
Change-Id: Ie92b9f74f4cc37b55f175cb828a88266c5167e0c
This commit is contained in:
Sourabh Betigeri
2023-08-17 15:34:53 -07:00
committed by Rakesh Roy
parent d517f0175c
commit 9cef1f0ac5
1 changed files with 16 additions and 3 deletions
Download Patch File Download Diff File Expand all files Collapse all files
catch/unit/cooperativeGrps/hipCGThreadBlockType.cc
+16 -3
View File
@@ -40,7 +40,8 @@ void kernel_cg_thread_block_type(int *sizeTestD,
int *thdRankTestD,
int *syncTestD,
dim3 *groupIndexTestD,
dim3 *thdIndexTestD)
dim3 *thdIndexTestD,
dim3 *groupDimTestD)
{
thread_block tb = this_thread_block();
int gIdx = (blockIdx.x * blockDim.x) + threadIdx.x;
@@ -64,6 +65,9 @@ void kernel_cg_thread_block_type(int *sizeTestD,
// Test thread_index
thdIndexTestD[gIdx] = tb.thread_index();
// Test group_dim aka number of threads in a block
groupDimTestD[gIdx] = tb.group_dim();
}
static void test_cg_thread_block_type(int blockSize)
@@ -74,7 +78,7 @@ static void test_cg_thread_block_type(int blockSize)
int *thdRankTestD, *thdRankTestH;
int *syncTestD, *syncTestH;
dim3 *groupIndexTestD, *groupIndexTestH;
dim3 *thdIndexTestD, *thdIndexTestH;
dim3 *thdIndexTestD, *thdIndexTestH, *groupDimTestD, *groupDimTestH;
// Allocate device memory
HIPCHECK(hipMalloc(&sizeTestD, nBytes));
@@ -82,6 +86,7 @@ static void test_cg_thread_block_type(int blockSize)
HIPCHECK(hipMalloc(&syncTestD, nBytes));
HIPCHECK(hipMalloc(&groupIndexTestD, nDim3Bytes));
HIPCHECK(hipMalloc(&thdIndexTestD, nDim3Bytes));
HIPCHECK(hipMalloc(&groupDimTestD, nDim3Bytes));
// Allocate host memory
HIPCHECK(hipHostMalloc(&sizeTestH, nBytes));
@@ -89,6 +94,7 @@ static void test_cg_thread_block_type(int blockSize)
HIPCHECK(hipHostMalloc(&syncTestH, nBytes));
HIPCHECK(hipHostMalloc(&groupIndexTestH, nDim3Bytes));
HIPCHECK(hipHostMalloc(&thdIndexTestH, nDim3Bytes));
HIPCHECK(hipHostMalloc(&groupDimTestH, nDim3Bytes));
// Launch Kernel
hipLaunchKernelGGL(kernel_cg_thread_block_type,
@@ -100,7 +106,8 @@ static void test_cg_thread_block_type(int blockSize)
thdRankTestD,
syncTestD,
groupIndexTestD,
thdIndexTestD);
thdIndexTestD,
groupDimTestD);
// Copy result from device to host
HIPCHECK(hipMemcpy(sizeTestH, sizeTestD, nBytes, hipMemcpyDeviceToHost));
@@ -108,6 +115,7 @@ static void test_cg_thread_block_type(int blockSize)
HIPCHECK(hipMemcpy(syncTestH, syncTestD, nBytes, hipMemcpyDeviceToHost));
HIPCHECK(hipMemcpy(groupIndexTestH, groupIndexTestD, nDim3Bytes, hipMemcpyDeviceToHost));
HIPCHECK(hipMemcpy(thdIndexTestH, thdIndexTestD, nDim3Bytes, hipMemcpyDeviceToHost));
HIPCHECK(hipMemcpy(groupDimTestH, groupDimTestD, nDim3Bytes, hipMemcpyDeviceToHost));
// Validate results for both blocks together
for (int i = 0; i < 2 * blockSize; ++i) {
@@ -120,6 +128,9 @@ static void test_cg_thread_block_type(int blockSize)
ASSERT_EQUAL(thdIndexTestH[i].x, (uint) i % blockSize);
ASSERT_EQUAL(thdIndexTestH[i].y, 0);
ASSERT_EQUAL(thdIndexTestH[i].z, 0);
ASSERT_EQUAL(groupDimTestH[i].x, blockSize);
ASSERT_EQUAL(groupDimTestH[i].y, 1);
ASSERT_EQUAL(groupDimTestH[i].z, 1);
}
// Free device memory
@@ -128,6 +139,7 @@ static void test_cg_thread_block_type(int blockSize)
HIPCHECK(hipFree(syncTestD));
HIPCHECK(hipFree(groupIndexTestD));
HIPCHECK(hipFree(thdIndexTestD));
HIPCHECK(hipFree(groupDimTestD));
//Free host memory
HIPCHECK(hipHostFree(sizeTestH));
@@ -135,6 +147,7 @@ static void test_cg_thread_block_type(int blockSize)
HIPCHECK(hipHostFree(syncTestH));
HIPCHECK(hipHostFree(groupIndexTestH));
HIPCHECK(hipHostFree(thdIndexTestH));
HIPCHECK(hipHostFree(groupDimTestH));
}
TEST_CASE("Unit_hipCGThreadBlockType") {