diff --git a/projects/clr/hipamd/tests/src/kernel/hipDynamicShared.cpp b/projects/clr/hipamd/tests/src/kernel/hipDynamicShared.cpp index 579432bf3d..98b9eacf1d 100644 --- a/projects/clr/hipamd/tests/src/kernel/hipDynamicShared.cpp +++ b/projects/clr/hipamd/tests/src/kernel/hipDynamicShared.cpp @@ -74,7 +74,7 @@ __global__ void testExternSharedKernel(const T* A_d, const T* B_d, T* C_d, } template -void testExternShared(size_t N, size_t groupElements) { +void testExternShared(size_t N, unsigned groupElements) { size_t Nbytes = N * sizeof(T); T *A_d, *B_d, *C_d; diff --git a/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGGridGroupTypeViaBaseType.cpp b/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGGridGroupTypeViaBaseType.cpp index 7c81cbe300..2338e8d0e1 100755 --- a/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGGridGroupTypeViaBaseType.cpp +++ b/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGGridGroupTypeViaBaseType.cpp @@ -44,7 +44,7 @@ void kernel_cg_grid_group_type_via_base_type(int *sizeTestD, int *isValidTestD, int *syncTestD) { - thread_group tg = this_grid(); + grid_group tg = this_grid(); int gIdx = (blockIdx.x * blockDim.x) + threadIdx.x; // Test size diff --git a/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGGridGroupTypeViaPublicApi.cpp b/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGGridGroupTypeViaPublicApi.cpp index be8b34e83b..5d5435f13a 100755 --- a/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGGridGroupTypeViaPublicApi.cpp +++ b/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGGridGroupTypeViaPublicApi.cpp @@ -54,7 +54,7 @@ void kernel_cg_grid_group_type_via_public_api(int *sizeTestD, thdRankTestD[gIdx] = thread_rank(gg); // Test is_valid api - isValidTestD[gIdx] = is_valid(gg); + isValidTestD[gIdx] = gg.is_valid(); // Test sync api if (blockIdx.x == 0 && threadIdx.x == 0) diff --git a/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupType.cpp b/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupType.cpp index 02be0a521b..89224620b6 100755 --- a/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupType.cpp +++ b/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupType.cpp @@ -22,8 +22,8 @@ THE SOFTWARE. /* HIT_START - * BUILD: %t %s ../../test_common.cpp - * TEST: %t + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 -rdc=true -gencode arch=compute_60,code=sm_60 + * TEST: %t EXCLUDE_HIP_PLATFORM nvidia * HIT_END */ diff --git a/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupTypeViaBaseType.cpp b/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupTypeViaBaseType.cpp index 0830e807c3..d20b80526b 100644 --- a/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupTypeViaBaseType.cpp +++ b/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupTypeViaBaseType.cpp @@ -22,8 +22,8 @@ THE SOFTWARE. /* HIT_START - * BUILD: %t %s ../../test_common.cpp - * TEST: %t + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 -rdc=true -gencode arch=compute_60,code=sm_60 + * TEST: %t EXCLUDE_HIP_PLATFORM nvidia * HIT_END */ @@ -47,7 +47,7 @@ void kernel_cg_multi_grid_group_type_via_base_type(int *sizeTestD, int *syncTestD, int *syncResultD) { - thread_group tg = this_multi_grid(); + multi_grid_group tg = this_multi_grid(); int gIdx = (blockIdx.x * blockDim.x) + threadIdx.x; // Test size diff --git a/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupTypeViaPublicApi.cpp b/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupTypeViaPublicApi.cpp index 5975ffa068..307e49459a 100644 --- a/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupTypeViaPublicApi.cpp +++ b/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupTypeViaPublicApi.cpp @@ -22,8 +22,8 @@ THE SOFTWARE. /* HIT_START - * BUILD: %t %s ../../test_common.cpp - * TEST: %t + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 -rdc=true -gencode arch=compute_60,code=sm_60 + * TEST: %t EXCLUDE_HIP_PLATFORM nvidia * HIT_END */ @@ -58,7 +58,7 @@ void kernel_cg_multi_grid_group_type_via_public_api(int *sizeTestD, thdRankTestD[gIdx] = thread_rank(mg); // Test is_valid api - isValidTestD[gIdx] = is_valid(mg); + isValidTestD[gIdx] = mg.is_valid(); // Test sync api // diff --git a/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockType.cpp b/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockType.cpp index dccac38bf3..0956fc94db 100755 --- a/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockType.cpp +++ b/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockType.cpp @@ -23,7 +23,7 @@ THE SOFTWARE. /* HIT_START * BUILD: %t %s ../../test_common.cpp - * TEST: %t + * TEST: %t EXCLUDE_HIP_PLATFORM nvidia * HIT_END */ @@ -39,7 +39,6 @@ using namespace cooperative_groups; static __global__ void kernel_cg_thread_block_type(int *sizeTestD, int *thdRankTestD, - int *isValidTestD, int *syncTestD, dim3 *groupIndexTestD, dim3 *thdIndexTestD) @@ -53,9 +52,6 @@ void kernel_cg_thread_block_type(int *sizeTestD, // Test thread_rank thdRankTestD[gIdx] = tb.thread_rank(); - // Test is_valid - isValidTestD[gIdx] = tb.is_valid(); - // Test sync __shared__ int sm[2]; if (threadIdx.x == 0) @@ -78,7 +74,6 @@ static void test_cg_thread_block_type(int blockSize) int nDim3Bytes = sizeof(dim3) * 2 * blockSize; int *sizeTestD, *sizeTestH; int *thdRankTestD, *thdRankTestH; - int *isValidTestD, *isValidTestH; int *syncTestD, *syncTestH; dim3 *groupIndexTestD, *groupIndexTestH; dim3 *thdIndexTestD, *thdIndexTestH; @@ -86,7 +81,6 @@ static void test_cg_thread_block_type(int blockSize) // Allocate device memory ASSERT_EQUAL(hipMalloc(&sizeTestD, nBytes), hipSuccess); ASSERT_EQUAL(hipMalloc(&thdRankTestD, nBytes), hipSuccess); - ASSERT_EQUAL(hipMalloc(&isValidTestD, nBytes), hipSuccess); ASSERT_EQUAL(hipMalloc(&syncTestD, nBytes), hipSuccess); ASSERT_EQUAL(hipMalloc(&groupIndexTestD, nDim3Bytes), hipSuccess); ASSERT_EQUAL(hipMalloc(&thdIndexTestD, nDim3Bytes), hipSuccess); @@ -94,7 +88,6 @@ static void test_cg_thread_block_type(int blockSize) // Allocate host memory ASSERT_EQUAL(hipHostMalloc(&sizeTestH, nBytes), hipSuccess); ASSERT_EQUAL(hipHostMalloc(&thdRankTestH, nBytes), hipSuccess); - ASSERT_EQUAL(hipHostMalloc(&isValidTestH, nBytes), hipSuccess); ASSERT_EQUAL(hipHostMalloc(&syncTestH, nBytes), hipSuccess); ASSERT_EQUAL(hipHostMalloc(&groupIndexTestH, nDim3Bytes), hipSuccess); ASSERT_EQUAL(hipHostMalloc(&thdIndexTestH, nDim3Bytes), hipSuccess); @@ -107,7 +100,6 @@ static void test_cg_thread_block_type(int blockSize) 0, sizeTestD, thdRankTestD, - isValidTestD, syncTestD, groupIndexTestD, thdIndexTestD); @@ -117,8 +109,6 @@ static void test_cg_thread_block_type(int blockSize) hipSuccess); ASSERT_EQUAL(hipMemcpy(thdRankTestH, thdRankTestD, nBytes, hipMemcpyDeviceToHost), hipSuccess); - ASSERT_EQUAL(hipMemcpy(isValidTestH, isValidTestD, nBytes, hipMemcpyDeviceToHost), - hipSuccess); ASSERT_EQUAL(hipMemcpy(syncTestH, syncTestD, nBytes, hipMemcpyDeviceToHost), hipSuccess); ASSERT_EQUAL(hipMemcpy(groupIndexTestH, groupIndexTestD, nDim3Bytes, hipMemcpyDeviceToHost), @@ -130,7 +120,6 @@ static void test_cg_thread_block_type(int blockSize) for (int i = 0; i < 2 * blockSize; ++i) { ASSERT_EQUAL(sizeTestH[i], blockSize); ASSERT_EQUAL(thdRankTestH[i], i % blockSize); - ASSERT_EQUAL(isValidTestH[i], 1); ASSERT_EQUAL(syncTestH[i], 200); ASSERT_EQUAL(groupIndexTestH[i].x, i / blockSize); ASSERT_EQUAL(groupIndexTestH[i].y, 0); @@ -143,7 +132,6 @@ static void test_cg_thread_block_type(int blockSize) // Free device memory ASSERT_EQUAL(hipFree(sizeTestD), hipSuccess); ASSERT_EQUAL(hipFree(thdRankTestD), hipSuccess); - ASSERT_EQUAL(hipFree(isValidTestD), hipSuccess); ASSERT_EQUAL(hipFree(syncTestD), hipSuccess); ASSERT_EQUAL(hipFree(groupIndexTestD), hipSuccess); ASSERT_EQUAL(hipFree(thdIndexTestD), hipSuccess); @@ -151,7 +139,6 @@ static void test_cg_thread_block_type(int blockSize) //Free host memory ASSERT_EQUAL(hipHostFree(sizeTestH), hipSuccess); ASSERT_EQUAL(hipHostFree(thdRankTestH), hipSuccess); - ASSERT_EQUAL(hipHostFree(isValidTestH), hipSuccess); ASSERT_EQUAL(hipHostFree(syncTestH), hipSuccess); ASSERT_EQUAL(hipHostFree(groupIndexTestH), hipSuccess); ASSERT_EQUAL(hipHostFree(thdIndexTestH), hipSuccess); diff --git a/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockTypeViaBaseType.cpp b/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockTypeViaBaseType.cpp index b0a42782c0..f9dba1cf89 100755 --- a/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockTypeViaBaseType.cpp +++ b/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockTypeViaBaseType.cpp @@ -23,7 +23,7 @@ THE SOFTWARE. /* HIT_START * BUILD: %t %s ../../test_common.cpp - * TEST: %t + * TEST: %t EXCLUDE_HIP_PLATFORM nvidia * HIT_END */ @@ -39,7 +39,6 @@ using namespace cooperative_groups; static __global__ void kernel_cg_thread_block_type_via_base_type(int *sizeTestD, int *thdRankTestD, - int *isValidTestD, int *syncTestD) { thread_group tg = this_thread_block(); @@ -51,9 +50,6 @@ void kernel_cg_thread_block_type_via_base_type(int *sizeTestD, // Test thread_rank thdRankTestD[gIdx] = tg.thread_rank(); - // Test is_valid - isValidTestD[gIdx] = tg.is_valid(); - // Test sync __shared__ int sm[2]; if (threadIdx.x == 0) @@ -69,19 +65,16 @@ static void test_cg_thread_block_type_via_base_type(int blockSize) int nBytes = sizeof(int) * 2 * blockSize; int *sizeTestD, *sizeTestH; int *thdRankTestD, *thdRankTestH; - int *isValidTestD, *isValidTestH; int *syncTestD, *syncTestH; // Allocate device memory ASSERT_EQUAL(hipMalloc(&sizeTestD, nBytes), hipSuccess); ASSERT_EQUAL(hipMalloc(&thdRankTestD, nBytes), hipSuccess); - ASSERT_EQUAL(hipMalloc(&isValidTestD, nBytes), hipSuccess); ASSERT_EQUAL(hipMalloc(&syncTestD, nBytes), hipSuccess); // Allocate host memory ASSERT_EQUAL(hipHostMalloc(&sizeTestH, nBytes), hipSuccess); ASSERT_EQUAL(hipHostMalloc(&thdRankTestH, nBytes), hipSuccess); - ASSERT_EQUAL(hipHostMalloc(&isValidTestH, nBytes), hipSuccess); ASSERT_EQUAL(hipHostMalloc(&syncTestH, nBytes), hipSuccess); // Launch Kernel @@ -92,7 +85,6 @@ static void test_cg_thread_block_type_via_base_type(int blockSize) 0, sizeTestD, thdRankTestD, - isValidTestD, syncTestD); // Copy result from device to host @@ -100,8 +92,6 @@ static void test_cg_thread_block_type_via_base_type(int blockSize) hipSuccess); ASSERT_EQUAL(hipMemcpy(thdRankTestH, thdRankTestD, nBytes, hipMemcpyDeviceToHost), hipSuccess); - ASSERT_EQUAL(hipMemcpy(isValidTestH, isValidTestD, nBytes, hipMemcpyDeviceToHost), - hipSuccess); ASSERT_EQUAL(hipMemcpy(syncTestH, syncTestD, nBytes, hipMemcpyDeviceToHost), hipSuccess); @@ -109,20 +99,17 @@ static void test_cg_thread_block_type_via_base_type(int blockSize) for (int i = 0; i < 2 * blockSize; ++i) { ASSERT_EQUAL(sizeTestH[i], blockSize); ASSERT_EQUAL(thdRankTestH[i], i % blockSize); - ASSERT_EQUAL(isValidTestH[i], 1); ASSERT_EQUAL(syncTestH[i], 200); } // Free device memory ASSERT_EQUAL(hipFree(sizeTestD), hipSuccess); ASSERT_EQUAL(hipFree(thdRankTestD), hipSuccess); - ASSERT_EQUAL(hipFree(isValidTestD), hipSuccess); ASSERT_EQUAL(hipFree(syncTestD), hipSuccess); //Free host memory ASSERT_EQUAL(hipHostFree(sizeTestH), hipSuccess); ASSERT_EQUAL(hipHostFree(thdRankTestH), hipSuccess); - ASSERT_EQUAL(hipHostFree(isValidTestH), hipSuccess); ASSERT_EQUAL(hipHostFree(syncTestH), hipSuccess); } diff --git a/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockTypeViaPublicApi.cpp b/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockTypeViaPublicApi.cpp index e4a6a6e330..9b626e9f31 100755 --- a/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockTypeViaPublicApi.cpp +++ b/projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockTypeViaPublicApi.cpp @@ -23,7 +23,7 @@ THE SOFTWARE. /* HIT_START * BUILD: %t %s ../../test_common.cpp - * TEST: %t + * TEST: %t EXCLUDE_HIP_PLATFORM nvidia * HIT_END */ @@ -39,7 +39,6 @@ using namespace cooperative_groups; static __global__ void kernel_cg_thread_block_type_via_public_api(int *sizeTestD, int *thdRankTestD, - int *isValidTestD, int *syncTestD) { thread_block tb = this_thread_block(); @@ -51,9 +50,6 @@ void kernel_cg_thread_block_type_via_public_api(int *sizeTestD, // Test thread_rank api thdRankTestD[gIdx] = thread_rank(tb); - // Test is_valid api - isValidTestD[gIdx] = is_valid(tb); - // Test sync api __shared__ int sm[2]; if (threadIdx.x == 0) @@ -69,19 +65,16 @@ static void test_cg_thread_block_type_via_public_api(int blockSize) int nBytes = sizeof(int) * 2 * blockSize; int *sizeTestD, *sizeTestH; int *thdRankTestD, *thdRankTestH; - int *isValidTestD, *isValidTestH; int *syncTestD, *syncTestH; // Allocate device memory ASSERT_EQUAL(hipMalloc(&sizeTestD, nBytes), hipSuccess); ASSERT_EQUAL(hipMalloc(&thdRankTestD, nBytes), hipSuccess); - ASSERT_EQUAL(hipMalloc(&isValidTestD, nBytes), hipSuccess); ASSERT_EQUAL(hipMalloc(&syncTestD, nBytes), hipSuccess); // Allocate host memory ASSERT_EQUAL(hipHostMalloc(&sizeTestH, nBytes), hipSuccess); ASSERT_EQUAL(hipHostMalloc(&thdRankTestH, nBytes), hipSuccess); - ASSERT_EQUAL(hipHostMalloc(&isValidTestH, nBytes), hipSuccess); ASSERT_EQUAL(hipHostMalloc(&syncTestH, nBytes), hipSuccess); // Launch Kernel @@ -92,7 +85,6 @@ static void test_cg_thread_block_type_via_public_api(int blockSize) 0, sizeTestD, thdRankTestD, - isValidTestD, syncTestD); // Copy result from device to host @@ -100,8 +92,6 @@ static void test_cg_thread_block_type_via_public_api(int blockSize) hipSuccess); ASSERT_EQUAL(hipMemcpy(thdRankTestH, thdRankTestD, nBytes, hipMemcpyDeviceToHost), hipSuccess); - ASSERT_EQUAL(hipMemcpy(isValidTestH, isValidTestD, nBytes, hipMemcpyDeviceToHost), - hipSuccess); ASSERT_EQUAL(hipMemcpy(syncTestH, syncTestD, nBytes, hipMemcpyDeviceToHost), hipSuccess); @@ -109,20 +99,17 @@ static void test_cg_thread_block_type_via_public_api(int blockSize) for (int i = 0; i < 2 * blockSize; ++i) { ASSERT_EQUAL(sizeTestH[i], blockSize); ASSERT_EQUAL(thdRankTestH[i], i % blockSize); - ASSERT_EQUAL(isValidTestH[i], 1); ASSERT_EQUAL(syncTestH[i], 200); } // Free device memory ASSERT_EQUAL(hipFree(sizeTestD), hipSuccess); ASSERT_EQUAL(hipFree(thdRankTestD), hipSuccess); - ASSERT_EQUAL(hipFree(isValidTestD), hipSuccess); ASSERT_EQUAL(hipFree(syncTestD), hipSuccess); //Free host memory ASSERT_EQUAL(hipHostFree(sizeTestH), hipSuccess); ASSERT_EQUAL(hipHostFree(thdRankTestH), hipSuccess); - ASSERT_EQUAL(hipHostFree(isValidTestH), hipSuccess); ASSERT_EQUAL(hipHostFree(syncTestH), hipSuccess); }