SWDEV-269439 - Fix dtests build failure on HIP CUDA

Change-Id: Ibdba5d364df1bcd3b9e158e22a602dd993044d7a


[ROCm/clr commit: 45e5d6aa67]

projects/clr/hipamd/tests/src/kernel/hipDynamicShared.cpp

@@ -74,7 +74,7 @@ __global__ void testExternSharedKernel(const T* A_d, const T* B_d, T* C_d,
 }
 
 template <typename T>
-void testExternShared(size_t N, size_t groupElements) {
+void testExternShared(size_t N, unsigned groupElements) {
     size_t Nbytes = N * sizeof(T);
 
     T *A_d, *B_d, *C_d;

projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGGridGroupTypeViaBaseType.cpp

@@ -44,7 +44,7 @@ void kernel_cg_grid_group_type_via_base_type(int *sizeTestD,
                                              int *isValidTestD,
                                              int *syncTestD)
 {
-  thread_group tg = this_grid();
+  grid_group tg = this_grid();
   int gIdx = (blockIdx.x * blockDim.x) + threadIdx.x;
 
   // Test size

projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGGridGroupTypeViaPublicApi.cpp

@@ -54,7 +54,7 @@ void kernel_cg_grid_group_type_via_public_api(int *sizeTestD,
   thdRankTestD[gIdx] = thread_rank(gg);
 
   // Test is_valid api
-  isValidTestD[gIdx] = is_valid(gg);
+  isValidTestD[gIdx] = gg.is_valid();
 
   // Test sync api
   if (blockIdx.x == 0 && threadIdx.x == 0)

projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupType.cpp

@@ -22,8 +22,8 @@ THE SOFTWARE.
 
 
 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp
- * TEST: %t
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 -rdc=true -gencode arch=compute_60,code=sm_60
+ * TEST: %t EXCLUDE_HIP_PLATFORM nvidia
  * HIT_END
  */
 

projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupTypeViaBaseType.cpp

@@ -22,8 +22,8 @@ THE SOFTWARE.
 
 
 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp
- * TEST: %t
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 -rdc=true -gencode arch=compute_60,code=sm_60
+ * TEST: %t EXCLUDE_HIP_PLATFORM nvidia
  * HIT_END
  */
 
@@ -47,7 +47,7 @@ void kernel_cg_multi_grid_group_type_via_base_type(int *sizeTestD,
                                                    int *syncTestD,
                                                    int *syncResultD)
 {
-  thread_group tg = this_multi_grid();
+  multi_grid_group tg = this_multi_grid();
   int gIdx = (blockIdx.x * blockDim.x) + threadIdx.x;
 
   // Test size

projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupTypeViaPublicApi.cpp

@@ -22,8 +22,8 @@ THE SOFTWARE.
 
 
 /* HIT_START
- * BUILD: %t %s ../../test_common.cpp
- * TEST: %t
+ * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 -rdc=true -gencode arch=compute_60,code=sm_60
+ * TEST: %t EXCLUDE_HIP_PLATFORM nvidia
  * HIT_END
  */
 
@@ -58,7 +58,7 @@ void kernel_cg_multi_grid_group_type_via_public_api(int *sizeTestD,
   thdRankTestD[gIdx] = thread_rank(mg);
 
   // Test is_valid api
-  isValidTestD[gIdx] = is_valid(mg);
+  isValidTestD[gIdx] = mg.is_valid();
 
   // Test sync api
   //

projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockType.cpp

@@ -23,7 +23,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * TEST: %t
+ * TEST: %t EXCLUDE_HIP_PLATFORM nvidia
  * HIT_END
  */
 
@@ -39,7 +39,6 @@ using namespace cooperative_groups;
 static __global__
 void kernel_cg_thread_block_type(int *sizeTestD,
                                  int *thdRankTestD,
-                                 int *isValidTestD,
                                  int *syncTestD,
                                  dim3 *groupIndexTestD,
                                  dim3 *thdIndexTestD)
@@ -53,9 +52,6 @@ void kernel_cg_thread_block_type(int *sizeTestD,
   // Test thread_rank
   thdRankTestD[gIdx] = tb.thread_rank();
 
-  // Test is_valid
-  isValidTestD[gIdx] = tb.is_valid();
-
   // Test sync
   __shared__ int sm[2];
   if (threadIdx.x == 0)
@@ -78,7 +74,6 @@ static void test_cg_thread_block_type(int blockSize)
   int nDim3Bytes = sizeof(dim3) * 2 * blockSize;
   int *sizeTestD, *sizeTestH;
   int *thdRankTestD, *thdRankTestH;
-  int *isValidTestD, *isValidTestH;
   int *syncTestD, *syncTestH;
   dim3 *groupIndexTestD, *groupIndexTestH;
   dim3 *thdIndexTestD, *thdIndexTestH;
@@ -86,7 +81,6 @@ static void test_cg_thread_block_type(int blockSize)
   // Allocate device memory
   ASSERT_EQUAL(hipMalloc(&sizeTestD, nBytes), hipSuccess);
   ASSERT_EQUAL(hipMalloc(&thdRankTestD, nBytes), hipSuccess);
-  ASSERT_EQUAL(hipMalloc(&isValidTestD, nBytes), hipSuccess);
   ASSERT_EQUAL(hipMalloc(&syncTestD, nBytes), hipSuccess);
   ASSERT_EQUAL(hipMalloc(&groupIndexTestD, nDim3Bytes), hipSuccess);
   ASSERT_EQUAL(hipMalloc(&thdIndexTestD, nDim3Bytes), hipSuccess);
@@ -94,7 +88,6 @@ static void test_cg_thread_block_type(int blockSize)
   // Allocate host memory
   ASSERT_EQUAL(hipHostMalloc(&sizeTestH, nBytes), hipSuccess);
   ASSERT_EQUAL(hipHostMalloc(&thdRankTestH, nBytes), hipSuccess);
-  ASSERT_EQUAL(hipHostMalloc(&isValidTestH, nBytes), hipSuccess);
   ASSERT_EQUAL(hipHostMalloc(&syncTestH, nBytes), hipSuccess);
   ASSERT_EQUAL(hipHostMalloc(&groupIndexTestH, nDim3Bytes), hipSuccess);
   ASSERT_EQUAL(hipHostMalloc(&thdIndexTestH, nDim3Bytes), hipSuccess);
@@ -107,7 +100,6 @@ static void test_cg_thread_block_type(int blockSize)
                      0,
                      sizeTestD,
                      thdRankTestD,
-                     isValidTestD,
                      syncTestD,
                      groupIndexTestD,
                      thdIndexTestD);
@@ -117,8 +109,6 @@ static void test_cg_thread_block_type(int blockSize)
                hipSuccess);
   ASSERT_EQUAL(hipMemcpy(thdRankTestH, thdRankTestD, nBytes, hipMemcpyDeviceToHost),
                hipSuccess);
-  ASSERT_EQUAL(hipMemcpy(isValidTestH, isValidTestD, nBytes, hipMemcpyDeviceToHost),
-               hipSuccess);
   ASSERT_EQUAL(hipMemcpy(syncTestH, syncTestD, nBytes, hipMemcpyDeviceToHost),
                hipSuccess);
   ASSERT_EQUAL(hipMemcpy(groupIndexTestH, groupIndexTestD, nDim3Bytes, hipMemcpyDeviceToHost),
@@ -130,7 +120,6 @@ static void test_cg_thread_block_type(int blockSize)
   for (int i = 0; i < 2 * blockSize; ++i) {
     ASSERT_EQUAL(sizeTestH[i], blockSize);
     ASSERT_EQUAL(thdRankTestH[i], i % blockSize);
-    ASSERT_EQUAL(isValidTestH[i], 1);
     ASSERT_EQUAL(syncTestH[i], 200);
     ASSERT_EQUAL(groupIndexTestH[i].x, i / blockSize);
     ASSERT_EQUAL(groupIndexTestH[i].y, 0);
@@ -143,7 +132,6 @@ static void test_cg_thread_block_type(int blockSize)
   // Free device memory
   ASSERT_EQUAL(hipFree(sizeTestD), hipSuccess);
   ASSERT_EQUAL(hipFree(thdRankTestD), hipSuccess);
-  ASSERT_EQUAL(hipFree(isValidTestD), hipSuccess);
   ASSERT_EQUAL(hipFree(syncTestD), hipSuccess);
   ASSERT_EQUAL(hipFree(groupIndexTestD), hipSuccess);
   ASSERT_EQUAL(hipFree(thdIndexTestD), hipSuccess);
@@ -151,7 +139,6 @@ static void test_cg_thread_block_type(int blockSize)
   //Free host memory
   ASSERT_EQUAL(hipHostFree(sizeTestH), hipSuccess);
   ASSERT_EQUAL(hipHostFree(thdRankTestH), hipSuccess);
-  ASSERT_EQUAL(hipHostFree(isValidTestH), hipSuccess);
   ASSERT_EQUAL(hipHostFree(syncTestH), hipSuccess);
   ASSERT_EQUAL(hipHostFree(groupIndexTestH), hipSuccess);
   ASSERT_EQUAL(hipHostFree(thdIndexTestH), hipSuccess);

projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockTypeViaBaseType.cpp

@@ -23,7 +23,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * TEST: %t
+ * TEST: %t EXCLUDE_HIP_PLATFORM nvidia
  * HIT_END
  */
 
@@ -39,7 +39,6 @@ using namespace cooperative_groups;
 static __global__
 void kernel_cg_thread_block_type_via_base_type(int *sizeTestD,
                                                int *thdRankTestD,
-                                               int *isValidTestD,
                                                int *syncTestD)
 {
   thread_group tg = this_thread_block();
@@ -51,9 +50,6 @@ void kernel_cg_thread_block_type_via_base_type(int *sizeTestD,
   // Test thread_rank
   thdRankTestD[gIdx] = tg.thread_rank();
 
-  // Test is_valid
-  isValidTestD[gIdx] = tg.is_valid();
-
   // Test sync
   __shared__ int sm[2];
   if (threadIdx.x == 0)
@@ -69,19 +65,16 @@ static void test_cg_thread_block_type_via_base_type(int blockSize)
   int nBytes = sizeof(int) * 2 * blockSize;
   int *sizeTestD, *sizeTestH;
   int *thdRankTestD, *thdRankTestH;
-  int *isValidTestD, *isValidTestH;
   int *syncTestD, *syncTestH;
 
   // Allocate device memory
   ASSERT_EQUAL(hipMalloc(&sizeTestD, nBytes), hipSuccess);
   ASSERT_EQUAL(hipMalloc(&thdRankTestD, nBytes), hipSuccess);
-  ASSERT_EQUAL(hipMalloc(&isValidTestD, nBytes), hipSuccess);
   ASSERT_EQUAL(hipMalloc(&syncTestD, nBytes), hipSuccess);
 
   // Allocate host memory
   ASSERT_EQUAL(hipHostMalloc(&sizeTestH, nBytes), hipSuccess);
   ASSERT_EQUAL(hipHostMalloc(&thdRankTestH, nBytes), hipSuccess);
-  ASSERT_EQUAL(hipHostMalloc(&isValidTestH, nBytes), hipSuccess);
   ASSERT_EQUAL(hipHostMalloc(&syncTestH, nBytes), hipSuccess);
 
   // Launch Kernel
@@ -92,7 +85,6 @@ static void test_cg_thread_block_type_via_base_type(int blockSize)
                      0,
                      sizeTestD,
                      thdRankTestD,
-                     isValidTestD,
                      syncTestD);
 
   // Copy result from device to host
@@ -100,8 +92,6 @@ static void test_cg_thread_block_type_via_base_type(int blockSize)
                hipSuccess);
   ASSERT_EQUAL(hipMemcpy(thdRankTestH, thdRankTestD, nBytes, hipMemcpyDeviceToHost),
                hipSuccess);
-  ASSERT_EQUAL(hipMemcpy(isValidTestH, isValidTestD, nBytes, hipMemcpyDeviceToHost),
-               hipSuccess);
   ASSERT_EQUAL(hipMemcpy(syncTestH, syncTestD, nBytes, hipMemcpyDeviceToHost),
                hipSuccess);
 
@@ -109,20 +99,17 @@ static void test_cg_thread_block_type_via_base_type(int blockSize)
   for (int i = 0; i < 2 * blockSize; ++i) {
     ASSERT_EQUAL(sizeTestH[i], blockSize);
     ASSERT_EQUAL(thdRankTestH[i], i % blockSize);
-    ASSERT_EQUAL(isValidTestH[i], 1);
     ASSERT_EQUAL(syncTestH[i], 200);
   }
 
   // Free device memory
   ASSERT_EQUAL(hipFree(sizeTestD), hipSuccess);
   ASSERT_EQUAL(hipFree(thdRankTestD), hipSuccess);
-  ASSERT_EQUAL(hipFree(isValidTestD), hipSuccess);
   ASSERT_EQUAL(hipFree(syncTestD), hipSuccess);
 
   //Free host memory
   ASSERT_EQUAL(hipHostFree(sizeTestH), hipSuccess);
   ASSERT_EQUAL(hipHostFree(thdRankTestH), hipSuccess);
-  ASSERT_EQUAL(hipHostFree(isValidTestH), hipSuccess);
   ASSERT_EQUAL(hipHostFree(syncTestH), hipSuccess);
 }
 

projects/clr/hipamd/tests/src/runtimeApi/cooperativeGrps/hipCGThreadBlockTypeViaPublicApi.cpp

@@ -23,7 +23,7 @@ THE SOFTWARE.
 
 /* HIT_START
  * BUILD: %t %s ../../test_common.cpp
- * TEST: %t
+ * TEST: %t EXCLUDE_HIP_PLATFORM nvidia
  * HIT_END
  */
 
@@ -39,7 +39,6 @@ using namespace cooperative_groups;
 static __global__
 void kernel_cg_thread_block_type_via_public_api(int *sizeTestD,
                                                 int *thdRankTestD,
-                                                int *isValidTestD,
                                                 int *syncTestD)
 {
   thread_block tb = this_thread_block();
@@ -51,9 +50,6 @@ void kernel_cg_thread_block_type_via_public_api(int *sizeTestD,
   // Test thread_rank api
   thdRankTestD[gIdx] = thread_rank(tb);
 
-  // Test is_valid api
-  isValidTestD[gIdx] = is_valid(tb);
-
   // Test sync api
   __shared__ int sm[2];
   if (threadIdx.x == 0)
@@ -69,19 +65,16 @@ static void test_cg_thread_block_type_via_public_api(int blockSize)
   int nBytes = sizeof(int) * 2 * blockSize;
   int *sizeTestD, *sizeTestH;
   int *thdRankTestD, *thdRankTestH;
-  int *isValidTestD, *isValidTestH;
   int *syncTestD, *syncTestH;
 
   // Allocate device memory
   ASSERT_EQUAL(hipMalloc(&sizeTestD, nBytes), hipSuccess);
   ASSERT_EQUAL(hipMalloc(&thdRankTestD, nBytes), hipSuccess);
-  ASSERT_EQUAL(hipMalloc(&isValidTestD, nBytes), hipSuccess);
   ASSERT_EQUAL(hipMalloc(&syncTestD, nBytes), hipSuccess);
 
   // Allocate host memory
   ASSERT_EQUAL(hipHostMalloc(&sizeTestH, nBytes), hipSuccess);
   ASSERT_EQUAL(hipHostMalloc(&thdRankTestH, nBytes), hipSuccess);
-  ASSERT_EQUAL(hipHostMalloc(&isValidTestH, nBytes), hipSuccess);
   ASSERT_EQUAL(hipHostMalloc(&syncTestH, nBytes), hipSuccess);
 
   // Launch Kernel
@@ -92,7 +85,6 @@ static void test_cg_thread_block_type_via_public_api(int blockSize)
                      0,
                      sizeTestD,
                      thdRankTestD,
-                     isValidTestD,
                      syncTestD);
 
   // Copy result from device to host
@@ -100,8 +92,6 @@ static void test_cg_thread_block_type_via_public_api(int blockSize)
                hipSuccess);
   ASSERT_EQUAL(hipMemcpy(thdRankTestH, thdRankTestD, nBytes, hipMemcpyDeviceToHost),
                hipSuccess);
-  ASSERT_EQUAL(hipMemcpy(isValidTestH, isValidTestD, nBytes, hipMemcpyDeviceToHost),
-               hipSuccess);
   ASSERT_EQUAL(hipMemcpy(syncTestH, syncTestD, nBytes, hipMemcpyDeviceToHost),
                hipSuccess);
 
@@ -109,20 +99,17 @@ static void test_cg_thread_block_type_via_public_api(int blockSize)
   for (int i = 0; i < 2 * blockSize; ++i) {
     ASSERT_EQUAL(sizeTestH[i], blockSize);
     ASSERT_EQUAL(thdRankTestH[i], i % blockSize);
-    ASSERT_EQUAL(isValidTestH[i], 1);
     ASSERT_EQUAL(syncTestH[i], 200);
   }
 
   // Free device memory
   ASSERT_EQUAL(hipFree(sizeTestD), hipSuccess);
   ASSERT_EQUAL(hipFree(thdRankTestD), hipSuccess);
-  ASSERT_EQUAL(hipFree(isValidTestD), hipSuccess);
   ASSERT_EQUAL(hipFree(syncTestD), hipSuccess);
 
   //Free host memory
   ASSERT_EQUAL(hipHostFree(sizeTestH), hipSuccess);
   ASSERT_EQUAL(hipHostFree(thdRankTestH), hipSuccess);
-  ASSERT_EQUAL(hipHostFree(isValidTestH), hipSuccess);
   ASSERT_EQUAL(hipHostFree(syncTestH), hipSuccess);
 }