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

SWDEV-294589 - Fix hipCGMultiGridGroupType* failure on NV (#2326)

Browse Source 
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
This commit is contained in:
TomSang
2021-08-26 23:21:15 -04:00
committed by GitHub
parent 67b3681d26
commit c57e0f8fe5
3 changed files with 57 additions and 103 deletions
Download Patch File Download Diff File Expand all files Collapse all files
tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupType.cpp
+18 -35
View File
@@ -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
tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupTypeViaBaseType.cpp
+23 -35
View File
@@ -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
tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupTypeViaPublicApi.cpp
+16 -33
View File
@@ -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