EXSWHTEC-256 - Implement tests for grid_group APIs (#153)

- Migrate basic Cooperative Groups tests and integrate to catch - Refactor basic Cooperative Groups tests - Rename tiled partition related files and fix minor bug - Add LaunchCooperativeKernal and LaunchCooperativeKernelMultiDevice tests - Refactor hipCGThreadBlockTileType to use common function - Fix updated file not added during merge - Add coalesced_group type tests - Add coalesced_group shuffle_up and shuffle_down tests - Add coalesced_group shuffle tests - test fails - Implement common code for cooperative group tests - Fixed compilation errror in cooperative_groups_common.hh - Implement busy wait device function - Reimplement tests for grid_group APIs - Add tests for grid_group member and non-member APIs - Refactor existing test for grid_group sync testing - Add thread and block dimensions generators - Add check of grid and block dimensions - Modify doxygen comments - Move cpu_grid.h and supporting functions to catch/include - Use warp_size from properties in grid/block dims generators - Fix condition for warp size 32 on AMD - Fix cpu_grid.h for warp function tests - Add missing include into cpu_grid.h - Code cleanup - Fix doxygen comments - Add missing include in utils header [ROCm/hip-tests commit: 22f3d9034b]
2023-07-18 09:25:00 +02:00
@@ -0,0 +1,154 @@
+/*
+Copyright (c) 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.
+*/
+
+#pragma once
+
+#include <optional>
+
+#include <hip_test_common.hh>
+#include <hip/hip_runtime_api.h>
+
+struct CPUGrid {
+  CPUGrid() = default;
+
+  CPUGrid(const dim3 grid_dim, const dim3 block_dim)
+      : grid_dim_{grid_dim},
+        block_dim_{block_dim},
+        block_count_{grid_dim.x * grid_dim.y * grid_dim.z},
+        threads_in_block_count_{block_dim.x * block_dim.y * block_dim.z},
+        thread_count_{block_count_ * threads_in_block_count_} {}
+
+  inline std::optional<unsigned int> thread_rank_in_block(
+      const unsigned int thread_rank_in_grid) const {
+    if (thread_rank_in_grid > thread_count_) {
+      return std::nullopt;
+    }
+
+    return thread_rank_in_grid % threads_in_block_count_;
+  }
+
+  inline std::optional<dim3> block_idx(const unsigned int thread_rank_in_grid) const {
+    if (thread_rank_in_grid > thread_count_) {
+      return std::nullopt;
+    }
+
+    dim3 block_idx;
+    const auto block_rank_in_grid = thread_rank_in_grid / threads_in_block_count_;
+    block_idx.x = block_rank_in_grid % grid_dim_.x;
+    block_idx.y = (block_rank_in_grid / grid_dim_.x) % grid_dim_.y;
+    block_idx.z = block_rank_in_grid / (grid_dim_.x * grid_dim_.y);
+
+    return block_idx;
+  }
+
+  inline std::optional<dim3> thread_idx(const unsigned int thread_rank_in_grid) const {
+    if (thread_rank_in_grid > thread_count_) {
+      return std::nullopt;
+    }
+
+    dim3 thread_idx;
+    const auto thread_rank_in_block = thread_rank_in_grid % threads_in_block_count_;
+    thread_idx.x = thread_rank_in_block % block_dim_.x;
+    thread_idx.y = (thread_rank_in_block / block_dim_.x) % block_dim_.y;
+    thread_idx.z = thread_rank_in_block / (block_dim_.x * block_dim_.y);
+
+    return thread_idx;
+  }
+
+  dim3 grid_dim_;
+  dim3 block_dim_;
+  unsigned int block_count_;
+  unsigned int threads_in_block_count_;
+  unsigned int thread_count_;
+};
+
+inline dim3 GenerateThreadDimensions() {
+  hipDeviceProp_t props;
+  HIP_CHECK(hipGetDeviceProperties(&props, 0));
+  const auto multipliers = {0.1, 0.2, 0.3, 0.4, 0.5, 0.6, 0.7, 0.8, 0.9, 1.0, 1.1, 1.2, 1.3,
+                            1.4, 1.5, 1.6, 1.7, 1.8, 1.9, 2.0, 2.1, 2.2, 2.3, 2.4, 2.5};
+  return GENERATE_COPY(
+      dim3(1, 1, 1), dim3(props.maxThreadsDim[0], 1, 1), dim3(1, props.maxThreadsDim[1], 1),
+      dim3(1, 1, props.maxThreadsDim[2]),
+      map([max = props.maxThreadsDim[0], warp_size = props.warpSize](
+              double i) { return dim3(std::min(static_cast<int>(i * warp_size), max), 1, 1); },
+          values(multipliers)),
+      map([max = props.maxThreadsDim[1], warp_size = props.warpSize](
+              double i) { return dim3(1, std::min(static_cast<int>(i * warp_size), max), 1); },
+          values(multipliers)),
+      map([max = props.maxThreadsDim[2], warp_size = props.warpSize](
+              double i) { return dim3(1, 1, std::min(static_cast<int>(i * warp_size), max)); },
+          values(multipliers)),
+      dim3(16, 8, 8), dim3(32, 32, 1), dim3(64, 8, 2), dim3(16, 16, 3), dim3(props.warpSize - 1, 3, 3),
+      dim3(props.warpSize + 1, 3, 3));
+}
+
+inline dim3 GenerateBlockDimensions() {
+  hipDeviceProp_t props;
+  HIP_CHECK(hipGetDeviceProperties(&props, 0));
+  const auto multipliers = {0.1, 0.5, 0.9, 1.0, 1.1, 1.5, 1.9, 2.0, 3.0, 4.0};
+  return GENERATE_COPY(dim3(1, 1, 1),
+                       map([sm = props.multiProcessorCount](
+                               double i) { return dim3(static_cast<int>(i * sm), 1, 1); },
+                           values(multipliers)),
+                       map([sm = props.multiProcessorCount](
+                               double i) { return dim3(1, static_cast<int>(i * sm), 1); },
+                           values(multipliers)),
+                       map([sm = props.multiProcessorCount](
+                               double i) { return dim3(1, 1, static_cast<int>(i * sm)); },
+                           values(multipliers)),
+                       dim3(5, 5, 5));
+}
+
+inline dim3 GenerateThreadDimensionsForShuffle() {
+  hipDeviceProp_t props;
+  HIP_CHECK(hipGetDeviceProperties(&props, 0));
+  const auto multipliers = {0.5, 0.9, 1.0, 1.5, 2.0};
+  return GENERATE_COPY(
+      dim3(1, 1, 1), dim3(props.maxThreadsDim[0], 1, 1), dim3(1, props.maxThreadsDim[1], 1),
+      dim3(1, 1, props.maxThreadsDim[2]),
+      map([max = props.maxThreadsDim[0], warp_size = props.warpSize](
+              double i) { return dim3(std::min(static_cast<int>(i * warp_size), max), 1, 1); },
+          values(multipliers)),
+      map([max = props.maxThreadsDim[1], warp_size = props.warpSize](
+              double i) { return dim3(1, std::min(static_cast<int>(i * warp_size), max), 1); },
+          values(multipliers)),
+      map([max = props.maxThreadsDim[2], warp_size = props.warpSize](
+              double i) { return dim3(1, 1, std::min(static_cast<int>(i * warp_size), max)); },
+          values(multipliers)),
+      dim3(16, 8, 8), dim3(32, 32, 1), dim3(64, 8, 2), dim3(16, 16, 3), dim3(props.warpSize - 1, 3, 3),
+      dim3(props.warpSize + 1, 3, 3));
+}
+
+inline dim3 GenerateBlockDimensionsForShuffle() {
+  hipDeviceProp_t props;
+  HIP_CHECK(hipGetDeviceProperties(&props, 0));
+  const auto multipliers = {0.5, 1.0};
+  return GENERATE_COPY(dim3(1, 1, 1),
+                       map([sm = props.multiProcessorCount](
+                               double i) { return dim3(static_cast<int>(i * sm), 1, 1); },
+                           values(multipliers)),
+                       map([sm = props.multiProcessorCount](
+                               double i) { return dim3(1, static_cast<int>(i * sm), 1); },
+                           values(multipliers)),
+                       map([sm = props.multiProcessorCount](
+                               double i) { return dim3(1, 1, static_cast<int>(i * sm)); },
+                           values(multipliers)),
+                       dim3(5, 5, 5));
+}
@@ -128,3 +128,10 @@ THE SOFTWARE.
 * This section describes the various kernel functions invocation.
 * @}
 */
