From a3cc0be869a9709d9d8b20dd66e113337e4571fe Mon Sep 17 00:00:00 2001 From: Nives Vukovic Date: Thu, 16 Nov 2023 18:25:33 +0000 Subject: [PATCH] EXSWHTEC-308 - Migrate and refactor cooperative groups tests from HIP repository (#238) Change-Id: Ib46b7c038a5bda9d05f5d55a7269a7c645b0d049 [ROCm/hip-tests commit: e0612f9346c7d16801b92f7d1ac8b1a7401f94d2] --- .../catch/unit/cooperativeGrps/CMakeLists.txt | 23 +- ...dGroups.cc => hipCGCoalescedGroups_old.cc} | 0 .../cooperativeGrps/hipCGGridGroupType_old.cc | 496 ++++++++++++++ .../hipCGMultiGridGroupType.cc | 240 ------- .../hipCGMultiGridGroupTypeViaBaseType.cc | 234 ------- .../hipCGMultiGridGroupTypeViaPublicApi.cc | 230 ------- .../hipCGMultiGridGroupType_old.cc | 638 ++++++++++++++++++ .../hipCGThreadBlockTileTypeShfl_old.cc | 198 ++++++ .../cooperativeGrps/hipCGThreadBlockType.cc | 177 ----- .../hipCGThreadBlockTypeViaBaseType.cc | 136 ---- .../hipCGThreadBlockTypeViaPublicApi.cc | 136 ---- .../hipCGThreadBlockType_old.cc | 225 ++++++ .../cooperativeGrps/hipCGTiledPartition.cc | 385 ----------- .../hipCGTiledPartitionType_old.cc | 279 ++++++++ ...pLaunchCooperativeKernelMultiDevice_old.cc | 606 +++++++++++++++++ .../hipLaunchCooperativeKernel_old.cc | 364 ++++++++++ .../unit/cooperativeGrps/hip_cg_common.hh | 68 ++ 17 files changed, 2885 insertions(+), 1550 deletions(-) rename projects/hip-tests/catch/unit/cooperativeGrps/{hipCGCoalescedGroups.cc => hipCGCoalescedGroups_old.cc} (100%) create mode 100644 projects/hip-tests/catch/unit/cooperativeGrps/hipCGGridGroupType_old.cc delete mode 100644 projects/hip-tests/catch/unit/cooperativeGrps/hipCGMultiGridGroupType.cc delete mode 100644 projects/hip-tests/catch/unit/cooperativeGrps/hipCGMultiGridGroupTypeViaBaseType.cc delete mode 100644 projects/hip-tests/catch/unit/cooperativeGrps/hipCGMultiGridGroupTypeViaPublicApi.cc create mode 100644 projects/hip-tests/catch/unit/cooperativeGrps/hipCGMultiGridGroupType_old.cc create mode 100644 projects/hip-tests/catch/unit/cooperativeGrps/hipCGThreadBlockTileTypeShfl_old.cc delete mode 100644 projects/hip-tests/catch/unit/cooperativeGrps/hipCGThreadBlockType.cc delete mode 100644 projects/hip-tests/catch/unit/cooperativeGrps/hipCGThreadBlockTypeViaBaseType.cc delete mode 100644 projects/hip-tests/catch/unit/cooperativeGrps/hipCGThreadBlockTypeViaPublicApi.cc create mode 100644 projects/hip-tests/catch/unit/cooperativeGrps/hipCGThreadBlockType_old.cc delete mode 100644 projects/hip-tests/catch/unit/cooperativeGrps/hipCGTiledPartition.cc create mode 100644 projects/hip-tests/catch/unit/cooperativeGrps/hipCGTiledPartitionType_old.cc create mode 100644 projects/hip-tests/catch/unit/cooperativeGrps/hipLaunchCooperativeKernelMultiDevice_old.cc create mode 100644 projects/hip-tests/catch/unit/cooperativeGrps/hipLaunchCooperativeKernel_old.cc create mode 100644 projects/hip-tests/catch/unit/cooperativeGrps/hip_cg_common.hh diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/CMakeLists.txt b/projects/hip-tests/catch/unit/cooperativeGrps/CMakeLists.txt index 829758b622..1090a68382 100644 --- a/projects/hip-tests/catch/unit/cooperativeGrps/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/cooperativeGrps/CMakeLists.txt @@ -1,26 +1,25 @@ # Common Tests - Test independent of all platforms set(TEST_SRC - hipCGThreadBlockType.cc - hipCGThreadBlockTypeViaBaseType.cc - hipCGThreadBlockTypeViaPublicApi.cc - hipCGMultiGridGroupType.cc - hipCGMultiGridGroupTypeViaBaseType.cc - hipCGMultiGridGroupTypeViaPublicApi.cc + hipCGThreadBlockType_old.cc + hipCGMultiGridGroupType_old.cc + hipCGGridGroupType_old.cc + hipCGTiledPartitionType_old.cc + hipCGThreadBlockTileTypeShfl_old.cc + hipCGCoalescedGroups_old.cc + hipLaunchCooperativeKernel_old.cc + hipLaunchCooperativeKernelMultiDevice_old.cc grid_group.cc coalesced_groups_shfl_down.cc coalesced_groups_shfl_up.cc - hipCGTiledPartition.cc - hipCGCoalescedGroups.cc coalesced_tiled_groups_metagrp.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") + set_source_files_properties(hipCGMultiGridGroupType_old.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(hipLaunchCooperativeKernelMultiDevice_old.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") 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") + 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, -gencode arch=compute_86,code=sm_86, -gencode=arch=compute_86,code=compute_86") else() hip_add_exe_to_target(NAME coopGrpTest TEST_SRC ${TEST_SRC} diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/hipCGCoalescedGroups.cc b/projects/hip-tests/catch/unit/cooperativeGrps/hipCGCoalescedGroups_old.cc similarity index 100% rename from projects/hip-tests/catch/unit/cooperativeGrps/hipCGCoalescedGroups.cc rename to projects/hip-tests/catch/unit/cooperativeGrps/hipCGCoalescedGroups_old.cc diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/hipCGGridGroupType_old.cc b/projects/hip-tests/catch/unit/cooperativeGrps/hipCGGridGroupType_old.cc new file mode 100644 index 0000000000..9f908a5be6 --- /dev/null +++ b/projects/hip-tests/catch/unit/cooperativeGrps/hipCGGridGroupType_old.cc @@ -0,0 +1,496 @@ +/* +Copyright (c) 2020 - 2022 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 +#include + +#include "hip_cg_common.hh" + +namespace cg = cooperative_groups; + +static __device__ int gm[2]; + +static __global__ void kernel_cg_grid_group_type(int* size_dev, int* thd_rank_dev, + int* is_valid_dev, int* sync_dev) { + cg::grid_group gg = cg::this_grid(); + int gIdx = (blockIdx.x * blockDim.x) + threadIdx.x; + + // Test size + size_dev[gIdx] = gg.size(); + + // Test thread_rank + thd_rank_dev[gIdx] = gg.thread_rank(); + + // Test is_valid + is_valid_dev[gIdx] = gg.is_valid(); + + // Test sync + if (blockIdx.x == 0 && threadIdx.x == 0) + gm[0] = 10; + else if (blockIdx.x == 1 && threadIdx.x == 0) + gm[1] = 20; + gg.sync(); + sync_dev[gIdx] = gm[1] * gm[0]; +} + +static __global__ void kernel_cg_grid_group_type_via_base_type(int* size_dev, int* thd_rank_dev, + int* is_valid_dev, int* sync_dev) { + cg::thread_group tg = cg::this_grid(); + int gIdx = (blockIdx.x * blockDim.x) + threadIdx.x; + + // Test size + size_dev[gIdx] = tg.size(); + + // Test thread_rank + thd_rank_dev[gIdx] = tg.thread_rank(); + + // Test is_valid +#ifdef __HIP_PLATFORM_AMD__ + is_valid_dev[gIdx] = tg.is_valid(); +#else + // Cuda has no thread_group.is_valid() + is_valid_dev[gIdx] = true; +#endif + + // Test sync + if (blockIdx.x == 0 && threadIdx.x == 0) + gm[0] = 10; + else if (blockIdx.x == 1 && threadIdx.x == 0) + gm[1] = 20; + tg.sync(); + sync_dev[gIdx] = gm[1] * gm[0]; +} + +static __global__ void kernel_cg_grid_group_type_via_public_api(int* size_dev, int* thd_rank_dev, + int* is_valid_dev, int* sync_dev) { + cg::grid_group gg = cg::this_grid(); + int gIdx = (blockIdx.x * blockDim.x) + threadIdx.x; + + // Test group_size api + size_dev[gIdx] = cg::group_size(gg); + + // Test thread_rank api + thd_rank_dev[gIdx] = cg::thread_rank(gg); + + // Test is_valid api + is_valid_dev[gIdx] = gg.is_valid(); + + // Test sync + if (blockIdx.x == 0 && threadIdx.x == 0) + gm[0] = 10; + else if (blockIdx.x == 1 && threadIdx.x == 0) + gm[1] = 20; + cg::sync(gg); + sync_dev[gIdx] = gm[1] * gm[0]; +} + +static __global__ void coop_kernel(unsigned int* first_array, unsigned int* second_array, + unsigned int loops, unsigned int array_len) { + cg::grid_group grid = cg::this_grid(); + unsigned int rank = grid.thread_rank(); + unsigned int grid_size = grid.size(); + + for (int i = 0; i < loops; i++) { + // The goal of this loop is to directly add in values from + // array one into array two, on a per-wave basis. + for (int offset = rank; offset < array_len; offset += grid_size) { + second_array[offset] += first_array[offset]; + } + + grid.sync(); + + // The goal of this loop is to pull data the "mirror" lane in + // array two and add it back into array one. This causes inter- + // thread swizzling. + for (int offset = rank; offset < array_len; offset += grid_size) { + unsigned int swizzle_offset = array_len - offset - 1; + first_array[offset] += second_array[swizzle_offset]; + } + + grid.sync(); + } +} + +static __global__ void test_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.x; + for (int i = 0; i < loops; i++) { + // Make the last thread run way behind everyone else. + // If the barrier below fails, then the other threads may hit the + // atomicInc instruction many times before the last thread ever gets to it. + // As such, without the barrier, the last array entry will eventually + // contain a very large value, defined by however many times the other + // wavefronts make it through this loop. + // If the barrier works, then it will likely contain some number + // near "total number of blocks". It will be the last wavefront to + // reach the atomicInc, but everyone will have only hit the atomic once. + if (rank == (grid.size() - 1)) { + long long time_diff = 0; + long long last_clock = clock64(); + do { + long long cur_clock = clock64(); + if (cur_clock > last_clock) { + time_diff += (cur_clock - last_clock); + } + // If it rolls over, we don't know how much to add to catch up. + // So just ignore those slipped cycles. + last_clock = cur_clock; + } while (time_diff < 1000000); + } + + if (threadIdx.x == 0) { + array[offset] = atomicInc(&atomic_val[0], UINT_MAX); + } + grid.sync(); + offset += gridDim.x; + } +} + +__global__ void test_kernel_gfx11(unsigned int* atomic_val, unsigned int* array, + unsigned int loops) { +#if HT_AMD + cg::grid_group grid = cg::this_grid(); + unsigned rank = grid.thread_rank(); + + int offset = blockIdx.x; + for (int i = 0; i < loops; i++) { + // Make the last thread run way behind everyone else. + // If the barrier below fails, then the other threads may hit the + // atomicInc instruction many times before the last thread ever gets + // to it. + // As such, without the barrier, the last array entry will eventually + // contain a very large value, defined by however many times the other + // wavefronts make it through this loop. + // If the barrier works, then it will likely contain some number + // near "total number of blocks". It will be the last wavefront to + // reach the atomicInc, but everyone will have only hit the atomic once. + if (rank == (grid.size() - 1)) { + long long time_diff = 0; + long long last_clock = wall_clock64(); + do { + long long cur_clock = wall_clock64(); + if (cur_clock > last_clock) { + time_diff += (cur_clock - last_clock); + } + // If it rolls over, we don't know how much to add to catch up. + // So just ignore those slipped cycles. + last_clock = cur_clock; + } while (time_diff < 1000000); + } + + if (threadIdx.x == 0) { + array[offset] = atomicInc(&atomic_val[0], UINT_MAX); + } + grid.sync(); + offset += gridDim.x; + } +#endif +} + +static void verify_coop_buffers(unsigned int* host_input, unsigned int* first_array, + unsigned int* second_array, unsigned int loops, + unsigned int array_len) { + unsigned int* expected_first_array = host_input; + unsigned int* expected_second_array = + reinterpret_cast(malloc(sizeof(unsigned int) * array_len)); + memset(expected_second_array, 0, sizeof(unsigned int) * array_len); + + for (int i = 0; i < loops; i++) { + for (int offset = 0; offset < array_len; offset++) { + expected_second_array[offset] += expected_first_array[offset]; + } + + for (int offset = 0; offset < array_len; offset++) { + unsigned int swizzle_offset = array_len - offset - 1; + expected_first_array[offset] += expected_second_array[swizzle_offset]; + } + } + + for (int i = 0; i < array_len; i++) { + REQUIRE(first_array[i] == expected_first_array[i]); + REQUIRE(second_array[i] == expected_second_array[i]); + } + + free(expected_second_array); +} + +static void verify_barrier_buffer(unsigned int loops, unsigned int warps, + unsigned int* host_buffer) { + unsigned int max_in_this_loop = 0; + for (unsigned int i = 0; i < loops; i++) { + max_in_this_loop += warps; + for (unsigned int j = 0; j < warps; j++) { + REQUIRE(host_buffer[i * warps + j] <= max_in_this_loop); + } + } +} + +template static void test_cg_grid_group_type(F kernel_func, int block_size) { + int num_bytes = sizeof(int) * 2 * block_size; + int *size_dev, *size_host; + int *thd_rank_dev, *thd_rank_host; + int *is_valid_dev, *is_valid_host; + int *sync_dev, *sync_host; + + // Allocate device memory + HIP_CHECK(hipMalloc(&size_dev, num_bytes)); + HIP_CHECK(hipMalloc(&thd_rank_dev, num_bytes)); + HIP_CHECK(hipMalloc(&is_valid_dev, num_bytes)); + HIP_CHECK(hipMalloc(&sync_dev, num_bytes)); + + // Allocate host memory + HIP_CHECK(hipHostMalloc(&size_host, num_bytes)); + HIP_CHECK(hipHostMalloc(&thd_rank_host, num_bytes)); + HIP_CHECK(hipHostMalloc(&is_valid_host, num_bytes)); + HIP_CHECK(hipHostMalloc(&sync_host, num_bytes)); + + // Launch Kernel + void* params[4]; + params[0] = &size_dev; + params[1] = &thd_rank_dev; + params[2] = &is_valid_dev; + params[3] = &sync_dev; + HIP_CHECK(hipLaunchCooperativeKernel(kernel_func, 2, block_size, params, 0, 0)); + + // Copy result from device to host + HIP_CHECK(hipMemcpy(size_host, size_dev, num_bytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(thd_rank_host, thd_rank_dev, num_bytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(is_valid_host, is_valid_dev, num_bytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(sync_host, sync_dev, num_bytes, hipMemcpyDeviceToHost)); + + // Validate results for both blocks together + for (int i = 0; i < 2 * block_size; ++i) { + ASSERT_EQUAL(size_host[i], 2 * block_size); + ASSERT_EQUAL(thd_rank_host[i], i); + ASSERT_EQUAL(is_valid_host[i], 1); + ASSERT_EQUAL(sync_host[i], 200); + } + + // Free device memory + HIP_CHECK(hipFree(size_dev)); + HIP_CHECK(hipFree(thd_rank_dev)); + HIP_CHECK(hipFree(is_valid_dev)); + HIP_CHECK(hipFree(sync_dev)); + + // Free host memory + HIP_CHECK(hipHostFree(size_host)); + HIP_CHECK(hipHostFree(thd_rank_host)); + HIP_CHECK(hipHostFree(is_valid_host)); + HIP_CHECK(hipHostFree(sync_host)); +} + +TEST_CASE("Unit_hipCGGridGroupType_Basic") { + // Use default device for validating the test + 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; + } + + void* (*kernel_func)(void); + + SECTION("Default grid group API test") { + kernel_func = reinterpret_cast(kernel_cg_grid_group_type); + } +#if HT_AMD + SECTION("Base type grid group API test") { + kernel_func = reinterpret_cast(kernel_cg_grid_group_type_via_base_type); + } +#endif + + SECTION("Public API grid group test") { + kernel_func = reinterpret_cast(kernel_cg_grid_group_type_via_public_api); + } + + // Test for block_size in powers of 2 + int max_threads_per_blk = device_properties.maxThreadsPerBlock; + for (int block_size = 2; block_size <= max_threads_per_blk; block_size = block_size * 2) { + test_cg_grid_group_type(kernel_func, block_size); + } + + // 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_grid_group_type(kernel_func, max(2, rand() % max_threads_per_blk)); + } +} + +TEST_CASE("Unit_hipCGGridGroupType_DataSharing") { + const auto device = GENERATE(range(0, HipTest::getDeviceCount())); + HIP_CHECK(hipSetDevice(device)); + + hipDeviceProp_t device_properties; + + HIP_CHECK(hipGetDeviceProperties(&device_properties, device)); + + if (!device_properties.cooperativeLaunch) { + HipTest::HIP_SKIP_TEST("Device doesn't support cooperative launch!"); + return; + } + + int loops = GENERATE(1, 2, 3, 4); + int width = GENERATE(512, 1024, 2048, 4096); + + // Launch enough waves to fill up all of the GPU + int warp_size = device_properties.warpSize; + int num_sms = device_properties.multiProcessorCount; + + // Calculate the device occupancy to know how many blocks can be run. + int max_blocks_per_sm; + HIP_CHECK( + hipOccupancyMaxActiveBlocksPerMultiprocessor(&max_blocks_per_sm, coop_kernel, warp_size, 0)); + + int num_blocks = max_blocks_per_sm * num_sms; + + // Create Streams + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + // Allocate and initialize data + + // Alocate the host input buffer, and two device buffers + unsigned int* input_buffer = + reinterpret_cast(malloc(sizeof(unsigned int) * width)); + for (int i = 0; i < width; i++) { + input_buffer[i] = i; + } + + unsigned int *dev_mem_1, *host_mem_1; + host_mem_1 = reinterpret_cast(malloc(sizeof(unsigned int) * width)); + HIP_CHECK(hipMalloc(&dev_mem_1, sizeof(unsigned int) * width)); + HIP_CHECK(hipMemcpyAsync(dev_mem_1, input_buffer, sizeof(unsigned int) * width, + hipMemcpyHostToDevice, stream)); + + unsigned int *dev_mem_2, *host_mem_2; + host_mem_2 = reinterpret_cast(malloc(sizeof(unsigned int) * width)); + HIP_CHECK(hipMalloc(&dev_mem_2, sizeof(unsigned int) * width)); + HIP_CHECK(hipMemsetAsync(dev_mem_2, 0, width * sizeof(unsigned int), stream)); + + // Launch the kernels + INFO("Launching a cooperative kernel with " << num_blocks << " blocks, each with " << warp_size + << " threads"); + + void* coop_params[4]; + coop_params[0] = reinterpret_cast(&dev_mem_1); + coop_params[1] = reinterpret_cast(&dev_mem_2); + coop_params[2] = reinterpret_cast(&loops); + coop_params[3] = reinterpret_cast(&width); + HIP_CHECK(hipLaunchCooperativeKernel(coop_kernel, num_blocks, warp_size, coop_params, 0, stream)); + + // Read back the buffers and print out their data + HIP_CHECK(hipMemcpyAsync(host_mem_1, dev_mem_1, sizeof(unsigned int) * width, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipMemcpyAsync(host_mem_2, dev_mem_2, sizeof(unsigned int) * width, + hipMemcpyDeviceToHost, stream)); + + HIP_CHECK(hipStreamSynchronize(stream)); + + verify_coop_buffers(input_buffer, host_mem_1, host_mem_2, loops, width); + + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipFree(dev_mem_1)); + HIP_CHECK(hipFree(dev_mem_2)); + free(input_buffer); + free(host_mem_1); + free(host_mem_2); +} + +TEST_CASE("Unit_hipCGGridGroupType_Barrier") { + const auto device = GENERATE(range(0, HipTest::getDeviceCount())); + HIP_CHECK(hipSetDevice(device)); + + hipDeviceProp_t device_properties; + + HIP_CHECK(hipGetDeviceProperties(&device_properties, device)); + + if (!device_properties.cooperativeLaunch) { + HipTest::HIP_SKIP_TEST("Device doesn't support cooperative launch!"); + return; + } + + uint32_t loops = GENERATE(1, 2, 3, 4); + uint32_t warps = GENERATE(4, 8, 16, 32); + uint32_t block_size = 1; + + // Test whether the requested size will fit on the GPU + int max_blocks_per_sm; + int warp_size = device_properties.warpSize; + int num_sms = device_properties.multiProcessorCount; + + int num_threads_in_block = block_size * warp_size; + + auto test_kernel_used = IsGfx11() ? test_kernel_gfx11 : test_kernel; + // Calculate the device occupancy to know how many blocks can be run. + HIP_CHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&max_blocks_per_sm, test_kernel_used, + num_threads_in_block, 0)); + + int requested_blocks = warps / block_size; + if (requested_blocks > max_blocks_per_sm * num_sms) { + INFO("Too many blocks requested!"); + REQUIRE(false); + } + + // Each block will output a single value per loop. + uint32_t total_buffer_len = requested_blocks * loops; + + // Alocate the buffer that will hold the kernel's output, and which will + // also be used to globally synchronize during GWS initialization + unsigned int* host_buffer = + reinterpret_cast(calloc(total_buffer_len, sizeof(unsigned int))); + + unsigned int* kernel_buffer; + HIP_CHECK(hipMalloc(&kernel_buffer, sizeof(unsigned int) * total_buffer_len)); + HIP_CHECK(hipMemcpy(kernel_buffer, host_buffer, sizeof(unsigned int) * total_buffer_len, + hipMemcpyHostToDevice)); + + unsigned int* kernel_atomic; + HIP_CHECK(hipMalloc(&kernel_atomic, sizeof(unsigned int))); + HIP_CHECK(hipMemset(kernel_atomic, 0, sizeof(unsigned int))); + + // Launch the kernel + INFO("Launching a cooperative kernel with " << warps << " warps in " << requested_blocks + << " thread blocks"); + + void* params[3]; + params[0] = reinterpret_cast(&kernel_atomic); + params[1] = reinterpret_cast(&kernel_buffer); + params[2] = reinterpret_cast(&loops); + HIP_CHECK(hipLaunchCooperativeKernel(test_kernel_used, requested_blocks, num_threads_in_block, + params, 0, 0)); + + // Read back the buffer to host + HIP_CHECK(hipMemcpy(host_buffer, kernel_buffer, sizeof(unsigned int) * total_buffer_len, + hipMemcpyDeviceToHost)); + + verify_barrier_buffer(loops, requested_blocks, host_buffer); + + HIP_CHECK(hipFree(kernel_buffer)); + HIP_CHECK(hipFree(kernel_atomic)); + free(host_buffer); +} diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/hipCGMultiGridGroupType.cc b/projects/hip-tests/catch/unit/cooperativeGrps/hipCGMultiGridGroupType.cc deleted file mode 100644 index 1dd2a8f3b5..0000000000 --- a/projects/hip-tests/catch/unit/cooperativeGrps/hipCGMultiGridGroupType.cc +++ /dev/null @@ -1,240 +0,0 @@ -/* -Copyright (c) 2020 - 2021 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. -*/ - - -/* HIT_START - * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 -rdc=true -gencode arch=compute_60,code=sm_60 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_80,code=sm_80 - * TEST: %t - * HIT_END - */ - -#include -#include - -#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, - int* gridRankTestD, - int *sizeTestD, - int *thdRankTestD, - int *isValidTestD, - int *syncTestD, - int *syncResultD) -{ - multi_grid_group mg = this_multi_grid(); - int gIdx = (blockIdx.x * blockDim.x) + threadIdx.x; - - // Test num_grids - numGridsTestD[gIdx] = mg.num_grids(); - - // Test grid_rank - gridRankTestD[gIdx] = mg.grid_rank(); - - // Test size - sizeTestD[gIdx] = mg.size(); - - // Test thread_rank - thdRankTestD[gIdx] = mg.thread_rank(); - - // Test is_valid - isValidTestD[gIdx] = mg.is_valid(); - - // Test sync - // - // Eech thread assign 1 to their respective location - syncTestD[gIdx] = 1; - // Grid level sync - 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 (uint i = 1; i < gridDim.x * blockDim.x; ++i) { - syncTestD[0] += syncTestD[i]; - } - syncResultD[mg.grid_rank() + 1] = syncTestD[0]; - } - // multi-grid level sync - mg.sync(); - // 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 (uint i = 1; i <= mg.num_grids(); ++i) { - syncResultD[0] += syncResultD[i]; - } - } -} - -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++) { - HIPCHECK(hipSetDevice(i)); - HIPCHECK(hipDeviceSynchronize()); // Make sure work is done on this device - HIPCHECK(hipStreamCreate(&stream[i])); - } - - // Allocate host and device memory - int nBytes = sizeof(int) * 2 * blockSize; - int *numGridsTestD[MaxGPUs], *numGridsTestH[MaxGPUs]; - int *gridRankTestD[MaxGPUs], *gridRankTestH[MaxGPUs]; - int *sizeTestD[MaxGPUs], *sizeTestH[MaxGPUs]; - int *thdRankTestD[MaxGPUs], *thdRankTestH[MaxGPUs]; - int *isValidTestD[MaxGPUs], *isValidTestH[MaxGPUs]; - int *syncTestD[MaxGPUs], *syncResultD; - for (int i = 0; i < nGpu; i++) { - HIPCHECK(hipSetDevice(i)); - - 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)); - - 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) { - HIPCHECK(hipHostMalloc(&syncResultD, sizeof(int) * (nGpu + 1), hipHostMallocCoherent)); - } - } - - // Launch Kernel - constexpr int NumKernelArgs = 7; - hipLaunchParams* launchParamsList = new hipLaunchParams[nGpu]; - void* args[MaxGPUs * NumKernelArgs]; - for (int i = 0; i < nGpu; i++) { - HIPCHECK(hipSetDevice(i)); - - args[i * NumKernelArgs] = &numGridsTestD[i]; - args[i * NumKernelArgs + 1] = &gridRankTestD[i]; - args[i * NumKernelArgs + 2] = &sizeTestD[i]; - args[i * NumKernelArgs + 3] = &thdRankTestD[i]; - args[i * NumKernelArgs + 4] = &isValidTestD[i]; - args[i * NumKernelArgs + 5] = &syncTestD[i]; - args[i * NumKernelArgs + 6] = &syncResultD; - - launchParamsList[i].func = reinterpret_cast(kernel_cg_multi_grid_group_type); - launchParamsList[i].gridDim = 2; - launchParamsList[i].blockDim = blockSize; - launchParamsList[i].sharedMem = 0; - launchParamsList[i].stream = stream[i]; - launchParamsList[i].args = &args[i * NumKernelArgs]; - } - HIPCHECK(hipLaunchCooperativeKernelMultiDevice(launchParamsList, nGpu, 0)); - - // Copy result from device to host - for (int i = 0; i < nGpu; i++) { - 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 - int gridsSeen[MaxGPUs]; - for (int i = 0; i < nGpu; ++i) { - for (int j = 0; j < 2 * blockSize; ++j) { - ASSERT_EQUAL(numGridsTestH[i][j], nGpu); - ASSERT_GE(gridRankTestH[i][j], 0); - ASSERT_LE(gridRankTestH[i][j], nGpu-1); - ASSERT_EQUAL(gridRankTestH[i][j], gridRankTestH[i][0]); - ASSERT_EQUAL(sizeTestH[i][j], nGpu * 2 * blockSize); - int gridRank = gridRankTestH[i][j]; - ASSERT_EQUAL(thdRankTestH[i][j], (gridRank * 2 * blockSize) + j); - ASSERT_EQUAL(isValidTestH[i][j], 1); - } - ASSERT_EQUAL(syncResultD[i+1], 2 * blockSize); - - // Validate uniqueness property of grid rank - 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_EQUAL(syncResultD[0], nGpu * 2 * blockSize); - - // Free host and device memory - delete [] launchParamsList; - for (int i = 0; i < nGpu; i++) { - HIPCHECK(hipSetDevice(i)); - - 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) { - HIPCHECK(hipHostFree(syncResultD)); - } - HIPCHECK(hipHostFree(numGridsTestH[i])); - HIPCHECK(hipHostFree(gridRankTestH[i])); - HIPCHECK(hipHostFree(sizeTestH[i])); - HIPCHECK(hipHostFree(thdRankTestH[i])); - HIPCHECK(hipHostFree(isValidTestH[i])); - } -} - -TEST_CASE("Unit_hipCGMultiGridGroupType") { - int nGpu = 0; - HIPCHECK(hipGetDeviceCount(&nGpu)); - nGpu = min(nGpu, MaxGPUs); - - // Set `maxThreadsPerBlock` by taking minimum among all available devices - int maxThreadsPerBlock = INT_MAX; - hipDeviceProp_t deviceProperties; - for (int i = 0; i < nGpu; i++) { - HIPCHECK(hipGetDeviceProperties(&deviceProperties, i)); - if (!deviceProperties.cooperativeMultiDeviceLaunch) { - HipTest::HIP_SKIP_TEST("Device doesn't support cooperative launch!"); - return; - } - maxThreadsPerBlock = min(maxThreadsPerBlock, deviceProperties.maxThreadsPerBlock); - } - - // 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 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); - } -} diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/hipCGMultiGridGroupTypeViaBaseType.cc b/projects/hip-tests/catch/unit/cooperativeGrps/hipCGMultiGridGroupTypeViaBaseType.cc deleted file mode 100644 index 408f3b0075..0000000000 --- a/projects/hip-tests/catch/unit/cooperativeGrps/hipCGMultiGridGroupTypeViaBaseType.cc +++ /dev/null @@ -1,234 +0,0 @@ -/* -Copyright (c) 2020 - 2021 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. -*/ - - -/* HIT_START - * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 -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 - * TEST: %t - * HIT_END - */ - -#include -#include -#include -#include -#include - -#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, - int* gridRankTestD, - int *thdRankTestD, - int *isValidTestD, - int *syncTestD, - int *syncResultD) -{ - thread_group tg = this_multi_grid(); // This can work if _CG_ABI_EXPERIMENTAL defined on Cuda - - int gIdx = (blockIdx.x * blockDim.x) + threadIdx.x; - - // Test size - sizeTestD[gIdx] = tg.size(); - - // Test thread_rank - gridRankTestD[gIdx] = this_multi_grid().grid_rank(); - thdRankTestD[gIdx] = tg.thread_rank(); - - // Test is_valid -#ifdef __HIP_PLATFORM_AMD__ - isValidTestD[gIdx] = tg.is_valid(); -#else - // Cuda has no thread_group.is_valid() - isValidTestD[gIdx] = true; -#endif - // Test sync - // - // Eech thread assign 1 to their respective location - syncTestD[gIdx] = 1; - // Grid level sync - 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 (uint i = 1; i < gridDim.x * blockDim.x; ++i) { - syncTestD[0] += syncTestD[i]; - } - syncResultD[this_multi_grid().grid_rank() + 1] = syncTestD[0]; - } - // multi-grid level sync - tg.sync(); - // 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 (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, int nGpu) -{ - // Create a stream each device - hipStream_t stream[MaxGPUs]; - for (int i = 0; i < nGpu; i++) { - HIPCHECK(hipSetDevice(i)); - HIPCHECK(hipDeviceSynchronize()); // Make sure work is done on this device - HIPCHECK(hipStreamCreate(&stream[i])); - } - - // Allocate host and device memory - int nBytes = sizeof(int) * 2 * blockSize; - int *sizeTestD[MaxGPUs], *sizeTestH[MaxGPUs]; - int *gridRankTestD[MaxGPUs], *gridRankTestH[MaxGPUs]; - int *thdRankTestD[MaxGPUs], *thdRankTestH[MaxGPUs]; - int *isValidTestD[MaxGPUs], *isValidTestH[MaxGPUs]; - int *syncTestD[MaxGPUs], *syncResultD; - for (int i = 0; i < nGpu; i++) { - HIPCHECK(hipSetDevice(i)); - - 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)); - - HIPCHECK(hipHostMalloc(&sizeTestH[i], nBytes)); - HIPCHECK(hipHostMalloc(&gridRankTestH[i], nBytes)); - HIPCHECK(hipHostMalloc(&thdRankTestH[i], nBytes)); - HIPCHECK(hipHostMalloc(&isValidTestH[i], nBytes)); - - if (i == 0) { - HIPCHECK(hipHostMalloc(&syncResultD, sizeof(int) * (nGpu + 1), hipHostMallocCoherent)); - } - } - - // Launch Kernel - constexpr int NumKernelArgs = 6; - hipLaunchParams* launchParamsList = new hipLaunchParams[nGpu]; - void* args[MaxGPUs * NumKernelArgs]; - for (int i = 0; i < nGpu; i++) { - HIPCHECK(hipSetDevice(i)); - - args[i * NumKernelArgs ] = &sizeTestD[i]; - args[i * NumKernelArgs + 1] = &gridRankTestD[i]; - args[i * NumKernelArgs + 2] = &thdRankTestD[i]; - args[i * NumKernelArgs + 3] = &isValidTestD[i]; - args[i * NumKernelArgs + 4] = &syncTestD[i]; - args[i * NumKernelArgs + 5] = &syncResultD; - - launchParamsList[i].func = reinterpret_cast(kernel_cg_multi_grid_group_type_via_base_type); - launchParamsList[i].gridDim = 2; - launchParamsList[i].blockDim = blockSize; - launchParamsList[i].sharedMem = 0; - launchParamsList[i].stream = stream[i]; - launchParamsList[i].args = &args[i * NumKernelArgs]; - } - HIPCHECK(hipLaunchCooperativeKernelMultiDevice(launchParamsList, nGpu, 0)); - - // Copy result from device to host - for (int i = 0; i < nGpu; i++) { - 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 - int gridsSeen[MaxGPUs]; - for (int i = 0; i < nGpu; ++i) { - for (int j = 0; j < 2 * blockSize; ++j) { - ASSERT_EQUAL(sizeTestH[i][j], nGpu * 2 * blockSize); - ASSERT_GE(gridRankTestH[i][j], 0); - ASSERT_LE(gridRankTestH[i][j], nGpu-1); - ASSERT_EQUAL(gridRankTestH[i][j], gridRankTestH[i][0]); - int gridRank = gridRankTestH[i][j]; - ASSERT_EQUAL(thdRankTestH[i][j], (gridRank * 2 * blockSize) + j); - ASSERT_EQUAL(isValidTestH[i][j], 1); - } - ASSERT_EQUAL(syncResultD[i+1], 2 * blockSize); - - // Validate uniqueness property of grid rank - 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_EQUAL(syncResultD[0], nGpu * 2 * blockSize); - - // Free host and device memory - delete [] launchParamsList; - for (int i = 0; i < nGpu; i++) { - HIPCHECK(hipSetDevice(i)); - - HIPCHECK(hipFree(sizeTestD[i])); - HIPCHECK(hipFree(gridRankTestD[i])); - HIPCHECK(hipFree(thdRankTestD[i])); - HIPCHECK(hipFree(isValidTestD[i])); - HIPCHECK(hipFree(syncTestD[i])); - - if (i == 0) - HIPCHECK(hipHostFree(syncResultD)); - - HIPCHECK(hipHostFree(sizeTestH[i])); - HIPCHECK(hipHostFree(gridRankTestH[i])); - HIPCHECK(hipHostFree(thdRankTestH[i])); - HIPCHECK(hipHostFree(isValidTestH[i])); - } -} - -TEST_CASE("Unit_hipCGMultiGridGroupType_BaseType") { - // Set `maxThreadsPerBlock` by taking minimum among all available devices - int nGpu = 0; - HIPCHECK(hipGetDeviceCount(&nGpu)); - nGpu = min(nGpu, MaxGPUs); - - int maxThreadsPerBlock = INT_MAX; - hipDeviceProp_t deviceProperties; - for (int i = 0; i < nGpu; i++) { - HIPCHECK(hipGetDeviceProperties(&deviceProperties, i)); - if (!deviceProperties.cooperativeMultiDeviceLaunch) { - HipTest::HIP_SKIP_TEST("Device doesn't support cooperative launch!"); - return; - } - maxThreadsPerBlock = min(maxThreadsPerBlock, deviceProperties.maxThreadsPerBlock); - } - - // 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 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); - } -} diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/hipCGMultiGridGroupTypeViaPublicApi.cc b/projects/hip-tests/catch/unit/cooperativeGrps/hipCGMultiGridGroupTypeViaPublicApi.cc deleted file mode 100644 index 3e5b97fe5a..0000000000 --- a/projects/hip-tests/catch/unit/cooperativeGrps/hipCGMultiGridGroupTypeViaPublicApi.cc +++ /dev/null @@ -1,230 +0,0 @@ -/* -Copyright (c) 2020 - 2021 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. -*/ - - -/* HIT_START - * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 -rdc=true -gencode arch=compute_60,code=sm_60 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_80,code=sm_80 - * TEST: %t - * HIT_END - */ - -#include -#include -#include -#include -#include - -#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, - int* gridRankTestD, - int *thdRankTestD, - int *isValidTestD, - int *syncTestD, - int *syncResultD) -{ - multi_grid_group mg = this_multi_grid(); - int gIdx = (blockIdx.x * blockDim.x) + threadIdx.x; - - // Test group_size api - sizeTestD[gIdx] = group_size(mg); - - // Test thread_rank api - gridRankTestD[gIdx] = this_multi_grid().grid_rank(); - thdRankTestD[gIdx] = thread_rank(mg); - - // Test is_valid api - isValidTestD[gIdx] = mg.is_valid(); - - // Test sync api - // - // Eech thread assign 1 to their respective location - syncTestD[gIdx] = 1; - // Grid level sync - 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 (uint i = 1; i < gridDim.x * blockDim.x; ++i) { - syncTestD[0] += syncTestD[i]; - } - syncResultD[this_multi_grid().grid_rank() + 1] = syncTestD[0]; - } - // multi-grid level sync via public api - sync(mg); - // 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 (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, int nGpu) -{ - // Create a stream each device - hipStream_t stream[MaxGPUs]; - for (int i = 0; i < nGpu; i++) { - HIPCHECK(hipSetDevice(i)); - HIPCHECK(hipDeviceSynchronize()); // Make sure work is done on this device - HIPCHECK(hipStreamCreate(&stream[i])); - } - - // Allocate host and device memory - int nBytes = sizeof(int) * 2 * blockSize; - int *sizeTestD[MaxGPUs], *sizeTestH[MaxGPUs]; - int *gridRankTestD[MaxGPUs], *gridRankTestH[MaxGPUs]; - int *thdRankTestD[MaxGPUs], *thdRankTestH[MaxGPUs]; - int *isValidTestD[MaxGPUs], *isValidTestH[MaxGPUs]; - int *syncTestD[MaxGPUs], *syncResultD; - for (int i = 0; i < nGpu; i++) { - HIPCHECK(hipSetDevice(i)); - - 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)); - - HIPCHECK(hipHostMalloc(&sizeTestH[i], nBytes)); - HIPCHECK(hipHostMalloc(&gridRankTestH[i], nBytes)); - HIPCHECK(hipHostMalloc(&thdRankTestH[i], nBytes)); - HIPCHECK(hipHostMalloc(&isValidTestH[i], nBytes)); - - if (i == 0) { - HIPCHECK(hipHostMalloc(&syncResultD, sizeof(int) * (nGpu + 1), hipHostMallocCoherent)); - } - } - - // Launch Kernel - constexpr int NumKernelArgs = 6; - hipLaunchParams* launchParamsList = new hipLaunchParams[nGpu]; - void* args[MaxGPUs * NumKernelArgs]; - for (int i = 0; i < nGpu; i++) { - HIPCHECK(hipSetDevice(i)); - - args[i * NumKernelArgs ] = &sizeTestD[i]; - args[i * NumKernelArgs + 1] = &gridRankTestD[i]; - args[i * NumKernelArgs + 2] = &thdRankTestD[i]; - args[i * NumKernelArgs + 3] = &isValidTestD[i]; - args[i * NumKernelArgs + 4] = &syncTestD[i]; - args[i * NumKernelArgs + 5] = &syncResultD; - - launchParamsList[i].func = reinterpret_cast(kernel_cg_multi_grid_group_type_via_public_api); - launchParamsList[i].gridDim = 2; - launchParamsList[i].blockDim = blockSize; - launchParamsList[i].sharedMem = 0; - launchParamsList[i].stream = stream[i]; - launchParamsList[i].args = &args[i * NumKernelArgs]; - } - HIPCHECK(hipLaunchCooperativeKernelMultiDevice(launchParamsList, nGpu, 0)); - - // Copy result from device to host - for (int i = 0; i < nGpu; i++) { - 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 - int gridsSeen[MaxGPUs]; - for (int i = 0; i < nGpu; ++i) { - for (int j = 0; j < 2 * blockSize; ++j) { - ASSERT_EQUAL(sizeTestH[i][j], nGpu * 2 * blockSize); - ASSERT_GE(gridRankTestH[i][j], 0); - ASSERT_LE(gridRankTestH[i][j], nGpu-1); - ASSERT_EQUAL(gridRankTestH[i][j], gridRankTestH[i][0]); - int gridRank = gridRankTestH[i][j]; - ASSERT_EQUAL(thdRankTestH[i][j], (gridRank * 2 * blockSize) + j); - ASSERT_EQUAL(isValidTestH[i][j], 1); - } - ASSERT_EQUAL(syncResultD[i+1], 2 * blockSize); - - // Validate uniqueness property of grid rank - 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_EQUAL(syncResultD[0], nGpu * 2 * blockSize); - - // Free host and device memory - delete [] launchParamsList; - for (int i = 0; i < nGpu; i++) { - HIPCHECK(hipSetDevice(i)); - - HIPCHECK(hipFree(sizeTestD[i])); - HIPCHECK(hipFree(gridRankTestD[i])); - HIPCHECK(hipFree(thdRankTestD[i])); - HIPCHECK(hipFree(isValidTestD[i])); - HIPCHECK(hipFree(syncTestD[i])); - - if (i == 0) - HIPCHECK(hipHostFree(syncResultD)); - - HIPCHECK(hipHostFree(sizeTestH[i])); - HIPCHECK(hipHostFree(gridRankTestH[i])); - HIPCHECK(hipHostFree(thdRankTestH[i])); - HIPCHECK(hipHostFree(isValidTestH[i])); - } -} - -TEST_CASE("Unit_hipCGMultiGridGroupType_PublicApi") { - // Set `maxThreadsPerBlock` by taking minimum among all available devices - int nGpu = 0; - HIPCHECK(hipGetDeviceCount(&nGpu)); - nGpu = min(nGpu, MaxGPUs); - - int maxThreadsPerBlock = INT_MAX; - hipDeviceProp_t deviceProperties; - for (int i = 0; i < nGpu; i++) { - HIPCHECK(hipGetDeviceProperties(&deviceProperties, i)); - if (!deviceProperties.cooperativeMultiDeviceLaunch) { - HipTest::HIP_SKIP_TEST("Device doesn't support cooperative launch!"); - return; - } - maxThreadsPerBlock = min(maxThreadsPerBlock, deviceProperties.maxThreadsPerBlock); - } - - // 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 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); - } -} diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/hipCGMultiGridGroupType_old.cc b/projects/hip-tests/catch/unit/cooperativeGrps/hipCGMultiGridGroupType_old.cc new file mode 100644 index 0000000000..4d980380ce --- /dev/null +++ b/projects/hip-tests/catch/unit/cooperativeGrps/hipCGMultiGridGroupType_old.cc @@ -0,0 +1,638 @@ +/* +Copyright (c) 2020 - 2022 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 +#include + +#include "hip_cg_common.hh" + +namespace cg = cooperative_groups; + +static __global__ void kernel_cg_multi_grid_group_type(int* grid_rank_dev, int* size_dev, + int* thd_rank_dev, int* is_valid_dev, + int* sync_dev, int* sync_result, + int* num_grids_dev) { + cg::multi_grid_group mg = cg::this_multi_grid(); + int gIdx = (blockIdx.x * blockDim.x) + threadIdx.x; + + // Test num_grids + num_grids_dev[gIdx] = mg.num_grids(); + + // Test grid_rank + grid_rank_dev[gIdx] = mg.grid_rank(); + + // Test size + size_dev[gIdx] = mg.size(); + + // Test thread_rank + thd_rank_dev[gIdx] = mg.thread_rank(); + + // Test is_valid + is_valid_dev[gIdx] = mg.is_valid(); + + // Test sync + // + // Eech thread assign 1 to their respective location + sync_dev[gIdx] = 1; + // Grid level sync + cg::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 (uint i = 1; i < gridDim.x * blockDim.x; ++i) { + sync_dev[0] += sync_dev[i]; + } + sync_result[mg.grid_rank() + 1] = sync_dev[0]; + } + // multi-grid level sync + mg.sync(); + // grid (gpu) 0 does final reduction across all grids (gpus) + if (mg.grid_rank() == 0 && blockIdx.x == 0 && threadIdx.x == 0) { + sync_result[0] = 0; + for (uint i = 1; i <= mg.num_grids(); ++i) { + sync_result[0] += sync_result[i]; + } + } +} + +static __global__ void kernel_cg_multi_grid_group_type_via_base_type( + int* grid_rank_dev, int* size_dev, int* thd_rank_dev, int* is_valid_dev, int* sync_dev, + int* sync_result) { + cg::thread_group tg = + cg::this_multi_grid(); // This can work if _CG_ABI_EXPERIMENTAL defined on Cuda + + int gIdx = (blockIdx.x * blockDim.x) + threadIdx.x; + + // Test size + size_dev[gIdx] = tg.size(); + + // Test thread_rank + grid_rank_dev[gIdx] = cg::this_multi_grid().grid_rank(); + thd_rank_dev[gIdx] = tg.thread_rank(); + + // Test is_valid +#ifdef __HIP_PLATFORM_AMD__ + is_valid_dev[gIdx] = tg.is_valid(); +#else + // Cuda has no thread_group.is_valid() + is_valid_dev[gIdx] = true; +#endif + // Test sync + // + // Eech thread assign 1 to their respective location + sync_dev[gIdx] = 1; + // Grid level sync + cg::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 (uint i = 1; i < gridDim.x * blockDim.x; ++i) { + sync_dev[0] += sync_dev[i]; + } + sync_result[cg::this_multi_grid().grid_rank() + 1] = sync_dev[0]; + } + // multi-grid level sync + tg.sync(); + // grid (gpu) 0 does final reduction across all grids (gpus) + if (cg::this_multi_grid().grid_rank() == 0 && blockIdx.x == 0 && threadIdx.x == 0) { + sync_result[0] = 0; + for (uint i = 1; i <= cg::this_multi_grid().num_grids(); ++i) { + sync_result[0] += sync_result[i]; + } + } +} + +static __global__ void kernel_cg_multi_grid_group_type_via_public_api( + int* grid_rank_dev, int* size_dev, int* thd_rank_dev, int* is_valid_dev, int* sync_dev, + int* sync_result) { + cg::multi_grid_group mg = cg::this_multi_grid(); + int gIdx = (blockIdx.x * blockDim.x) + threadIdx.x; + + // Test group_size api + size_dev[gIdx] = cg::group_size(mg); + + // Test thread_rank api + grid_rank_dev[gIdx] = cg::this_multi_grid().grid_rank(); + thd_rank_dev[gIdx] = cg::thread_rank(mg); + + // Test is_valid api + is_valid_dev[gIdx] = mg.is_valid(); + + // Test sync api + // + // Eech thread assign 1 to their respective location + sync_dev[gIdx] = 1; + // Grid level sync + cg::sync(cg::this_grid()); + // Thread 0 from work-group 0 of current grid (gpu) does grid level reduction + if (blockIdx.x == 0 && threadIdx.x == 0) { + for (uint i = 1; i < gridDim.x * blockDim.x; ++i) { + sync_dev[0] += sync_dev[i]; + } + sync_result[cg::this_multi_grid().grid_rank() + 1] = sync_dev[0]; + } + // multi-grid level sync via public api + cg::sync(mg); + // grid (gpu) 0 does final reduction across all grids (gpus) + if (cg::this_multi_grid().grid_rank() == 0 && blockIdx.x == 0 && threadIdx.x == 0) { + sync_result[0] = 0; + for (uint i = 1; i <= cg::this_multi_grid().num_grids(); ++i) { + sync_result[0] += sync_result[i]; + } + } +} + +static __global__ void test_kernel(unsigned int* atomic_val, unsigned int* global_array, + unsigned int* array, uint32_t loops) { + cg::grid_group grid = cg::this_grid(); + cg::multi_grid_group mgrid = cg::this_multi_grid(); + unsigned rank = grid.thread_rank(); + unsigned global_rank = mgrid.thread_rank(); + + int offset = blockIdx.x; + for (int i = 0; i < loops; i++) { + // Make the last thread run way behind everyone else. + // If the grid barrier below fails, then the other threads may hit the + // atomicInc instruction many times before the last thread ever gets + // to it. + // As such, without the barrier, the last array entry will eventually + // contain a very large value, defined by however many times the other + // wavefronts make it through this loop. + // If the barrier works, then it will likely contain some number + // near "total number of blocks". It will be the last wavefront to + // reach the atomicInc, but everyone will have only hit the atomic once. + if (rank == (grid.size() - 1)) { + long long time_diff = 0; + long long last_clock = clock64(); + do { + long long cur_clock = clock64(); + if (cur_clock > last_clock) { + time_diff += (cur_clock - last_clock); + } + // If it rolls over, we don't know how much to add to catch up. + // So just ignore those slipped cycles. + last_clock = cur_clock; + } while (time_diff < 1000000); + } + if (threadIdx.x == 0) { + array[offset] = atomicInc(atomic_val, UINT_MAX); + } + grid.sync(); + + // Make the last thread in the entire multi-grid run way behind + // everyone else. + // If the mgrid barrier below fails, then the two global_array entries + // will end up being out of sync, because the intermingling of adds + // and multiplies will not be aligned between to the two GPUs. + if (global_rank == (mgrid.size() - 1)) { + long long time_diff = 0; + long long last_clock = clock64(); + do { + long long cur_clock = clock64(); + if (cur_clock > last_clock) { + time_diff += (cur_clock - last_clock); + } + // If it rolls over, we don't know how much to add to catch up. + // So just ignore those slipped cycles. + last_clock = cur_clock; + } while (time_diff < 1000000); + } + // During even iterations, add into your own array entry + // During odd iterations, add into your partner's array entry + unsigned grid_rank = mgrid.grid_rank(); + unsigned inter_gpu_offset = (grid_rank + i) % mgrid.num_grids(); + if (rank == (grid.size() - 1)) { + if (i % mgrid.num_grids() == 0) { + global_array[grid_rank] += 2; + } else { + global_array[inter_gpu_offset] *= 2; + } + } + mgrid.sync(); + offset += gridDim.x; + } +} + +__global__ void test_kernel_gfx11(unsigned int* atomic_val, unsigned int* global_array, + unsigned int* array, uint32_t loops) { +#if HT_AMD + cg::grid_group grid = cg::this_grid(); + cg::multi_grid_group mgrid = cg::this_multi_grid(); + unsigned rank = grid.thread_rank(); + unsigned global_rank = mgrid.thread_rank(); + + int offset = blockIdx.x; + for (int i = 0; i < loops; i++) { + // Make the last thread run way behind everyone else. + // If the grid barrier below fails, then the other threads may hit the + // atomicInc instruction many times before the last thread ever gets + // to it. + // As such, without the barrier, the last array entry will eventually + // contain a very large value, defined by however many times the other + // wavefronts make it through this loop. + // If the barrier works, then it will likely contain some number + // near "total number of blocks". It will be the last wavefront to + // reach the atomicInc, but everyone will have only hit the atomic once. + if (rank == (grid.size() - 1)) { + long long time_diff = 0; + long long last_clock = wall_clock64(); + do { + long long cur_clock = wall_clock64(); + if (cur_clock > last_clock) { + time_diff += (cur_clock - last_clock); + } + // If it rolls over, we don't know how much to add to catch up. + // So just ignore those slipped cycles. + last_clock = cur_clock; + } while (time_diff < 1000000); + } + if (threadIdx.x == 0) { + array[offset] = atomicInc(atomic_val, UINT_MAX); + } + grid.sync(); + + // Make the last thread in the entire multi-grid run way behind + // everyone else. + // If the mgrid barrier below fails, then the two global_array entries + // will end up being out of sync, because the intermingling of adds + // and multiplies will not be aligned between to the two GPUs. + if (global_rank == (mgrid.size() - 1)) { + long long time_diff = 0; + long long last_clock = wall_clock64(); + do { + long long cur_clock = wall_clock64(); + if (cur_clock > last_clock) { + time_diff += (cur_clock - last_clock); + } + // If it rolls over, we don't know how much to add to catch up. + // So just ignore those slipped cycles. + last_clock = cur_clock; + } while (time_diff < 1000000); + } + // During even iterations, add into your own array entry + // During odd iterations, add into your partner's array entry + unsigned grid_rank = mgrid.grid_rank(); + unsigned inter_gpu_offset = (grid_rank + i) % mgrid.num_grids(); + if (rank == (grid.size() - 1)) { + if (i % mgrid.num_grids() == 0) { + global_array[grid_rank] += 2; + } else { + global_array[inter_gpu_offset] *= 2; + } + } + mgrid.sync(); + offset += gridDim.x; + } +#endif +} + +static void verify_barrier_buffer(unsigned int loops, unsigned int warps, unsigned int* host_buffer, + unsigned int num_devs) { + unsigned int max_in_this_loop = 0; + for (unsigned int i = 0; i < loops; i++) { + max_in_this_loop += (warps * num_devs); + for (unsigned int j = 0; j < warps; j++) { + REQUIRE(host_buffer[i * warps + j] <= max_in_this_loop); + } + } +} + +static void verify_multi_gpu_buffer(unsigned int loops, unsigned int array_val) { + unsigned int desired_val = 0; + for (int i = 0; i < loops; i++) { + if (i % 2 == 0) { + desired_val += 2; + } else { + desired_val *= 2; + } + } + + REQUIRE(array_val == desired_val); +} + +template +static void test_cg_multi_grid_group_type(F kernel_func, int num_devices, int block_size, + bool specific_api_test) { + // Create a stream each device + hipStream_t stream[MaxGPUs]; + for (int i = 0; i < num_devices; i++) { + HIP_CHECK(hipSetDevice(i)); + HIP_CHECK(hipDeviceSynchronize()); // Make sure work is done on this device + HIP_CHECK(hipStreamCreate(&stream[i])); + } + + // Allocate host and device memory + int num_bytes = sizeof(int) * 2 * block_size; + int *num_grids_dev[MaxGPUs], *num_grids_host[MaxGPUs]; + int *grid_rank_dev[MaxGPUs], *grid_rank_host[MaxGPUs]; + int *size_dev[MaxGPUs], *size_host[MaxGPUs]; + int *thd_rank_dev[MaxGPUs], *thd_rank_host[MaxGPUs]; + int *is_valid_dev[MaxGPUs], *is_valid_host[MaxGPUs]; + int *sync_dev[MaxGPUs], *sync_result; + for (int i = 0; i < num_devices; i++) { + HIP_CHECK(hipSetDevice(i)); + + if (specific_api_test) { + HIP_CHECK(hipMalloc(&num_grids_dev[i], num_bytes)); + HIP_CHECK(hipHostMalloc(&num_grids_host[i], num_bytes)); + } + + HIP_CHECK(hipMalloc(&grid_rank_dev[i], num_bytes)); + HIP_CHECK(hipMalloc(&size_dev[i], num_bytes)); + HIP_CHECK(hipMalloc(&thd_rank_dev[i], num_bytes)); + HIP_CHECK(hipMalloc(&is_valid_dev[i], num_bytes)); + HIP_CHECK(hipMalloc(&sync_dev[i], num_bytes)); + + HIP_CHECK(hipHostMalloc(&grid_rank_host[i], num_bytes)); + HIP_CHECK(hipHostMalloc(&size_host[i], num_bytes)); + HIP_CHECK(hipHostMalloc(&thd_rank_host[i], num_bytes)); + HIP_CHECK(hipHostMalloc(&is_valid_host[i], num_bytes)); + + if (i == 0) { + HIP_CHECK( + hipHostMalloc(&sync_result, sizeof(int) * (num_devices + 1), hipHostMallocCoherent)); + } + } + + // Launch Kernel + int NumKernelArgs = 6; + if (specific_api_test) { + NumKernelArgs = 7; + } + hipLaunchParams* launchParamsList = new hipLaunchParams[num_devices]; + void* args[MaxGPUs * NumKernelArgs]; + for (int i = 0; i < num_devices; i++) { + HIP_CHECK(hipSetDevice(i)); + + args[i * NumKernelArgs] = &grid_rank_dev[i]; + args[i * NumKernelArgs + 1] = &size_dev[i]; + args[i * NumKernelArgs + 2] = &thd_rank_dev[i]; + args[i * NumKernelArgs + 3] = &is_valid_dev[i]; + args[i * NumKernelArgs + 4] = &sync_dev[i]; + args[i * NumKernelArgs + 5] = &sync_result; + if (specific_api_test) { + args[i * NumKernelArgs + 6] = &num_grids_dev[i]; + } + + launchParamsList[i].func = reinterpret_cast(kernel_func); + launchParamsList[i].gridDim = 2; + launchParamsList[i].blockDim = block_size; + launchParamsList[i].sharedMem = 0; + launchParamsList[i].stream = stream[i]; + launchParamsList[i].args = &args[i * NumKernelArgs]; + } + HIP_CHECK(hipLaunchCooperativeKernelMultiDevice(launchParamsList, num_devices, 0)); + + // Copy result from device to host + for (int i = 0; i < num_devices; i++) { + HIP_CHECK(hipSetDevice(i)); + if (specific_api_test) { + HIP_CHECK(hipMemcpy(num_grids_host[i], num_grids_dev[i], num_bytes, hipMemcpyDeviceToHost)); + } + HIP_CHECK(hipMemcpy(grid_rank_host[i], grid_rank_dev[i], num_bytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(size_host[i], size_dev[i], num_bytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(thd_rank_host[i], thd_rank_dev[i], num_bytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(is_valid_host[i], is_valid_dev[i], num_bytes, hipMemcpyDeviceToHost)); + } + + // Validate results + int grids_seen[MaxGPUs]; + for (int i = 0; i < num_devices; ++i) { + for (int j = 0; j < 2 * block_size; ++j) { + if (specific_api_test) { + ASSERT_EQUAL(num_grids_host[i][j], num_devices); + } + ASSERT_GE(grid_rank_host[i][j], 0); + ASSERT_LE(grid_rank_host[i][j], num_devices - 1); + ASSERT_EQUAL(grid_rank_host[i][j], grid_rank_host[i][0]); + ASSERT_EQUAL(size_host[i][j], num_devices * 2 * block_size); + int gridRank = grid_rank_host[i][j]; + ASSERT_EQUAL(thd_rank_host[i][j], (gridRank * 2 * block_size) + j); + ASSERT_EQUAL(is_valid_host[i][j], 1); + } + ASSERT_EQUAL(sync_result[i + 1], 2 * block_size); + + // Validate uniqueness property of grid rank + grids_seen[i] = grid_rank_host[i][0]; + for (int k = 0; k < i; ++k) { + INFO("Grid rank in multi-gpu setup should be unique"); + REQUIRE(grids_seen[k] != grids_seen[i]); + } + } + ASSERT_EQUAL(sync_result[0], num_devices * 2 * block_size); + + // Free host and device memory + delete[] launchParamsList; + for (int i = 0; i < num_devices; i++) { + HIP_CHECK(hipSetDevice(i)); + + if (specific_api_test) { + HIP_CHECK(hipFree(num_grids_dev[i])); + HIP_CHECK(hipHostFree(num_grids_host[i])); + } + + HIP_CHECK(hipFree(grid_rank_dev[i])); + HIP_CHECK(hipFree(size_dev[i])); + HIP_CHECK(hipFree(thd_rank_dev[i])); + HIP_CHECK(hipFree(is_valid_dev[i])); + HIP_CHECK(hipFree(sync_dev[i])); + + if (i == 0) { + HIP_CHECK(hipHostFree(sync_result)); + } + HIP_CHECK(hipHostFree(grid_rank_host[i])); + HIP_CHECK(hipHostFree(size_host[i])); + HIP_CHECK(hipHostFree(thd_rank_host[i])); + HIP_CHECK(hipHostFree(is_valid_host[i])); + } +} + +TEST_CASE("Unit_hipCGMultiGridGroupType_Basic") { + int num_devices = 0; + HIP_CHECK(hipGetDeviceCount(&num_devices)); + num_devices = min(num_devices, MaxGPUs); + + // Set `max_threads_per_blk` by taking minimum among all available devices + int max_threads_per_blk = INT_MAX; + hipDeviceProp_t device_properties; + for (int i = 0; i < num_devices; i++) { + HIP_CHECK(hipGetDeviceProperties(&device_properties, i)); + if (!device_properties.cooperativeMultiDeviceLaunch) { + HipTest::HIP_SKIP_TEST("Device doesn't support cooperative launch!"); + return; + } + max_threads_per_blk = min(max_threads_per_blk, device_properties.maxThreadsPerBlock); + } + + void* (*kernel_func)(void); + bool specific_api_test = false; + + SECTION("Default multi grid group API test") { + kernel_func = reinterpret_cast(kernel_cg_multi_grid_group_type); + specific_api_test = true; + } + + SECTION("Base type multi grid group API test") { + kernel_func = reinterpret_cast(kernel_cg_multi_grid_group_type_via_base_type); + } + + SECTION("Public API multi grid group test") { + kernel_func = reinterpret_cast(kernel_cg_multi_grid_group_type_via_public_api); + } + + // Test for blockSizes in powers of 2 + for (int block_size = 2; block_size <= max_threads_per_blk; block_size = block_size * 2) { + test_cg_multi_grid_group_type(kernel_func, num_devices, block_size, specific_api_test); + } + + // 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(kernel_func, num_devices, max(2, rand() % max_threads_per_blk), + specific_api_test); + } +} + +TEST_CASE("Unit_hipCGMultiGridGroupType_Barrier") { + int num_devices = 0; + uint32_t loops = GENERATE(1, 2, 3, 4); + uint32_t warps = GENERATE(4, 8, 16, 32); + uint32_t block_size = 1; + + HIP_CHECK(hipGetDeviceCount(&num_devices)); + if (num_devices < 2) { + HipTest::HIP_SKIP_TEST("Device number is < 2"); + return; + } + + hipDeviceProp_t device_properties[num_devices]; + for (int i = 0; i < num_devices; i++) { + HIP_CHECK(hipGetDeviceProperties(&device_properties[i], i)); + if (!device_properties[i].cooperativeMultiDeviceLaunch) { + HipTest::HIP_SKIP_TEST("Device doesn't support cooperative launch!"); + return; + } + } + + // Test whether the requested size will fit on the GPU + int warp_sizes[num_devices]; + int num_sms[num_devices]; + int warp_size = INT_MAX; + int num_sm = INT_MAX; + for (int i = 0; i < num_devices; i++) { + warp_sizes[i] = device_properties[i].warpSize; + if (warp_sizes[i] < warp_size) { + warp_size = warp_sizes[i]; + } + num_sms[i] = device_properties[i].multiProcessorCount; + if (num_sms[i] < num_sm) { + num_sm = num_sms[i]; + } + } + + int num_threads_in_block = block_size * warp_size; + + // Calculate the device occupancy to know how many blocks can be run. + int max_blocks_per_sm_arr[num_devices]; + int max_blocks_per_sm = INT_MAX; + for (int i = 0; i < num_devices; i++) { + HIP_CHECK(hipSetDevice(i)); + auto test_kernel_used = IsGfx11() ? test_kernel_gfx11 : test_kernel; + HIP_CHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor( + &max_blocks_per_sm_arr[i], test_kernel_used, num_threads_in_block, 0)); + if (max_blocks_per_sm_arr[i] < max_blocks_per_sm) { + max_blocks_per_sm = max_blocks_per_sm_arr[i]; + } + } + + int requested_blocks = warps / block_size; + + // Each block will output a single value per loop. + uint32_t total_buffer_len = requested_blocks * loops; + + // Alocate the buffer that will hold the kernel's output, and which will + // also be used to globally synchronize during GWS initialization + unsigned int* host_buffer[num_devices]; + unsigned int* kernel_buffer[num_devices]; + unsigned int* kernel_atomic[num_devices]; + hipStream_t streams[num_devices]; + for (int i = 0; i < num_devices; i++) { + host_buffer[i] = + reinterpret_cast(calloc(total_buffer_len, sizeof(unsigned int))); + HIP_CHECK(hipSetDevice(i)); + HIP_CHECK(hipMalloc(&kernel_buffer[i], sizeof(unsigned int) * total_buffer_len)); + HIP_CHECK(hipMemcpy(kernel_buffer[i], host_buffer[i], sizeof(unsigned int) * total_buffer_len, + hipMemcpyHostToDevice)); + HIP_CHECK(hipMalloc(&kernel_atomic[i], sizeof(unsigned int))); + HIP_CHECK(hipMemset(kernel_atomic[i], 0, sizeof(unsigned int))); + HIP_CHECK(hipStreamCreate(&streams[i])); + } + + // Single kernel atomic shared between both devices; put it on the host + unsigned int* global_array; + HIP_CHECK(hipHostMalloc(&global_array, sizeof(unsigned int) * num_devices)); + HIP_CHECK(hipMemset(global_array, 0, num_devices * sizeof(unsigned int))); + + // Launch the kernels + INFO("Launching a cooperative kernel with " << warps << " warps in " << requested_blocks + << " thread blocks"); + + void* dev_params[num_devices][4]; + hipLaunchParams md_params[num_devices]; + for (int i = 0; i < num_devices; i++) { + HIP_CHECK(hipSetDevice(i)); + auto test_kernel_used = IsGfx11() ? test_kernel_gfx11 : test_kernel; + dev_params[i][0] = reinterpret_cast(&kernel_atomic[i]); + dev_params[i][1] = reinterpret_cast(&global_array); + dev_params[i][2] = reinterpret_cast(&kernel_buffer[i]); + dev_params[i][3] = reinterpret_cast(&loops); + md_params[i].func = reinterpret_cast(test_kernel_used); + md_params[i].gridDim = requested_blocks; + md_params[i].blockDim = num_threads_in_block; + md_params[i].sharedMem = 0; + md_params[i].stream = streams[i]; + md_params[i].args = dev_params[i]; + } + + HIP_CHECK(hipLaunchCooperativeKernelMultiDevice(md_params, num_devices, 0)); + HIP_CHECK(hipDeviceSynchronize()); + + // Read back the buffer to host + for (int dev = 0; dev < num_devices; dev++) { + HIP_CHECK(hipMemcpy(host_buffer[dev], kernel_buffer[dev], + sizeof(unsigned int) * total_buffer_len, hipMemcpyDeviceToHost)); + } + + for (unsigned int dev = 0; dev < num_devices; dev++) { + verify_barrier_buffer(loops, requested_blocks, host_buffer[dev], num_devices); + } + + for (int dev = 0; dev < num_devices; dev++) { + verify_multi_gpu_buffer(loops, global_array[dev]); + } + + HIP_CHECK(hipHostFree(global_array)); + for (int k = 0; k < num_devices; ++k) { + HIP_CHECK(hipFree(kernel_buffer[k])); + HIP_CHECK(hipFree(kernel_atomic[k])); + HIP_CHECK(hipStreamDestroy(streams[k])); + free(host_buffer[k]); + } +} diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/hipCGThreadBlockTileTypeShfl_old.cc b/projects/hip-tests/catch/unit/cooperativeGrps/hipCGThreadBlockTileTypeShfl_old.cc new file mode 100644 index 0000000000..14dd68116f --- /dev/null +++ b/projects/hip-tests/catch/unit/cooperativeGrps/hipCGThreadBlockTileTypeShfl_old.cc @@ -0,0 +1,198 @@ +/* +Copyright (c) 2020 - 2022 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 +#include + +#include "hip_cg_common.hh" + +namespace cg = cooperative_groups; + +enum class TiledGroupShflTests { shflDown, shflXor, shflUp }; + +template +__device__ int reduction_kernel_shfl_down(cg::thread_block_tile const& g, + volatile int val) { + int sz = g.size(); + + for (int i = sz / 2; i > 0; i >>= 1) { + val += g.shfl_down(val, i); + } + + // 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 +__device__ int reduction_kernel_shfl_xor(cg::thread_block_tile const& g, int val) { + int sz = g.size(); + + for (int i = sz / 2; i > 0; i >>= 1) { + val += g.shfl_xor(val, i); + } + + // 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 +__device__ int prefix_sum_kernel(cg::thread_block_tile const& g, volatile int val) { + int sz = g.size(); +#pragma unroll + for (int i = 1; i < sz; i <<= 1) { + int temp = g.shfl_up(val, i); + + if (g.thread_rank() >= i) { + val += temp; + } + } + return val; +} + +template +static __global__ void kernel_cg_group_partition_static(int* result, + TiledGroupShflTests shfl_test) { + cg::thread_block thread_block_CG_ty = cg::this_thread_block(); + int input, output_sum; + + // Choose a leader thread to print the results + if (thread_block_CG_ty.thread_rank() == 0) { + printf(" Creating %d groups, of tile size %d threads:\n\n", + (int)thread_block_CG_ty.size() / tile_size, tile_size); + } + + thread_block_CG_ty.sync(); + + cg::thread_block_tile tiled_part = cg::tiled_partition(thread_block_CG_ty); + + input = tiled_part.thread_rank(); + + switch (shfl_test) { + case (TiledGroupShflTests::shflDown): + output_sum = reduction_kernel_shfl_down(tiled_part, input); + break; + case (TiledGroupShflTests::shflXor): + output_sum = reduction_kernel_shfl_xor(tiled_part, input); + break; + case (TiledGroupShflTests::shflUp): + output_sum = prefix_sum_kernel(tiled_part, input); + result[thread_block_CG_ty.thread_rank()] = output_sum; + } + + if (tiled_part.thread_rank() == 0 && shfl_test != TiledGroupShflTests::shflUp) { + printf(" Sum of all ranks 0..%d in this tiled_part group is %d\n", tiled_part.size() - 1, + output_sum); + result[thread_block_CG_ty.thread_rank() / (tile_size)] = output_sum; + } +} + +static void expected_result_calc(int* expected_result, int tile_size, int size, + TiledGroupShflTests shfl_test) { + switch (shfl_test) { + case (TiledGroupShflTests::shflDown): + case (TiledGroupShflTests::shflXor): { + int expected_sum = ((tile_size - 1) * tile_size / 2); + for (int i = 0; i < size; i++) { + expected_result[i] = expected_sum; + } + break; + } + case (TiledGroupShflTests::shflUp): { + for (int i = 0; i < size / tile_size; i++) { + int acc = 0; + for (int j = 0; j < tile_size; j++) { + acc += j; + expected_result[i * tile_size + j] = acc; + } + } + break; + } + } +} + +template static void test_group_partition(TiledGroupShflTests shfl_test) { + int block_size = 1; + int threads_per_blk = 64; + + int num_elem = (block_size * threads_per_blk) / tile_size; + if (shfl_test == TiledGroupShflTests::shflUp) { + num_elem = block_size * threads_per_blk; + } + + int* expected_result = new int[num_elem]; + + int* result_dev = NULL; + int* result_host = NULL; + + HIP_CHECK(hipHostMalloc(&result_host, num_elem * sizeof(int), hipHostMallocDefault)); + memset(result_host, 0, num_elem * sizeof(int)); + + HIP_CHECK(hipMalloc(&result_dev, num_elem * sizeof(int))); + + // Launch Kernel + hipLaunchKernelGGL(kernel_cg_group_partition_static, block_size, threads_per_blk, + threads_per_blk * sizeof(int), 0, result_dev, shfl_test); + HIP_CHECK(hipDeviceSynchronize()); + + + HIP_CHECK(hipMemcpy(result_host, result_dev, sizeof(int) * num_elem, hipMemcpyDeviceToHost)); + + expected_result_calc(expected_result, tile_size, num_elem, shfl_test); + compareResults(expected_result, result_host, num_elem * sizeof(int)); + + // Free all allocated memory on host and device + HIP_CHECK(hipFree(result_dev)); + HIP_CHECK(hipHostFree(result_host)); + delete[] expected_result; +} + +TEST_CASE("Unit_hipCGThreadBlockTileType_Shfl") { + // Use default device for validating the test + 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; + } + + TiledGroupShflTests shfl_test = GENERATE( + TiledGroupShflTests::shflDown, TiledGroupShflTests::shflXor, TiledGroupShflTests::shflUp); + test_group_partition<2>(shfl_test); + test_group_partition<4>(shfl_test); + test_group_partition<8>(shfl_test); + test_group_partition<16>(shfl_test); + test_group_partition<32>(shfl_test); +} diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/hipCGThreadBlockType.cc b/projects/hip-tests/catch/unit/cooperativeGrps/hipCGThreadBlockType.cc deleted file mode 100644 index 98f611aa8a..0000000000 --- a/projects/hip-tests/catch/unit/cooperativeGrps/hipCGThreadBlockType.cc +++ /dev/null @@ -1,177 +0,0 @@ -/* -Copyright (c) 2020 - 2021 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. -*/ - - -/* HIT_START - * BUILD: %t %s ../../test_common.cpp - * TEST: %t - * HIT_END - */ - -#include -#include -#include - -#define ASSERT_EQUAL(lhs, rhs) HIPASSERT(lhs == rhs) - -using namespace cooperative_groups; - -static __global__ -void kernel_cg_thread_block_type(int *sizeTestD, - int *thdRankTestD, - int *syncTestD, - dim3 *groupIndexTestD, - dim3 *thdIndexTestD, - dim3 *groupDimTestD) -{ - thread_block tb = this_thread_block(); - int gIdx = (blockIdx.x * blockDim.x) + threadIdx.x; - // Test size - sizeTestD[gIdx] = tb.size(); - - // Test thread_rank - thdRankTestD[gIdx] = tb.thread_rank(); - - // Test sync - __shared__ int sm[2]; - if (threadIdx.x == 0) - sm[0] = 10; - else if (threadIdx.x == 1) - sm[1] = 20; - tb.sync(); - syncTestD[gIdx] = sm[1] * sm[0]; - - // Test group_index - groupIndexTestD[gIdx] = tb.group_index(); - - // Test thread_index - thdIndexTestD[gIdx] = tb.thread_index(); - - // Test group_dim aka number of threads in a block - groupDimTestD[gIdx] = tb.group_dim(); -} - -static void test_cg_thread_block_type(int blockSize) -{ - int nBytes = sizeof(int) * 2 * blockSize; - int nDim3Bytes = sizeof(dim3) * 2 * blockSize; - int *sizeTestD, *sizeTestH; - int *thdRankTestD, *thdRankTestH; - int *syncTestD, *syncTestH; - dim3 *groupIndexTestD, *groupIndexTestH; - dim3 *thdIndexTestD, *thdIndexTestH, *groupDimTestD, *groupDimTestH; - - // Allocate device memory - HIPCHECK(hipMalloc(&sizeTestD, nBytes)); - HIPCHECK(hipMalloc(&thdRankTestD, nBytes)); - HIPCHECK(hipMalloc(&syncTestD, nBytes)); - HIPCHECK(hipMalloc(&groupIndexTestD, nDim3Bytes)); - HIPCHECK(hipMalloc(&thdIndexTestD, nDim3Bytes)); - HIPCHECK(hipMalloc(&groupDimTestD, nDim3Bytes)); - - // Allocate host memory - HIPCHECK(hipHostMalloc(&sizeTestH, nBytes)); - HIPCHECK(hipHostMalloc(&thdRankTestH, nBytes)); - HIPCHECK(hipHostMalloc(&syncTestH, nBytes)); - HIPCHECK(hipHostMalloc(&groupIndexTestH, nDim3Bytes)); - HIPCHECK(hipHostMalloc(&thdIndexTestH, nDim3Bytes)); - HIPCHECK(hipHostMalloc(&groupDimTestH, nDim3Bytes)); - - // Launch Kernel - hipLaunchKernelGGL(kernel_cg_thread_block_type, - 2, - blockSize, - 0, - 0, - sizeTestD, - thdRankTestD, - syncTestD, - groupIndexTestD, - thdIndexTestD, - groupDimTestD); - - // Copy result from device to host - 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)); - HIPCHECK(hipMemcpy(groupDimTestH, groupDimTestD, 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, (uint) i / blockSize); - ASSERT_EQUAL(groupIndexTestH[i].y, 0); - ASSERT_EQUAL(groupIndexTestH[i].z, 0); - ASSERT_EQUAL(thdIndexTestH[i].x, (uint) i % blockSize); - ASSERT_EQUAL(thdIndexTestH[i].y, 0); - ASSERT_EQUAL(thdIndexTestH[i].z, 0); - ASSERT_EQUAL(groupDimTestH[i].x, blockSize); - ASSERT_EQUAL(groupDimTestH[i].y, 1); - ASSERT_EQUAL(groupDimTestH[i].z, 1); - } - - // Free device memory - HIPCHECK(hipFree(sizeTestD)); - HIPCHECK(hipFree(thdRankTestD)); - HIPCHECK(hipFree(syncTestD)); - HIPCHECK(hipFree(groupIndexTestD)); - HIPCHECK(hipFree(thdIndexTestD)); - HIPCHECK(hipFree(groupDimTestD)); - - //Free host memory - HIPCHECK(hipHostFree(sizeTestH)); - HIPCHECK(hipHostFree(thdRankTestH)); - HIPCHECK(hipHostFree(syncTestH)); - HIPCHECK(hipHostFree(groupIndexTestH)); - HIPCHECK(hipHostFree(thdIndexTestH)); - HIPCHECK(hipHostFree(groupDimTestH)); -} - -TEST_CASE("Unit_hipCGThreadBlockType") { - // Use default device for validating the test - int deviceId; - hipDeviceProp_t deviceProperties; - HIPCHECK(hipGetDevice(&deviceId)); - HIPCHECK(hipGetDeviceProperties(&deviceProperties, deviceId)); - - if (!deviceProperties.cooperativeLaunch) { - HipTest::HIP_SKIP_TEST("Device doesn't support cooperative launch!"); - return; - } - - // 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); - } - - // 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)); - } -} diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/hipCGThreadBlockTypeViaBaseType.cc b/projects/hip-tests/catch/unit/cooperativeGrps/hipCGThreadBlockTypeViaBaseType.cc deleted file mode 100644 index 69f5e91ad2..0000000000 --- a/projects/hip-tests/catch/unit/cooperativeGrps/hipCGThreadBlockTypeViaBaseType.cc +++ /dev/null @@ -1,136 +0,0 @@ -/* -Copyright (c) 2020 - 2021 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. -*/ - - -/* HIT_START - * BUILD: %t %s ../../test_common.cpp - * TEST: %t - * HIT_END - */ - -#include -#include "hip/hip_cooperative_groups.h" -#include - -#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs) - -using namespace cooperative_groups; - -static __global__ -void kernel_cg_thread_block_type_via_base_type(int *sizeTestD, - int *thdRankTestD, - int *syncTestD) -{ - thread_group tg = this_thread_block(); - int gIdx = (blockIdx.x * blockDim.x) + threadIdx.x; - - // Test size - sizeTestD[gIdx] = tg.size(); - - // Test thread_rank - thdRankTestD[gIdx] = tg.thread_rank(); - - // Test sync - __shared__ int sm[2]; - if (threadIdx.x == 0) - sm[0] = 10; - else if (threadIdx.x == 1) - sm[1] = 20; - tg.sync(); - syncTestD[gIdx] = sm[1] * sm[0]; -} - -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 *syncTestD, *syncTestH; - - // Allocate device memory - HIPCHECK(hipMalloc(&sizeTestD, nBytes)); - HIPCHECK(hipMalloc(&thdRankTestD, nBytes)); - HIPCHECK(hipMalloc(&syncTestD, nBytes)); - - // Allocate host memory - HIPCHECK(hipHostMalloc(&sizeTestH, nBytes)); - HIPCHECK(hipHostMalloc(&thdRankTestH, nBytes)); - HIPCHECK(hipHostMalloc(&syncTestH, nBytes)); - - // Launch Kernel - hipLaunchKernelGGL(kernel_cg_thread_block_type_via_base_type, - 2, - blockSize, - 0, - 0, - sizeTestD, - thdRankTestD, - syncTestD); - - // Copy result from device to host - 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) { - ASSERT_EQUAL(sizeTestH[i], blockSize); - ASSERT_EQUAL(thdRankTestH[i], i % blockSize); - ASSERT_EQUAL(syncTestH[i], 200); - } - - // Free device memory - HIPCHECK(hipFree(sizeTestD)); - HIPCHECK(hipFree(thdRankTestD)); - HIPCHECK(hipFree(syncTestD)); - - //Free host memory - HIPCHECK(hipHostFree(sizeTestH)); - HIPCHECK(hipHostFree(thdRankTestH)); - HIPCHECK(hipHostFree(syncTestH)); -} - -TEST_CASE("Unit_hipCGThreadBlockType_BaseType") { - // Use default device for validating the test - int deviceId; - hipDeviceProp_t deviceProperties; - HIPCHECK(hipGetDevice(&deviceId)); - HIPCHECK(hipGetDeviceProperties(&deviceProperties, deviceId)); - - if (!deviceProperties.cooperativeLaunch) { - HipTest::HIP_SKIP_TEST("Device doesn't support cooperative launch!"); - return; - } - - // 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); - } - - // 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)); - } -} diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/hipCGThreadBlockTypeViaPublicApi.cc b/projects/hip-tests/catch/unit/cooperativeGrps/hipCGThreadBlockTypeViaPublicApi.cc deleted file mode 100644 index f4913ad2c7..0000000000 --- a/projects/hip-tests/catch/unit/cooperativeGrps/hipCGThreadBlockTypeViaPublicApi.cc +++ /dev/null @@ -1,136 +0,0 @@ -/* -Copyright (c) 2020 - 2021 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. -*/ - - -/* HIT_START - * BUILD: %t %s ../../test_common.cpp - * TEST: %t - * HIT_END - */ - -#include -#include "hip/hip_cooperative_groups.h" -#include - -#define ASSERT_EQUAL(lhs, rhs) assert(lhs == rhs) - -using namespace cooperative_groups; - -static __global__ -void kernel_cg_thread_block_type_via_public_api(int *sizeTestD, - int *thdRankTestD, - int *syncTestD) -{ - thread_block tb = this_thread_block(); - int gIdx = (blockIdx.x * blockDim.x) + threadIdx.x; - - // Test group_size api - sizeTestD[gIdx] = group_size(tb); - - // Test thread_rank api - thdRankTestD[gIdx] = thread_rank(tb); - - // Test sync api - __shared__ int sm[2]; - if (threadIdx.x == 0) - sm[0] = 10; - else if (threadIdx.x == 1) - sm[1] = 20; - sync(tb); - syncTestD[gIdx] = sm[1] * sm[0]; -} - -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 *syncTestD, *syncTestH; - - // Allocate device memory - HIPCHECK(hipMalloc(&sizeTestD, nBytes)); - HIPCHECK(hipMalloc(&thdRankTestD, nBytes)); - HIPCHECK(hipMalloc(&syncTestD, nBytes)); - - // Allocate host memory - HIPCHECK(hipHostMalloc(&sizeTestH, nBytes)); - HIPCHECK(hipHostMalloc(&thdRankTestH, nBytes)); - HIPCHECK(hipHostMalloc(&syncTestH, nBytes)); - - // Launch Kernel - hipLaunchKernelGGL(kernel_cg_thread_block_type_via_public_api, - 2, - blockSize, - 0, - 0, - sizeTestD, - thdRankTestD, - syncTestD); - - // Copy result from device to host - 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) { - ASSERT_EQUAL(sizeTestH[i], blockSize); - ASSERT_EQUAL(thdRankTestH[i], i % blockSize); - ASSERT_EQUAL(syncTestH[i], 200); - } - - // Free device memory - HIPCHECK(hipFree(sizeTestD)); - HIPCHECK(hipFree(thdRankTestD)); - HIPCHECK(hipFree(syncTestD)); - - //Free host memory - HIPCHECK(hipHostFree(sizeTestH)); - HIPCHECK(hipHostFree(thdRankTestH)); - HIPCHECK(hipHostFree(syncTestH)); -} - -TEST_CASE("Unit_hipCGThreadBlockType_PublicApi") { - // Use default device for validating the test - int deviceId; - hipDeviceProp_t deviceProperties; - HIPCHECK(hipGetDevice(&deviceId)); - HIPCHECK(hipGetDeviceProperties(&deviceProperties, deviceId)); - - if (!deviceProperties.cooperativeLaunch) { - HipTest::HIP_SKIP_TEST("Device doesn't support cooperative launch!"); - return; - } - - // 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); - } - - // 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)); - } -} diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/hipCGThreadBlockType_old.cc b/projects/hip-tests/catch/unit/cooperativeGrps/hipCGThreadBlockType_old.cc new file mode 100644 index 0000000000..87ec21d748 --- /dev/null +++ b/projects/hip-tests/catch/unit/cooperativeGrps/hipCGThreadBlockType_old.cc @@ -0,0 +1,225 @@ +/* +Copyright (c) 2020 - 2022 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 +#include + +#include "hip_cg_common.hh" + +namespace cg = cooperative_groups; + +enum class ThreadBlockTypeTests { basicApi, baseType, publicApi }; + +static __global__ void kernel_cg_thread_block_type(int* size_dev, int* thd_rank_dev, int* sync_dev, + dim3* group_index_dev, dim3* thd_index_dev, + dim3* group_dim_dev) { + cg::thread_block tb = cg::this_thread_block(); + int gIdx = (blockIdx.x * blockDim.x) + threadIdx.x; + // Test size + size_dev[gIdx] = tb.size(); + + // Test thread_rank + thd_rank_dev[gIdx] = tb.thread_rank(); + + // Test sync + __shared__ int sm[2]; + if (threadIdx.x == 0) + sm[0] = 10; + else if (threadIdx.x == 1) + sm[1] = 20; + tb.sync(); + sync_dev[gIdx] = sm[1] * sm[0]; + + // Test group_index + group_index_dev[gIdx] = tb.group_index(); + + // Test thread_index + thd_index_dev[gIdx] = tb.thread_index(); + + // Test group_dim aka number of threads in a block + group_dim_dev[gIdx] = tb.group_dim(); +} + +static __global__ void kernel_cg_thread_block_type_via_base_type(int* size_dev, int* thd_rank_dev, + int* sync_dev) { + cg::thread_group tg = cg::this_thread_block(); + int gIdx = (blockIdx.x * blockDim.x) + threadIdx.x; + + // Test size + size_dev[gIdx] = tg.size(); + + // Test thread_rank + thd_rank_dev[gIdx] = tg.thread_rank(); + + // Test sync + __shared__ int sm[2]; + if (threadIdx.x == 0) + sm[0] = 10; + else if (threadIdx.x == 1) + sm[1] = 20; + tg.sync(); + sync_dev[gIdx] = sm[1] * sm[0]; +} + +static __global__ void kernel_cg_thread_block_type_via_public_api(int* size_dev, int* thd_rank_dev, + int* sync_dev) { + cg::thread_block tb = cg::this_thread_block(); + int gIdx = (blockIdx.x * blockDim.x) + threadIdx.x; + + // Test group_size api + size_dev[gIdx] = cg::group_size(tb); + + // Test thread_rank api + thd_rank_dev[gIdx] = cg::thread_rank(tb); + + // Test sync api + __shared__ int sm[2]; + if (threadIdx.x == 0) + sm[0] = 10; + else if (threadIdx.x == 1) + sm[1] = 20; + cg::sync(tb); + sync_dev[gIdx] = sm[1] * sm[0]; +} + +static void test_cg_thread_block_type(ThreadBlockTypeTests test_type, int block_size) { + int num_bytes = sizeof(int) * 2 * block_size; + int num_dim3_bytes = sizeof(dim3) * 2 * block_size; + int *size_dev, *size_host; + int *thd_rank_dev, *thd_rank_host; + int *sync_dev, *sync_host; + dim3 *group_index_dev, *group_index_host; + dim3 *thd_index_dev, *thd_index_host; + dim3 *group_dim_dev, *group_dim_host; + + // Allocate device memory + HIP_CHECK(hipMalloc(&size_dev, num_bytes)); + HIP_CHECK(hipMalloc(&thd_rank_dev, num_bytes)); + HIP_CHECK(hipMalloc(&sync_dev, num_bytes)); + + // Allocate host memory + HIP_CHECK(hipHostMalloc(&size_host, num_bytes)); + HIP_CHECK(hipHostMalloc(&thd_rank_host, num_bytes)); + HIP_CHECK(hipHostMalloc(&sync_host, num_bytes)); + + switch (test_type) { + case (ThreadBlockTypeTests::basicApi): + HIP_CHECK(hipMalloc(&group_index_dev, num_dim3_bytes)); + HIP_CHECK(hipMalloc(&thd_index_dev, num_dim3_bytes)); + HIP_CHECK(hipMalloc(&group_dim_dev, num_dim3_bytes)); + HIP_CHECK(hipHostMalloc(&group_index_host, num_dim3_bytes)); + HIP_CHECK(hipHostMalloc(&thd_index_host, num_dim3_bytes)); + HIP_CHECK(hipHostMalloc(&group_dim_host, num_dim3_bytes)); + + hipLaunchKernelGGL(kernel_cg_thread_block_type, 2, block_size, 0, 0, size_dev, thd_rank_dev, + sync_dev, group_index_dev, thd_index_dev, group_dim_dev); + break; + case (ThreadBlockTypeTests::baseType): + hipLaunchKernelGGL(kernel_cg_thread_block_type_via_base_type, 2, block_size, 0, 0, size_dev, + thd_rank_dev, sync_dev); + break; + case (ThreadBlockTypeTests::publicApi): + hipLaunchKernelGGL(kernel_cg_thread_block_type_via_public_api, 2, block_size, 0, 0, size_dev, + thd_rank_dev, sync_dev); + } + + // Copy result from device to host + HIP_CHECK(hipMemcpy(size_host, size_dev, num_bytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(thd_rank_host, thd_rank_dev, num_bytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(sync_host, sync_dev, num_bytes, hipMemcpyDeviceToHost)); + if (test_type == ThreadBlockTypeTests::basicApi) { + HIP_CHECK(hipMemcpy(group_index_host, group_index_dev, num_dim3_bytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(thd_index_host, thd_index_dev, num_dim3_bytes, hipMemcpyDeviceToHost)); + HIP_CHECK(hipMemcpy(group_dim_host, group_dim_dev, num_dim3_bytes, hipMemcpyDeviceToHost)); + } + + // Validate results for both blocks together + for (int i = 0; i < 2 * block_size; ++i) { + ASSERT_EQUAL(size_host[i], block_size); + ASSERT_EQUAL(thd_rank_host[i], i % block_size); + ASSERT_EQUAL(sync_host[i], 200); + if (test_type == ThreadBlockTypeTests::basicApi) { + ASSERT_EQUAL(group_index_host[i].x, (uint)i / block_size); + ASSERT_EQUAL(group_index_host[i].y, 0); + ASSERT_EQUAL(group_index_host[i].z, 0); + ASSERT_EQUAL(thd_index_host[i].x, (uint)i % block_size); + ASSERT_EQUAL(thd_index_host[i].y, 0); + ASSERT_EQUAL(thd_index_host[i].z, 0); + ASSERT_EQUAL(group_dim_host[i].x, block_size); + ASSERT_EQUAL(group_dim_host[i].y, 1); + ASSERT_EQUAL(group_dim_host[i].z, 1); + } + } + + // Free device memory + HIP_CHECK(hipFree(size_dev)); + HIP_CHECK(hipFree(thd_rank_dev)); + HIP_CHECK(hipFree(sync_dev)); + + // Free host memory + HIP_CHECK(hipHostFree(size_host)); + HIP_CHECK(hipHostFree(thd_rank_host)); + HIP_CHECK(hipHostFree(sync_host)); + + if (test_type == ThreadBlockTypeTests::basicApi) { + HIP_CHECK(hipFree(group_index_dev)); + HIP_CHECK(hipFree(thd_index_dev)); + HIP_CHECK(hipFree(group_dim_dev)); + HIP_CHECK(hipHostFree(group_index_host)); + HIP_CHECK(hipHostFree(thd_index_host)); + HIP_CHECK(hipHostFree(group_dim_host)); + } +} + + +TEST_CASE("Unit_hipCGThreadBlockType") { + // Use default device for validating the test + 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; + } + + ThreadBlockTypeTests test_type = ThreadBlockTypeTests::basicApi; + + SECTION("Default thread block API test") { test_type = ThreadBlockTypeTests::basicApi; } + + SECTION("Base type thread block API test") { test_type = ThreadBlockTypeTests::baseType; } + + SECTION("Public API thread block test") { test_type = ThreadBlockTypeTests::publicApi; } + + // Test for blockSizes in powers of 2 + int max_threads_per_blk = device_properties.maxThreadsPerBlock; + for (int block_size = 2; block_size <= max_threads_per_blk; block_size = block_size * 2) { + test_cg_thread_block_type(test_type, block_size); + } + + // Test for random block_size, 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(test_type, max(2, rand() % max_threads_per_blk)); + } +} diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/hipCGTiledPartition.cc b/projects/hip-tests/catch/unit/cooperativeGrps/hipCGTiledPartition.cc deleted file mode 100644 index 783d7c8036..0000000000 --- a/projects/hip-tests/catch/unit/cooperativeGrps/hipCGTiledPartition.cc +++ /dev/null @@ -1,385 +0,0 @@ -/* -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, - 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 -#include -#include -#include - -using namespace cooperative_groups; - -/* 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(); - - 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 -__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 tiledPartition = tiled_partition(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* workspace = NULL; - - if (isGlobalMem) { - workspace = globalMem; - } else { - // Declare a shared memory - extern __shared__ int sharedMem[]; - workspace = sharedMem; - } - - int input, outputSum; - - // input to reduction, for each thread, is its' rank in the group - input = threadBlockCGTy.thread_rank(); - - 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 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; - HIPCHECK(hipMalloc((void**)&dResult, numTiles * sizeof(int))); - - int* globalMem = NULL; - if (useGlobalMem) { - HIPCHECK(hipMalloc((void**)&globalMem, threadsPerBlock * sizeof(int))); - } - - int* hResult = NULL; - HIPCHECK(hipHostMalloc(&hResult, numTiles * sizeof(int), hipHostMallocDefault)); - memset(hResult, 0, numTiles * sizeof(int)); - - if (useGlobalMem) { - // Launch Kernel - hipLaunchKernelGGL(kernel_cg_group_partition_static, 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, 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)); - } - } - - HIPCHECK(hipMemcpy(hResult, dResult, numTiles * sizeof(int), hipMemcpyDeviceToHost)); - - verifyResults(expectedSum, hResult, numTiles); - - // Free all allocated memory on host and device - HIPCHECK(hipFree(dResult)); - HIPCHECK(hipFree(hResult)); - if (useGlobalMem) { - HIPCHECK(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; - HIPCHECK(hipMalloc(&dResult, sizeof(int) * numTiles)); - - int* globalMem = NULL; - if (useGlobalMem) { - HIPCHECK(hipMalloc((void**)&globalMem, threadsPerBlock * sizeof(int))); - } - - int* hResult = NULL; - HIPCHECK(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)); - } - } - - HIPCHECK(hipMemcpy(hResult, dResult, numTiles * sizeof(int), hipMemcpyDeviceToHost)); - - verifyResults(expectedSum, hResult, numTiles); - - // Free all allocated memory on host and device - HIPCHECK(hipFree(dResult)); - HIPCHECK(hipFree(hResult)); - if (useGlobalMem) { - HIPCHECK(hipFree(globalMem)); - } - delete[] expectedSum; - - printf("\n...PASSED.\n\n"); -} - -TEST_CASE("Unit_tiled_partition") { - // Use default device for validating the test - int deviceId; - HIP_CHECK_ERROR(hipGetDevice(&deviceId), hipSuccess); - hipDeviceProp_t deviceProperties; - HIP_CHECK_ERROR(hipGetDeviceProperties(&deviceProperties, deviceId), hipSuccess); - - 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 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; -} diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/hipCGTiledPartitionType_old.cc b/projects/hip-tests/catch/unit/cooperativeGrps/hipCGTiledPartitionType_old.cc new file mode 100644 index 0000000000..af43709ac4 --- /dev/null +++ b/projects/hip-tests/catch/unit/cooperativeGrps/hipCGTiledPartitionType_old.cc @@ -0,0 +1,279 @@ +/* +Copyright (c) 2020 - 2022 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, + 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 +#include +#include + +#include "hip_cg_common.hh" + +namespace cg = cooperative_groups; + +/* 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(cg::thread_group g, int* x, int val) { + int lane = g.thread_rank(); + + 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 +__global__ void kernel_cg_group_partition_static(int* result, bool is_global_mem, int* global_mem) { + cg::thread_block thread_block_CG_ty = cg::this_thread_block(); + + int* workspace = NULL; + + if (is_global_mem) { + workspace = global_mem; + } else { + // Declare a shared memory + extern __shared__ int shared_mem[]; + workspace = shared_mem; + } + + int input, output_sum, expected_output; + + // input to reduction, for each thread, is its' rank in the group + input = thread_block_CG_ty.thread_rank(); + + expected_output = (thread_block_CG_ty.size() - 1) * thread_block_CG_ty.size() / 2; + + output_sum = reduction_kernel(thread_block_CG_ty, workspace, input); + + if (thread_block_CG_ty.thread_rank() == 0) { + printf(" Sum of all ranks 0..%d in threadBlockCooperativeGroup is %d (expected %d)\n\n", + (int)thread_block_CG_ty.size() - 1, output_sum, expected_output); + printf(" Creating %d groups, of tile size %d threads:\n\n", + (int)thread_block_CG_ty.size() / tile_size, tile_size); + } + + thread_block_CG_ty.sync(); + + cg::thread_block_tile tiled_part = cg::tiled_partition(thread_block_CG_ty); + + // This offset allows each group to have its own unique area in the workspace array + int workspace_offset = thread_block_CG_ty.thread_rank() - tiled_part.thread_rank(); + + output_sum = reduction_kernel(tiled_part, workspace + workspace_offset, input); + + if (tiled_part.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", + tiled_part.size() - 1, output_sum, tiled_part.meta_group_rank(), + tiled_part.meta_group_size()); + result[input / (tile_size)] = output_sum; + } + return; +} + +__global__ void kernel_cg_group_partition_dynamic(unsigned int tile_size, int* result, + bool is_global_mem, int* global_mem) { + cg::thread_block thread_block_CG_ty = cg::this_thread_block(); + + int* workspace = NULL; + + if (is_global_mem) { + workspace = global_mem; + } else { + // Declare a shared memory + extern __shared__ int shared_mem[]; + workspace = shared_mem; + } + + int input, output_sum; + + // input to reduction, for each thread, is its' rank in the group + input = thread_block_CG_ty.thread_rank(); + + output_sum = reduction_kernel(thread_block_CG_ty, workspace, input); + + if (thread_block_CG_ty.thread_rank() == 0) { + printf("\n\n\n Sum of all ranks 0..%d in threadBlockCooperativeGroup is %d\n\n", + (int)thread_block_CG_ty.size() - 1, output_sum); + printf(" Creating %d groups, of tile size %d threads:\n\n", + (int)thread_block_CG_ty.size() / tile_size, tile_size); + } + + thread_block_CG_ty.sync(); + + cg::thread_group tiled_part = cg::tiled_partition(thread_block_CG_ty, tile_size); + + // This offset allows each group to have its own unique area in the workspace array + int workspace_offset = thread_block_CG_ty.thread_rank() - tiled_part.thread_rank(); + + output_sum = reduction_kernel(tiled_part, workspace + workspace_offset, input); + + if (tiled_part.thread_rank() == 0) { + printf( + " Sum of all ranks 0..%d in this tiledPartition group is %d. Corresponding parent thread " + "rank: %d\n", + static_cast(tiled_part.size()) - 1, output_sum, input); + result[input / (tile_size)] = output_sum; + } + return; +} + +template +static void common_group_partition(F kernel_func, unsigned int tile_size, void** params, + size_t num_params, bool use_global_mem) { + int block_size = 1; + int threads_per_blk = 64; + + int num_tiles = (block_size * threads_per_blk) / tile_size; + + // 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* expected_sum = new int[num_tiles]; + int temp = 0, sum = 0; + + for (int i = 1; i <= num_tiles; i++) { + sum = temp; + temp = (((tile_size * i) - 1) * (tile_size * i)) / 2; + expected_sum[i - 1] = temp - sum; + } + + int* result_dev = NULL; + HIP_CHECK(hipMalloc((void**)&result_dev, num_tiles * sizeof(int))); + + int* global_mem = NULL; + if (use_global_mem) { + HIP_CHECK(hipMalloc((void**)&global_mem, threads_per_blk * sizeof(int))); + } + + int* result_host = NULL; + HIP_CHECK(hipHostMalloc(&result_host, num_tiles * sizeof(int), hipHostMallocDefault)); + memset(result_host, 0, num_tiles * sizeof(int)); + + params[num_params + 0] = &result_dev; + params[num_params + 1] = &use_global_mem; + params[num_params + 2] = &global_mem; + + if (use_global_mem) { + // Launch Kernel + HIP_CHECK(hipLaunchCooperativeKernel(kernel_func, block_size, threads_per_blk, params, 0, 0)); + HIP_CHECK(hipDeviceSynchronize()); + } else { + // Launch Kernel + HIP_CHECK(hipLaunchCooperativeKernel(kernel_func, block_size, threads_per_blk, params, + threads_per_blk * sizeof(int), 0)); + HIP_CHECK(hipDeviceSynchronize()); + } + + HIP_CHECK(hipMemcpy(result_host, result_dev, num_tiles * sizeof(int), hipMemcpyDeviceToHost)); + + verifyResults(expected_sum, result_host, num_tiles); + + // Free all allocated memory on host and device + HIP_CHECK(hipFree(result_dev)); + HIP_CHECK(hipHostFree(result_host)); + if (use_global_mem) { + HIP_CHECK(hipFree(global_mem)); + } + delete[] expected_sum; +} + +template static void test_group_partition(bool use_global_mem) { + void* params[3]; + size_t num_params = 0; + common_group_partition(kernel_cg_group_partition_static, tile_size, params, num_params, + use_global_mem); +} + +static void test_group_partition(unsigned int tile_size, bool use_global_mem) { + void* params[4]; + params[0] = &tile_size; + size_t num_params = 1; + common_group_partition(kernel_cg_group_partition_dynamic, tile_size, params, num_params, + use_global_mem); +} + +TEST_CASE("Unit_hipCGThreadBlockTileType") { + // Use default device for validating the test + 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; + } + + bool use_global_mem = GENERATE(true, false); + + SECTION("Static tile partition") { + test_group_partition<2>(use_global_mem); + test_group_partition<4>(use_global_mem); + test_group_partition<8>(use_global_mem); + test_group_partition<16>(use_global_mem); + test_group_partition<32>(use_global_mem); + } + + SECTION("Dynamic tile partition") { + unsigned int tile_size = GENERATE(2, 4, 8, 16, 32); + test_group_partition(tile_size, use_global_mem); + } +} diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/hipLaunchCooperativeKernelMultiDevice_old.cc b/projects/hip-tests/catch/unit/cooperativeGrps/hipLaunchCooperativeKernelMultiDevice_old.cc new file mode 100644 index 0000000000..6af1a6f4a9 --- /dev/null +++ b/projects/hip-tests/catch/unit/cooperativeGrps/hipLaunchCooperativeKernelMultiDevice_old.cc @@ -0,0 +1,606 @@ +/* +Copyright (c) 2020 - 2022 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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +// Test Description: +/*The general idea of the application is to test how multi-GPU Cooperative +Groups kernel launches to a stream interact with other things that may be +simultaneously running in the same streams. + +The HIP specification says that a multi-GPU cooperative launch will wait +until all of the streams it's using finish their work. Only then will the +cooperative kernel be launched to all of the devices. Then no other work +can take part in the any of the streams until all of the multi-GPU +cooperative work is done. + +However, there are flags that allow you to disable each of these +serialization points: hipCooperativeLaunchMultiDeviceNoPreSync and +hipCooperativeLaunchMultiDeviceNoPostSync. + +As such, this benchmark tests the following five situations launching +to two GPUs (and thus two streams): + + 1. Normal multi-GPU cooperative kernel: + This should result in the following pattern: + Stream 0: Cooperative + Stream 1: Cooperative + 2. Regular kernel launches and multi-GPU cooperative kernel launches + with the default flags, resulting in the following pattern: + Stream 0: Regular --> Cooperative + Stream 1: --> Cooperative --> Regular + + 3. Regular kernel launches and multi-GPU cooperative kernel launches + that turn off "pre-sync". This should allow a cooperative kernel + to launch even if work is already in a stream pointing to + another GPU. + This should result in the following pattern: + Stream 0: Regular --> Cooperative + Stream 1: Cooperative --> Regular + + 4. Regular kernel launches and multi-GPU cooperative kernel launches + that turn off "post-sync". This should allow a new kernel to enter + a GPU even if another GPU still has a cooperative kernel on it. + This should result in the following pattern: + Stream 0: Regular --> Cooperative + Stream 1: --> Cooperative--> Regular + + 5. Regular kernel launches and multi-GPU cooperative kernel launches + that turn off both pre- and post-sync. This should allow any of + the kernels to launch to their GPU regardless of the status of + other kernels in other multi-GPU stream groups. + This should result in the following pattern: + Stream 0: Regular --> Cooperative + Stream 1: Cooperative --> Regular + +We time how long it takes to run each of these benchmarks and print it as +the output of the benchmark. The kernels themselves are just useless time- +wasting code so that the kernel takes a meaningful amount of time on the +GPU before it exits. We only launch a single wavefront for each kernel, so +any serialization should not be because of GPU occupancy concerns. + +If tests 2, 3, and 4 take roughly 3x as long as #1, that implies that +cooperative kernels are serialized as expected. + +If test #5 takes roughly twice as long as #1, that implies that the +overlap-allowing flags work as expected. +*/ + +#include +#include + +namespace cg = cooperative_groups; + +static constexpr size_t kBufferLen = 1024 * 1024; + +__global__ void test_gws(uint* buf, uint buf_size, long* tmp_buf, long* result) { + extern __shared__ long tmp[]; + uint groups = gridDim.x; + uint group_id = blockIdx.x; + uint local_id = threadIdx.x; + uint chunk = gridDim.x * blockDim.x; + + uint i = group_id * blockDim.x + local_id; + long sum = 0; + while (i < buf_size) { + sum += buf[i]; + i += chunk; + } + tmp[local_id] = sum; + __syncthreads(); + i = 0; + if (local_id == 0) { + sum = 0; + while (i < blockDim.x) { + sum += tmp[i]; + i++; + } + tmp_buf[group_id] = sum; + } + // wait + cg::this_grid().sync(); + + if (((blockIdx.x * blockDim.x) + threadIdx.x) == 0) { + for (uint i = 1; i < groups; ++i) { + sum += tmp_buf[i]; + } + //*result = sum; + result[1 + cg::this_multi_grid().grid_rank()] = sum; + } + cg::this_multi_grid().sync(); + if (cg::this_multi_grid().grid_rank() == 0) { + sum = 0; + for (uint i = 1; i <= cg::this_multi_grid().num_grids(); ++i) { + sum += result[i]; + } + *result = sum; + } +} + +__global__ void test_coop_kernel(unsigned int loops, long long* array, int fast_gpu) { + cg::multi_grid_group mgrid = cg::this_multi_grid(); + unsigned int rank = blockIdx.x * blockDim.x + threadIdx.x; + + if (mgrid.grid_rank() == fast_gpu) { + return; + } + + for (int i = 0; i < loops; i++) { + long long time_diff = 0; + long long last_clock = clock64(); + do { + long long cur_clock = clock64(); + if (cur_clock > last_clock) { + time_diff += (cur_clock - last_clock); + } + // If it rolls over, we don't know how much to add to catch up. + // So just ignore those slipped cycles. + last_clock = cur_clock; + } while (time_diff < 1000000); + array[rank] += clock64(); + } +} + +__global__ void test_coop_kernel_gfx11(unsigned int loops, long long* array, int fast_gpu) { +#if HT_AMD + cg::multi_grid_group mgrid = cg::this_multi_grid(); + unsigned int rank = blockIdx.x * blockDim.x + threadIdx.x; + + if (mgrid.grid_rank() == fast_gpu) { + return; + } + + for (int i = 0; i < loops; i++) { + long long time_diff = 0; + long long last_clock = wall_clock64(); + do { + long long cur_clock = wall_clock64(); + if (cur_clock > last_clock) { + time_diff += (cur_clock - last_clock); + } + // If it rolls over, we don't know how much to add to catch up. + // So just ignore those slipped cycles. + last_clock = cur_clock; + } while (time_diff < 1000000); + array[rank] += wall_clock64(); + } +#endif +} + +__global__ void test_kernel(uint32_t loops, unsigned long long* array) { + unsigned int rank = blockIdx.x * blockDim.x + threadIdx.x; + + for (int i = 0; i < loops; i++) { + long long time_diff = 0; + long long last_clock = clock64(); + do { + long long cur_clock = clock64(); + if (cur_clock > last_clock) { + time_diff += (cur_clock - last_clock); + } + // If it rolls over, we don't know how much to add to catch up. + // So just ignore those slipped cycles. + last_clock = cur_clock; + } while (time_diff < 1000000); + array[rank] += clock64(); + } +} + +__global__ void test_kernel_gfx11(uint32_t loops, unsigned long long* array) { +#if HT_AMD + unsigned int rank = blockIdx.x * blockDim.x + threadIdx.x; + + for (int i = 0; i < loops; i++) { + long long time_diff = 0; + long long last_clock = wall_clock64(); + do { + long long cur_clock = wall_clock64(); + if (cur_clock > last_clock) { + time_diff += (cur_clock - last_clock); + } + // If it rolls over, we don't know how much to add to catch up. + // So just ignore those slipped cycles. + last_clock = cur_clock; + } while (time_diff < 1000000); + array[rank] += wall_clock64(); + } +#endif +} + +static void verify_time(double single_kernel_time, double multi_kernel_time, float low_bound, + float high_bound) { + // Test that multiple kernel times are inside expected boundaries + REQUIRE(multi_kernel_time >= low_bound * single_kernel_time); + REQUIRE(multi_kernel_time <= high_bound * single_kernel_time); +} + +void test_multigrid_streams(int device_num) { + uint32_t loops = 2000; + int32_t fast_gpu = -1; + + // We will launch enough waves to fill up all of the GPU + int warp_sizes[2]; + int num_sms[2]; + hipDeviceProp_t device_properties[2]; + int warp_size = INT_MAX; + int num_sm = INT_MAX; + for (int dev = 0; dev < (device_num - 1); ++dev) { + for (int i = 0; i < 2; i++) { + HIP_CHECK(hipGetDeviceProperties(&device_properties[i], (dev + i))); + warp_sizes[i] = device_properties[i].warpSize; + if (warp_sizes[i] < warp_size) { + warp_size = warp_sizes[i]; + } + num_sms[i] = device_properties[i].multiProcessorCount; + if (num_sms[i] < num_sm) { + num_sm = num_sms[i]; + } + } + + // Calculate the device occupancy to know how many blocks can be run. + int max_blocks_per_sm_arr[2]; + int max_blocks_per_sm = INT_MAX; + for (int i = 0; i < 2; i++) { + HIP_CHECK(hipSetDevice(dev + i)); + auto test_kernel_used = IsGfx11() ? test_kernel_gfx11 : test_kernel; + HIP_CHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&max_blocks_per_sm_arr[i], + test_kernel_used, warp_size, 0)); + if (max_blocks_per_sm_arr[i] < max_blocks_per_sm) { + max_blocks_per_sm = max_blocks_per_sm_arr[i]; + } + } + int desired_blocks = 1; + + if (desired_blocks > max_blocks_per_sm * num_sm) { + INFO("The requested number of blocks will not fit on the GPU"); + REQUIRE(desired_blocks < max_blocks_per_sm * num_sm); + return; + } + + // Create the streams we will use in this test + hipStream_t streams[2]; + for (int i = 0; i < 2; i++) { + HIP_CHECK(hipSetDevice(dev + i)); + HIP_CHECK(hipStreamCreate(&streams[i])); + } + + // Set up data to pass into the kernel + // Alocate the host input buffer, and two device-focused buffers that we + // will use for our test. + unsigned long long* dev_array[2]; + for (int i = 0; i < 2; i++) { + int good_size = desired_blocks * warp_size * sizeof(long long); + HIP_CHECK(hipSetDevice(dev + i)); + HIP_CHECK(hipMalloc(reinterpret_cast(&dev_array[i]), good_size)); + HIP_CHECK(hipMemsetAsync(dev_array[i], 0, good_size, streams[i])); + } + for (int i = 0; i < 2; i++) { + HIP_CHECK(hipSetDevice(dev + i)); + HIP_CHECK(hipDeviceSynchronize()); + } + + /* Launch the kernels ****************************************************/ + void* dev_params[2][3]; + hipLaunchParams md_params[2]; + std::chrono::time_point start_time[2]; + std::chrono::time_point end_time[2]; + + // Test 0: Launching a multi-GPU cooperative kernel + // Both GPUs launch a long cooperative kernel + INFO("GPU " << dev << ": Long Coop Kernel"); + INFO("GPU " << (dev + 1) << ": Long Coop Kernel"); + + auto test_coop_kernel_used = IsGfx11() ? test_coop_kernel_gfx11 : test_coop_kernel; + for (int i = 0; i < 2; i++) { + dev_params[i][0] = reinterpret_cast(&loops); + dev_params[i][1] = reinterpret_cast(&dev_array[i]); + dev_params[i][2] = reinterpret_cast(&fast_gpu); + md_params[i].func = reinterpret_cast(test_coop_kernel_used); + md_params[i].gridDim = desired_blocks; + md_params[i].blockDim = warp_size; + md_params[i].sharedMem = 0; + md_params[i].stream = streams[i]; + md_params[i].args = dev_params[i]; + } + + start_time[0] = std::chrono::system_clock::now(); + HIP_CHECK(hipLaunchCooperativeKernelMultiDevice(md_params, 2, 0)); + for (int i = 0; i < 2; i++) { + HIP_CHECK(hipSetDevice(dev + i)); + HIP_CHECK(hipDeviceSynchronize()); + } + end_time[0] = std::chrono::system_clock::now(); + + std::chrono::duration single_kernel_time = (end_time[0] - start_time[0]); + INFO("A single kernel on both GPUs took: " << single_kernel_time.count() << " seconds"); + + SECTION("GPU1 - Standard/ Long Coop, GPU2 - Coop/Standard") { + INFO("GPU " << dev << ": Standard/Long Coop"); + INFO("GPU " << (dev + 1) << ": Coop/Standard"); + fast_gpu = 1; + start_time[1] = std::chrono::system_clock::now(); + HIP_CHECK(hipSetDevice(dev)); + auto test_kernel_used = IsGfx11() ? test_kernel_gfx11 : test_kernel; + hipLaunchKernelGGL(test_kernel_used, dim3(desired_blocks), dim3(warp_size), 0, streams[0], + loops, dev_array[0]); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipLaunchCooperativeKernelMultiDevice(md_params, 2, 0)); + HIP_CHECK(hipSetDevice(dev + 1)); + test_kernel_used = IsGfx11() ? test_kernel_gfx11 : test_kernel; + hipLaunchKernelGGL(test_kernel_used, dim3(desired_blocks), dim3(warp_size), 0, streams[1], + loops, dev_array[1]); + HIP_CHECK(hipGetLastError()); + for (int i = 0; i < 2; i++) { + HIP_CHECK(hipSetDevice(dev + i)); + HIP_CHECK(hipDeviceSynchronize()); + } + end_time[1] = std::chrono::system_clock::now(); + std::chrono::duration serialized_gpu0_time = (end_time[1] - start_time[1]); + INFO("Serialized set of three kernels with GPU0 being long took: " + << serialized_gpu0_time.count() << " seconds"); + + verify_time(single_kernel_time.count(), serialized_gpu0_time.count(), 2.7f, 3.3f); + } + + SECTION("GPU1 - Standard/Coop, GPU2 - Long Coop/Standard") { + INFO("GPU " << dev << ": Standard/Coop"); + INFO("GPU " << (dev + 1) << ": Long Coop/Standard"); + fast_gpu = 0; + start_time[1] = std::chrono::system_clock::now(); + HIP_CHECK(hipSetDevice(dev)); + auto test_kernel_used = IsGfx11() ? test_kernel_gfx11 : test_kernel; + hipLaunchKernelGGL(test_kernel_used, dim3(desired_blocks), dim3(warp_size), 0, streams[0], + loops, dev_array[0]); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipLaunchCooperativeKernelMultiDevice(md_params, 2, 0)); + HIP_CHECK(hipSetDevice(dev + 1)); + test_kernel_used = IsGfx11() ? test_kernel_gfx11 : test_kernel; + hipLaunchKernelGGL(test_kernel_used, dim3(desired_blocks), dim3(warp_size), 0, streams[1], + loops, dev_array[1]); + HIP_CHECK(hipGetLastError()); + for (int i = 0; i < 2; i++) { + HIP_CHECK(hipSetDevice(dev + i)); + HIP_CHECK(hipDeviceSynchronize()); + } + end_time[1] = std::chrono::system_clock::now(); + std::chrono::duration serialized_gpu1_time = (end_time[1] - start_time[1]); + INFO("Serialized set of three kernels with GPU1 being long took: " + << serialized_gpu1_time.count() << " seconds"); + + verify_time(single_kernel_time.count(), serialized_gpu1_time.count(), 2.7f, 3.3f); + } + + SECTION( + "GPU1 - Standard/Coop, GPU2 - Long Coop/Standard - regular and coop kernel overlap at " + "beginning") { + INFO("GPU " << dev << ": Standard/Coop with multi device no pre sync"); + INFO("GPU " << (dev + 1) << ": Long Coop/Standard with multi device no pre sync"); + fast_gpu = 0; + start_time[1] = std::chrono::system_clock::now(); + HIP_CHECK(hipSetDevice(dev)); + auto test_kernel_used = IsGfx11() ? test_kernel_gfx11 : test_kernel; + hipLaunchKernelGGL(test_kernel_used, dim3(desired_blocks), dim3(warp_size), 0, streams[0], + loops, dev_array[0]); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipLaunchCooperativeKernelMultiDevice(md_params, 2, + hipCooperativeLaunchMultiDeviceNoPreSync)); + HIP_CHECK(hipSetDevice(dev + 1)); + test_kernel_used = IsGfx11() ? test_kernel_gfx11 : test_kernel; + hipLaunchKernelGGL(test_kernel_used, dim3(desired_blocks), dim3(warp_size), 0, streams[1], + loops, dev_array[1]); + HIP_CHECK(hipGetLastError()); + for (int i = 0; i < 2; i++) { + HIP_CHECK(hipSetDevice(dev + i)); + HIP_CHECK(hipDeviceSynchronize()); + } + end_time[1] = std::chrono::system_clock::now(); + std::chrono::duration pre_overlapped_time = (end_time[1] - start_time[1]); + INFO("Multiple kernels with pre-overlap allowed took: " << pre_overlapped_time.count() + << " seconds"); + + verify_time(single_kernel_time.count(), pre_overlapped_time.count(), 1.7f, 2.3f); + } + + SECTION( + "GPU1 - Standard/Long Coop, GPU2 - Coop/Standard - regular and coop kernel overlap at " + "end") { + INFO("GPU " << dev << ": Standard/Long Coop with multi device no post sync"); + INFO("GPU " << (dev + 1) << ": Coop/Standard with multi device no post sync"); + fast_gpu = 1; + start_time[1] = std::chrono::system_clock::now(); + HIP_CHECK(hipSetDevice(dev)); + auto test_kernel_used = IsGfx11() ? test_kernel_gfx11 : test_kernel; + hipLaunchKernelGGL(test_kernel_used, dim3(desired_blocks), dim3(warp_size), 0, streams[0], + loops, dev_array[0]); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipLaunchCooperativeKernelMultiDevice(md_params, 2, + hipCooperativeLaunchMultiDeviceNoPostSync)); + HIP_CHECK(hipSetDevice(dev + 1)); + test_kernel_used = IsGfx11() ? test_kernel_gfx11 : test_kernel; + hipLaunchKernelGGL(test_kernel_used, dim3(desired_blocks), dim3(warp_size), 0, streams[1], + loops, dev_array[1]); + for (int i = 0; i < 2; i++) { + HIP_CHECK(hipSetDevice(dev + i)); + HIP_CHECK(hipDeviceSynchronize()); + } + end_time[1] = std::chrono::system_clock::now(); + std::chrono::duration post_overlapped_time = (end_time[1] - start_time[1]); + INFO("Multiple kernels with post-overlap allowed took: " << post_overlapped_time.count() + << " seconds"); + + verify_time(single_kernel_time.count(), post_overlapped_time.count(), 1.7f, 2.3f); + } + + SECTION( + "GPU1 - Standard/Long Coop, GPU2 - Long Coop/Standard - regular and coop kernel overlap") { + INFO("GPU " << dev << ": Standard/Long Coop with multi device no pre or post sync"); + INFO("GPU " << (dev + 1) << ": Long Coop/Standard with multi device no pre or post sync"); + start_time[1] = std::chrono::system_clock::now(); + HIP_CHECK(hipSetDevice(dev)); + auto test_kernel_used = IsGfx11() ? test_kernel_gfx11 : test_kernel; + hipLaunchKernelGGL(test_kernel_used, dim3(desired_blocks), dim3(warp_size), 0, streams[0], + loops, dev_array[0]); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipLaunchCooperativeKernelMultiDevice( + md_params, 2, + hipCooperativeLaunchMultiDeviceNoPreSync | hipCooperativeLaunchMultiDeviceNoPostSync)); + HIP_CHECK(hipSetDevice(dev + 1)); + test_kernel_used = IsGfx11() ? test_kernel_gfx11 : test_kernel; + hipLaunchKernelGGL(test_kernel_used, dim3(desired_blocks), dim3(warp_size), 0, streams[1], + loops, dev_array[1]); + HIP_CHECK(hipGetLastError()); + for (int i = 0; i < 2; i++) { + HIP_CHECK(hipSetDevice(dev + i)); + HIP_CHECK(hipDeviceSynchronize()); + } + end_time[1] = std::chrono::system_clock::now(); + std::chrono::duration overlapped_time = (end_time[1] - start_time[1]); + INFO("Multiple kernels with overlap allowed took: " << overlapped_time.count() << " seconds"); + + verify_time(single_kernel_time.count(), overlapped_time.count(), 1.8f, 2.2f); + } + + for (int k = 0; k < 2; ++k) { + HIP_CHECK(hipFree(dev_array[k])); + HIP_CHECK(hipStreamDestroy(streams[k])); + } + } +} + +TEST_CASE("Unit_hipLaunchCooperativeKernelMultiDevice_Basic") { + constexpr uint num_kernel_args = 4; + + int device_num = 0; + HIP_CHECK(hipGetDeviceCount(&device_num)); + + size_t buffer_size = kBufferLen * sizeof(int); + + int* A_h = reinterpret_cast(malloc(buffer_size * device_num)); + for (uint32_t i = 0; i < kBufferLen * device_num; ++i) { + A_h[i] = static_cast(i); + } + + int* A_d[device_num]; + long* B_d[device_num]; + long* C_d; + hipStream_t stream[device_num]; + + hipDeviceProp_t device_properties[device_num]; + + for (int i = 0; i < device_num; i++) { + HIP_CHECK(hipSetDevice(i)); + + // Calculate the device occupancy to know how many blocks can be run concurrently + HIP_CHECK(hipGetDeviceProperties(&device_properties[i], 0)); + if (!device_properties[i].cooperativeMultiDeviceLaunch) { + HipTest::HIP_SKIP_TEST("Device doesn't support cooperative launch!"); + return; + } + + HIP_CHECK(hipMalloc(&A_d[i], buffer_size)); + HIP_CHECK(hipMemcpy(A_d[i], &A_h[i * kBufferLen], buffer_size, hipMemcpyHostToDevice)); + if (i == 0) { + HIP_CHECK(hipHostMalloc(&C_d, (device_num + 1) * sizeof(long))); + } + + HIP_CHECK(hipStreamCreate(&stream[i])); + HIP_CHECK(hipDeviceSynchronize()); + } + + dim3 dimBlock; + dim3 dimGrid; + dimGrid.x = 1; + dimGrid.y = 1; + dimGrid.z = 1; + dimBlock.x = 64; + dimBlock.y = 1; + dimBlock.z = 1; + + int num_blocks = 0; + uint workgroup = GENERATE(64, 128, 256); + + hipLaunchParams* launch_params_list = new hipLaunchParams[device_num]; + void* args[device_num * num_kernel_args]; + + for (int i = 0; i < device_num; i++) { + HIP_CHECK(hipSetDevice(i)); + + dimBlock.x = workgroup; + HIP_CHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor( + &num_blocks, test_gws, dimBlock.x * dimBlock.y * dimBlock.z, dimBlock.x * sizeof(long))); + + INFO("GPU" << i << " has block size = " << dimBlock.x << " and num blocks per CU " << num_blocks + << "\n"); + + dimGrid.x = device_properties[i].multiProcessorCount * std::min(num_blocks, 32); + + HIP_CHECK(hipMalloc(&B_d[i], dimGrid.x * sizeof(long))); + + args[i * num_kernel_args] = (void*)&A_d[i]; + args[i * num_kernel_args + 1] = (void*)&kBufferLen; + args[i * num_kernel_args + 2] = (void*)&B_d[i]; + args[i * num_kernel_args + 3] = (void*)&C_d; + + launch_params_list[i].func = reinterpret_cast(test_gws); + launch_params_list[i].gridDim = dimGrid; + launch_params_list[i].blockDim = dimBlock; + launch_params_list[i].sharedMem = dimBlock.x * sizeof(long); + launch_params_list[i].stream = stream[i]; + launch_params_list[i].args = &args[i * num_kernel_args]; + } + + HIP_CHECK(hipLaunchCooperativeKernelMultiDevice(launch_params_list, device_num, 0)); + for (int i = 0; i < device_num; i++) { + HIP_CHECK(hipStreamSynchronize(stream[i])); + } + + size_t processed_Dwords = kBufferLen * device_num; + REQUIRE(*C_d == (((long)(processed_Dwords) * (processed_Dwords - 1)) / 2)); + + delete[] launch_params_list; + + HIP_CHECK(hipSetDevice(0)); + HIP_CHECK(hipHostFree(C_d)); + for (int i = 0; i < device_num; i++) { + HIP_CHECK(hipSetDevice(i)); + HIP_CHECK(hipFree(A_d[i])); + HIP_CHECK(hipFree(B_d[i])); + HIP_CHECK(hipStreamDestroy(stream[i])); + } + + free(A_h); +} + +TEST_CASE("Unit_hipLaunchCooperativeKernelMultiDevice_Streams") { + int device_num = 0; + HIP_CHECK(hipGetDeviceCount(&device_num)); + + if (device_num < 2) { + HipTest::HIP_SKIP_TEST("Skipping because devices < 2"); + return; + } + + hipDeviceProp_t device_properties; + for (int i = 0; i < device_num; i++) { + HIP_CHECK(hipGetDeviceProperties(&device_properties, i)); + if (!device_properties.cooperativeMultiDeviceLaunch) { + HipTest::HIP_SKIP_TEST("Device doesn't support cooperative launch!"); + return; + } + } + + test_multigrid_streams(device_num); +} diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/hipLaunchCooperativeKernel_old.cc b/projects/hip-tests/catch/unit/cooperativeGrps/hipLaunchCooperativeKernel_old.cc new file mode 100644 index 0000000000..02ed362b55 --- /dev/null +++ b/projects/hip-tests/catch/unit/cooperativeGrps/hipLaunchCooperativeKernel_old.cc @@ -0,0 +1,364 @@ +/* +Copyright (c) 2020 - 2021 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 +#include + +namespace cg = cooperative_groups; + +static constexpr size_t kBufferLen = 1024 * 1024; + +__global__ void test_gws(int* buf, size_t buf_size, long* tmp_buf, long* result) { + extern __shared__ long tmp[]; + uint offset = blockIdx.x * blockDim.x + threadIdx.x; + uint stride = gridDim.x * blockDim.x; + cg::grid_group gg = cg::this_grid(); + + long sum = 0; + for (uint i = offset; i < buf_size; i += stride) { + sum += buf[i]; + } + tmp[threadIdx.x] = sum; + + __syncthreads(); + + if (threadIdx.x == 0) { + sum = 0; + for (uint i = 0; i < blockDim.x; i++) { + sum += tmp[i]; + } + tmp_buf[blockIdx.x] = sum; + } + + gg.sync(); + + if (offset == 0) { + for (uint i = 1; i < gridDim.x; ++i) { + sum += tmp_buf[i]; + } + *result = sum; + } +} + +__global__ void test_kernel(uint32_t loops, unsigned long long* array, long long totalTicks) { + cg::thread_block tb = cg::this_thread_block(); + unsigned int rank = blockIdx.x * blockDim.x + threadIdx.x; + + for (int i = 0; i < loops; i++) { + long long time_diff = 0; + long long last_clock = clock64(); + do { + long long cur_clock = clock64(); + if (cur_clock > last_clock) { + time_diff += (cur_clock - last_clock); + } + // If it rolls over, we don't know how much to add to catch up. + // So just ignore those slipped cycles. + last_clock = cur_clock; + } while (time_diff < totalTicks); + tb.sync(); + array[rank] += clock64(); + } +} + +__global__ void test_kernel_gfx11(uint32_t loops, unsigned long long* array, long long totalTicks) { +#if HT_AMD + cg::thread_block tb = cg::this_thread_block(); + unsigned int rank = blockIdx.x * blockDim.x + threadIdx.x; + + for (int i = 0; i < loops; i++) { + long long time_diff = 0; + long long last_clock = wall_clock64(); + do { + long long cur_clock = wall_clock64(); + if (cur_clock > last_clock) { + time_diff += (cur_clock - last_clock); + } + // If it rolls over, we don't know how much to add to catch up. + // So just ignore those slipped cycles. + last_clock = cur_clock; + } while (time_diff < totalTicks); + tb.sync(); + array[rank] += wall_clock64(); + } +#endif +} + +template +static void verifyLeastCapacity(T& single_kernel_time, T& double_kernel_time, + T& triple_kernel_time) { +#if HT_AMD + // hipLaunchCooperativeKernel() follows serialization policy on AMD devices + // Test that the two cooperative kernels took roughly twice as long as the one + REQUIRE(double_kernel_time.count() >= 1.8 * single_kernel_time.count()); + REQUIRE(double_kernel_time.count() <= 2.2 * single_kernel_time.count()); +#else + // hipLaunchCooperativeKernel() doesn't follow serialization policy on NV devices + // Test that the two cooperative kernels took roughly as long as the one + REQUIRE(double_kernel_time.count() >= 0.8 * single_kernel_time.count()); + REQUIRE(double_kernel_time.count() <= 1.2 * single_kernel_time.count()); +#endif + + // Test that the three kernels together took roughly as long as the two + // cooperative kernels. + REQUIRE(triple_kernel_time.count() <= 1.1 * double_kernel_time.count()); +} + +template +static void verifyHalfCapacity(T& single_kernel_time, T& double_kernel_time, + T& triple_kernel_time) { + // Test that the two cooperative kernels took roughly twice as long as the one + REQUIRE(double_kernel_time.count() >= 1.8 * single_kernel_time.count()); + REQUIRE(double_kernel_time.count() <= 2.2 * single_kernel_time.count()); + + // Test that the three kernels together took roughly as long as the two + // cooperative kernels. + REQUIRE(triple_kernel_time.count() <= 1.1 * double_kernel_time.count()); +} + +template +static void verifyFullCapacity(T& single_kernel_time, T& double_kernel_time, + T& triple_kernel_time) { + // Test that the two cooperative kernels took roughly twice as long as the one + REQUIRE(double_kernel_time.count() >= 1.8 * single_kernel_time.count()); + REQUIRE(double_kernel_time.count() <= 2.2 * single_kernel_time.count()); + + // Test that the three kernels together took roughly 1.6 times as long as the two + // cooperative kernels. If the first 2 kernels run very fast, the third + // won't share much time with the second kernel. + REQUIRE(triple_kernel_time.count() <= 1.7 * double_kernel_time.count()); +} + +template +static void verify(int tests, T& single_kernel_time, T& double_kernel_time, T& triple_kernel_time) { + switch (tests) { + case 0: + verifyLeastCapacity(single_kernel_time, double_kernel_time, triple_kernel_time); + break; + case 1: + verifyHalfCapacity(single_kernel_time, double_kernel_time, triple_kernel_time); + break; + case 2: + verifyFullCapacity(single_kernel_time, double_kernel_time, triple_kernel_time); + break; + default: + break; + } +} + +static void test_cooperative_streams(int dev, int p_tests) { + hipStream_t streams[3]; + unsigned long long* dev_array[3]; + int loops = 1000; + + HIP_CHECK(hipSetDevice(dev)); + hipDeviceProp_t device_properties; + HIP_CHECK(hipGetDeviceProperties(&device_properties, dev)); + + // Test whether target device supports cooperative groups + if (device_properties.cooperativeLaunch == 0) { + std::cout << "Cooperative group support not available in device " << dev << std::endl; + return; + } + + // We will launch enough waves to fill up all of the GPU + int warp_size = device_properties.warpSize; + int num_sms = device_properties.multiProcessorCount; + long long totalTicks = device_properties.clockRate; + int max_blocks_per_sm = 0; + // Calculate the device occupancy to know how many blocks can be run. + auto test_kernel_used = IsGfx11() ? test_kernel_gfx11 : test_kernel; + HIP_CHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&max_blocks_per_sm, test_kernel_used, + warp_size, 0)); + int max_active_blocks = max_blocks_per_sm * num_sms; + int coop_blocks = 0; + int reg_blocks = 0; + + switch (p_tests) { + case 0: + // 1 block + coop_blocks = 1; + reg_blocks = 1; + break; + case 1: + // Half capacity + // To make sure the second kernel launched by hipLaunchCooperativeKernel + // is invoked after the first kernel finished + coop_blocks = max_active_blocks / 2 + 1; + // To make sure the third kernel launched by hipLaunchKernelGGL is invoked + // concurrently with the second kernel + reg_blocks = max_active_blocks - coop_blocks; + break; + case 2: + // Full capacity + coop_blocks = max_active_blocks; + reg_blocks = max_active_blocks; + break; + default: + break; + } + + for (int i = 0; i < 3; i++) { + HIP_CHECK(hipStreamCreate(&streams[i])); + } + + // Set up data to pass into the kernel + + for (int i = 0; i < 3; i++) { + HIP_CHECK(hipMalloc(reinterpret_cast(&dev_array[i]), warp_size * sizeof(long long))); + HIP_CHECK(hipMemsetAsync(dev_array[i], 0, warp_size * sizeof(long long), streams[i])); + } + + HIP_CHECK(hipDeviceSynchronize()); + + // Launch the kernels + void* coop_params[3][3]; + for (int i = 0; i < 3; i++) { + coop_params[i][0] = reinterpret_cast(&loops); + coop_params[i][1] = reinterpret_cast(&dev_array[i]); + coop_params[i][2] = reinterpret_cast(&totalTicks); + } + + // We need exclude the the initial launching as it will need time to load code obj. + HIP_CHECK(hipLaunchCooperativeKernel(reinterpret_cast(test_kernel_used), max_active_blocks, + warp_size, coop_params[0], 0, streams[0])); + HIP_CHECK(hipDeviceSynchronize()); + + // Launching a single cooperative kernel + auto single_start = std::chrono::system_clock::now(); + HIP_CHECK(hipLaunchCooperativeKernel(reinterpret_cast(test_kernel_used), max_active_blocks, + warp_size, coop_params[0], 0, streams[0])); + HIP_CHECK(hipDeviceSynchronize()); + auto single_end = std::chrono::system_clock::now(); + + std::chrono::duration single_kernel_time = (single_end - single_start); + + // Launching 2 cooperative kernels to different streams + auto double_start = std::chrono::system_clock::now(); + HIP_CHECK(hipLaunchCooperativeKernel(reinterpret_cast(test_kernel_used), coop_blocks, + warp_size, coop_params[0], 0, streams[0])); + HIP_CHECK(hipLaunchCooperativeKernel(reinterpret_cast(test_kernel_used), coop_blocks, + warp_size, coop_params[1], 0, streams[1])); + + HIP_CHECK(hipDeviceSynchronize()); + auto double_end = std::chrono::system_clock::now(); + + // Launching 2 cooperative kernels and 1 normal kernel + std::chrono::duration double_kernel_time = (double_end - double_start); + + auto triple_start = std::chrono::system_clock::now(); + HIP_CHECK(hipLaunchCooperativeKernel(reinterpret_cast(test_kernel_used), coop_blocks, + warp_size, coop_params[0], 0, streams[0])); + HIP_CHECK(hipLaunchCooperativeKernel(reinterpret_cast(test_kernel_used), coop_blocks, + warp_size, coop_params[1], 0, streams[1])); + hipLaunchKernelGGL(test_kernel_used, dim3(reg_blocks), dim3(warp_size), 0, streams[2], loops, + dev_array[2], totalTicks); + + HIP_CHECK(hipDeviceSynchronize()); + auto triple_end = std::chrono::system_clock::now(); + std::chrono::duration triple_kernel_time = (triple_end - triple_start); + + for (int k = 0; k < 3; ++k) { + HIP_CHECK(hipFree(dev_array[k])); + HIP_CHECK(hipStreamDestroy(streams[k])); + } + + + INFO("A single kernel took : " << single_kernel_time.count() << " seconds"); + INFO("Two cooperative kernels took: " << double_kernel_time.count() << " seconds"); + INFO("Two coop kernels and a third regular kernel took: " << triple_kernel_time.count() + << " seconds"); + + verify(p_tests, single_kernel_time, double_kernel_time, triple_kernel_time); +} + +TEST_CASE("Unit_hipLaunchCooperativeKernel_Basic") { + // Use default device for validating the test + int device; + int *A_h, *A_d; + long* B_d; + long* C_d; + 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; + } + + size_t buffer_size = kBufferLen * sizeof(int); + + A_h = reinterpret_cast(malloc(buffer_size)); + for (uint32_t i = 0; i < kBufferLen; ++i) { + A_h[i] = static_cast(i); + } + + HIP_CHECK(hipMalloc(&A_d, buffer_size)); + HIP_CHECK(hipMemcpy(A_d, A_h, buffer_size, hipMemcpyHostToDevice)); + HIP_CHECK(hipHostMalloc(&C_d, sizeof(long))); + + hipStream_t stream; + HIPCHECK(hipStreamCreate(&stream)); + + dim3 dimBlock = dim3(1); + dim3 dimGrid = dim3(1); + int numBlocks = 0; + + uint32_t workgroup = GENERATE(32, 64, 128, 256); + + dimBlock.x = workgroup; + + // Calculate the device occupancy to know how many blocks can be run concurrently + HIP_CHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor( + &numBlocks, test_gws, dimBlock.x * dimBlock.y * dimBlock.z, dimBlock.x * sizeof(long))); + + dimGrid.x = device_properties.multiProcessorCount * std::min(numBlocks, 32); + HIP_CHECK(hipMalloc(&B_d, dimGrid.x * sizeof(long))); + + void* params[4]; + params[0] = (void*)&A_d; + params[1] = (void*)&kBufferLen; + params[2] = (void*)&B_d; + params[3] = (void*)&C_d; + + INFO("Testing with grid size = " << dimGrid.x << " and block size = " << dimBlock.x << "\n"); + HIP_CHECK(hipLaunchCooperativeKernel(reinterpret_cast(test_gws), dimGrid, dimBlock, params, + dimBlock.x * sizeof(long), stream)); + + HIP_CHECK(hipStreamSynchronize(stream)); + + REQUIRE(*C_d == (((long)(kBufferLen) * (kBufferLen - 1)) / 2)); + + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipHostFree(C_d)); + HIP_CHECK(hipFree(B_d)); + HIP_CHECK(hipFree(A_d)); + free(A_h); +} + +TEST_CASE("Unit_hipLaunchCooperativeKernel_Streams") { + const auto device = GENERATE(range(0, HipTest::getDeviceCount())); + int p_tests = GENERATE(0, 1, 2); + + test_cooperative_streams(device, p_tests); +} diff --git a/projects/hip-tests/catch/unit/cooperativeGrps/hip_cg_common.hh b/projects/hip-tests/catch/unit/cooperativeGrps/hip_cg_common.hh new file mode 100644 index 0000000000..a041c3d673 --- /dev/null +++ b/projects/hip-tests/catch/unit/cooperativeGrps/hip_cg_common.hh @@ -0,0 +1,68 @@ +/* +Copyright (c) 2020 - 2021 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 + +#define ASSERT_EQUAL(lhs, rhs) HIP_ASSERT(lhs == rhs) +#define ASSERT_LE(lhs, rhs) HIPASSERT(lhs <= rhs) +#define ASSERT_GE(lhs, rhs) HIPASSERT(lhs >= rhs) + +constexpr int MaxGPUs = 8; + +template +void printResults(T* ptr, int size) { + for (int i = 0; i < size; i++) { + std::cout << ptr[i] << " "; + } + std::cout << '\n'; +} + +template +void compareResults(T* cpu, T* gpu, int size) { + for (unsigned int i = 0; i < size / sizeof(T); i++) { + if (cpu[i] != gpu[i]) { + INFO("Results do not match at index " << i); + REQUIRE(cpu[i] == gpu[i]); + } + } +} + + +// Search if the sum exists in the expected results array +template +void verifyResults(T* hPtr, T* 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) { + INFO("Result verification failed!"); + REQUIRE(j != size); + } + } +} \ No newline at end of file