diff --git a/projects/hip-tests/catch/include/cpu_grid.h b/projects/hip-tests/catch/include/cpu_grid.h new file mode 100644 index 0000000000..98e5521840 --- /dev/null +++ b/projects/hip-tests/catch/include/cpu_grid.h @@ -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 + +#include +#include + +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 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 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 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(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(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(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(i * sm), 1, 1); }, + values(multipliers)), + map([sm = props.multiProcessorCount]( + double i) { return dim3(1, static_cast(i * sm), 1); }, + values(multipliers)), + map([sm = props.multiProcessorCount]( + double i) { return dim3(1, 1, static_cast(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(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(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(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(i * sm), 1, 1); }, + values(multipliers)), + map([sm = props.multiProcessorCount]( + double i) { return dim3(1, static_cast(i * sm), 1); }, + values(multipliers)), + map([sm = props.multiProcessorCount]( + double i) { return dim3(1, 1, static_cast(i * sm)); }, + values(multipliers)), + dim3(5, 5, 5)); +} \ No newline at end of file diff --git a/projects/hip-tests/catch/include/hip_test_defgroups.hh b/projects/hip-tests/catch/include/hip_test_defgroups.hh index 590f680880..f0966ae93b 100644 --- a/projects/hip-tests/catch/include/hip_test_defgroups.hh +++ b/projects/hip-tests/catch/include/hip_test_defgroups.hh @@ -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. + * @} + */ diff --git a/projects/hip-tests/catch/include/utils.hh b/projects/hip-tests/catch/include/utils.hh index 98f6676d6d..6dd8070b9c 100644 --- a/projects/hip-tests/catch/include/utils.hh +++ b/projects/hip-tests/catch/include/utils.hh @@ -20,6 +20,7 @@ THE SOFTWARE. #pragma once #include +#include #include #include @@ -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 +static inline void ArrayAllOf(const T* arr, uint32_t count, F value_gen) { + for (auto i = 0u; i < count; ++i) { + const std::optional 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 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) { diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/CMakeLists.txt b/projects/hip-tests/catch/unit/cooperativeGrps/CMakeLists.txt index ce7632fb8b..ad31c0a22a 100644 --- a/projects/hip-tests/catch/unit/cooperativeGrps/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/cooperativeGrps/CMakeLists.txt @@ -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 diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/cooperative_groups_common.hh b/projects/hip-tests/catch/unit/cooperativeGrps/cooperative_groups_common.hh new file mode 100644 index 0000000000..b0c8598581 --- /dev/null +++ b/projects/hip-tests/catch/unit/cooperativeGrps/cooperative_groups_common.hh @@ -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 +#include + +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 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; +} \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/grid_group.cc b/projects/hip-tests/catch/unit/cooperativeGrps/grid_group.cc new file mode 100644 index 0000000000..ee7402341a --- /dev/null +++ b/projects/hip-tests/catch/unit/cooperativeGrps/grid_group.cc @@ -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 +#include +#include + +/** + * @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 uint_arr_dev(LinearAllocs::hipMalloc, + grid.thread_count_ * sizeof(unsigned int)); + LinearAllocGuard 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 uint_arr_dev(LinearAllocs::hipMalloc, + grid.thread_count_ * sizeof(unsigned int)); + LinearAllocGuard 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 uint_arr_dev(LinearAllocs::hipMalloc, + array_len * sizeof(unsigned int)); + LinearAllocGuard uint_arr(LinearAllocs::hipHostMalloc, + array_len * sizeof(unsigned int)); + LinearAllocGuard 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(&atomic_val_ptr); + params[1] = reinterpret_cast(&uint_arr_dev_ptr); + params[2] = reinterpret_cast(&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); + } +} \ No newline at end of file