diff --git a/projects/hip/tests/catch/unit/CMakeLists.txt b/projects/hip/tests/catch/unit/CMakeLists.txt index 6e97f24663..d856006312 100644 --- a/projects/hip/tests/catch/unit/CMakeLists.txt +++ b/projects/hip/tests/catch/unit/CMakeLists.txt @@ -33,6 +33,7 @@ add_subdirectory(kernel) add_subdirectory(multiThread) add_subdirectory(compiler) add_subdirectory(errorHandling) +add_subdirectory(cooperativeGrps) if(HIP_PLATFORM STREQUAL "amd") add_subdirectory(clock) endif() diff --git a/projects/hip/tests/catch/unit/cooperativeGrps/CMakeLists.txt b/projects/hip/tests/catch/unit/cooperativeGrps/CMakeLists.txt new file mode 100644 index 0000000000..61c567aee4 --- /dev/null +++ b/projects/hip/tests/catch/unit/cooperativeGrps/CMakeLists.txt @@ -0,0 +1,22 @@ +# Common Tests - Test independent of all platforms +set(TEST_SRC + hipCGThreadBlockType.cc + hipCGThreadBlockTypeViaBaseType.cc + hipCGThreadBlockTypeViaPublicApi.cc + hipCGMultiGridGroupType.cc + hipCGMultiGridGroupTypeViaBaseType.cc + hipCGMultiGridGroupTypeViaPublicApi.cc +) +if(HIP_PLATFORM STREQUAL "nvidia") + set_source_files_properties(hipCGMultiGridGroupType.cc PROPERTIES COMPILE_FLAGS "-rdc=true -gencode arch=compute_60,code=sm_60 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_80,code=sm_80") + set_source_files_properties(hipCGMultiGridGroupTypeViaBaseType.cc PROPERTIES COMPILE_FLAGS "-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") + set_source_files_properties(hipCGMultiGridGroupTypeViaPublicApi.cc PROPERTIES COMPILE_FLAGS "-rdc=true -gencode arch=compute_60,code=sm_60 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_80,code=sm_80") + hip_add_exe_to_target(NAME coopGrpTest + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME build_tests + LINKER_LIBS "-rdc=true -gencode arch=compute_60,code=sm_60 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_80,code=sm_80") +else() + hip_add_exe_to_target(NAME coopGrpTest + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME build_tests) +endif() diff --git a/projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupType.cpp b/projects/hip/tests/catch/unit/cooperativeGrps/hipCGMultiGridGroupType.cc similarity index 58% rename from projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupType.cpp rename to projects/hip/tests/catch/unit/cooperativeGrps/hipCGMultiGridGroupType.cc index 0b170de1f7..1dd2a8f3b5 100644 --- a/projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupType.cpp +++ b/projects/hip/tests/catch/unit/cooperativeGrps/hipCGMultiGridGroupType.cc @@ -27,20 +27,15 @@ THE SOFTWARE. * HIT_END */ -#include "test_common.h" -#include "hip/hip_cooperative_groups.h" -#include -#include -#include +#include +#include -#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs) -#define ASSERT_LE(lhs, rhs) assert(lhs <= rhs) -#define ASSERT_GE(lhs, rhs) assert(lhs >= rhs) - -constexpr int MaxGPUs = 8; -int nGpu = 0; +#define ASSERT_EQUAL(lhs, rhs) HIPASSERT(lhs == rhs) +#define ASSERT_LE(lhs, rhs) HIPASSERT(lhs <= rhs) +#define ASSERT_GE(lhs, rhs) HIPASSERT(lhs >= rhs) using namespace cooperative_groups; +constexpr int MaxGPUs = 8; static __global__ void kernel_cg_multi_grid_group_type(int* numGridsTestD, @@ -77,7 +72,7 @@ void kernel_cg_multi_grid_group_type(int* numGridsTestD, this_grid().sync(); // Thread 0 from work-group 0 of current grid (gpu) does grid level reduction if (blockIdx.x == 0 && threadIdx.x == 0) { - for (int i = 1; i < gridDim.x * blockDim.x; ++i) { + for (uint i = 1; i < gridDim.x * blockDim.x; ++i) { syncTestD[0] += syncTestD[i]; } syncResultD[mg.grid_rank() + 1] = syncTestD[0]; @@ -87,20 +82,20 @@ void kernel_cg_multi_grid_group_type(int* numGridsTestD, // grid (gpu) 0 does final reduction across all grids (gpus) if (mg.grid_rank() == 0 && blockIdx.x == 0 && threadIdx.x == 0) { syncResultD[0] = 0; - for (int i = 1; i <= mg.num_grids(); ++i) { + for (uint i = 1; i <= mg.num_grids(); ++i) { syncResultD[0] += syncResultD[i]; } } } -static void test_cg_multi_grid_group_type(int blockSize) +static void test_cg_multi_grid_group_type(int blockSize, int nGpu) { // 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); + HIPCHECK(hipSetDevice(i)); + HIPCHECK(hipDeviceSynchronize()); // Make sure work is done on this device + HIPCHECK(hipStreamCreate(&stream[i])); } // Allocate host and device memory @@ -112,25 +107,23 @@ static void test_cg_multi_grid_group_type(int blockSize) int *isValidTestD[MaxGPUs], *isValidTestH[MaxGPUs]; int *syncTestD[MaxGPUs], *syncResultD; for (int i = 0; i < nGpu; i++) { - ASSERT_EQUAL(hipSetDevice(i), hipSuccess); + HIPCHECK(hipSetDevice(i)); - ASSERT_EQUAL(hipMalloc(&numGridsTestD[i], nBytes), hipSuccess); - ASSERT_EQUAL(hipMalloc(&gridRankTestD[i], nBytes), hipSuccess); - ASSERT_EQUAL(hipMalloc(&sizeTestD[i], nBytes), hipSuccess); - ASSERT_EQUAL(hipMalloc(&thdRankTestD[i], nBytes), hipSuccess); - ASSERT_EQUAL(hipMalloc(&isValidTestD[i], nBytes), hipSuccess); - ASSERT_EQUAL(hipMalloc(&syncTestD[i], nBytes), hipSuccess); + HIPCHECK(hipMalloc(&numGridsTestD[i], nBytes)); + HIPCHECK(hipMalloc(&gridRankTestD[i], nBytes)); + HIPCHECK(hipMalloc(&sizeTestD[i], nBytes)); + HIPCHECK(hipMalloc(&thdRankTestD[i], nBytes)); + HIPCHECK(hipMalloc(&isValidTestD[i], nBytes)); + HIPCHECK(hipMalloc(&syncTestD[i], nBytes)); - ASSERT_EQUAL(hipHostMalloc(&numGridsTestH[i], nBytes), hipSuccess); - ASSERT_EQUAL(hipHostMalloc(&gridRankTestH[i], nBytes), hipSuccess); - ASSERT_EQUAL(hipHostMalloc(&sizeTestH[i], nBytes), hipSuccess); - ASSERT_EQUAL(hipHostMalloc(&thdRankTestH[i], nBytes), hipSuccess); - ASSERT_EQUAL(hipHostMalloc(&isValidTestH[i], nBytes), hipSuccess); + HIPCHECK(hipHostMalloc(&numGridsTestH[i], nBytes)); + HIPCHECK(hipHostMalloc(&gridRankTestH[i], nBytes)); + HIPCHECK(hipHostMalloc(&sizeTestH[i], nBytes)); + HIPCHECK(hipHostMalloc(&thdRankTestH[i], nBytes)); + HIPCHECK(hipHostMalloc(&isValidTestH[i], nBytes)); if (i == 0) { - ASSERT_EQUAL( - hipHostMalloc(&syncResultD, sizeof(int) * (nGpu + 1), hipHostMallocCoherent), - hipSuccess); + HIPCHECK(hipHostMalloc(&syncResultD, sizeof(int) * (nGpu + 1), hipHostMallocCoherent)); } } @@ -139,7 +132,7 @@ static void test_cg_multi_grid_group_type(int blockSize) hipLaunchParams* launchParamsList = new hipLaunchParams[nGpu]; void* args[MaxGPUs * NumKernelArgs]; for (int i = 0; i < nGpu; i++) { - ASSERT_EQUAL(hipSetDevice(i), hipSuccess); + HIPCHECK(hipSetDevice(i)); args[i * NumKernelArgs] = &numGridsTestD[i]; args[i * NumKernelArgs + 1] = &gridRankTestD[i]; @@ -160,18 +153,12 @@ static void test_cg_multi_grid_group_type(int blockSize) // 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), - hipSuccess); - ASSERT_EQUAL(hipMemcpy(sizeTestH[i], sizeTestD[i], nBytes, hipMemcpyDeviceToHost), - hipSuccess); - ASSERT_EQUAL(hipMemcpy(thdRankTestH[i], thdRankTestD[i], nBytes, hipMemcpyDeviceToHost), - hipSuccess); - ASSERT_EQUAL(hipMemcpy(isValidTestH[i], isValidTestD[i], nBytes, hipMemcpyDeviceToHost), - hipSuccess); - + HIPCHECK(hipSetDevice(i)); + HIPCHECK(hipMemcpy(numGridsTestH[i], numGridsTestD[i], nBytes, hipMemcpyDeviceToHost)); + HIPCHECK(hipMemcpy(gridRankTestH[i], gridRankTestD[i], nBytes, hipMemcpyDeviceToHost)); + HIPCHECK(hipMemcpy(sizeTestH[i], sizeTestD[i], nBytes, hipMemcpyDeviceToHost)); + HIPCHECK(hipMemcpy(thdRankTestH[i], thdRankTestD[i], nBytes, hipMemcpyDeviceToHost)); + HIPCHECK(hipMemcpy(isValidTestH[i], isValidTestD[i], nBytes, hipMemcpyDeviceToHost)); } // Validate results @@ -193,7 +180,7 @@ static void test_cg_multi_grid_group_type(int blockSize) gridsSeen[i] = gridRankTestH[i][0]; for (int k = 0; k < i; ++k) { if (gridsSeen[k] == gridsSeen[i]) { - assert (false && "Grid rank in multi-gpu setup should be unique"); + assert(false && "Grid rank in multi-gpu setup should be unique"); } } } @@ -202,60 +189,52 @@ static void test_cg_multi_grid_group_type(int blockSize) // Free host and device memory delete [] launchParamsList; for (int i = 0; i < nGpu; i++) { - ASSERT_EQUAL(hipSetDevice(i), hipSuccess); + HIPCHECK(hipSetDevice(i)); - ASSERT_EQUAL(hipFree(numGridsTestD[i]), hipSuccess); - ASSERT_EQUAL(hipFree(gridRankTestD[i]), hipSuccess); - ASSERT_EQUAL(hipFree(sizeTestD[i]), hipSuccess); - ASSERT_EQUAL(hipFree(thdRankTestD[i]), hipSuccess); - ASSERT_EQUAL(hipFree(isValidTestD[i]), hipSuccess); - ASSERT_EQUAL(hipFree(syncTestD[i]), hipSuccess); + HIPCHECK(hipFree(numGridsTestD[i])); + HIPCHECK(hipFree(gridRankTestD[i])); + HIPCHECK(hipFree(sizeTestD[i])); + HIPCHECK(hipFree(thdRankTestD[i])); + HIPCHECK(hipFree(isValidTestD[i])); + HIPCHECK(hipFree(syncTestD[i])); if (i == 0) { - ASSERT_EQUAL(hipHostFree(syncResultD), hipSuccess); + HIPCHECK(hipHostFree(syncResultD)); } - 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); + HIPCHECK(hipHostFree(numGridsTestH[i])); + HIPCHECK(hipHostFree(gridRankTestH[i])); + HIPCHECK(hipHostFree(sizeTestH[i])); + HIPCHECK(hipHostFree(thdRankTestH[i])); + HIPCHECK(hipHostFree(isValidTestH[i])); } } -int main() -{ +TEST_CASE("Unit_hipCGMultiGridGroupType") { + int nGpu = 0; + HIPCHECK(hipGetDeviceCount(&nGpu)); + nGpu = min(nGpu, MaxGPUs); + // Set `maxThreadsPerBlock` by taking minimum among all available devices - ASSERT_EQUAL(hipGetDeviceCount(&nGpu), hipSuccess); - if (nGpu > MaxGPUs) { - nGpu = MaxGPUs; - } int maxThreadsPerBlock = INT_MAX; + hipDeviceProp_t deviceProperties; for (int i = 0; i < nGpu; i++) { - hipDeviceProp_t deviceProperties; - ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, i), hipSuccess); + HIPCHECK(hipGetDeviceProperties(&deviceProperties, i)); if (!deviceProperties.cooperativeMultiDeviceLaunch) { - printf("Device doesn't support cooperative launch!"); - passed(); + HipTest::HIP_SKIP_TEST("Device doesn't support cooperative launch!"); + return; } - int curDeviceMaxThreadsPerBlock = deviceProperties.maxThreadsPerBlock; - maxThreadsPerBlock = min(maxThreadsPerBlock, curDeviceMaxThreadsPerBlock); + maxThreadsPerBlock = min(maxThreadsPerBlock, deviceProperties.maxThreadsPerBlock); } - // Test block sizes which are powers of 2 - int i = 0; - while (true) { - int blockSize = pow(2, i); - if (blockSize > maxThreadsPerBlock) - break; - test_cg_multi_grid_group_type(blockSize); - ++i; + // Test for blockSizes in powers of 2 + for (int blockSize = 2; blockSize <= maxThreadsPerBlock; blockSize = blockSize*2) { + test_cg_multi_grid_group_type(blockSize, nGpu); } - // Test some random block sizes - for(int j = 0; j < 10 ; ++j) { - int blockSize = rand() % maxThreadsPerBlock; - test_cg_multi_grid_group_type(blockSize); + // Test for random blockSizes, but the sequence is the same every execution + srand(0); + for (int i = 0; i < 10; i++) { + // Test fails for 0 thread per block + test_cg_multi_grid_group_type(max(2, rand() % maxThreadsPerBlock), nGpu); } - - passed(); } diff --git a/projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupTypeViaBaseType.cpp b/projects/hip/tests/catch/unit/cooperativeGrps/hipCGMultiGridGroupTypeViaBaseType.cc similarity index 63% rename from projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupTypeViaBaseType.cpp rename to projects/hip/tests/catch/unit/cooperativeGrps/hipCGMultiGridGroupTypeViaBaseType.cc index 0c5dc7b104..408f3b0075 100644 --- a/projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupTypeViaBaseType.cpp +++ b/projects/hip/tests/catch/unit/cooperativeGrps/hipCGMultiGridGroupTypeViaBaseType.cc @@ -27,20 +27,18 @@ THE SOFTWARE. * HIT_END */ -#include "test_common.h" -#include "hip/hip_cooperative_groups.h" +#include +#include #include #include #include -#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs) -#define ASSERT_LE(lhs, rhs) assert(lhs <= rhs) -#define ASSERT_GE(lhs, rhs) assert(lhs >= rhs) - -constexpr int MaxGPUs = 8; -int nGpu = 0; +#define ASSERT_EQUAL(lhs, rhs) HIPASSERT(lhs == rhs) +#define ASSERT_LE(lhs, rhs) HIPASSERT(lhs <= rhs) +#define ASSERT_GE(lhs, rhs) HIPASSERT(lhs >= rhs) using namespace cooperative_groups; +constexpr int MaxGPUs = 8; static __global__ void kernel_cg_multi_grid_group_type_via_base_type(int *sizeTestD, @@ -76,7 +74,7 @@ void kernel_cg_multi_grid_group_type_via_base_type(int *sizeTestD, this_grid().sync(); // Thread 0 from work-group 0 of current grid (gpu) does grid level reduction if (blockIdx.x == 0 && threadIdx.x == 0) { - for (int i = 1; i < gridDim.x * blockDim.x; ++i) { + for (uint i = 1; i < gridDim.x * blockDim.x; ++i) { syncTestD[0] += syncTestD[i]; } syncResultD[this_multi_grid().grid_rank() + 1] = syncTestD[0]; @@ -86,20 +84,20 @@ void kernel_cg_multi_grid_group_type_via_base_type(int *sizeTestD, // grid (gpu) 0 does final reduction across all grids (gpus) 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) { + for (uint i = 1; i <= this_multi_grid().num_grids(); ++i) { syncResultD[0] += syncResultD[i]; } } } -static void test_cg_multi_grid_group_type_via_base_type(int blockSize) +static void test_cg_multi_grid_group_type_via_base_type(int blockSize, int nGpu) { // 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); + HIPCHECK(hipSetDevice(i)); + HIPCHECK(hipDeviceSynchronize()); // Make sure work is done on this device + HIPCHECK(hipStreamCreate(&stream[i])); } // Allocate host and device memory @@ -110,23 +108,21 @@ static void test_cg_multi_grid_group_type_via_base_type(int blockSize) int *isValidTestD[MaxGPUs], *isValidTestH[MaxGPUs]; int *syncTestD[MaxGPUs], *syncResultD; for (int i = 0; i < nGpu; i++) { - ASSERT_EQUAL(hipSetDevice(i), hipSuccess); + HIPCHECK(hipSetDevice(i)); - ASSERT_EQUAL(hipMalloc(&sizeTestD[i], nBytes), hipSuccess); - ASSERT_EQUAL(hipMalloc(&gridRankTestD[i], nBytes), hipSuccess); - ASSERT_EQUAL(hipMalloc(&thdRankTestD[i], nBytes), hipSuccess); - ASSERT_EQUAL(hipMalloc(&isValidTestD[i], nBytes), hipSuccess); - ASSERT_EQUAL(hipMalloc(&syncTestD[i], nBytes), hipSuccess); + HIPCHECK(hipMalloc(&sizeTestD[i], nBytes)); + HIPCHECK(hipMalloc(&gridRankTestD[i], nBytes)); + HIPCHECK(hipMalloc(&thdRankTestD[i], nBytes)); + HIPCHECK(hipMalloc(&isValidTestD[i], nBytes)); + HIPCHECK(hipMalloc(&syncTestD[i], nBytes)); - ASSERT_EQUAL(hipHostMalloc(&sizeTestH[i], nBytes), hipSuccess); - ASSERT_EQUAL(hipHostMalloc(&gridRankTestH[i], nBytes), hipSuccess); - ASSERT_EQUAL(hipHostMalloc(&thdRankTestH[i], nBytes), hipSuccess); - ASSERT_EQUAL(hipHostMalloc(&isValidTestH[i], nBytes), hipSuccess); + HIPCHECK(hipHostMalloc(&sizeTestH[i], nBytes)); + HIPCHECK(hipHostMalloc(&gridRankTestH[i], nBytes)); + HIPCHECK(hipHostMalloc(&thdRankTestH[i], nBytes)); + HIPCHECK(hipHostMalloc(&isValidTestH[i], nBytes)); if (i == 0) { - ASSERT_EQUAL( - hipHostMalloc(&syncResultD, sizeof(int) * (nGpu + 1), hipHostMallocCoherent), - hipSuccess); + HIPCHECK(hipHostMalloc(&syncResultD, sizeof(int) * (nGpu + 1), hipHostMallocCoherent)); } } @@ -135,7 +131,7 @@ static void test_cg_multi_grid_group_type_via_base_type(int blockSize) hipLaunchParams* launchParamsList = new hipLaunchParams[nGpu]; void* args[MaxGPUs * NumKernelArgs]; for (int i = 0; i < nGpu; i++) { - ASSERT_EQUAL(hipSetDevice(i), hipSuccess); + HIPCHECK(hipSetDevice(i)); args[i * NumKernelArgs ] = &sizeTestD[i]; args[i * NumKernelArgs + 1] = &gridRankTestD[i]; @@ -155,16 +151,11 @@ static void test_cg_multi_grid_group_type_via_base_type(int blockSize) // Copy result from device to host for (int i = 0; i < nGpu; i++) { - ASSERT_EQUAL(hipSetDevice(i), hipSuccess); - - ASSERT_EQUAL(hipMemcpy(sizeTestH[i], sizeTestD[i], nBytes, hipMemcpyDeviceToHost), - hipSuccess); - ASSERT_EQUAL(hipMemcpy(gridRankTestH[i], gridRankTestD[i], nBytes, hipMemcpyDeviceToHost), - hipSuccess); - ASSERT_EQUAL(hipMemcpy(thdRankTestH[i], thdRankTestD[i], nBytes, hipMemcpyDeviceToHost), - hipSuccess); - ASSERT_EQUAL(hipMemcpy(isValidTestH[i], isValidTestD[i], nBytes, hipMemcpyDeviceToHost), - hipSuccess); + HIPCHECK(hipSetDevice(i)); + HIPCHECK(hipMemcpy(sizeTestH[i], sizeTestD[i], nBytes, hipMemcpyDeviceToHost)); + HIPCHECK(hipMemcpy(gridRankTestH[i], gridRankTestD[i], nBytes, hipMemcpyDeviceToHost)); + HIPCHECK(hipMemcpy(thdRankTestH[i], thdRankTestD[i], nBytes, hipMemcpyDeviceToHost)); + HIPCHECK(hipMemcpy(isValidTestH[i], isValidTestD[i], nBytes, hipMemcpyDeviceToHost)); } // Validate results @@ -194,58 +185,50 @@ static void test_cg_multi_grid_group_type_via_base_type(int blockSize) // Free host and device memory delete [] launchParamsList; for (int i = 0; i < nGpu; i++) { - ASSERT_EQUAL(hipSetDevice(i), hipSuccess); + HIPCHECK(hipSetDevice(i)); - ASSERT_EQUAL(hipFree(sizeTestD[i]), hipSuccess); - ASSERT_EQUAL(hipFree(gridRankTestD[i]), hipSuccess); - ASSERT_EQUAL(hipFree(thdRankTestD[i]), hipSuccess); - ASSERT_EQUAL(hipFree(isValidTestD[i]), hipSuccess); - ASSERT_EQUAL(hipFree(syncTestD[i]), hipSuccess); + HIPCHECK(hipFree(sizeTestD[i])); + HIPCHECK(hipFree(gridRankTestD[i])); + HIPCHECK(hipFree(thdRankTestD[i])); + HIPCHECK(hipFree(isValidTestD[i])); + HIPCHECK(hipFree(syncTestD[i])); if (i == 0) - ASSERT_EQUAL(hipHostFree(syncResultD), hipSuccess); + HIPCHECK(hipHostFree(syncResultD)); - ASSERT_EQUAL(hipHostFree(sizeTestH[i]), hipSuccess); - ASSERT_EQUAL(hipHostFree(gridRankTestH[i]), hipSuccess); - ASSERT_EQUAL(hipHostFree(thdRankTestH[i]), hipSuccess); - ASSERT_EQUAL(hipHostFree(isValidTestH[i]), hipSuccess); + HIPCHECK(hipHostFree(sizeTestH[i])); + HIPCHECK(hipHostFree(gridRankTestH[i])); + HIPCHECK(hipHostFree(thdRankTestH[i])); + HIPCHECK(hipHostFree(isValidTestH[i])); } } -int main() -{ +TEST_CASE("Unit_hipCGMultiGridGroupType_BaseType") { // Set `maxThreadsPerBlock` by taking minimum among all available devices - ASSERT_EQUAL(hipGetDeviceCount(&nGpu), hipSuccess); - if (nGpu > MaxGPUs) { - nGpu = MaxGPUs; - } + int nGpu = 0; + HIPCHECK(hipGetDeviceCount(&nGpu)); + nGpu = min(nGpu, MaxGPUs); + int maxThreadsPerBlock = INT_MAX; + hipDeviceProp_t deviceProperties; for (int i = 0; i < nGpu; i++) { - hipDeviceProp_t deviceProperties; - ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, i), hipSuccess); + HIPCHECK(hipGetDeviceProperties(&deviceProperties, i)); if (!deviceProperties.cooperativeMultiDeviceLaunch) { - printf("Device doesn't support cooperative launch!"); - passed(); + HipTest::HIP_SKIP_TEST("Device doesn't support cooperative launch!"); + return; } - int curDeviceMaxThreadsPerBlock = deviceProperties.maxThreadsPerBlock; - maxThreadsPerBlock = min(maxThreadsPerBlock, curDeviceMaxThreadsPerBlock); + maxThreadsPerBlock = min(maxThreadsPerBlock, deviceProperties.maxThreadsPerBlock); } - // Test block sizes which are powers of 2 - int i = 0; - while (true) { - int blockSize = pow(2, i); - if (blockSize > maxThreadsPerBlock) - break; - test_cg_multi_grid_group_type_via_base_type(blockSize); - ++i; + // Test for blockSizes in powers of 2 + for (int blockSize = 2; blockSize <= maxThreadsPerBlock; blockSize = blockSize*2) { + test_cg_multi_grid_group_type_via_base_type(blockSize, nGpu); } - // Test some random block sizes - for(int j = 0; j < 10 ; ++j) { - int blockSize = rand() % maxThreadsPerBlock; - test_cg_multi_grid_group_type_via_base_type(blockSize); + // Test for random blockSizes, but the sequence is the same every execution + srand(0); + for (int i = 0; i < 10; i++) { + // Test fails for 0 thread per block + test_cg_multi_grid_group_type_via_base_type(max(2, rand() % maxThreadsPerBlock), nGpu); } - - passed(); } diff --git a/projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupTypeViaPublicApi.cpp b/projects/hip/tests/catch/unit/cooperativeGrps/hipCGMultiGridGroupTypeViaPublicApi.cc similarity index 62% rename from projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupTypeViaPublicApi.cpp rename to projects/hip/tests/catch/unit/cooperativeGrps/hipCGMultiGridGroupTypeViaPublicApi.cc index a00a40345c..3e5b97fe5a 100644 --- a/projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupTypeViaPublicApi.cpp +++ b/projects/hip/tests/catch/unit/cooperativeGrps/hipCGMultiGridGroupTypeViaPublicApi.cc @@ -27,20 +27,18 @@ THE SOFTWARE. * HIT_END */ -#include "test_common.h" -#include "hip/hip_cooperative_groups.h" +#include +#include #include #include #include -#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs) -#define ASSERT_LE(lhs, rhs) assert(lhs <= rhs) -#define ASSERT_GE(lhs, rhs) assert(lhs >= rhs) - -constexpr int MaxGPUs = 8; -int nGpu = 0; +#define ASSERT_EQUAL(lhs, rhs) HIPASSERT(lhs == rhs) +#define ASSERT_LE(lhs, rhs) HIPASSERT(lhs <= rhs) +#define ASSERT_GE(lhs, rhs) HIPASSERT(lhs >= rhs) using namespace cooperative_groups; +constexpr int MaxGPUs = 8; static __global__ void kernel_cg_multi_grid_group_type_via_public_api(int *sizeTestD, @@ -71,7 +69,7 @@ void kernel_cg_multi_grid_group_type_via_public_api(int *sizeTestD, sync(this_grid()); // Thread 0 from work-group 0 of current grid (gpu) does grid level reduction if (blockIdx.x == 0 && threadIdx.x == 0) { - for (int i = 1; i < gridDim.x * blockDim.x; ++i) { + for (uint i = 1; i < gridDim.x * blockDim.x; ++i) { syncTestD[0] += syncTestD[i]; } syncResultD[this_multi_grid().grid_rank() + 1] = syncTestD[0]; @@ -81,20 +79,20 @@ void kernel_cg_multi_grid_group_type_via_public_api(int *sizeTestD, // grid (gpu) 0 does final reduction across all grids (gpus) 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) { + for (uint i = 1; i <= this_multi_grid().num_grids(); ++i) { syncResultD[0] += syncResultD[i]; } } } -static void test_cg_multi_grid_group_type_via_public_api(int blockSize) +static void test_cg_multi_grid_group_type_via_public_api(int blockSize, int nGpu) { // 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); + HIPCHECK(hipSetDevice(i)); + HIPCHECK(hipDeviceSynchronize()); // Make sure work is done on this device + HIPCHECK(hipStreamCreate(&stream[i])); } // Allocate host and device memory @@ -105,23 +103,21 @@ static void test_cg_multi_grid_group_type_via_public_api(int blockSize) int *isValidTestD[MaxGPUs], *isValidTestH[MaxGPUs]; int *syncTestD[MaxGPUs], *syncResultD; for (int i = 0; i < nGpu; i++) { - ASSERT_EQUAL(hipSetDevice(i), hipSuccess); + HIPCHECK(hipSetDevice(i)); - ASSERT_EQUAL(hipMalloc(&sizeTestD[i], nBytes), hipSuccess); - ASSERT_EQUAL(hipMalloc(&gridRankTestD[i], nBytes), hipSuccess); - ASSERT_EQUAL(hipMalloc(&thdRankTestD[i], nBytes), hipSuccess); - ASSERT_EQUAL(hipMalloc(&isValidTestD[i], nBytes), hipSuccess); - ASSERT_EQUAL(hipMalloc(&syncTestD[i], nBytes), hipSuccess); + HIPCHECK(hipMalloc(&sizeTestD[i], nBytes)); + HIPCHECK(hipMalloc(&gridRankTestD[i], nBytes)); + HIPCHECK(hipMalloc(&thdRankTestD[i], nBytes)); + HIPCHECK(hipMalloc(&isValidTestD[i], nBytes)); + HIPCHECK(hipMalloc(&syncTestD[i], nBytes)); - ASSERT_EQUAL(hipHostMalloc(&sizeTestH[i], nBytes), hipSuccess); - ASSERT_EQUAL(hipHostMalloc(&gridRankTestH[i], nBytes), hipSuccess); - ASSERT_EQUAL(hipHostMalloc(&thdRankTestH[i], nBytes), hipSuccess); - ASSERT_EQUAL(hipHostMalloc(&isValidTestH[i], nBytes), hipSuccess); + HIPCHECK(hipHostMalloc(&sizeTestH[i], nBytes)); + HIPCHECK(hipHostMalloc(&gridRankTestH[i], nBytes)); + HIPCHECK(hipHostMalloc(&thdRankTestH[i], nBytes)); + HIPCHECK(hipHostMalloc(&isValidTestH[i], nBytes)); if (i == 0) { - ASSERT_EQUAL( - hipHostMalloc(&syncResultD, sizeof(int) * (nGpu + 1), hipHostMallocCoherent), - hipSuccess); + HIPCHECK(hipHostMalloc(&syncResultD, sizeof(int) * (nGpu + 1), hipHostMallocCoherent)); } } @@ -130,7 +126,7 @@ static void test_cg_multi_grid_group_type_via_public_api(int blockSize) hipLaunchParams* launchParamsList = new hipLaunchParams[nGpu]; void* args[MaxGPUs * NumKernelArgs]; for (int i = 0; i < nGpu; i++) { - ASSERT_EQUAL(hipSetDevice(i), hipSuccess); + HIPCHECK(hipSetDevice(i)); args[i * NumKernelArgs ] = &sizeTestD[i]; args[i * NumKernelArgs + 1] = &gridRankTestD[i]; @@ -150,16 +146,12 @@ static void test_cg_multi_grid_group_type_via_public_api(int blockSize) // Copy result from device to host for (int i = 0; i < nGpu; i++) { - ASSERT_EQUAL(hipSetDevice(i), hipSuccess); + HIPCHECK(hipSetDevice(i)); - ASSERT_EQUAL(hipMemcpy(sizeTestH[i], sizeTestD[i], nBytes, hipMemcpyDeviceToHost), - hipSuccess); - ASSERT_EQUAL(hipMemcpy(gridRankTestH[i], gridRankTestD[i], nBytes, hipMemcpyDeviceToHost), - hipSuccess); - ASSERT_EQUAL(hipMemcpy(thdRankTestH[i], thdRankTestD[i], nBytes, hipMemcpyDeviceToHost), - hipSuccess); - ASSERT_EQUAL(hipMemcpy(isValidTestH[i], isValidTestD[i], nBytes, hipMemcpyDeviceToHost), - hipSuccess); + HIPCHECK(hipMemcpy(sizeTestH[i], sizeTestD[i], nBytes, hipMemcpyDeviceToHost)); + HIPCHECK(hipMemcpy(gridRankTestH[i], gridRankTestD[i], nBytes, hipMemcpyDeviceToHost)); + HIPCHECK(hipMemcpy(thdRankTestH[i], thdRankTestD[i], nBytes, hipMemcpyDeviceToHost)); + HIPCHECK(hipMemcpy(isValidTestH[i], isValidTestD[i], nBytes, hipMemcpyDeviceToHost)); } // Validate results @@ -189,58 +181,50 @@ static void test_cg_multi_grid_group_type_via_public_api(int blockSize) // Free host and device memory delete [] launchParamsList; for (int i = 0; i < nGpu; i++) { - ASSERT_EQUAL(hipSetDevice(i), hipSuccess); + HIPCHECK(hipSetDevice(i)); - ASSERT_EQUAL(hipFree(sizeTestD[i]), hipSuccess); - ASSERT_EQUAL(hipFree(gridRankTestD[i]), hipSuccess); - ASSERT_EQUAL(hipFree(thdRankTestD[i]), hipSuccess); - ASSERT_EQUAL(hipFree(isValidTestD[i]), hipSuccess); - ASSERT_EQUAL(hipFree(syncTestD[i]), hipSuccess); + HIPCHECK(hipFree(sizeTestD[i])); + HIPCHECK(hipFree(gridRankTestD[i])); + HIPCHECK(hipFree(thdRankTestD[i])); + HIPCHECK(hipFree(isValidTestD[i])); + HIPCHECK(hipFree(syncTestD[i])); if (i == 0) - ASSERT_EQUAL(hipHostFree(syncResultD), hipSuccess); + HIPCHECK(hipHostFree(syncResultD)); - ASSERT_EQUAL(hipHostFree(sizeTestH[i]), hipSuccess); - ASSERT_EQUAL(hipHostFree(gridRankTestH[i]), hipSuccess); - ASSERT_EQUAL(hipHostFree(thdRankTestH[i]), hipSuccess); - ASSERT_EQUAL(hipHostFree(isValidTestH[i]), hipSuccess); + HIPCHECK(hipHostFree(sizeTestH[i])); + HIPCHECK(hipHostFree(gridRankTestH[i])); + HIPCHECK(hipHostFree(thdRankTestH[i])); + HIPCHECK(hipHostFree(isValidTestH[i])); } } -int main() -{ +TEST_CASE("Unit_hipCGMultiGridGroupType_PublicApi") { // Set `maxThreadsPerBlock` by taking minimum among all available devices - ASSERT_EQUAL(hipGetDeviceCount(&nGpu), hipSuccess); - if (nGpu > MaxGPUs) { - nGpu = MaxGPUs; - } + int nGpu = 0; + HIPCHECK(hipGetDeviceCount(&nGpu)); + nGpu = min(nGpu, MaxGPUs); + int maxThreadsPerBlock = INT_MAX; + hipDeviceProp_t deviceProperties; for (int i = 0; i < nGpu; i++) { - hipDeviceProp_t deviceProperties; - ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, i), hipSuccess); + HIPCHECK(hipGetDeviceProperties(&deviceProperties, i)); if (!deviceProperties.cooperativeMultiDeviceLaunch) { - printf("Device doesn't support cooperative launch!"); - passed(); + HipTest::HIP_SKIP_TEST("Device doesn't support cooperative launch!"); + return; } - int curDeviceMaxThreadsPerBlock = deviceProperties.maxThreadsPerBlock; - maxThreadsPerBlock = min(maxThreadsPerBlock, curDeviceMaxThreadsPerBlock); + maxThreadsPerBlock = min(maxThreadsPerBlock, deviceProperties.maxThreadsPerBlock); } - // Test block sizes which are powers of 2 - int i = 0; - while (true) { - int blockSize = pow(2, i); - if (blockSize > maxThreadsPerBlock) - break; - test_cg_multi_grid_group_type_via_public_api(blockSize); - ++i; + // Test for blockSizes in powers of 2 + for (int blockSize = 2; blockSize <= maxThreadsPerBlock; blockSize = blockSize*2) { + test_cg_multi_grid_group_type_via_public_api(blockSize, nGpu); } - // Test some random block sizes - for(int j = 0; j < 10 ; ++j) { - int blockSize = rand() % maxThreadsPerBlock; - test_cg_multi_grid_group_type_via_public_api(blockSize); + // Test for random blockSizes, but the sequence is the same every execution + srand(0); + for (int i = 0; i < 10; i++) { + // Test fails for 0 thread per block + test_cg_multi_grid_group_type_via_public_api(max(2, rand() % maxThreadsPerBlock), nGpu); } - - passed(); } diff --git a/projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockType.cpp b/projects/hip/tests/catch/unit/cooperativeGrps/hipCGThreadBlockType.cc similarity index 56% rename from projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockType.cpp rename to projects/hip/tests/catch/unit/cooperativeGrps/hipCGThreadBlockType.cc index 0c7dea861f..bc4d671da0 100644 --- a/projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockType.cpp +++ b/projects/hip/tests/catch/unit/cooperativeGrps/hipCGThreadBlockType.cc @@ -27,12 +27,11 @@ THE SOFTWARE. * HIT_END */ -#include "test_common.h" -#include "hip/hip_cooperative_groups.h" -#include +#include +#include #include -#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs) +#define ASSERT_EQUAL(lhs, rhs) HIPASSERT(lhs == rhs) using namespace cooperative_groups; @@ -45,7 +44,6 @@ void kernel_cg_thread_block_type(int *sizeTestD, { thread_block tb = this_thread_block(); int gIdx = (blockIdx.x * blockDim.x) + threadIdx.x; - // Test size sizeTestD[gIdx] = tb.size(); @@ -79,18 +77,18 @@ static void test_cg_thread_block_type(int blockSize) dim3 *thdIndexTestD, *thdIndexTestH; // Allocate device memory - ASSERT_EQUAL(hipMalloc(&sizeTestD, nBytes), hipSuccess); - ASSERT_EQUAL(hipMalloc(&thdRankTestD, nBytes), hipSuccess); - ASSERT_EQUAL(hipMalloc(&syncTestD, nBytes), hipSuccess); - ASSERT_EQUAL(hipMalloc(&groupIndexTestD, nDim3Bytes), hipSuccess); - ASSERT_EQUAL(hipMalloc(&thdIndexTestD, nDim3Bytes), hipSuccess); + HIPCHECK(hipMalloc(&sizeTestD, nBytes)); + HIPCHECK(hipMalloc(&thdRankTestD, nBytes)); + HIPCHECK(hipMalloc(&syncTestD, nBytes)); + HIPCHECK(hipMalloc(&groupIndexTestD, nDim3Bytes)); + HIPCHECK(hipMalloc(&thdIndexTestD, nDim3Bytes)); // Allocate host memory - ASSERT_EQUAL(hipHostMalloc(&sizeTestH, nBytes), hipSuccess); - ASSERT_EQUAL(hipHostMalloc(&thdRankTestH, nBytes), hipSuccess); - ASSERT_EQUAL(hipHostMalloc(&syncTestH, nBytes), hipSuccess); - ASSERT_EQUAL(hipHostMalloc(&groupIndexTestH, nDim3Bytes), hipSuccess); - ASSERT_EQUAL(hipHostMalloc(&thdIndexTestH, nDim3Bytes), hipSuccess); + HIPCHECK(hipHostMalloc(&sizeTestH, nBytes)); + HIPCHECK(hipHostMalloc(&thdRankTestH, nBytes)); + HIPCHECK(hipHostMalloc(&syncTestH, nBytes)); + HIPCHECK(hipHostMalloc(&groupIndexTestH, nDim3Bytes)); + HIPCHECK(hipHostMalloc(&thdIndexTestH, nDim3Bytes)); // Launch Kernel hipLaunchKernelGGL(kernel_cg_thread_block_type, @@ -105,79 +103,62 @@ static void test_cg_thread_block_type(int blockSize) thdIndexTestD); // Copy result from device to host - ASSERT_EQUAL(hipMemcpy(sizeTestH, sizeTestD, nBytes, hipMemcpyDeviceToHost), - hipSuccess); - ASSERT_EQUAL(hipMemcpy(thdRankTestH, thdRankTestD, nBytes, hipMemcpyDeviceToHost), - hipSuccess); - ASSERT_EQUAL(hipMemcpy(syncTestH, syncTestD, nBytes, hipMemcpyDeviceToHost), - hipSuccess); - ASSERT_EQUAL(hipMemcpy(groupIndexTestH, groupIndexTestD, nDim3Bytes, hipMemcpyDeviceToHost), - hipSuccess); - ASSERT_EQUAL(hipMemcpy(thdIndexTestH, thdIndexTestD, nDim3Bytes, hipMemcpyDeviceToHost), - hipSuccess); + HIPCHECK(hipMemcpy(sizeTestH, sizeTestD, nBytes, hipMemcpyDeviceToHost)); + HIPCHECK(hipMemcpy(thdRankTestH, thdRankTestD, nBytes, hipMemcpyDeviceToHost)); + HIPCHECK(hipMemcpy(syncTestH, syncTestD, nBytes, hipMemcpyDeviceToHost)); + HIPCHECK(hipMemcpy(groupIndexTestH, groupIndexTestD, nDim3Bytes, hipMemcpyDeviceToHost)); + HIPCHECK(hipMemcpy(thdIndexTestH, thdIndexTestD, nDim3Bytes, hipMemcpyDeviceToHost)); // Validate results for both blocks together for (int i = 0; i < 2 * blockSize; ++i) { ASSERT_EQUAL(sizeTestH[i], blockSize); ASSERT_EQUAL(thdRankTestH[i], i % blockSize); ASSERT_EQUAL(syncTestH[i], 200); - ASSERT_EQUAL(groupIndexTestH[i].x, i / blockSize); + ASSERT_EQUAL(groupIndexTestH[i].x, (uint) i / blockSize); ASSERT_EQUAL(groupIndexTestH[i].y, 0); ASSERT_EQUAL(groupIndexTestH[i].z, 0); - ASSERT_EQUAL(thdIndexTestH[i].x, i % blockSize); + ASSERT_EQUAL(thdIndexTestH[i].x, (uint) i % blockSize); ASSERT_EQUAL(thdIndexTestH[i].y, 0); ASSERT_EQUAL(thdIndexTestH[i].z, 0); } // Free device memory - ASSERT_EQUAL(hipFree(sizeTestD), hipSuccess); - ASSERT_EQUAL(hipFree(thdRankTestD), hipSuccess); - ASSERT_EQUAL(hipFree(syncTestD), hipSuccess); - ASSERT_EQUAL(hipFree(groupIndexTestD), hipSuccess); - ASSERT_EQUAL(hipFree(thdIndexTestD), hipSuccess); + HIPCHECK(hipFree(sizeTestD)); + HIPCHECK(hipFree(thdRankTestD)); + HIPCHECK(hipFree(syncTestD)); + HIPCHECK(hipFree(groupIndexTestD)); + HIPCHECK(hipFree(thdIndexTestD)); //Free host memory - ASSERT_EQUAL(hipHostFree(sizeTestH), hipSuccess); - ASSERT_EQUAL(hipHostFree(thdRankTestH), hipSuccess); - ASSERT_EQUAL(hipHostFree(syncTestH), hipSuccess); - ASSERT_EQUAL(hipHostFree(groupIndexTestH), hipSuccess); - ASSERT_EQUAL(hipHostFree(thdIndexTestH), hipSuccess); + HIPCHECK(hipHostFree(sizeTestH)); + HIPCHECK(hipHostFree(thdRankTestH)); + HIPCHECK(hipHostFree(syncTestH)); + HIPCHECK(hipHostFree(groupIndexTestH)); + HIPCHECK(hipHostFree(thdIndexTestH)); } -int main() -{ +TEST_CASE("Unit_hipCGThreadBlockType") { // Use default device for validating the test int deviceId; - ASSERT_EQUAL(hipGetDevice(&deviceId), hipSuccess); hipDeviceProp_t deviceProperties; - ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, deviceId), hipSuccess); - int maxThreadsPerBlock = deviceProperties.maxThreadsPerBlock; + HIPCHECK(hipGetDevice(&deviceId)); + HIPCHECK(hipGetDeviceProperties(&deviceProperties, deviceId)); if (!deviceProperties.cooperativeLaunch) { - std::cout << "info: Device doesn't support cooperative launch! skipping the test!\n"; - if (hip_skip_tests_enabled()) { - return hip_skip_retcode(); - } else { - passed(); - } - return 0; + HipTest::HIP_SKIP_TEST("Device doesn't support cooperative launch!"); + return; } - // Test block sizes which are powers of 2 - int i = 1; - while (true) { - int blockSize = pow(2, i); - if (blockSize > maxThreadsPerBlock) - break; - test_cg_thread_block_type(blockSize); - ++i; - } - - // Test some random block sizes - for(int j = 0; j < 10 ; ++j) { - int blockSize = rand() % maxThreadsPerBlock; + // Test for blockSizes in powers of 2 + int maxThreadsPerBlock = deviceProperties.maxThreadsPerBlock; + for (int blockSize = 2; blockSize <= maxThreadsPerBlock; blockSize = blockSize*2) { test_cg_thread_block_type(blockSize); } - passed(); + // Test for random blockSizes, but the sequence is the same every execution + srand(0); + for (int i = 0; i < 10; i++) { + // Test fails for only 1 thread per block + test_cg_thread_block_type(max(2, rand() % maxThreadsPerBlock)); + } } diff --git a/projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockTypeViaBaseType.cpp b/projects/hip/tests/catch/unit/cooperativeGrps/hipCGThreadBlockTypeViaBaseType.cc similarity index 63% rename from projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockTypeViaBaseType.cpp rename to projects/hip/tests/catch/unit/cooperativeGrps/hipCGThreadBlockTypeViaBaseType.cc index 8b51c2c3fc..69f5e91ad2 100644 --- a/projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockTypeViaBaseType.cpp +++ b/projects/hip/tests/catch/unit/cooperativeGrps/hipCGThreadBlockTypeViaBaseType.cc @@ -27,9 +27,8 @@ THE SOFTWARE. * HIT_END */ -#include "test_common.h" +#include #include "hip/hip_cooperative_groups.h" -#include #include #define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs) @@ -68,14 +67,14 @@ static void test_cg_thread_block_type_via_base_type(int blockSize) int *syncTestD, *syncTestH; // Allocate device memory - ASSERT_EQUAL(hipMalloc(&sizeTestD, nBytes), hipSuccess); - ASSERT_EQUAL(hipMalloc(&thdRankTestD, nBytes), hipSuccess); - ASSERT_EQUAL(hipMalloc(&syncTestD, nBytes), hipSuccess); + HIPCHECK(hipMalloc(&sizeTestD, nBytes)); + HIPCHECK(hipMalloc(&thdRankTestD, nBytes)); + HIPCHECK(hipMalloc(&syncTestD, nBytes)); // Allocate host memory - ASSERT_EQUAL(hipHostMalloc(&sizeTestH, nBytes), hipSuccess); - ASSERT_EQUAL(hipHostMalloc(&thdRankTestH, nBytes), hipSuccess); - ASSERT_EQUAL(hipHostMalloc(&syncTestH, nBytes), hipSuccess); + HIPCHECK(hipHostMalloc(&sizeTestH, nBytes)); + HIPCHECK(hipHostMalloc(&thdRankTestH, nBytes)); + HIPCHECK(hipHostMalloc(&syncTestH, nBytes)); // Launch Kernel hipLaunchKernelGGL(kernel_cg_thread_block_type_via_base_type, @@ -88,12 +87,9 @@ static void test_cg_thread_block_type_via_base_type(int blockSize) syncTestD); // Copy result from device to host - ASSERT_EQUAL(hipMemcpy(sizeTestH, sizeTestD, nBytes, hipMemcpyDeviceToHost), - hipSuccess); - ASSERT_EQUAL(hipMemcpy(thdRankTestH, thdRankTestD, nBytes, hipMemcpyDeviceToHost), - hipSuccess); - ASSERT_EQUAL(hipMemcpy(syncTestH, syncTestD, nBytes, hipMemcpyDeviceToHost), - hipSuccess); + HIPCHECK(hipMemcpy(sizeTestH, sizeTestD, nBytes, hipMemcpyDeviceToHost)); + HIPCHECK(hipMemcpy(thdRankTestH, thdRankTestD, nBytes, hipMemcpyDeviceToHost)); + HIPCHECK(hipMemcpy(syncTestH, syncTestD, nBytes, hipMemcpyDeviceToHost)); // Validate results for both blocks together for (int i = 0; i < 2 * blockSize; ++i) { @@ -103,50 +99,38 @@ static void test_cg_thread_block_type_via_base_type(int blockSize) } // Free device memory - ASSERT_EQUAL(hipFree(sizeTestD), hipSuccess); - ASSERT_EQUAL(hipFree(thdRankTestD), hipSuccess); - ASSERT_EQUAL(hipFree(syncTestD), hipSuccess); + HIPCHECK(hipFree(sizeTestD)); + HIPCHECK(hipFree(thdRankTestD)); + HIPCHECK(hipFree(syncTestD)); //Free host memory - ASSERT_EQUAL(hipHostFree(sizeTestH), hipSuccess); - ASSERT_EQUAL(hipHostFree(thdRankTestH), hipSuccess); - ASSERT_EQUAL(hipHostFree(syncTestH), hipSuccess); + HIPCHECK(hipHostFree(sizeTestH)); + HIPCHECK(hipHostFree(thdRankTestH)); + HIPCHECK(hipHostFree(syncTestH)); } -int main() -{ +TEST_CASE("Unit_hipCGThreadBlockType_BaseType") { // Use default device for validating the test int deviceId; - ASSERT_EQUAL(hipGetDevice(&deviceId), hipSuccess); hipDeviceProp_t deviceProperties; - ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, deviceId), hipSuccess); - int maxThreadsPerBlock = deviceProperties.maxThreadsPerBlock; + HIPCHECK(hipGetDevice(&deviceId)); + HIPCHECK(hipGetDeviceProperties(&deviceProperties, deviceId)); if (!deviceProperties.cooperativeLaunch) { - std::cout << "info: Device doesn't support cooperative launch! skipping the test!\n"; - if (hip_skip_tests_enabled()) { - return hip_skip_retcode(); - } else { - passed(); - } - return 0; + HipTest::HIP_SKIP_TEST("Device doesn't support cooperative launch!"); + return; } - // Test block sizes which are powers of 2 - int i = 1; - while (true) { - int blockSize = pow(2, i); - if (blockSize > maxThreadsPerBlock) - break; - test_cg_thread_block_type_via_base_type(blockSize); - ++i; - } - - // Test some random block sizes - for(int j = 0; j < 10 ; ++j) { - int blockSize = rand() % maxThreadsPerBlock; + // Test for blockSizes in powers of 2 + int maxThreadsPerBlock = deviceProperties.maxThreadsPerBlock; + for (int blockSize = 2; blockSize <= maxThreadsPerBlock; blockSize = blockSize*2) { test_cg_thread_block_type_via_base_type(blockSize); } - passed(); + // Test for random blockSizes, but the sequence is the same every execution + srand(0); + for (int i = 0; i < 10; i++) { + // Test fails for only 1 thread per block + test_cg_thread_block_type_via_base_type(max(2, rand() % maxThreadsPerBlock)); + } } diff --git a/projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockTypeViaPublicApi.cpp b/projects/hip/tests/catch/unit/cooperativeGrps/hipCGThreadBlockTypeViaPublicApi.cc similarity index 63% rename from projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockTypeViaPublicApi.cpp rename to projects/hip/tests/catch/unit/cooperativeGrps/hipCGThreadBlockTypeViaPublicApi.cc index 3042526c49..f4913ad2c7 100644 --- a/projects/hip/tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockTypeViaPublicApi.cpp +++ b/projects/hip/tests/catch/unit/cooperativeGrps/hipCGThreadBlockTypeViaPublicApi.cc @@ -27,9 +27,8 @@ THE SOFTWARE. * HIT_END */ -#include "test_common.h" +#include #include "hip/hip_cooperative_groups.h" -#include #include #define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs) @@ -68,14 +67,14 @@ static void test_cg_thread_block_type_via_public_api(int blockSize) int *syncTestD, *syncTestH; // Allocate device memory - ASSERT_EQUAL(hipMalloc(&sizeTestD, nBytes), hipSuccess); - ASSERT_EQUAL(hipMalloc(&thdRankTestD, nBytes), hipSuccess); - ASSERT_EQUAL(hipMalloc(&syncTestD, nBytes), hipSuccess); + HIPCHECK(hipMalloc(&sizeTestD, nBytes)); + HIPCHECK(hipMalloc(&thdRankTestD, nBytes)); + HIPCHECK(hipMalloc(&syncTestD, nBytes)); // Allocate host memory - ASSERT_EQUAL(hipHostMalloc(&sizeTestH, nBytes), hipSuccess); - ASSERT_EQUAL(hipHostMalloc(&thdRankTestH, nBytes), hipSuccess); - ASSERT_EQUAL(hipHostMalloc(&syncTestH, nBytes), hipSuccess); + HIPCHECK(hipHostMalloc(&sizeTestH, nBytes)); + HIPCHECK(hipHostMalloc(&thdRankTestH, nBytes)); + HIPCHECK(hipHostMalloc(&syncTestH, nBytes)); // Launch Kernel hipLaunchKernelGGL(kernel_cg_thread_block_type_via_public_api, @@ -88,12 +87,9 @@ static void test_cg_thread_block_type_via_public_api(int blockSize) syncTestD); // Copy result from device to host - ASSERT_EQUAL(hipMemcpy(sizeTestH, sizeTestD, nBytes, hipMemcpyDeviceToHost), - hipSuccess); - ASSERT_EQUAL(hipMemcpy(thdRankTestH, thdRankTestD, nBytes, hipMemcpyDeviceToHost), - hipSuccess); - ASSERT_EQUAL(hipMemcpy(syncTestH, syncTestD, nBytes, hipMemcpyDeviceToHost), - hipSuccess); + HIPCHECK(hipMemcpy(sizeTestH, sizeTestD, nBytes, hipMemcpyDeviceToHost)); + HIPCHECK(hipMemcpy(thdRankTestH, thdRankTestD, nBytes, hipMemcpyDeviceToHost)); + HIPCHECK(hipMemcpy(syncTestH, syncTestD, nBytes, hipMemcpyDeviceToHost)); // Validate results for both blocks together for (int i = 0; i < 2 * blockSize; ++i) { @@ -103,50 +99,38 @@ static void test_cg_thread_block_type_via_public_api(int blockSize) } // Free device memory - ASSERT_EQUAL(hipFree(sizeTestD), hipSuccess); - ASSERT_EQUAL(hipFree(thdRankTestD), hipSuccess); - ASSERT_EQUAL(hipFree(syncTestD), hipSuccess); + HIPCHECK(hipFree(sizeTestD)); + HIPCHECK(hipFree(thdRankTestD)); + HIPCHECK(hipFree(syncTestD)); //Free host memory - ASSERT_EQUAL(hipHostFree(sizeTestH), hipSuccess); - ASSERT_EQUAL(hipHostFree(thdRankTestH), hipSuccess); - ASSERT_EQUAL(hipHostFree(syncTestH), hipSuccess); + HIPCHECK(hipHostFree(sizeTestH)); + HIPCHECK(hipHostFree(thdRankTestH)); + HIPCHECK(hipHostFree(syncTestH)); } -int main() -{ +TEST_CASE("Unit_hipCGThreadBlockType_PublicApi") { // Use default device for validating the test int deviceId; - ASSERT_EQUAL(hipGetDevice(&deviceId), hipSuccess); hipDeviceProp_t deviceProperties; - ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, deviceId), hipSuccess); - int maxThreadsPerBlock = deviceProperties.maxThreadsPerBlock; + HIPCHECK(hipGetDevice(&deviceId)); + HIPCHECK(hipGetDeviceProperties(&deviceProperties, deviceId)); if (!deviceProperties.cooperativeLaunch) { - std::cout << "info: Device doesn't support cooperative launch! skipping the test!\n"; - if (hip_skip_tests_enabled()) { - return hip_skip_retcode(); - } else { - passed(); - } - return 0; + HipTest::HIP_SKIP_TEST("Device doesn't support cooperative launch!"); + return; } - // Test block sizes which are powers of 2 - int i = 1; - while (true) { - int blockSize = pow(2, i); - if (blockSize > maxThreadsPerBlock) - break; - test_cg_thread_block_type_via_public_api(blockSize); - ++i; - } - - // Test some random block sizes - for(int j = 0; j < 10 ; ++j) { - int blockSize = rand() % maxThreadsPerBlock; + // Test for blockSizes in powers of 2 + int maxThreadsPerBlock = deviceProperties.maxThreadsPerBlock; + for (int blockSize = 2; blockSize <= maxThreadsPerBlock; blockSize = blockSize*2) { test_cg_thread_block_type_via_public_api(blockSize); } - passed(); + // Test for random blockSizes, but the sequence is the same every execution + srand(0); + for (int i = 0; i < 10; i++) { + // Test fails for only 1 thread per block + test_cg_thread_block_type_via_public_api(max(2, rand() % maxThreadsPerBlock)); + } }