+
+/**
+ * @defgroup DeviceLanguageTest Device Language
+ * @{
+ * This section describes tests for the Device Language API.
+ * @}
+ */
@@ -20,6 +20,7 @@ THE SOFTWARE.
 #pragma once

 #include <chrono>
+#include <optional>

 #include <hip_test_common.hh>
 #include <hip/hip_runtime_api.h>
@@ -54,6 +55,20 @@ void ArrayFindIfNot(T* const array, const T expected_value, const size_t num_ele
  ArrayFindIfNot(array, array + num_elements, expected_value);
 }

+template <typename T, typename F>
+static inline void ArrayAllOf(const T* arr, uint32_t count, F value_gen) {
+  for (auto i = 0u; i < count; ++i) {
+    const std::optional<T> expected_val = value_gen(i);
+    if (!expected_val.has_value()) continue;
+    // Using require on every iteration leads to a noticeable performance loss on large arrays,
+    // even when the require passes.
+    if (arr[i] != expected_val.value()) {
+      INFO("Mismatch at index: " << i);
+      REQUIRE(arr[i] == expected_val.value());
+    }
+  }
+}
+
 template <typename T, typename F>
 void PitchedMemoryVerify(T* const ptr, const size_t pitch, const size_t width, const size_t height,
                         const size_t depth, F expected_value_generator) {
@@ -6,6 +6,7 @@ set(TEST_SRC
  hipCGMultiGridGroupType.cc
  hipCGMultiGridGroupTypeViaBaseType.cc
  hipCGMultiGridGroupTypeViaPublicApi.cc
+  grid_group.cc
  coalesced_groups_shfl_down.cc
  coalesced_groups_shfl_up.cc
  simple_coalesced_groups.cc
@@ -0,0 +1,71 @@
+/*
+Copyright (c) 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.
+*/
+
+#pragma once
+
+#include <hip_test_common.hh>
+#include <hip/hip_cooperative_groups.h>
+
+namespace {
+#if (!__GFX8__ && !__GFX9__) || HT_NVIDIA
+constexpr size_t kWarpSize = 32;
+#else
+constexpr size_t kWarpSize = 64;
+#endif
+}  // namespace
+
+constexpr int MaxGPUs = 8;
+
+__device__ inline unsigned int thread_rank_in_grid() {
+  const auto block_size = blockDim.x * blockDim.y * blockDim.z;
+  const auto block_rank_in_grid = (blockIdx.z * gridDim.y + blockIdx.y) * gridDim.x + blockIdx.x;
+  const auto thread_rank_in_block =
+      (threadIdx.z * blockDim.y + threadIdx.y) * blockDim.x + threadIdx.x;
+  return block_rank_in_grid * block_size + thread_rank_in_block;
+}
+
+static __device__ void busy_wait(unsigned long long wait_period) {
+  unsigned long long time_diff = 0;
+  unsigned long long last_clock = clock64();
+  while (time_diff < wait_period) {
+    unsigned long long cur_clock = clock64();
+    if (cur_clock > last_clock) {
+      time_diff += (cur_clock - last_clock);
+    }
+    last_clock = cur_clock;
+  }
+}
+
+template <class T> bool CheckDimensions(unsigned int device, T kernel, dim3 blocks, dim3 threads) {
+  hipDeviceProp_t props;
+  int max_blocks_per_sm = 0;
+  int num_sm = 0;
+  HIP_CHECK(hipSetDevice(device));
+  HIP_CHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&max_blocks_per_sm, kernel,
+                                                         threads.x * threads.y * threads.z, 0));
+
+  HIP_CHECK(hipGetDeviceProperties(&props, device));
+  num_sm = props.multiProcessorCount;
+
+  if ((blocks.x * blocks.y * blocks.z) > max_blocks_per_sm * num_sm) {
+    return false;
+  }
+
+  return true;
+}
@@ -0,0 +1,285 @@
+/*
+Copyright (c) 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.
+*/
+#include "cooperative_groups_common.hh"
+
+#include <cpu_grid.h>
+#include <resource_guards.hh>
+#include <utils.hh>
+
+/**
+ * @addtogroup grid_group grid_group
+ * @{
+ * @ingroup DeviceLanguageTest
+ * Contains unit tests for all grid_group APIs
+ */
+
+namespace cg = cooperative_groups;
+
+static __global__ void grid_group_size_getter(unsigned int* sizes) {
+  sizes[thread_rank_in_grid()] = cg::this_grid().size();
+}
+
+static __global__ void grid_group_thread_rank_getter(unsigned int* thread_ranks) {
+  thread_ranks[thread_rank_in_grid()] = cg::this_grid().thread_rank();
+}
+
+static __global__ void grid_group_is_valid_getter(unsigned int* is_valid_flags) {
+  is_valid_flags[thread_rank_in_grid()] = cg::this_grid().is_valid();
+}
+
+static __global__ void grid_group_non_member_size_getter(unsigned int* sizes) {
+  sizes[thread_rank_in_grid()] = cg::group_size(cg::this_grid());
+}
+
+static __global__ void grid_group_non_member_thread_rank_getter(unsigned int* thread_ranks) {
+  thread_ranks[thread_rank_in_grid()] = cg::thread_rank(cg::this_grid());
+}
+
+static __global__ void sync_kernel(unsigned int* atomic_val, unsigned int* array,
+                                   unsigned int loops) {
+  cg::grid_group grid = cg::this_grid();
+  unsigned rank = grid.thread_rank();
+
+  int offset = (blockIdx.z * gridDim.y + blockIdx.y) * gridDim.x + blockIdx.x;
+  for (int i = 0; i < loops; i++) {
+    // Make the last thread run way behind everyone else.
+    // If the sync below fails, then the other threads may hit the
+    // atomicInc instruction many times before the last thread ever gets to it.
+    // If the sync works, then it will likely contain "total number of blocks"*i
+    if (rank == (grid.size() - 1)) {
+      busy_wait(100000);
+    }
+    if (threadIdx.x == blockDim.x - 1 && threadIdx.y == blockDim.y - 1 &&
+        threadIdx.z == blockDim.z - 1) {
+      array[offset] = atomicInc(&atomic_val[0], UINT_MAX);
+    }
+    grid.sync();
+    offset += gridDim.x * gridDim.y * gridDim.z;
+  }
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Launches kernels that write the return values of size, thread_rank and is_valid member
+ * functions to an output array that is validated on the host side. The kernels are run
+ * sequentially, reusing the output array, to avoid running out of device memory for large kernel
+ * launches.
+ * Test source
+ * ------------------------
+ *  - unit/cooperativeGrps/grid_group.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ *  - Device supports cooperative launch
+ */
+TEST_CASE("Unit_Grid_Group_Getters_Positive_Basic") {
+  int device;
+  hipDeviceProp_t device_properties;
+  HIP_CHECK(hipGetDevice(&device));
+  HIP_CHECK(hipGetDeviceProperties(&device_properties, device));
+
+  if (!device_properties.cooperativeLaunch) {
+    HipTest::HIP_SKIP_TEST("Device doesn't support cooperative launch!");
+    return;
+  }
+
+  const auto blocks = GenerateBlockDimensions();
+  const auto threads = GenerateThreadDimensions();
+  if (!CheckDimensions(device, grid_group_size_getter, blocks, threads)) return;
+  INFO("Grid dimensions: x " << blocks.x << ", y " << blocks.y << ", z " << blocks.z);
+  INFO("Block dimensions: x " << threads.x << ", y " << threads.y << ", z " << threads.z);
+
+  const CPUGrid grid(blocks, threads);
+
+  LinearAllocGuard<unsigned int> uint_arr_dev(LinearAllocs::hipMalloc,
+                                              grid.thread_count_ * sizeof(unsigned int));
+  LinearAllocGuard<unsigned int> uint_arr(LinearAllocs::hipHostMalloc,
+                                          grid.thread_count_ * sizeof(unsigned int));
+
+  // Launch Kernel
+  unsigned int* uint_arr_dev_ptr = uint_arr_dev.ptr();
+  void* params[1];
+  params[0] = &uint_arr_dev_ptr;
+
+  HIP_CHECK(hipLaunchCooperativeKernel(grid_group_size_getter, blocks, threads, params, 0, 0));
+
+  HIP_CHECK(hipMemcpy(uint_arr.ptr(), uint_arr_dev.ptr(),
+                      grid.thread_count_ * sizeof(*uint_arr.ptr()), hipMemcpyDeviceToHost));
+  HIP_CHECK(hipDeviceSynchronize());
+  HIP_CHECK(
+      hipLaunchCooperativeKernel(grid_group_thread_rank_getter, blocks, threads, params, 0, 0));
+
+  // Verify grid_group.size() values
+  ArrayAllOf(uint_arr.ptr(), grid.thread_count_,
+             [size = grid.thread_count_](uint32_t) { return size; });
+
+  HIP_CHECK(hipMemcpy(uint_arr.ptr(), uint_arr_dev.ptr(),
+                      grid.thread_count_ * sizeof(*uint_arr.ptr()), hipMemcpyDeviceToHost));
+  HIP_CHECK(hipDeviceSynchronize());
+  HIP_CHECK(hipLaunchCooperativeKernel(grid_group_is_valid_getter, blocks, threads, params, 0, 0));
+
+  // Verify grid_group.thread_rank() values
+  ArrayAllOf(uint_arr.ptr(), grid.thread_count_, [](uint32_t i) { return i; });
+
+  HIP_CHECK(hipMemcpy(uint_arr.ptr(), uint_arr_dev.ptr(),
+                      grid.thread_count_ * sizeof(*uint_arr.ptr()), hipMemcpyDeviceToHost));
+  HIP_CHECK(hipDeviceSynchronize());
+
+  // Verify grid_group.is_valid() values
+  ArrayAllOf(uint_arr.ptr(), grid.thread_count_, [](uint32_t i) { return 1; });
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Launches kernels that write the return values of size and thread_rank non-member functions
+ * to an output array that is validated on the host side. The kernels are run sequentially, reusing
+ * the output array, to avoid running out of device memory for large kernel launches.
+ * Test source
+ * ------------------------
+ *  - unit/cooperativeGrps/grid_group.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ *  - Device supports cooperative launch
+ */
+TEST_CASE("Unit_Grid_Group_Getters_Via_Non_Member_Functions_Positive_Basic") {
+  int device;
+  hipDeviceProp_t device_properties;
+  HIP_CHECK(hipGetDevice(&device));
+  HIP_CHECK(hipGetDeviceProperties(&device_properties, device));
+
+  if (!device_properties.cooperativeLaunch) {
+    HipTest::HIP_SKIP_TEST("Device doesn't support cooperative launch!");
+    return;
+  }
+
+  const auto blocks = GenerateBlockDimensions();
+  const auto threads = GenerateThreadDimensions();
+  if (!CheckDimensions(device, grid_group_non_member_size_getter, blocks, threads)) return;
+  INFO("Grid dimensions: x " << blocks.x << ", y " << blocks.y << ", z " << blocks.z);
+  INFO("Block dimensions: x " << threads.x << ", y " << threads.y << ", z " << threads.z);
+
+  const CPUGrid grid(blocks, threads);
+
+  LinearAllocGuard<unsigned int> uint_arr_dev(LinearAllocs::hipMalloc,
+                                              grid.thread_count_ * sizeof(unsigned int));
+  LinearAllocGuard<unsigned int> uint_arr(LinearAllocs::hipHostMalloc,
+                                          grid.thread_count_ * sizeof(unsigned int));
+
+  // Launch Kernel
+  unsigned int* uint_arr_dev_ptr = uint_arr_dev.ptr();
+  void* params[1];
+  params[0] = &uint_arr_dev_ptr;
+
+  HIP_CHECK(
+      hipLaunchCooperativeKernel(grid_group_non_member_size_getter, blocks, threads, params, 0, 0));
+
+  HIP_CHECK(hipMemcpy(uint_arr.ptr(), uint_arr_dev.ptr(),
+                      grid.thread_count_ * sizeof(*uint_arr.ptr()), hipMemcpyDeviceToHost));
+  HIP_CHECK(hipDeviceSynchronize());
+  HIP_CHECK(hipLaunchCooperativeKernel(grid_group_non_member_thread_rank_getter, blocks, threads,
+                                       params, 0, 0));
+
+  // Verify grid_group.size() values
+  ArrayAllOf(uint_arr.ptr(), grid.thread_count_,
+             [size = grid.thread_count_](uint32_t) { return size; });
+
+  HIP_CHECK(hipMemcpy(uint_arr.ptr(), uint_arr_dev.ptr(),
+                      grid.thread_count_ * sizeof(*uint_arr.ptr()), hipMemcpyDeviceToHost));
+  HIP_CHECK(hipDeviceSynchronize());
+
+  // Verify grid_group.thread_rank() values
+  ArrayAllOf(uint_arr.ptr(), grid.thread_count_, [](uint32_t i) { return i; });
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *    - Launches a kernel where the last thread in a block atomically increments a global variable
+ * within a work loop. The value returned from this atomic increment entirely depends on the order
+ * the threads arrive at the atomic instruction. Each thread then stores the result in the global
+ * array based on its block id. A wait loop is inserted into the last thread so that it runs behind
+ * all other threads. If the sync doesn't work, the other threads will increment the atomic variable
+ * many times before the last thread gets to it and it will read a very large value. If the sync
+ * works, each thread will increment the variable once per loop iteration and the last thread will
+ * contain total number of blocks * loop iteration.
+ * Test source
+ * ------------------------
+ *  - unit/cooperativeGrps/grid_group.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 5.2
+ *  - Device supports cooperative launch
+ */
+TEST_CASE("Unit_Grid_Group_Sync_Positive_Basic") {
+  int device;
+  hipDeviceProp_t device_properties;
+  HIP_CHECK(hipGetDevice(&device));
+  HIP_CHECK(hipGetDeviceProperties(&device_properties, device));
+
+  if (!device_properties.cooperativeLaunch) {
+    HipTest::HIP_SKIP_TEST("Device doesn't support cooperative launch!");
+    return;
+  }
+
+  auto loops = GENERATE(2, 4, 8, 16);
+  const auto blocks = GenerateBlockDimensions();
+  const auto threads = GenerateThreadDimensions();
+  if (!CheckDimensions(device, sync_kernel, blocks, threads)) return;
+  INFO("Grid dimensions: x " << blocks.x << ", y " << blocks.y << ", z " << blocks.z);
+  INFO("Block dimensions: x " << threads.x << ", y " << threads.y << ", z " << threads.z);
+
+  const CPUGrid grid(blocks, threads);
+  unsigned int array_len = grid.block_count_ * loops;
+
+  LinearAllocGuard<unsigned int> uint_arr_dev(LinearAllocs::hipMalloc,
+                                              array_len * sizeof(unsigned int));
+  LinearAllocGuard<unsigned int> uint_arr(LinearAllocs::hipHostMalloc,
+                                          array_len * sizeof(unsigned int));
+  LinearAllocGuard<unsigned int> atomic_val(LinearAllocs::hipMalloc, sizeof(unsigned int));
+  HIP_CHECK(hipMemset(atomic_val.ptr(), 0, sizeof(unsigned int)));
+
+  // Launch Kernel
+  unsigned int* uint_arr_dev_ptr = uint_arr_dev.ptr();
+  unsigned int* atomic_val_ptr = atomic_val.ptr();
+  void* params[3];
+  params[0] = reinterpret_cast<void*>(&atomic_val_ptr);
+  params[1] = reinterpret_cast<void*>(&uint_arr_dev_ptr);
+  params[2] = reinterpret_cast<void*>(&loops);
+
+  HIP_CHECK(hipLaunchCooperativeKernel(sync_kernel, blocks, threads, params, 0, 0));
+
+  HIP_CHECK(hipMemcpy(uint_arr.ptr(), uint_arr_dev.ptr(), array_len * sizeof(*uint_arr.ptr()),
+                      hipMemcpyDeviceToHost));
+
+  HIP_CHECK(hipDeviceSynchronize());
+
+  // Verify host buffer values
+  unsigned int max_in_this_loop = 0;
+  for (unsigned int i = 0; i < loops; i++) {
+    max_in_this_loop += grid.block_count_;
+    unsigned int j = 0;
+    for (j = 0; j < grid.block_count_ - 1; j++) {
+      REQUIRE(uint_arr.ptr()[i * grid.block_count_ + j] < max_in_this_loop);
+    }
+    REQUIRE(uint_arr.ptr()[i * grid.block_count_ + j] == max_in_this_loop - 1);
+  }
+}