EXSWHTEC-308 - Migrate and refactor cooperative groups tests from HIP repository (#238)
Change-Id: Ib46b7c038a5bda9d05f5d55a7269a7c645b0d049
[ROCm/hip-tests commit: e0612f9346]
Este cometimento está contido em:
Nives Vukovic
cometido por
Maneesh Gupta
Maneesh Gupta
ascendente
6250954456
cometimento
a3cc0be869
@@ -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}
|
||||
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <hip/hip_cooperative_groups.h>
|
||||
|
||||
#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<unsigned int*>(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 <typename F> 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<void* (*)()>(kernel_cg_grid_group_type);
|
||||
}
|
||||
#if HT_AMD
|
||||
SECTION("Base type grid group API test") {
|
||||
kernel_func = reinterpret_cast<void* (*)()>(kernel_cg_grid_group_type_via_base_type);
|
||||
}
|
||||
#endif
|
||||
|
||||
SECTION("Public API grid group test") {
|
||||
kernel_func = reinterpret_cast<void* (*)()>(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<unsigned int*>(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<unsigned int*>(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<unsigned int*>(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<void*>(&dev_mem_1);
|
||||
coop_params[1] = reinterpret_cast<void*>(&dev_mem_2);
|
||||
coop_params[2] = reinterpret_cast<void*>(&loops);
|
||||
coop_params[3] = reinterpret_cast<void*>(&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<unsigned int*>(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<void*>(&kernel_atomic);
|
||||
params[1] = reinterpret_cast<void*>(&kernel_buffer);
|
||||
params[2] = reinterpret_cast<void*>(&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);
|
||||
}
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <hip/hip_cooperative_groups.h>
|
||||
|
||||
#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<void*>(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);
|
||||
}
|
||||
}
|
||||
-234
@@ -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 <hip_test_common.hh>
|
||||
#include <hip/hip_cooperative_groups.h>
|
||||
#include <cmath>
|
||||
#include <cstdlib>
|
||||
#include <climits>
|
||||
|
||||
#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<void*>(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);
|
||||
}
|
||||
}
|
||||
-230
@@ -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 <hip_test_common.hh>
|
||||
#include <hip/hip_cooperative_groups.h>
|
||||
#include <cmath>
|
||||
#include <cstdlib>
|
||||
#include <climits>
|
||||
|
||||
#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<void*>(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);
|
||||
}
|
||||
}
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <hip/hip_cooperative_groups.h>
|
||||
|
||||
#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 <typename F>
|
||||
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<void*>(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<void* (*)()>(kernel_cg_multi_grid_group_type);
|
||||
specific_api_test = true;
|
||||
}
|
||||
|
||||
SECTION("Base type multi grid group API test") {
|
||||
kernel_func = reinterpret_cast<void* (*)()>(kernel_cg_multi_grid_group_type_via_base_type);
|
||||
}
|
||||
|
||||
SECTION("Public API multi grid group test") {
|
||||
kernel_func = reinterpret_cast<void* (*)()>(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<unsigned int*>(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<void*>(&kernel_atomic[i]);
|
||||
dev_params[i][1] = reinterpret_cast<void*>(&global_array);
|
||||
dev_params[i][2] = reinterpret_cast<void*>(&kernel_buffer[i]);
|
||||
dev_params[i][3] = reinterpret_cast<void*>(&loops);
|
||||
md_params[i].func = reinterpret_cast<void*>(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]);
|
||||
}
|
||||
}
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <hip/hip_cooperative_groups.h>
|
||||
|
||||
#include "hip_cg_common.hh"
|
||||
|
||||
namespace cg = cooperative_groups;
|
||||
|
||||
enum class TiledGroupShflTests { shflDown, shflXor, shflUp };
|
||||
|
||||
template <unsigned int tileSz>
|
||||
__device__ int reduction_kernel_shfl_down(cg::thread_block_tile<tileSz> 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 <unsigned int tileSz>
|
||||
__device__ int reduction_kernel_shfl_xor(cg::thread_block_tile<tileSz> 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 <unsigned int tileSz>
|
||||
__device__ int prefix_sum_kernel(cg::thread_block_tile<tileSz> 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 <unsigned int tile_size>
|
||||
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<tile_size> tiled_part = cg::tiled_partition<tile_size>(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 <unsigned int tile_size> 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<tile_size>, 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);
|
||||
}
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <hip/hip_cooperative_groups.h>
|
||||
#include <cstdlib>
|
||||
|
||||
#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));
|
||||
}
|
||||
}
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include "hip/hip_cooperative_groups.h"
|
||||
#include <cstdlib>
|
||||
|
||||
#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));
|
||||
}
|
||||
}
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include "hip/hip_cooperative_groups.h"
|
||||
#include <cstdlib>
|
||||
|
||||
#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));
|
||||
}
|
||||
}
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <hip/hip_cooperative_groups.h>
|
||||
|
||||
#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));
|
||||
}
|
||||
}
|
||||
@@ -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<tileSz>,
|
||||
or in runtime as tiled_partition(thread_group parent, tileSz). This test covers both these
|
||||
cases.
|
||||
This test tests functionality of cg group partitioning, (static and dynamic) and its respective
|
||||
API's size(), thread_rank(), and sync().
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip/hip_cooperative_groups.h>
|
||||
#include <stdio.h>
|
||||
#include <vector>
|
||||
|
||||
using namespace cooperative_groups;
|
||||
|
||||
/* 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 <unsigned int tileSz>
|
||||
__global__ void kernel_cg_group_partition_static(int* result, bool isGlobalMem, int* globalMem) {
|
||||
thread_block threadBlockCGTy = this_thread_block();
|
||||
int threadBlockGroupSize = threadBlockCGTy.size();
|
||||
|
||||
int* workspace = NULL;
|
||||
|
||||
if (isGlobalMem) {
|
||||
workspace = globalMem;
|
||||
} else {
|
||||
// Declare a shared memory
|
||||
extern __shared__ int sharedMem[];
|
||||
workspace = sharedMem;
|
||||
}
|
||||
|
||||
int input, outputSum, expectedOutput;
|
||||
|
||||
// we pass its own thread rank as inputs
|
||||
input = threadBlockCGTy.thread_rank();
|
||||
|
||||
expectedOutput = (threadBlockGroupSize - 1) * threadBlockGroupSize / 2;
|
||||
|
||||
outputSum = reduction_kernel(threadBlockCGTy, workspace, input);
|
||||
|
||||
// Choose a leader thread to print the results
|
||||
if (threadBlockCGTy.thread_rank() == 0) {
|
||||
printf(" Sum of all ranks 0..%d in threadBlockCooperativeGroup is %d (expected %d)\n\n",
|
||||
(int)threadBlockCGTy.size() - 1, outputSum, expectedOutput);
|
||||
printf(" Creating %d groups, of tile size %d threads:\n\n",
|
||||
(int)threadBlockCGTy.size() / tileSz, tileSz);
|
||||
}
|
||||
|
||||
threadBlockCGTy.sync();
|
||||
|
||||
thread_block_tile<tileSz> tiledPartition = tiled_partition<tileSz>(threadBlockCGTy);
|
||||
|
||||
// This offset allows each group to have its own unique area in the workspace array
|
||||
int workspaceOffset = threadBlockCGTy.thread_rank() - tiledPartition.thread_rank();
|
||||
|
||||
outputSum = reduction_kernel(tiledPartition, workspace + workspaceOffset, input);
|
||||
|
||||
if (tiledPartition.thread_rank() == 0) {
|
||||
printf(
|
||||
" Sum of all ranks 0..%d in this tiledPartition group is %d. Corresponding parent thread "
|
||||
"rank via meta_group_rank : %d and the total number of groups created when partitioned : %d\n",
|
||||
tiledPartition.size() - 1, outputSum, tiledPartition.meta_group_rank(), tiledPartition.meta_group_size());
|
||||
result[input / (tileSz)] = outputSum;
|
||||
}
|
||||
return;
|
||||
}
|
||||
|
||||
|
||||
__global__ void kernel_cg_group_partition_dynamic(unsigned int tileSz, int* result,
|
||||
bool isGlobalMem, int* globalMem) {
|
||||
thread_block threadBlockCGTy = this_thread_block();
|
||||
|
||||
int* 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 <unsigned int tileSz> static void test_group_partition(bool useGlobalMem) {
|
||||
hipError_t err;
|
||||
int blockSize = 1;
|
||||
int threadsPerBlock = 64;
|
||||
|
||||
int numTiles = (blockSize * threadsPerBlock) / tileSz;
|
||||
|
||||
// Build an array of expected reduction sum output on the host
|
||||
// based on the sum of their respective thread ranks for verification.
|
||||
// eg: parent group has 64threads.
|
||||
// child thread ranks: 0-15, 16-31, 32-47, 48-63
|
||||
// expected sum: 120, 376, 632, 888
|
||||
int* expectedSum = new int[numTiles];
|
||||
int temp = 0, sum = 0;
|
||||
|
||||
for (int i = 1; i <= numTiles; i++) {
|
||||
sum = temp;
|
||||
temp = (((tileSz * i) - 1) * (tileSz * i)) / 2;
|
||||
expectedSum[i-1] = temp - sum;
|
||||
}
|
||||
|
||||
int* dResult = NULL;
|
||||
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<tileSz>, blockSize, threadsPerBlock, 0, 0,
|
||||
dResult, useGlobalMem, globalMem);
|
||||
err = hipDeviceSynchronize();
|
||||
if (err != hipSuccess) {
|
||||
fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));
|
||||
}
|
||||
} else {
|
||||
// Launch Kernel
|
||||
hipLaunchKernelGGL(kernel_cg_group_partition_static<tileSz>, blockSize, threadsPerBlock,
|
||||
threadsPerBlock * sizeof(int), 0, dResult, useGlobalMem, globalMem);
|
||||
err = hipDeviceSynchronize();
|
||||
if (err != hipSuccess) {
|
||||
fprintf(stderr, "Failed to launch kernel (error code %s)!\n", hipGetErrorString(err));
|
||||
}
|
||||
}
|
||||
|
||||
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<unsigned int> tileSizes = {2, 4, 8, 16, 32};
|
||||
std::cout << "\nUsing global memory for computation\n";
|
||||
for (auto i : tileSizes) {
|
||||
std::cout << "TEST " << testNo << ":" << '\n' << std::endl;
|
||||
test_group_partition(i, useGlobalMem);
|
||||
testNo++;
|
||||
}
|
||||
|
||||
useGlobalMem = false;
|
||||
testNo = 1;
|
||||
std::cout << "\nUsing shared memory for computation\n";
|
||||
for (auto i : tileSizes) {
|
||||
std::cout << "TEST " << testNo << ":" << '\n' << std::endl;
|
||||
test_group_partition(i, useGlobalMem);
|
||||
testNo++;
|
||||
}
|
||||
printf("\n...PASSED.\n\n");
|
||||
return;
|
||||
}
|
||||
@@ -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<tileSz>,
|
||||
or in runtime as tiled_partition(thread_group parent, tileSz). This test covers both these
|
||||
cases.
|
||||
This test tests functionality of cg group partitioning, (static and dynamic) and its respective
|
||||
API's size(), thread_rank(), and sync().
|
||||
*/
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
#include <hip/hip_cooperative_groups.h>
|
||||
#include <cstdlib>
|
||||
|
||||
#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 <unsigned int tile_size>
|
||||
__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<tile_size> tiled_part = cg::tiled_partition<tile_size>(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<int>(tiled_part.size()) - 1, output_sum, input);
|
||||
result[input / (tile_size)] = output_sum;
|
||||
}
|
||||
return;
|
||||
}
|
||||
|
||||
template <typename F>
|
||||
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 <unsigned int tile_size> 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>, 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);
|
||||
}
|
||||
}
|
||||
+606
@@ -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 <hip_test_common.hh>
|
||||
#include <hip/hip_cooperative_groups.h>
|
||||
|
||||
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<void**>(&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<std::chrono::system_clock> start_time[2];
|
||||
std::chrono::time_point<std::chrono::system_clock> 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<void*>(&loops);
|
||||
dev_params[i][1] = reinterpret_cast<void*>(&dev_array[i]);
|
||||
dev_params[i][2] = reinterpret_cast<void*>(&fast_gpu);
|
||||
md_params[i].func = reinterpret_cast<void*>(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<double> 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<double> 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<double> 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<double> 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<double> 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<double> 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<int*>(malloc(buffer_size * device_num));
|
||||
for (uint32_t i = 0; i < kBufferLen * device_num; ++i) {
|
||||
A_h[i] = static_cast<int>(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<void*>(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);
|
||||
}
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <hip/hip_cooperative_groups.h>
|
||||
|
||||
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 <typename T>
|
||||
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 <typename T>
|
||||
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 <typename T>
|
||||
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 <typename T>
|
||||
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<void**>(&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<void*>(&loops);
|
||||
coop_params[i][1] = reinterpret_cast<void*>(&dev_array[i]);
|
||||
coop_params[i][2] = reinterpret_cast<void*>(&totalTicks);
|
||||
}
|
||||
|
||||
// We need exclude the the initial launching as it will need time to load code obj.
|
||||
HIP_CHECK(hipLaunchCooperativeKernel(reinterpret_cast<void*>(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<void*>(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<double> 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<void*>(test_kernel_used), coop_blocks,
|
||||
warp_size, coop_params[0], 0, streams[0]));
|
||||
HIP_CHECK(hipLaunchCooperativeKernel(reinterpret_cast<void*>(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> double_kernel_time = (double_end - double_start);
|
||||
|
||||
auto triple_start = std::chrono::system_clock::now();
|
||||
HIP_CHECK(hipLaunchCooperativeKernel(reinterpret_cast<void*>(test_kernel_used), coop_blocks,
|
||||
warp_size, coop_params[0], 0, streams[0]));
|
||||
HIP_CHECK(hipLaunchCooperativeKernel(reinterpret_cast<void*>(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<double> 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<int*>(malloc(buffer_size));
|
||||
for (uint32_t i = 0; i < kBufferLen; ++i) {
|
||||
A_h[i] = static_cast<int>(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<void*>(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);
|
||||
}
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <hip/hip_cooperative_groups.h>
|
||||
|
||||
#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 <typename T>
|
||||
void printResults(T* ptr, int size) {
|
||||
for (int i = 0; i < size; i++) {
|
||||
std::cout << ptr[i] << " ";
|
||||
}
|
||||
std::cout << '\n';
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
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 <typename T>
|
||||
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);
|
||||
}
|
||||
}
|
||||
}
|
||||
Criar uma nova questão referindo esta
Bloquear um utilizador