SWDEV-396656 - Adds CG tiled_partition and coalesced_groups tests

- Tests meta_group_rank/size() in static tile
               partitioning

Change-Id: I1a0d83224145a2356aa12b5e83b094834b53d5a4

catch/unit/cooperativeGrps/CMakeLists.txt

@@ -9,7 +9,8 @@ set(TEST_SRC
   grid_group.cc
   coalesced_groups_shfl_down.cc
   coalesced_groups_shfl_up.cc
-  simple_coalesced_groups.cc
+  hipCGTiledPartition.cc
+  hipCGCoalescedGroups.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")

catch/unit/cooperativeGrps/hipCGCoalescedGroups.cc

@@ -1,5 +1,5 @@
 /*
-Copyright (c) 2020 - 2021 Advanced Micro Devices, Inc. All rights reserved.
+Copyright (c) 2020 - 2023 Advanced Micro Devices, Inc. All rights reserved.
 
 Permission is hereby granted, free of charge, to any person obtaining a copy
 of this software and associated documentation files (the "Software"), to deal
@@ -20,18 +20,6 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */
 
-// Test Description:
-/* This test implements sum reduction kernel, first with each threads own rank
-   as input and comparing the sum with expected sum output derieved from n(n-1)/2
-   formula. The second part, partitions this parent group into child subgroups
-   a.k.a tiles using using tiled_partition() collective operation. This can be called
-   with a static tile size, passed in templated non-type variable-tiled_partition<tileSz>,
-   or in runtime as tiled_partition(thread_group parent, tileSz). This test covers both these
-   cases.
-   This test tests functionality of cg group partitioning, (static and dynamic) and its respective
-   API's size(), thread_rank(), and sync().
-*/
-
 #include <hip_test_common.hh>
 #include <hip/hip_cooperative_groups.h>
 #include <stdio.h>
@@ -173,10 +161,10 @@ __global__ void kernel_cg_coalesced_group_partition(unsigned int tileSz, int* re
 
     if (tiledPartition.thread_rank() == 0) {
       printf(
-          "   Sum of all ranks 0..%d in this tiledPartition group is %d. Corresponding parent thread "
-          "rank: %d\n",
-          tiledPartition.size() - 1, outputSum, input);
-
+          "   Sum of all ranks 0..%d in this tiledPartition group is %d. Corresponding parent thread rank"
+          " obtained from meta_group_rank : %d and number of tiles created : %d\n",
+          tiledPartition.size() - 1, outputSum, tiledPartition.meta_group_rank(),
+          tiledPartition.meta_group_size());
         result[input / (tileSz)] = outputSum;
     }
     return;
@@ -226,7 +214,7 @@ void printResultsSimpleCoalescedGroups(int* ptr, int size) {
 void compareResultsSimpleCoalescedGroups(int* cpu, int* gpu, int size) {
   for (unsigned int i = 0; i < size / sizeof(int); i++) {
     if (cpu[i] != gpu[i]) {
-      INFO(" results do not match.");
+      REQUIRE(" results do not match.");
     }
   }
 }
@@ -238,7 +226,7 @@ static void test_active_threads_grouping() {
 
   // Launch Kernel
     hipLaunchKernelGGL(kernel_coalesced_active_groups, blockSize, threadsPerBlock, 0, 0);
-    HIP_CHECK(hipGetLastError()); 
+    HIP_CHECK(hipGetLastError());
 
     err = hipDeviceSynchronize();
     if (err != hipSuccess) {
@@ -257,7 +245,7 @@ void verifyResultsSimpleCoalescedGroups(int* hPtr, int* dPtr, int size) {
       }
     }
     if (j == size) {
-      INFO(" Result verification failed!");
+      REQUIRE(" Result verification failed!");
     }
   }
 }
@@ -303,7 +291,7 @@ static void test_group_partition(unsigned int tileSz, bool useGlobalMem) {
     if (useGlobalMem) {
       hipLaunchKernelGGL(kernel_cg_coalesced_group_partition, blockSize, threadsPerBlock, 0, 0, tileSz,
                          dResult, useGlobalMem, globalMem, i);
-      HIP_CHECK(hipGetLastError()); 
+      HIP_CHECK(hipGetLastError());
 
       err = hipDeviceSynchronize();
       if (err != hipSuccess) {
@@ -312,7 +300,7 @@ static void test_group_partition(unsigned int tileSz, bool useGlobalMem) {
     } else {
       hipLaunchKernelGGL(kernel_cg_coalesced_group_partition, blockSize, threadsPerBlock,
                          threadsPerBlock * sizeof(int), 0, tileSz, dResult, useGlobalMem, globalMem, i);
-      HIP_CHECK(hipGetLastError()); 
+      HIP_CHECK(hipGetLastError());
 
       err = hipDeviceSynchronize();
       if (err != hipSuccess) {
@@ -387,7 +375,7 @@ static void test_shfl_any_to_any() {
     // Launch Kernel
     hipLaunchKernelGGL(kernel_shfl_any_to_any, blockSize, threadsPerBlock,
                        threadsPerBlock * sizeof(int), 0 , dPtr, dsrcArr, dResults, i);
-    HIP_CHECK(hipGetLastError()); 
+    HIP_CHECK(hipGetLastError());
     HIPCHECK(hipMemcpy(hPtr, dResults, group_size_in_bytes, hipMemcpyDeviceToHost));
     err = hipDeviceSynchronize();
     if (err != hipSuccess) {
@@ -457,7 +445,7 @@ static void test_shfl_broadcast() {
     // Launch Kernel
     hipLaunchKernelGGL(kernel_shfl, blockSize, threadsPerBlock,
                        threadsPerBlock * sizeof(int), 0, dPtr, dResults, srcLane, i);
-    HIP_CHECK(hipGetLastError()); 
+    HIP_CHECK(hipGetLastError());
     HIPCHECK(hipMemcpy(hPtr, dResults, group_size_in_bytes, hipMemcpyDeviceToHost));
     err = hipDeviceSynchronize();
     if (err != hipSuccess) {

catch/unit/cooperativeGrps/hipCGTiledPartition.cc

@@ -0,0 +1,392 @@
+/*
+Copyright (c) 2020 - 2023 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+// Test Description:
+/* This test implements sum reduction kernel, first with each threads own rank
+   as input and comparing the sum with expected sum output derieved from n(n-1)/2
+   formula. The second part, partitions this parent group into child subgroups
+   a.k.a tiles using using tiled_partition() collective operation. This can be called
+   with a static tile size, passed in templated non-type variable-tiled_partition<tileSz>,
+   or in runtime as tiled_partition(thread_group parent, tileSz). This test covers both these
+   cases.
+   This test tests functionality of cg group partitioning, (static and dynamic) and its respective
+   API's size(), thread_rank(), and sync().
+*/
+
+#include <hip_test_common.hh>
+#include <hip/hip_cooperative_groups.h>
+#include <stdio.h>
+#include <vector>
+
+using namespace cooperative_groups;
+
+#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs)
+
+/* Parallel reduce kernel.
+ *
+ * Step complexity: O(log n)
+ * Work complexity: O(n)
+ *
+ * Note: This kernel works only with power of 2 input arrays.
+ */
+__device__ int reduction_kernel(thread_group g, int* x, int val) {
+  int lane = g.thread_rank();
+  int sz = g.size();
+
+  for (int i = g.size() / 2; i > 0; i /= 2) {
+    // use lds to store the temporary result
+    x[lane] = val;
+    // Ensure all the stores are completed.
+    g.sync();
+
+    if (lane < i) {
+      val += x[lane + i];
+    }
+    // It must work on one tiled thread group at a time,
+    // and it must make sure all memory operations are
+    // completed before moving to the next stride.
+    // sync() here just does that.
+    g.sync();
+  }
+
+  // Choose the 0'th indexed thread that holds the reduction value to return
+  if (g.thread_rank() == 0) {
+    return val;
+  }
+  // Rest of the threads return no useful values
+  else {
+    return -1;
+  }
+}
+
+template <unsigned int tileSz>
+__global__ void kernel_cg_group_partition_static(int* result, bool isGlobalMem, int* globalMem) {
+  thread_block threadBlockCGTy = this_thread_block();
+  int threadBlockGroupSize = threadBlockCGTy.size();
+
+  int* workspace = NULL;
+
+  if (isGlobalMem) {
+    workspace = globalMem;
+  } else {
+    // Declare a shared memory
+    extern __shared__ int sharedMem[];
+    workspace = sharedMem;
+  }
+
+  int input, outputSum, expectedOutput;
+
+  // we pass its own thread rank as inputs
+  input = threadBlockCGTy.thread_rank();
+
+  expectedOutput = (threadBlockGroupSize - 1) * threadBlockGroupSize / 2;
+
+  outputSum = reduction_kernel(threadBlockCGTy, workspace, input);
+
+  // Choose a leader thread to print the results
+  if (threadBlockCGTy.thread_rank() == 0) {
+    printf(" Sum of all ranks 0..%d in threadBlockCooperativeGroup is %d (expected %d)\n\n",
+           (int)threadBlockCGTy.size() - 1, outputSum, expectedOutput);
+    printf(" Creating %d groups, of tile size %d threads:\n\n",
+           (int)threadBlockCGTy.size() / tileSz, tileSz);
+  }
+
+  threadBlockCGTy.sync();
+
+  thread_block_tile<tileSz> tiledPartition = tiled_partition<tileSz>(threadBlockCGTy);
+
+  // This offset allows each group to have its own unique area in the workspace array
+  int workspaceOffset = threadBlockCGTy.thread_rank() - tiledPartition.thread_rank();
+
+  outputSum = reduction_kernel(tiledPartition, workspace + workspaceOffset, input);
+
+  if (tiledPartition.thread_rank() == 0) {
+    printf(
+        "   Sum of all ranks 0..%d in this tiledPartition group is %d. Corresponding parent thread "
+        "rank via meta_group_rank : %d and the total number of groups created when partitioned : %d\n",
+        tiledPartition.size() - 1, outputSum, tiledPartition.meta_group_rank(), tiledPartition.meta_group_size());
+    result[input / (tileSz)] = outputSum;
+  }
+  return;
+}
+
+
+__global__ void kernel_cg_group_partition_dynamic(unsigned int tileSz, int* result,
+                                                  bool isGlobalMem, int* globalMem) {
+  thread_block threadBlockCGTy = this_thread_block();
+  int threadBlockGroupSize = threadBlockCGTy.size();
+
+  int* workspace = NULL;
+
+  if (isGlobalMem) {
+    workspace = globalMem;
+  } else {
+    // Declare a shared memory
+    extern __shared__ int sharedMem[];
+    workspace = sharedMem;
+  }
+
+  int input, outputSum, expectedOutput;
+
+  // input to reduction, for each thread, is its' rank in the group
+  input = threadBlockCGTy.thread_rank();
+
+  expectedOutput = (threadBlockGroupSize - 1) * threadBlockGroupSize / 2;
+
+  outputSum = reduction_kernel(threadBlockCGTy, workspace, input);
+
+  if (threadBlockCGTy.thread_rank() == 0) {
+    printf(" Sum of all ranks 0..%d in threadBlockCooperativeGroup is %d\n\n",
+           (int)threadBlockCGTy.size() - 1, outputSum);
+    printf(" Creating %d groups, of tile size %d threads:\n\n",
+           (int)threadBlockCGTy.size() / tileSz, tileSz);
+  }
+
+  threadBlockCGTy.sync();
+
+  thread_group tiledPartition = tiled_partition(threadBlockCGTy, tileSz);
+
+  // This offset allows each group to have its own unique area in the workspace array
+  int workspaceOffset = threadBlockCGTy.thread_rank() - tiledPartition.thread_rank();
+
+  outputSum = reduction_kernel(tiledPartition, workspace + workspaceOffset, input);
+
+  if (tiledPartition.thread_rank() == 0) {
+     printf(
+        "   Sum of all ranks 0..%d in this tiledPartition group is %d. Corresponding parent thread "
+        " %d\n", tiledPartition.size() - 1, outputSum, input);
+    result[input / (tileSz)] = outputSum;
+  }
+  return;
+}
+
+// Search if the sum exists in the expected results array
+void verifyResults(int* hPtr, int* dPtr, int size) {
+  int i = 0, j = 0;
+  for (i = 0; i < size; i++) {
+    for (j = 0; j < size; j++) {
+      if (hPtr[i] == dPtr[j]) {
+        break;
+      }
+    }
+    if (j == size) {
+      REQUIRE(" Result verification failed!");
+    }
+  }
+}
+
+
+template <unsigned int tileSz> static void test_group_partition(bool useGlobalMem) {
+  hipError_t err;
+  int blockSize = 1;
+  int threadsPerBlock = 64;
+
+  int numTiles = (blockSize * threadsPerBlock) / tileSz;
+
+  // Build an array of expected reduction sum output on the host
+  // based on the sum of their respective thread ranks for verification.
+  // eg: parent group has 64threads.
+  // child thread ranks: 0-15, 16-31, 32-47, 48-63
+  // expected sum:       120,   376,  632,  888
+  int* expectedSum = new int[numTiles];
+  int temp = 0, sum = 0;
+
+  for (int i = 1; i <= numTiles; i++) {
+    sum = temp;
+    temp = (((tileSz * i) - 1) * (tileSz * i)) / 2;
+    expectedSum[i-1] = temp - sum;
+  }
+
+  int* dResult = NULL;
+  hipMalloc((void**)&dResult, numTiles * sizeof(int));
+
+  int* globalMem = NULL;
+  if (useGlobalMem) {
+    hipMalloc((void**)&globalMem, threadsPerBlock * sizeof(int));
+  }
+
+  int* hResult = NULL;
+  hipHostMalloc(&hResult, numTiles * sizeof(int), hipHostMallocDefault);
+  memset(hResult, 0, numTiles * sizeof(int));
+
+  if (useGlobalMem) {
+    // Launch Kernel
+    hipLaunchKernelGGL(kernel_cg_group_partition_static<tileSz>, blockSize, threadsPerBlock, 0, 0,
+                       dResult, useGlobalMem, globalMem);
+    err = hipDeviceSynchronize();
+    if (err != hipSuccess) {
+      fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));
+    }
+  } else {
+    // Launch Kernel
+    hipLaunchKernelGGL(kernel_cg_group_partition_static<tileSz>, blockSize, threadsPerBlock,
+                       threadsPerBlock * sizeof(int), 0, dResult, useGlobalMem, globalMem);
+    err = hipDeviceSynchronize();
+    if (err != hipSuccess) {
+      fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));
+    }
+  }
+
+  hipMemcpy(hResult, dResult, numTiles * sizeof(int), hipMemcpyDeviceToHost);
+
+  verifyResults(expectedSum, hResult, numTiles);
+
+  // Free all allocated memory on host and device
+  hipFree(dResult);
+  hipFree(hResult);
+  if (useGlobalMem) {
+    hipFree(globalMem);
+  }
+  delete[] expectedSum;
+
+  printf("\n...PASSED.\n\n");
+}
+
+static void test_group_partition(unsigned int tileSz, bool useGlobalMem) {
+  hipError_t err;
+  int blockSize = 1;
+  int threadsPerBlock = 64;
+
+  int numTiles = (blockSize * threadsPerBlock) / tileSz;
+  // Build an array of expected reduction sum output on the host
+  // based on the sum of their respective thread ranks to use for verification
+  int* expectedSum = new int[numTiles];
+  int temp = 0, sum = 0;
+  for (int i = 1; i <= numTiles; i++) {
+    sum = temp;
+    temp = (((tileSz * i) - 1) * (tileSz * i)) / 2;
+    expectedSum[i-1] = temp - sum;
+  }
+
+  int* dResult = NULL;
+  hipMalloc(&dResult, sizeof(int) * numTiles);
+
+  int* globalMem = NULL;
+  if (useGlobalMem) {
+    hipMalloc((void**)&globalMem, threadsPerBlock * sizeof(int));
+  }
+
+  int* hResult = NULL;
+  hipHostMalloc(&hResult, numTiles * sizeof(int), hipHostMallocDefault);
+  memset(hResult, 0, numTiles * sizeof(int));
+
+  // Launch Kernel
+  if (useGlobalMem) {
+    hipLaunchKernelGGL(kernel_cg_group_partition_dynamic, blockSize, threadsPerBlock, 0, 0, tileSz,
+                       dResult, useGlobalMem, globalMem);
+
+    err = hipDeviceSynchronize();
+    if (err != hipSuccess) {
+      fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));
+    }
+  } else {
+    hipLaunchKernelGGL(kernel_cg_group_partition_dynamic, blockSize, threadsPerBlock,
+                       threadsPerBlock * sizeof(int), 0, tileSz, dResult, useGlobalMem, globalMem);
+
+    err = hipDeviceSynchronize();
+    if (err != hipSuccess) {
+      fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));
+    }
+  }
+
+  hipMemcpy(hResult, dResult, numTiles * sizeof(int), hipMemcpyDeviceToHost);
+
+  verifyResults(expectedSum, hResult, numTiles);
+
+  // Free all allocated memory on host and device
+  hipFree(dResult);
+  hipFree(hResult);
+  if (useGlobalMem) {
+    hipFree(globalMem);
+  }
+  delete[] expectedSum;
+
+  printf("\n...PASSED.\n\n");
+}
+
+TEST_CASE("Unit_tiled_partition") {
+  // 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;
+
+  if (!deviceProperties.cooperativeLaunch) {
+    HipTest::HIP_SKIP_TEST("Device doesn't support cooperative launch!");
+  }
+
+  bool useGlobalMem = true;
+  std::cout << "Testing static tiled_partition for different tile sizes" << std::endl;
+  std::cout << "\nUsing global memory for computation\n";
+  /* Test static tile_partition */
+  std::cout << "TEST 1:" << '\n' << std::endl;
+  test_group_partition<2>(useGlobalMem);
+  std::cout << "TEST 2:" << '\n' << std::endl;
+  test_group_partition<4>(useGlobalMem);
+  std::cout << "TEST 3:" << '\n' << std::endl;
+  test_group_partition<8>(useGlobalMem);
+  std::cout << "TEST 4:" << '\n' << std::endl;
+  test_group_partition<16>(useGlobalMem);
+  std::cout << "TEST 5:" << '\n' << std::endl;
+  test_group_partition<32>(useGlobalMem);
+
+  useGlobalMem = false;
+  std::cout << "Testing static tiled_partition for different tile sizes" << std::endl;
+  std::cout << "\nUsing shared memory for computation\n";
+  /* Test static tile_partition */
+  std::cout << "TEST 1:" << '\n' << std::endl;
+  test_group_partition<2>(useGlobalMem);
+  std::cout << "TEST 2:" << '\n' << std::endl;
+  test_group_partition<4>(useGlobalMem);
+  std::cout << "TEST 3:" << '\n' << std::endl;
+  test_group_partition<8>(useGlobalMem);
+  std::cout << "TEST 4:" << '\n' << std::endl;
+  test_group_partition<16>(useGlobalMem);
+  std::cout << "TEST 5:" << '\n' << std::endl;
+  test_group_partition<32>(useGlobalMem);
+
+
+  std::cout << "Now testing dynamic tiled_partition for different tile sizes" << '\n' << std::endl;
+
+  /* Test dynamic group partition*/
+  useGlobalMem = true;
+  int testNo = 1;
+  std::vector<unsigned int> tileSizes = {2, 4, 8, 16, 32};
+  std::cout << "\nUsing global memory for computation\n";
+  for (auto i : tileSizes) {
+    std::cout << "TEST " << testNo << ":" << '\n' << std::endl;
+    test_group_partition(i, useGlobalMem);
+    testNo++;
+  }
+
+  useGlobalMem = false;
+  testNo = 1;
+  std::cout << "\nUsing shared memory for computation\n";
+  for (auto i : tileSizes) {
+    std::cout << "TEST " << testNo << ":" << '\n' << std::endl;
+    test_group_partition(i, useGlobalMem);
+    testNo++;
+  }
+  printf("\n...PASSED.\n\n");
+  return;
+}