SWDEV-299127 - Merge 'develop' into 'amd-staging'

Change-Id: If2461c35ec8281d56859996ea07ec83665420f09


[ROCm/hip commit: 6bc3306f95]

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()

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()

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 <cmath>
-#include <cstdlib>
-#include <climits>
+#include <hip_test_common.hh>
+#include <hip/hip_cooperative_groups.h>
 
-#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();
 }

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 <hip_test_common.hh>
+#include <hip/hip_cooperative_groups.h>
 #include <cmath>
 #include <cstdlib>
 #include <climits>
 
-#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();
 }

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 <hip_test_common.hh>
+#include <hip/hip_cooperative_groups.h>
 #include <cmath>
 #include <cstdlib>
 #include <climits>
 
-#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();
 }

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 <cmath>
+#include <hip_test_common.hh>
+#include <hip/hip_cooperative_groups.h>
 #include <cstdlib>
 
-#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));
+  }
 }

projects/hip/tests/catch/unit/cooperativeGrps/hipCGThreadBlockTypeViaBaseType.cc

@@ -27,9 +27,8 @@ THE SOFTWARE.
  * HIT_END
  */
 
-#include "test_common.h"
+#include <hip_test_common.hh>
 #include "hip/hip_cooperative_groups.h"
-#include <cmath>
 #include <cstdlib>
 
 #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));
+  }
 }

projects/hip/tests/catch/unit/cooperativeGrps/hipCGThreadBlockTypeViaPublicApi.cc

@@ -27,9 +27,8 @@ THE SOFTWARE.
  * HIT_END
  */
 
-#include "test_common.h"
+#include <hip_test_common.hh>
 #include "hip/hip_cooperative_groups.h"
-#include <cmath>
 #include <cstdlib>
 
 #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));
+  }
 }