From c57e0f8fe568f962654f555c64e6ee592d081f2a Mon Sep 17 00:00:00 2001 From: TomSang Date: Thu, 26 Aug 2021 23:21:15 -0400 Subject: [PATCH] SWDEV-294589 - Fix hipCGMultiGridGroupType* failure on NV (#2326) Fix the following failed tests on NV, hipCGMultiGridGroupType hipCGMultiGridGroupTypeViaBaseType hipCGMultiGridGroupTypeViaPublicApi 1. Fix wrong logic in kernel for both AMD and NV. 2. Remove unnecessary hipDeviceSynchronize(). 3. In hipCGMultiGridGroupTypeViaBaseType.cpp, change multi_grid_group as thread_group which is originally expected. 4. hipFree(syncResultD) is fixed as hipHostFree(syncResultD) 5. Optimize some host codes. Change-Id: I3fe6dac35a7b14bab12adf397b7885df83d28059 --- .../hipCGMultiGridGroupType.cpp | 53 ++++++----------- .../hipCGMultiGridGroupTypeViaBaseType.cpp | 58 ++++++++----------- .../hipCGMultiGridGroupTypeViaPublicApi.cpp | 49 +++++----------- 3 files changed, 57 insertions(+), 103 deletions(-) diff --git a/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupType.cpp b/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupType.cpp index f3af976d35..0b170de1f7 100755 --- a/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupType.cpp +++ b/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupType.cpp @@ -22,8 +22,8 @@ THE SOFTWARE. /* HIT_START - * 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 + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 -rdc=true -gencode arch=compute_60,code=sm_60 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_80,code=sm_80 + * TEST: %t * HIT_END */ @@ -37,6 +37,9 @@ THE SOFTWARE. #define ASSERT_LE(lhs, rhs) assert(lhs <= rhs) #define ASSERT_GE(lhs, rhs) assert(lhs >= rhs) +constexpr int MaxGPUs = 8; +int nGpu = 0; + using namespace cooperative_groups; static __global__ @@ -82,7 +85,7 @@ void kernel_cg_multi_grid_group_type(int* numGridsTestD, // multi-grid level sync mg.sync(); // grid (gpu) 0 does final reduction across all grids (gpus) - if (mg.grid_rank() == 0) { + if (mg.grid_rank() == 0 && blockIdx.x == 0 && threadIdx.x == 0) { syncResultD[0] = 0; for (int i = 1; i <= mg.num_grids(); ++i) { syncResultD[0] += syncResultD[i]; @@ -92,29 +95,12 @@ void kernel_cg_multi_grid_group_type(int* numGridsTestD, static void test_cg_multi_grid_group_type(int blockSize) { - // Get device count - constexpr int MaxGPUs = 8; - int nGpu = 0; - ASSERT_EQUAL(hipGetDeviceCount(&nGpu), hipSuccess); - - // Check if device suppurts multi gpu cooperative group support - hipDeviceProp_t deviceProp[MaxGPUs]; - for (int i = 0; i < nGpu; i++) { - ASSERT_EQUAL(hipSetDevice(i), hipSuccess); - hipGetDeviceProperties(&deviceProp[i], 0); - if (!deviceProp[i].cooperativeMultiDeviceLaunch) { - printf("Device doesn't support multi gpu cooperative launch"); - return; - } - hipDeviceSynchronize(); - } - // Create a stream each device hipStream_t stream[MaxGPUs]; for (int i = 0; i < nGpu; i++) { ASSERT_EQUAL(hipSetDevice(i), hipSuccess); + hipDeviceSynchronize(); // Make sure work is done on this device ASSERT_EQUAL(hipStreamCreate(&stream[i]), hipSuccess); - hipDeviceSynchronize(); } // Allocate host and device memory @@ -146,8 +132,6 @@ static void test_cg_multi_grid_group_type(int blockSize) hipHostMalloc(&syncResultD, sizeof(int) * (nGpu + 1), hipHostMallocCoherent), hipSuccess); } - - hipDeviceSynchronize(); } // Launch Kernel @@ -171,15 +155,12 @@ static void test_cg_multi_grid_group_type(int blockSize) launchParamsList[i].sharedMem = 0; launchParamsList[i].stream = stream[i]; launchParamsList[i].args = &args[i * NumKernelArgs]; - - hipDeviceSynchronize(); } - hipLaunchCooperativeKernelMultiDevice(launchParamsList, nGpu, 0); + HIPCHECK(hipLaunchCooperativeKernelMultiDevice(launchParamsList, nGpu, 0)); // Copy result from device to host for (int i = 0; i < nGpu; i++) { ASSERT_EQUAL(hipSetDevice(i), hipSuccess); - ASSERT_EQUAL(hipMemcpy(numGridsTestH[i], numGridsTestD[i], nBytes, hipMemcpyDeviceToHost), hipSuccess); ASSERT_EQUAL(hipMemcpy(gridRankTestH[i], gridRankTestD[i], nBytes, hipMemcpyDeviceToHost), @@ -191,7 +172,6 @@ static void test_cg_multi_grid_group_type(int blockSize) ASSERT_EQUAL(hipMemcpy(isValidTestH[i], isValidTestD[i], nBytes, hipMemcpyDeviceToHost), hipSuccess); - hipDeviceSynchronize(); } // Validate results @@ -231,31 +211,34 @@ static void test_cg_multi_grid_group_type(int blockSize) ASSERT_EQUAL(hipFree(isValidTestD[i]), hipSuccess); ASSERT_EQUAL(hipFree(syncTestD[i]), hipSuccess); - if (i == 0) - ASSERT_EQUAL(hipFree(syncResultD), hipSuccess); - + if (i == 0) { + ASSERT_EQUAL(hipHostFree(syncResultD), hipSuccess); + } ASSERT_EQUAL(hipHostFree(numGridsTestH[i]), hipSuccess); ASSERT_EQUAL(hipHostFree(gridRankTestH[i]), hipSuccess); ASSERT_EQUAL(hipHostFree(sizeTestH[i]), hipSuccess); ASSERT_EQUAL(hipHostFree(thdRankTestH[i]), hipSuccess); ASSERT_EQUAL(hipHostFree(isValidTestH[i]), hipSuccess); - - hipDeviceSynchronize(); } } int main() { // Set `maxThreadsPerBlock` by taking minimum among all available devices - int nGpu = 0; ASSERT_EQUAL(hipGetDeviceCount(&nGpu), hipSuccess); + if (nGpu > MaxGPUs) { + nGpu = MaxGPUs; + } int maxThreadsPerBlock = INT_MAX; for (int i = 0; i < nGpu; i++) { hipDeviceProp_t deviceProperties; ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, i), hipSuccess); + if (!deviceProperties.cooperativeMultiDeviceLaunch) { + printf("Device doesn't support cooperative launch!"); + passed(); + } int curDeviceMaxThreadsPerBlock = deviceProperties.maxThreadsPerBlock; maxThreadsPerBlock = min(maxThreadsPerBlock, curDeviceMaxThreadsPerBlock); - hipDeviceSynchronize(); } // Test block sizes which are powers of 2 diff --git a/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupTypeViaBaseType.cpp b/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupTypeViaBaseType.cpp index 4431a8a18f..0c5dc7b104 100644 --- a/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupTypeViaBaseType.cpp +++ b/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupTypeViaBaseType.cpp @@ -22,8 +22,8 @@ THE SOFTWARE. /* HIT_START - * 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 + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 -D_CG_ABI_EXPERIMENTAL -rdc=true -gencode arch=compute_60,code=sm_60 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_80,code=sm_80 + * TEST: %t * HIT_END */ @@ -37,6 +37,9 @@ THE SOFTWARE. #define ASSERT_LE(lhs, rhs) assert(lhs <= rhs) #define ASSERT_GE(lhs, rhs) assert(lhs >= rhs) +constexpr int MaxGPUs = 8; +int nGpu = 0; + using namespace cooperative_groups; static __global__ @@ -47,7 +50,8 @@ void kernel_cg_multi_grid_group_type_via_base_type(int *sizeTestD, int *syncTestD, int *syncResultD) { - multi_grid_group tg = this_multi_grid(); + thread_group tg = this_multi_grid(); // This can work if _CG_ABI_EXPERIMENTAL defined on Cuda + int gIdx = (blockIdx.x * blockDim.x) + threadIdx.x; // Test size @@ -58,8 +62,12 @@ void kernel_cg_multi_grid_group_type_via_base_type(int *sizeTestD, thdRankTestD[gIdx] = tg.thread_rank(); // Test is_valid +#ifdef __HIP_PLATFORM_AMD__ isValidTestD[gIdx] = tg.is_valid(); - +#else + // Cuda has no thread_group.is_valid() + isValidTestD[gIdx] = true; +#endif // Test sync // // Eech thread assign 1 to their respective location @@ -76,7 +84,7 @@ void kernel_cg_multi_grid_group_type_via_base_type(int *sizeTestD, // multi-grid level sync tg.sync(); // grid (gpu) 0 does final reduction across all grids (gpus) - if (this_multi_grid().grid_rank() == 0) { + if (this_multi_grid().grid_rank() == 0 && blockIdx.x == 0 && threadIdx.x == 0) { syncResultD[0] = 0; for (int i = 1; i <= this_multi_grid().num_grids(); ++i) { syncResultD[0] += syncResultD[i]; @@ -86,29 +94,12 @@ void kernel_cg_multi_grid_group_type_via_base_type(int *sizeTestD, static void test_cg_multi_grid_group_type_via_base_type(int blockSize) { - // Get device count - constexpr int MaxGPUs = 8; - int nGpu = 0; - ASSERT_EQUAL(hipGetDeviceCount(&nGpu), hipSuccess); - - // Check if device suppurts multi gpu cooperative group support - hipDeviceProp_t deviceProp[MaxGPUs]; - for (int i = 0; i < nGpu; i++) { - ASSERT_EQUAL(hipSetDevice(i), hipSuccess); - hipGetDeviceProperties(&deviceProp[i], 0); - if (!deviceProp[i].cooperativeMultiDeviceLaunch) { - printf("Device doesn't support multi gpu cooperative launch"); - return; - } - hipDeviceSynchronize(); - } - // Create a stream each device hipStream_t stream[MaxGPUs]; for (int i = 0; i < nGpu; i++) { ASSERT_EQUAL(hipSetDevice(i), hipSuccess); + hipDeviceSynchronize(); // Make sure work is done on this device ASSERT_EQUAL(hipStreamCreate(&stream[i]), hipSuccess); - hipDeviceSynchronize(); } // Allocate host and device memory @@ -137,8 +128,6 @@ static void test_cg_multi_grid_group_type_via_base_type(int blockSize) hipHostMalloc(&syncResultD, sizeof(int) * (nGpu + 1), hipHostMallocCoherent), hipSuccess); } - - hipDeviceSynchronize(); } // Launch Kernel @@ -161,10 +150,8 @@ static void test_cg_multi_grid_group_type_via_base_type(int blockSize) launchParamsList[i].sharedMem = 0; launchParamsList[i].stream = stream[i]; launchParamsList[i].args = &args[i * NumKernelArgs]; - - hipDeviceSynchronize(); } - hipLaunchCooperativeKernelMultiDevice(launchParamsList, nGpu, 0); + HIPCHECK(hipLaunchCooperativeKernelMultiDevice(launchParamsList, nGpu, 0)); // Copy result from device to host for (int i = 0; i < nGpu; i++) { @@ -178,8 +165,6 @@ static void test_cg_multi_grid_group_type_via_base_type(int blockSize) hipSuccess); ASSERT_EQUAL(hipMemcpy(isValidTestH[i], isValidTestD[i], nBytes, hipMemcpyDeviceToHost), hipSuccess); - - hipDeviceSynchronize(); } // Validate results @@ -218,29 +203,32 @@ static void test_cg_multi_grid_group_type_via_base_type(int blockSize) ASSERT_EQUAL(hipFree(syncTestD[i]), hipSuccess); if (i == 0) - ASSERT_EQUAL(hipFree(syncResultD), hipSuccess); + ASSERT_EQUAL(hipHostFree(syncResultD), hipSuccess); ASSERT_EQUAL(hipHostFree(sizeTestH[i]), hipSuccess); ASSERT_EQUAL(hipHostFree(gridRankTestH[i]), hipSuccess); ASSERT_EQUAL(hipHostFree(thdRankTestH[i]), hipSuccess); ASSERT_EQUAL(hipHostFree(isValidTestH[i]), hipSuccess); - - hipDeviceSynchronize(); } } int main() { // Set `maxThreadsPerBlock` by taking minimum among all available devices - int nGpu = 0; ASSERT_EQUAL(hipGetDeviceCount(&nGpu), hipSuccess); + if (nGpu > MaxGPUs) { + nGpu = MaxGPUs; + } int maxThreadsPerBlock = INT_MAX; for (int i = 0; i < nGpu; i++) { hipDeviceProp_t deviceProperties; ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, i), hipSuccess); + if (!deviceProperties.cooperativeMultiDeviceLaunch) { + printf("Device doesn't support cooperative launch!"); + passed(); + } int curDeviceMaxThreadsPerBlock = deviceProperties.maxThreadsPerBlock; maxThreadsPerBlock = min(maxThreadsPerBlock, curDeviceMaxThreadsPerBlock); - hipDeviceSynchronize(); } // Test block sizes which are powers of 2 diff --git a/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupTypeViaPublicApi.cpp b/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupTypeViaPublicApi.cpp index 60014e4959..a00a40345c 100644 --- a/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupTypeViaPublicApi.cpp +++ b/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupTypeViaPublicApi.cpp @@ -22,8 +22,8 @@ THE SOFTWARE. /* HIT_START - * 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 all + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 -rdc=true -gencode arch=compute_60,code=sm_60 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_80,code=sm_80 + * TEST: %t * HIT_END */ @@ -37,6 +37,9 @@ THE SOFTWARE. #define ASSERT_LE(lhs, rhs) assert(lhs <= rhs) #define ASSERT_GE(lhs, rhs) assert(lhs >= rhs) +constexpr int MaxGPUs = 8; +int nGpu = 0; + using namespace cooperative_groups; static __global__ @@ -76,7 +79,7 @@ void kernel_cg_multi_grid_group_type_via_public_api(int *sizeTestD, // multi-grid level sync via public api sync(mg); // grid (gpu) 0 does final reduction across all grids (gpus) - if (this_multi_grid().grid_rank() == 0) { + if (this_multi_grid().grid_rank() == 0 && blockIdx.x == 0 && threadIdx.x == 0) { syncResultD[0] = 0; for (int i = 1; i <= this_multi_grid().num_grids(); ++i) { syncResultD[0] += syncResultD[i]; @@ -86,29 +89,12 @@ void kernel_cg_multi_grid_group_type_via_public_api(int *sizeTestD, static void test_cg_multi_grid_group_type_via_public_api(int blockSize) { - // Get device count - constexpr int MaxGPUs = 8; - int nGpu = 0; - ASSERT_EQUAL(hipGetDeviceCount(&nGpu), hipSuccess); - - // Check if device suppurts multi gpu cooperative group support - hipDeviceProp_t deviceProp[MaxGPUs]; - for (int i = 0; i < nGpu; i++) { - ASSERT_EQUAL(hipSetDevice(i), hipSuccess); - hipGetDeviceProperties(&deviceProp[i], 0); - if (!deviceProp[i].cooperativeMultiDeviceLaunch) { - printf("Device doesn't support multi gpu cooperative launch"); - return; - } - hipDeviceSynchronize(); - } - // Create a stream each device hipStream_t stream[MaxGPUs]; for (int i = 0; i < nGpu; i++) { ASSERT_EQUAL(hipSetDevice(i), hipSuccess); + hipDeviceSynchronize(); // Make sure work is done on this device ASSERT_EQUAL(hipStreamCreate(&stream[i]), hipSuccess); - hipDeviceSynchronize(); } // Allocate host and device memory @@ -137,8 +123,6 @@ static void test_cg_multi_grid_group_type_via_public_api(int blockSize) hipHostMalloc(&syncResultD, sizeof(int) * (nGpu + 1), hipHostMallocCoherent), hipSuccess); } - - hipDeviceSynchronize(); } // Launch Kernel @@ -161,10 +145,8 @@ static void test_cg_multi_grid_group_type_via_public_api(int blockSize) launchParamsList[i].sharedMem = 0; launchParamsList[i].stream = stream[i]; launchParamsList[i].args = &args[i * NumKernelArgs]; - - hipDeviceSynchronize(); } - hipLaunchCooperativeKernelMultiDevice(launchParamsList, nGpu, 0); + HIPCHECK(hipLaunchCooperativeKernelMultiDevice(launchParamsList, nGpu, 0)); // Copy result from device to host for (int i = 0; i < nGpu; i++) { @@ -178,8 +160,6 @@ static void test_cg_multi_grid_group_type_via_public_api(int blockSize) hipSuccess); ASSERT_EQUAL(hipMemcpy(isValidTestH[i], isValidTestD[i], nBytes, hipMemcpyDeviceToHost), hipSuccess); - - hipDeviceSynchronize(); } // Validate results @@ -218,29 +198,32 @@ static void test_cg_multi_grid_group_type_via_public_api(int blockSize) ASSERT_EQUAL(hipFree(syncTestD[i]), hipSuccess); if (i == 0) - ASSERT_EQUAL(hipFree(syncResultD), hipSuccess); + ASSERT_EQUAL(hipHostFree(syncResultD), hipSuccess); ASSERT_EQUAL(hipHostFree(sizeTestH[i]), hipSuccess); ASSERT_EQUAL(hipHostFree(gridRankTestH[i]), hipSuccess); ASSERT_EQUAL(hipHostFree(thdRankTestH[i]), hipSuccess); ASSERT_EQUAL(hipHostFree(isValidTestH[i]), hipSuccess); - - hipDeviceSynchronize(); } } int main() { // Set `maxThreadsPerBlock` by taking minimum among all available devices - int nGpu = 0; ASSERT_EQUAL(hipGetDeviceCount(&nGpu), hipSuccess); + if (nGpu > MaxGPUs) { + nGpu = MaxGPUs; + } int maxThreadsPerBlock = INT_MAX; for (int i = 0; i < nGpu; i++) { hipDeviceProp_t deviceProperties; ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, i), hipSuccess); + if (!deviceProperties.cooperativeMultiDeviceLaunch) { + printf("Device doesn't support cooperative launch!"); + passed(); + } int curDeviceMaxThreadsPerBlock = deviceProperties.maxThreadsPerBlock; maxThreadsPerBlock = min(maxThreadsPerBlock, curDeviceMaxThreadsPerBlock); - hipDeviceSynchronize(); } // Test block sizes which are powers of 2