[dtest] Cooperative Groups related tests
Converted tests from below git to hipdirected_tests http://gitlab1.amd.com/jgreatho/cuda_cooperative_groups_test Modified to cover multi-GPUs Disabled tests for cuda because of some header file issues SWDEV-238517 for enhancing hip unit tests Change-Id: If35fd710e8ab61debcf66bca5b6503539c567ec1
This commit is contained in:
Lakhan singh Thakur
committed by
Mohan Kumar Mithur
Mohan Kumar Mithur
parent
449848cb5a
commit
40ca4a5ea8
@@ -0,0 +1,280 @@
|
||||
/*
|
||||
Copyright (c) 2020-present 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:
|
||||
/*The general idea of the application is to test how Cooperative Groups kernel
|
||||
launches work when launching too many warps to the target device. This test
|
||||
first queries the nominal warp size of the target device. It then walks through
|
||||
block sizes from 1 thread, 1 warp, 2 warps, ... `maximum_warps_in_a_block`. For
|
||||
each of these, it queries the maximum number of blocks that can fit in each SM.
|
||||
It then queries the number of SMs on the target device. This will yield a
|
||||
calculation for the maximum number of blocks that can be co-scheduled on this
|
||||
device.
|
||||
|
||||
The Cooperative Groups API says that users should not launch more than this
|
||||
many warps (or blocks, etc.) to the target device. This test first tires to
|
||||
launch 2x as many blcoks, to confirm that the runtime prevents such a launch
|
||||
by returning a proper error value (`hipErrorCooperativeLaunchTooLarge`).
|
||||
|
||||
It then ensures that trying to launch too large of a kernel invocation does
|
||||
not break the GPU by launching a kernel with exactly the maximum number of
|
||||
blocks.
|
||||
|
||||
Finally, we run the same test for a block size that is larger than the maximum
|
||||
allowed by the device, to ensure that this case is properly detected by the
|
||||
runtime and that nothing breaks.*/
|
||||
|
||||
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <hip/hip_cooperative_groups.h>
|
||||
#include "test_common.h"
|
||||
|
||||
|
||||
static inline void hipCheckAndFail(hipError_t errval,
|
||||
const char *file, int line) {
|
||||
hipError_t last_err = hipGetLastError();
|
||||
if (errval != hipSuccess) {
|
||||
std::cerr << "hip error: " << hipGetErrorString(errval);
|
||||
std::cerr << std::endl;
|
||||
std::cerr << " Location: " << file << ":" << line << std::endl;
|
||||
failed("");
|
||||
}
|
||||
if (last_err != errval) {
|
||||
std::cerr << "Error: the return value of a function was not the same ";
|
||||
std::cerr << "as the value returned by hipGetLastError()" << std::endl;
|
||||
std::cerr << " Location: " << file << ":" << line << std::endl;
|
||||
std::cerr << " Function returned: " << hipGetErrorString(errval);
|
||||
std::cerr << " (" << errval << ")" << std::endl;
|
||||
std::cerr << "hipGetLastError() returned: " << hipGetErrorString(last_err);
|
||||
std::cerr << " (" << last_err << ")" << std::endl;
|
||||
failed("");
|
||||
}
|
||||
}
|
||||
#define hipCheckErr(errval) \
|
||||
do { hipCheckAndFail((errval), __FILE__, __LINE__); } while (0)
|
||||
|
||||
static inline bool hipCheckExpected(hipError_t errval,
|
||||
hipError_t expected_err, const char *file, int line) {
|
||||
hipError_t last_err = hipGetLastError();
|
||||
if (errval != expected_err) {
|
||||
std::cerr << "hip error: " << hipGetErrorString(errval);
|
||||
std::cerr << std::endl;
|
||||
std::cerr << " Location: " << file << ":" << line << std::endl;
|
||||
return false;
|
||||
}
|
||||
if (last_err != errval) {
|
||||
std::cerr << "Error: the return value of a function was not the same ";
|
||||
std::cerr << "as the value returned by hipGetLastError()" << std::endl;
|
||||
std::cerr << " Location: " << file << ":" << line << std::endl;
|
||||
std::cerr << " Function returned: " << hipGetErrorString(errval);
|
||||
std::cerr << " (" << errval << ")" << std::endl;
|
||||
std::cerr << "hipGetLastError() returned: " << hipGetErrorString(last_err);
|
||||
std::cerr << " (" << last_err << ")" << std::endl;
|
||||
return false;
|
||||
}
|
||||
return true;
|
||||
}
|
||||
|
||||
static bool cooperative_groups_support(int device_id) {
|
||||
hipError_t err;
|
||||
int cooperative_attribute;
|
||||
HIPCHECK(hipDeviceGetAttribute(&cooperative_attribute,
|
||||
hipDeviceAttributeCooperativeLaunch, device_id));
|
||||
if (!cooperative_attribute) {
|
||||
std::cerr << "Cooperative launch support not available in ";
|
||||
std::cerr << "the device attribute for device " << device_id;
|
||||
std::cerr << std::endl;
|
||||
return false;
|
||||
}
|
||||
hipDeviceProp_t device_properties;
|
||||
HIPCHECK(hipGetDeviceProperties(&device_properties, device_id));
|
||||
if (device_properties.cooperativeLaunch == 0) {
|
||||
std::cerr << "Cooperative group support not available in ";
|
||||
std::cerr << "device properties." << std::endl;
|
||||
return false;
|
||||
}
|
||||
return true;
|
||||
}
|
||||
|
||||
__global__ void test_kernel(long long *array) {
|
||||
unsigned int rank = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
array[rank] += clock64();
|
||||
}
|
||||
|
||||
int main(int argc, char** argv) {
|
||||
hipError_t err;
|
||||
int device_num, FailFlag = 0;
|
||||
// Alocate the host input buffer, and two device-focused buffers that we
|
||||
// will use for our test.
|
||||
unsigned int *dev_array[2];
|
||||
HIPCHECK(hipGetDeviceCount(&device_num));
|
||||
for (int dev = 0; dev < device_num; ++dev) {
|
||||
/*************************************************************************/
|
||||
/* Test whether target device supports cooperative groups ****************/
|
||||
HIPCHECK(hipSetDevice(dev));
|
||||
if (!cooperative_groups_support(dev)) {
|
||||
std::cout << "Skipping the test with Pass result.\n";
|
||||
passed();
|
||||
}
|
||||
|
||||
/*************************************************************************/
|
||||
/* Create the streams we will use in this test. **************************/
|
||||
hipStream_t streams[2];
|
||||
for (int i = 0; i < 2; i++) {
|
||||
HIPCHECK(hipStreamCreate(&streams[i]));
|
||||
}
|
||||
|
||||
/*************************************************************************/
|
||||
/* We will try to launch more waves than the GPU can fit. ***************/
|
||||
hipDeviceProp_t device_properties;
|
||||
HIPCHECK(hipGetDeviceProperties(&device_properties, dev));
|
||||
int warp_size = device_properties.warpSize;
|
||||
int num_sms = device_properties.multiProcessorCount;
|
||||
int max_num_threads = device_properties.maxThreadsPerBlock;
|
||||
|
||||
// Check single-thread block, all numbers of warps, then too-large block
|
||||
for (int block_size = 0; block_size <= (max_num_threads + warp_size);
|
||||
block_size += warp_size) {
|
||||
if (block_size == 0) {
|
||||
block_size = 1;
|
||||
}
|
||||
int max_blocks_per_sm;
|
||||
// Calculate the device occupancy to know how many blocks can be run.
|
||||
HIPCHECK(hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
|
||||
&max_blocks_per_sm, test_kernel, block_size, 0,
|
||||
hipOccupancyDefault));
|
||||
|
||||
if ((block_size > max_num_threads) && (max_blocks_per_sm != 0)) {
|
||||
std::cerr << "ERROR! Occupancy API indicated that we can have >0 ";
|
||||
std::cerr << "blocks in a kernel when the block size is too large ";
|
||||
std::cerr << "to work on the device." << std::endl;
|
||||
std::cerr << "This is incorrect, and could possibly lead users ";
|
||||
std::cerr << "to try to launch kernels that will fail." << std::endl;
|
||||
//failed("");
|
||||
FailFlag = 1;
|
||||
break;
|
||||
}
|
||||
|
||||
int desired_blocks = max_blocks_per_sm * num_sms;
|
||||
bool expect_fail = false;
|
||||
if (desired_blocks == 0) {
|
||||
desired_blocks = 1;
|
||||
expect_fail = true;
|
||||
}
|
||||
|
||||
/**********************************************************************/
|
||||
/* Set up data to pass into the kernel ********************************/
|
||||
|
||||
for (int i = 0; i < 2; i++) {
|
||||
int test_size;
|
||||
// Case where we expect to fail at launch.
|
||||
if (i == 0) {
|
||||
test_size = 2 * desired_blocks;
|
||||
} else {
|
||||
test_size = desired_blocks;
|
||||
}
|
||||
HIPCHECK(hipMalloc(reinterpret_cast<void**>(&dev_array[i]),
|
||||
test_size * block_size * sizeof(long long)));
|
||||
HIPCHECK(hipMemsetAsync(dev_array[i], 0,
|
||||
test_size * block_size * sizeof(long long),
|
||||
streams[i]));
|
||||
}
|
||||
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
/***********************************************************************/
|
||||
/* Launch the kernels **************************************************/
|
||||
void *coop_params[2][1];
|
||||
for (int i = 0; i < 2; i++) {
|
||||
coop_params[i][0] = reinterpret_cast<void*>(&dev_array[i]);
|
||||
}
|
||||
|
||||
err = hipLaunchCooperativeKernel(reinterpret_cast<void*>(test_kernel),
|
||||
2 * desired_blocks, block_size,
|
||||
coop_params[0], 0, streams[0]);
|
||||
|
||||
hipError_t expect_to_see;
|
||||
if (expect_fail) {
|
||||
expect_to_see = hipErrorInvalidConfiguration;
|
||||
} else {
|
||||
expect_to_see = hipErrorCooperativeLaunchTooLarge;
|
||||
}
|
||||
if (!hipCheckExpected(err, expect_to_see, __FILE__, __LINE__)) {
|
||||
std::cerr << "ERROR! Tried to launch a cooperative kernel with ";
|
||||
std::cerr << "too many warps." << std::endl;
|
||||
std::cerr << "This SHOULD have failed with the error ";
|
||||
std::cerr << hipGetErrorString(expect_to_see);
|
||||
std::cerr << " (" << expect_to_see << ")." << std::endl;
|
||||
std::cerr << "Instead, the launch returned " << hipGetErrorName(err);
|
||||
std::cerr << " (" << err << ")" << std::endl;
|
||||
FailFlag = 1;
|
||||
break;
|
||||
}
|
||||
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
err = hipLaunchCooperativeKernel(reinterpret_cast<void*>(test_kernel),
|
||||
desired_blocks, block_size,
|
||||
coop_params[1], 0, streams[1]);
|
||||
|
||||
if (expect_fail) {
|
||||
expect_to_see = hipErrorInvalidConfiguration;
|
||||
} else {
|
||||
expect_to_see = hipSuccess;
|
||||
}
|
||||
if (!hipCheckExpected(err, expect_to_see, __FILE__, __LINE__)) {
|
||||
std::cerr << "ERROR! Tried to launch a cooperative kernel ";
|
||||
std::cerr << "with a normal size, but a block size of ";
|
||||
std::cerr << desired_blocks << std::endl;
|
||||
std::cerr << "This SHOULD have returned ";
|
||||
std::cerr << hipGetErrorString(expect_to_see);
|
||||
std::cerr << " (" << expect_to_see << ")." << std::endl;
|
||||
std::cerr << "Instead, the launch returned " << hipGetErrorName(err);
|
||||
std::cerr << " (" << err << ")" << std::endl;
|
||||
FailFlag = 1;
|
||||
break;
|
||||
}
|
||||
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
if (block_size == 1) {
|
||||
block_size = 0;
|
||||
}
|
||||
for (int m = 0; m < 2; ++m) {
|
||||
HIPCHECK(hipFree(dev_array[m]));
|
||||
}
|
||||
}
|
||||
for (int m = 0; m < 2; ++m) {
|
||||
HIPCHECK(hipStreamDestroy(streams[m]));
|
||||
}
|
||||
if (FailFlag == 1) {
|
||||
for (int m = 0; m < 2; ++m) {
|
||||
HIPCHECK(hipFree(dev_array[m]));
|
||||
}
|
||||
failed("");
|
||||
}
|
||||
}
|
||||
passed();
|
||||
}
|
||||
@@ -0,0 +1,283 @@
|
||||
/*
|
||||
Copyright (c) 2020-present 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 Cooperative Groups kernel
|
||||
launches to a stream interact with other kernels being launched to different
|
||||
streams.
|
||||
|
||||
For example: the HIP runtime will force cooperative kernel launches to run
|
||||
serially, even if they are launched to different streams. However,
|
||||
cooperative kernel launches can run in parallel with regular kernels that
|
||||
are launched to other streams. This limitation is so that the cooperative
|
||||
kernels do not conflict with one another for resources and potentially
|
||||
deadlock the system.
|
||||
|
||||
As such, this benchmark tests three situations:
|
||||
|
||||
1. Launching a cooperative kernel by itself to stream[0]
|
||||
2. Launching two cooperative kernels in parallel to stream[0] and stream[1]
|
||||
3. Launching two cooperative kernels in parallel to stream[0] and stream[1]
|
||||
and launching a third non-cooperative kernel to stream[2]
|
||||
|
||||
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 test #2 takes roughly twice as long as #1, that implies that cooperative
|
||||
kernels are properly serialized with each other by the runtime.
|
||||
|
||||
If test #3 takes the same amount of time as test #2, that implies that
|
||||
regular kernels can properly run in parallel with cooperative kernels.
|
||||
*/
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include <chrono>
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <hip/hip_cooperative_groups.h>
|
||||
#include "test_common.h"
|
||||
|
||||
static inline void hipCheckAndFail(hipError_t errval,
|
||||
const char *file, int line) {
|
||||
hipError_t last_err = hipGetLastError();
|
||||
if (errval != hipSuccess) {
|
||||
std::cerr << "hip error: " << hipGetErrorString(errval);
|
||||
std::cerr << std::endl;
|
||||
std::cerr << "Location: " << file << ":" << line << std::endl;
|
||||
failed("");
|
||||
}
|
||||
if (last_err != errval) {
|
||||
std::cerr << "Error: the return value of a function was not the same ";
|
||||
std::cerr << "as the value returned by hipGetLastError()" << std::endl;
|
||||
std::cerr << "Location: " << file << ":" << line << std::endl;
|
||||
std::cerr << "Function returned: " << hipGetErrorString(errval);
|
||||
std::cerr << " (" << errval << ")" << std::endl;
|
||||
std::cerr << "hipGetLastError() returned: " << hipGetErrorString(last_err);
|
||||
std::cerr << " (" << last_err << ")" << std::endl;
|
||||
failed("");
|
||||
}
|
||||
}
|
||||
#define hipCheckErr(errval) \
|
||||
do { hipCheckAndFail((errval), __FILE__, __LINE__); } while (0)
|
||||
|
||||
static int cooperative_groups_support(int device_id) {
|
||||
hipError_t err;
|
||||
int cooperative_attribute;
|
||||
HIPCHECK(hipDeviceGetAttribute(&cooperative_attribute,
|
||||
hipDeviceAttributeCooperativeLaunch, device_id));
|
||||
if (!cooperative_attribute) {
|
||||
std::cerr << "Cooperative launch support not available in ";
|
||||
std::cerr << "the device attribute for device " << device_id;
|
||||
std::cerr << std::endl;
|
||||
return 0;
|
||||
}
|
||||
|
||||
hipDeviceProp_t device_properties;
|
||||
HIPCHECK(hipGetDeviceProperties(&device_properties, device_id));
|
||||
if (device_properties.cooperativeLaunch == 0) {
|
||||
std::cerr << "Cooperative group support not available in ";
|
||||
std::cerr << "device properties." << std::endl;
|
||||
return 0;
|
||||
}
|
||||
return 1;
|
||||
}
|
||||
|
||||
__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 start_clock = clock64();
|
||||
while (clock64() < (start_clock+1000000)) {}
|
||||
array[rank] += clock64();
|
||||
}
|
||||
}
|
||||
|
||||
int main(int argc, char** argv) {
|
||||
hipError_t err;
|
||||
/*************************************************************************/
|
||||
int device_num = 0, loops = 1000, FailFlag = 0;
|
||||
/* Create the streams we will use in this test. **************************/
|
||||
hipStream_t streams[3];
|
||||
// Alocate the host input buffer, and two device-focused buffers that we
|
||||
// will use for our test.
|
||||
unsigned long long *dev_array[3];
|
||||
HIPCHECK(hipGetDeviceCount(&device_num));
|
||||
for (int dev = 0; dev < device_num; ++dev) {
|
||||
/*************************************************************************/
|
||||
/* Test whether target device supports cooperative groups ****************/
|
||||
HIPCHECK(hipSetDevice(dev));
|
||||
if (!cooperative_groups_support(dev)) {
|
||||
std::cout << "Skipping the test with Pass result.\n";
|
||||
passed();
|
||||
}
|
||||
|
||||
/*************************************************************************/
|
||||
/* We will launch enough waves to fill up all of the GPU *****************/
|
||||
hipDeviceProp_t device_properties;
|
||||
HIPCHECK(hipGetDeviceProperties(&device_properties, dev));
|
||||
int warp_size = device_properties.warpSize;
|
||||
int num_sms = device_properties.multiProcessorCount;
|
||||
int desired_blocks = 1;
|
||||
std::cout << "Device: " << dev << std::endl;
|
||||
std::cout << "Device name: " << device_properties.name << std::endl;
|
||||
|
||||
int max_blocks_per_sm;
|
||||
// Calculate the device occupancy to know how many blocks can be run.
|
||||
HIPCHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&max_blocks_per_sm,
|
||||
test_kernel,
|
||||
warp_size, 0));
|
||||
|
||||
if (desired_blocks > max_blocks_per_sm * num_sms) {
|
||||
std::cerr << "The requested number of blocks will not fit on the GPU";
|
||||
std::cerr << std::endl;
|
||||
std::cerr << "You requested " << desired_blocks << " but we can only ";
|
||||
std::cerr << "fit " << (max_blocks_per_sm * num_sms) << std::endl;
|
||||
failed("");
|
||||
}
|
||||
|
||||
/*************************************************************************/
|
||||
for (int i = 0; i < 3; i++) {
|
||||
HIPCHECK(hipStreamCreate(&streams[i]));
|
||||
}
|
||||
|
||||
/*************************************************************************/
|
||||
/* Set up data to pass into the kernel ***********************************/
|
||||
|
||||
for (int i = 0; i < 3; i++) {
|
||||
HIPCHECK(hipMalloc(reinterpret_cast<void**>(&dev_array[i]),
|
||||
warp_size * sizeof(long long)));
|
||||
HIPCHECK(hipMemsetAsync(dev_array[i], 0, warp_size * sizeof(long long),
|
||||
streams[i]));
|
||||
}
|
||||
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
/*************************************************************************/
|
||||
/* Launch the kernels ****************************************************/
|
||||
void *coop_params[3][2];
|
||||
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]);
|
||||
}
|
||||
|
||||
std::cout << "Launching a single cooperative kernel..." << std::endl;
|
||||
auto single_start = std::chrono::system_clock::now();
|
||||
HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast<void*>(test_kernel),
|
||||
desired_blocks, warp_size,
|
||||
coop_params[0], 0, streams[0]));
|
||||
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
auto single_end = std::chrono::system_clock::now();
|
||||
std::cout << "Launching 2 cooperative kernels to different streams...";
|
||||
std::cout << std::endl;
|
||||
|
||||
auto double_start = std::chrono::system_clock::now();
|
||||
HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast<void*>(test_kernel),
|
||||
desired_blocks, warp_size,
|
||||
coop_params[0], 0, streams[0]));
|
||||
HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast<void*>(test_kernel),
|
||||
desired_blocks, warp_size,
|
||||
coop_params[1], 0, streams[1]));
|
||||
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
auto double_end = std::chrono::system_clock::now();
|
||||
std::cout << "Launching 2 cooperative kernels and 1 normal kernel...";
|
||||
std::cout << std::endl;
|
||||
|
||||
auto triple_start = std::chrono::system_clock::now();
|
||||
HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast<void*>(test_kernel),
|
||||
desired_blocks, warp_size,
|
||||
coop_params[0], 0, streams[0]));
|
||||
HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast<void*>(test_kernel),
|
||||
desired_blocks, warp_size,
|
||||
coop_params[1], 0, streams[1]));
|
||||
hipLaunchKernelGGL(test_kernel, dim3(desired_blocks), dim3(warp_size),
|
||||
0, streams[2], loops, dev_array[2]);
|
||||
err = hipGetLastError();
|
||||
hipCheckErr(err);
|
||||
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
auto triple_end = std::chrono::system_clock::now();
|
||||
std::chrono::duration<double> single_kernel_time =
|
||||
(single_end - single_start);
|
||||
std::chrono::duration<double> double_kernel_time =
|
||||
(double_end - double_start);
|
||||
std::chrono::duration<double> triple_kernel_time =
|
||||
(triple_end - triple_start);
|
||||
|
||||
std::cout << "A single kernel took:" << std::endl;
|
||||
std::cout << " " << single_kernel_time.count();
|
||||
std::cout << " seconds" << std::endl;
|
||||
std::cout << std::endl;
|
||||
std::cout << "Two cooperative kernels that could run together took:";
|
||||
std::cout << std::endl;
|
||||
std::cout << " " << double_kernel_time.count();
|
||||
std::cout << " seconds" << std::endl;
|
||||
std::cout << std::endl;
|
||||
std::cout << "Two coop kernels and a third regular kernel took:";
|
||||
std::cout << std::endl << " ";
|
||||
std::cout << triple_kernel_time.count();
|
||||
std::cout << " seconds" << std::endl;
|
||||
|
||||
std::cout << "Testing whether these times make sense.." << std::endl;
|
||||
// Test that two cooperative kernels is roughly twice as long as one
|
||||
if (double_kernel_time < 1.8 * single_kernel_time) {
|
||||
std::cerr << "ERROR!" << std::endl;
|
||||
std::cerr << "Two cooperative kernels launched at the same ";
|
||||
std::cerr << "time did not take roughly twice as long as a single ";
|
||||
std::cerr << "cooperative kernel." << std::endl;
|
||||
std::cerr << "Were they truly serialized?" << std::endl;
|
||||
FailFlag = 1;
|
||||
break;
|
||||
}
|
||||
|
||||
// Test that the three kernels together took roughly as long as two
|
||||
// cooperative kernels.
|
||||
if (triple_kernel_time > 1.1 * double_kernel_time) {
|
||||
std::cerr << "ERROR!" << std::endl;
|
||||
std::cerr << "Launching a normal kernel in parallel with two ";
|
||||
std::cerr << "back-to-back cooperative kernels still ended up taking ";
|
||||
std::cerr << "more than 10% longer than the two cooperative kernels ";
|
||||
std::cerr << "alone." << std::endl;
|
||||
std::cerr << "Is the normal kernel being serialized with the ";
|
||||
std::cerr << "cooperative kernels on different streams?" << std::endl;
|
||||
FailFlag = 1;
|
||||
break;
|
||||
}
|
||||
for (int k = 0; k < 3; ++k) {
|
||||
HIPCHECK(hipFree(dev_array[k]));
|
||||
HIPCHECK(hipStreamDestroy(streams[k]));
|
||||
}
|
||||
}
|
||||
if (FailFlag == 1) {
|
||||
for (int k = 0; k < 3; ++k) {
|
||||
HIPCHECK(hipFree(dev_array[k]));
|
||||
HIPCHECK(hipStreamDestroy(streams[k]));
|
||||
}
|
||||
failed("");
|
||||
}
|
||||
passed();
|
||||
}
|
||||
@@ -0,0 +1,303 @@
|
||||
/*
|
||||
Copyright (c) 2020-present 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 create a buffer of width N. N is a
|
||||
command line parameter, and the user will need to make sure that we can fit
|
||||
two buffers of N unsigned integers onto the target GPU at the same time.
|
||||
|
||||
We then launch a fixed number of warps to the GPU. This number is calculated
|
||||
to fill the GPU with as many warps as can simultaneously run on the GPU.
|
||||
The threads in these warps then walk over two arrays. First, values from
|
||||
A[offset] are added into B[offset]. After all of A is added into all of B
|
||||
in this element-wise manner, all of the waves barrier with one another.
|
||||
|
||||
After the barrier, the waves start adding values from B[mirror_offset] into
|
||||
A[offset]. Mirror offset means that the wave that is writing into A[7] is
|
||||
reading from B[7 before the last value]. This was probably written by a
|
||||
different thread before the barrier.
|
||||
|
||||
After going through this loop a certain number of times, the kernel ends and
|
||||
we read the arrays back out and recalculate this algorithm serially on the
|
||||
CPU. We compare the serial version to the version that has inter-thread data
|
||||
sharing and barriers and ensure they result in the same answer.
|
||||
|
||||
If they do have the same answer, then we can pretty confidently say that
|
||||
writing from thread X and then hitting a barrier allows thread Y to see the
|
||||
values.*/
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <hip/hip_cooperative_groups.h>
|
||||
#include "test_common.h"
|
||||
|
||||
static inline void hipCheckAndFail(hipError_t errval,
|
||||
const char *file, int line) {
|
||||
hipError_t last_err = hipGetLastError();
|
||||
if (errval != hipSuccess) {
|
||||
std::cerr << "hip error: " << hipGetErrorString(errval);
|
||||
std::cerr << std::endl;
|
||||
std::cerr << " Location: " << file << ":" << line << std::endl;
|
||||
exit(errval);
|
||||
}
|
||||
if (last_err != errval) {
|
||||
std::cerr << "Error: the return value of a function was not the same ";
|
||||
std::cerr << "as the value returned by hipGetLastError()" << std::endl;
|
||||
std::cerr << " Location: " << file << ":" << line << std::endl;
|
||||
std::cerr << " Function returned: " << hipGetErrorString(errval);
|
||||
std::cerr << " (" << errval << ")" << std::endl;
|
||||
std::cerr << "hipGetLastError() returned: " << hipGetErrorString(last_err);
|
||||
std::cerr << " (" << last_err << ")" << std::endl;
|
||||
failed("");
|
||||
}
|
||||
}
|
||||
#define hipCheckErr(errval)\
|
||||
do { hipCheckAndFail((errval), __FILE__, __LINE__); } while (0)
|
||||
|
||||
static int cooperative_groups_support(int device_id) {
|
||||
hipError_t err;
|
||||
|
||||
int cooperative_attribute;
|
||||
HIPCHECK(hipDeviceGetAttribute(&cooperative_attribute,
|
||||
hipDeviceAttributeCooperativeLaunch, device_id));
|
||||
if (!cooperative_attribute) {
|
||||
std::cerr << "Cooperative launch support not available in ";
|
||||
std::cerr << "the device attribute for device " << device_id;
|
||||
std::cerr << std::endl;
|
||||
return 0;
|
||||
}
|
||||
|
||||
hipDeviceProp_t device_properties;
|
||||
HIPCHECK(hipGetDeviceProperties(&device_properties, device_id));
|
||||
if (device_properties.cooperativeLaunch == 0) {
|
||||
std::cerr << "Cooperative group support not available in ";
|
||||
std::cerr << "device properties." << std::endl;
|
||||
return 0;
|
||||
}
|
||||
return 1;
|
||||
}
|
||||
|
||||
static int verify_coop_arrays(unsigned int loops, unsigned int *host_input,
|
||||
unsigned int *first_array,
|
||||
unsigned int *second_array,
|
||||
unsigned int array_len) {
|
||||
unsigned int *host_first_array = host_input;
|
||||
unsigned int *host_second_array = (unsigned int*)calloc(array_len,
|
||||
sizeof(int));
|
||||
|
||||
for (int i = 0; i < loops; i++) {
|
||||
for (int offset = 0; offset < array_len; offset++) {
|
||||
host_second_array[offset] += host_first_array[offset];
|
||||
}
|
||||
|
||||
for (int offset = 0; offset < array_len; offset++) {
|
||||
unsigned int swizzle_offset = array_len - offset - 1;
|
||||
host_first_array[offset] += host_second_array[swizzle_offset];
|
||||
}
|
||||
}
|
||||
|
||||
for (int i = 0; i < array_len; i++) {
|
||||
if (host_first_array[i] != first_array[i]) {
|
||||
std::cerr << "Test failure!" << std::endl;
|
||||
std::cerr << " host_first_array[" << i << "] contains the ";
|
||||
std::cerr << "value " << host_first_array[i] << std::endl;
|
||||
std::cerr << " GPU first_array[" << i << "] contains the ";
|
||||
std::cerr << "value " << first_array[i] << std::endl;
|
||||
return -1;
|
||||
}
|
||||
if (host_second_array[i] != second_array[i]) {
|
||||
std::cerr << "Test failure!" << std::endl;
|
||||
std::cerr << " host_second_array[" << i << "] contains the ";
|
||||
std::cerr << "value " << host_second_array[i] << std::endl;
|
||||
std::cerr << " GPU second_array[" << i << "] contains the ";
|
||||
std::cerr << "value " << second_array[i] << std::endl;
|
||||
return -1;
|
||||
}
|
||||
}
|
||||
|
||||
std::cout << "Coop test appears to work properly!" << std::endl;
|
||||
free(host_second_array);
|
||||
return 0;
|
||||
}
|
||||
|
||||
__global__ void
|
||||
coop_kernel(unsigned int *first_array, unsigned int *second_array,
|
||||
unsigned int loops, unsigned int array_len) {
|
||||
cooperative_groups::grid_group grid = cooperative_groups::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();
|
||||
}
|
||||
}
|
||||
|
||||
int main(int argc, char** argv) {
|
||||
hipError_t err;
|
||||
/*************************************************************************/
|
||||
/* Parse the command line parameters *************************************/
|
||||
// Arguments to pull out of the command line.
|
||||
int device_num = 0, loops = 2, width = 4096, flag = 0;
|
||||
HIPCHECK(hipGetDeviceCount(&device_num));
|
||||
for (int dev = 0; dev < device_num; ++dev) {
|
||||
std::cout << "Device number: " << dev << std::endl;
|
||||
std::cout << "Loops: " << loops << std::endl;
|
||||
std::cout << "Width: " << width << std::endl;
|
||||
|
||||
/*************************************************************************/
|
||||
/* Test whether target device supports cooperative groups ****************/
|
||||
HIPCHECK(hipSetDevice(dev));
|
||||
|
||||
if (!cooperative_groups_support(dev)) {
|
||||
std::cout << "Skipping the test with Pass result.\n";
|
||||
passed();
|
||||
}
|
||||
|
||||
/*************************************************************************/
|
||||
/* We will launch enough waves to fill up all of the GPU *****************/
|
||||
hipDeviceProp_t device_properties;
|
||||
HIPCHECK(hipGetDeviceProperties(&device_properties, dev));
|
||||
|
||||
int warp_size = device_properties.warpSize;
|
||||
int num_sms = device_properties.multiProcessorCount;
|
||||
|
||||
std::cout << "Device name: " << device_properties.name << std::endl;
|
||||
std::cout << std::endl;
|
||||
|
||||
// Calculate the device occupancy to know how many blocks can be run.
|
||||
int max_blocks_per_sm;
|
||||
HIPCHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&max_blocks_per_sm,
|
||||
coop_kernel,
|
||||
warp_size, 0));
|
||||
|
||||
int total_blocks = max_blocks_per_sm * num_sms;
|
||||
|
||||
/*************************************************************************/
|
||||
/* Create the streams we will use in this test. **************************/
|
||||
hipStream_t streams[2];
|
||||
for (int i = 0; i < 2; i++) {
|
||||
HIPCHECK(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 int *input_buffer = (unsigned int*)calloc(width,
|
||||
sizeof(unsigned int));
|
||||
for (int i = 0; i < width; i++) {
|
||||
input_buffer[i] = i;
|
||||
}
|
||||
|
||||
unsigned int *first_dev_array;
|
||||
HIPCHECK(hipMalloc(reinterpret_cast<void**>(&first_dev_array),
|
||||
width * sizeof(unsigned int)));
|
||||
|
||||
HIPCHECK(hipMemcpyAsync(first_dev_array, input_buffer,
|
||||
width * sizeof(unsigned int),
|
||||
hipMemcpyHostToDevice, streams[0]));
|
||||
|
||||
unsigned int *second_dev_array;
|
||||
HIPCHECK(hipMalloc(reinterpret_cast<void**>(&second_dev_array),
|
||||
width * sizeof(unsigned int)));
|
||||
HIPCHECK(hipMemsetAsync(second_dev_array, 0, width * sizeof(unsigned int),
|
||||
streams[0]));
|
||||
|
||||
/*************************************************************************/
|
||||
/* Launch the kernels ****************************************************/
|
||||
std::cout << "Launching a cooperative kernel with " << total_blocks;
|
||||
std::cout << " thread blocks, each with " << warp_size << " threads";
|
||||
std::cout << std::endl;
|
||||
|
||||
void *coop_params[4];
|
||||
coop_params[0] = reinterpret_cast<void*>(&first_dev_array);
|
||||
coop_params[1] = reinterpret_cast<void*>(&second_dev_array);
|
||||
coop_params[2] = reinterpret_cast<void*>(&loops);
|
||||
coop_params[3] = reinterpret_cast<void*>(&width);
|
||||
HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast<void*>(coop_kernel),
|
||||
total_blocks, warp_size, coop_params,
|
||||
0, streams[0]));
|
||||
|
||||
/*************************************************************************/
|
||||
/* Read back the buffers and print out their data ************************/
|
||||
unsigned int *first_array = (unsigned int*)calloc(width,
|
||||
sizeof(unsigned int));
|
||||
unsigned int *second_array = (unsigned int*)calloc(width,
|
||||
sizeof(unsigned int));
|
||||
HIPCHECK(hipMemcpyAsync(first_array, first_dev_array,
|
||||
width * sizeof(unsigned int),
|
||||
hipMemcpyDeviceToHost, streams[0]));
|
||||
|
||||
HIPCHECK(hipMemcpyAsync(second_array, second_dev_array,
|
||||
width * sizeof(unsigned int),
|
||||
hipMemcpyDeviceToHost, streams[0]));
|
||||
|
||||
std::cout << "Waiting for cooperative work to finish..." << std::endl;
|
||||
std::cout << std::flush;
|
||||
|
||||
HIPCHECK(hipStreamSynchronize(streams[0]));
|
||||
|
||||
|
||||
int ret_val = 0;
|
||||
|
||||
std::cout << "Attemping to verify buffers." << std::endl;
|
||||
std::cout << std::flush;
|
||||
ret_val = verify_coop_arrays(loops, input_buffer, first_array,
|
||||
second_array, width);
|
||||
if (!ret_val) {
|
||||
std::cout << "It appears that inter-thread data sharing at ";
|
||||
std::cout << "grid_group sync points works properly!" << std::endl;
|
||||
} else {
|
||||
flag = 1;
|
||||
}
|
||||
for (int k = 0; k < 2; ++k) {
|
||||
HIPCHECK(hipStreamDestroy(streams[k]));
|
||||
}
|
||||
HIPCHECK(hipFree(first_dev_array));
|
||||
HIPCHECK(hipFree(second_dev_array));
|
||||
free(input_buffer);
|
||||
free(first_array);
|
||||
free(second_array);
|
||||
}
|
||||
if (!flag) {
|
||||
passed();
|
||||
} else {
|
||||
failed("");
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,568 @@
|
||||
/*
|
||||
Copyright (c) 2020-present 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 Cooperative Groups kernel
|
||||
launches work when launching too many warps to multiple target devices. This
|
||||
tests the following failure modes for hipLaunchCooperativeKernelMultiDevice:
|
||||
1) Do not launch more warps to any device than can fit on that device
|
||||
2) All device targets for the multi-device launch function must be different
|
||||
3) All streams must be explicit (non-NULL)
|
||||
4) The kernels sent in must be identical between devices
|
||||
5) The grid and block sizes must be identical between devices
|
||||
6) The block dimensions must be non-zero
|
||||
7) The dynamic shared memory size must be identical between devices.
|
||||
|
||||
This test ensures that the proper error conditions are returned, even if the
|
||||
target kernel does not actually use any fo the cooperative groups features.
|
||||
|
||||
Note that tests 4, 5, and 7 only hold on Nvidia GPUs. AMD GPUs running ROCm
|
||||
do not have these constraints. As such, the test checks to see whether they
|
||||
should fail or succeed and compares this to what actually happens.
|
||||
*/
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <hip/hip_cooperative_groups.h>
|
||||
#include "test_common.h"
|
||||
|
||||
static inline void hipCheckAndFail(hipError_t errval,
|
||||
const char *file, int line) {
|
||||
hipError_t last_err = hipGetLastError();
|
||||
if (errval != hipSuccess) {
|
||||
std::cerr << "hip error: " << hipGetErrorString(errval);
|
||||
std::cerr << std::endl;
|
||||
std::cerr << " Location: " << file << ":" << line << std::endl;
|
||||
failed("");
|
||||
}
|
||||
if (last_err != errval) {
|
||||
std::cerr << "Error: the return value of a function was not the same ";
|
||||
std::cerr << "as the value returned by hipGetLastError()" << std::endl;
|
||||
std::cerr << " Location: " << file << ":" << line << std::endl;
|
||||
std::cerr << " Function returned: " << hipGetErrorString(errval);
|
||||
std::cerr << " (" << errval << ")" << std::endl;
|
||||
std::cerr << "hipGetLastError() returned: " << hipGetErrorString(last_err);
|
||||
std::cerr << " (" << last_err << ")" << std::endl;
|
||||
failed("");
|
||||
}
|
||||
}
|
||||
#define hipCheckErr(errval) \
|
||||
do { hipCheckAndFail((errval), __FILE__, __LINE__); } while (0)
|
||||
|
||||
static int cooperative_groups_support(int device_id) {
|
||||
hipError_t err;
|
||||
|
||||
int cooperative_attribute;
|
||||
HIPCHECK(hipDeviceGetAttribute(&cooperative_attribute,
|
||||
hipDeviceAttributeCooperativeLaunch, device_id));
|
||||
if (!cooperative_attribute) {
|
||||
std::cerr << "Cooperative launch support not available in ";
|
||||
std::cerr << "the device attribute for device " << device_id;
|
||||
std::cerr << std::endl;
|
||||
return 0;
|
||||
}
|
||||
|
||||
int multi_gpu_cooperative_attribute;
|
||||
HIPCHECK(hipDeviceGetAttribute(&multi_gpu_cooperative_attribute,
|
||||
hipDeviceAttributeCooperativeMultiDeviceLaunch, device_id));
|
||||
|
||||
if (!multi_gpu_cooperative_attribute) {
|
||||
std::cerr << "Multi-GPU cooperative launch support not available in ";
|
||||
std::cerr << "the device attribute for device " << device_id;
|
||||
std::cerr << std::endl;
|
||||
return 0;
|
||||
}
|
||||
|
||||
hipDeviceProp_t device_properties;
|
||||
HIPCHECK(hipGetDeviceProperties(&device_properties, device_id));
|
||||
if (device_properties.cooperativeLaunch == 0) {
|
||||
std::cerr << "Cooperative group support not available in ";
|
||||
std::cerr << "device properties." << std::endl;
|
||||
return 0;
|
||||
}
|
||||
if (device_properties.cooperativeMultiDeviceLaunch == 0) {
|
||||
std::cerr << "Multi-GPU cooperative group support not available in ";
|
||||
std::cerr << "device properties." << std::endl;
|
||||
return 0;
|
||||
}
|
||||
return 1;
|
||||
}
|
||||
|
||||
static int support_for_separate_kernels(int device_id) {
|
||||
hipError_t err;
|
||||
|
||||
int separate_kernel_supported;
|
||||
HIPCHECK(hipDeviceGetAttribute(&separate_kernel_supported,
|
||||
hipDeviceAttributeCooperativeMultiDeviceUnmatchedFunc,
|
||||
device_id));
|
||||
if (!separate_kernel_supported) {
|
||||
return 0;
|
||||
}
|
||||
|
||||
hipDeviceProp_t device_properties;
|
||||
HIPCHECK(hipGetDeviceProperties(&device_properties, device_id));
|
||||
if (device_properties.cooperativeMultiDeviceUnmatchedFunc == 0) {
|
||||
return 0;
|
||||
}
|
||||
return 1;
|
||||
}
|
||||
|
||||
static int support_for_separate_grid_sizes(int device_id) {
|
||||
hipError_t err;
|
||||
int separate_sizes_supported;
|
||||
HIPCHECK(hipDeviceGetAttribute(&separate_sizes_supported,
|
||||
hipDeviceAttributeCooperativeMultiDeviceUnmatchedGridDim,
|
||||
device_id));
|
||||
if (!separate_sizes_supported) {
|
||||
return 0;
|
||||
}
|
||||
|
||||
hipDeviceProp_t device_properties;
|
||||
HIPCHECK(hipGetDeviceProperties(&device_properties, device_id));
|
||||
if (device_properties.cooperativeMultiDeviceUnmatchedGridDim == 0) {
|
||||
return 0;
|
||||
}
|
||||
return 1;
|
||||
}
|
||||
|
||||
static int support_for_separate_block_dims(int device_id) {
|
||||
hipError_t err;
|
||||
int separate_sizes_supported;
|
||||
HIPCHECK(hipDeviceGetAttribute(&separate_sizes_supported,
|
||||
hipDeviceAttributeCooperativeMultiDeviceUnmatchedBlockDim,
|
||||
device_id));
|
||||
if (!separate_sizes_supported) {
|
||||
return 0;
|
||||
}
|
||||
|
||||
hipDeviceProp_t device_properties;
|
||||
HIPCHECK(hipGetDeviceProperties(&device_properties, device_id));
|
||||
if (device_properties.cooperativeMultiDeviceUnmatchedBlockDim == 0) {
|
||||
return 0;
|
||||
}
|
||||
return 1;
|
||||
}
|
||||
|
||||
static int support_for_separate_shared_sizes(int device_id) {
|
||||
hipError_t err;
|
||||
int separate_sizes_supported;
|
||||
HIPCHECK(hipDeviceGetAttribute(&separate_sizes_supported,
|
||||
hipDeviceAttributeCooperativeMultiDeviceUnmatchedSharedMem,
|
||||
device_id));
|
||||
if (!separate_sizes_supported) {
|
||||
return 0;
|
||||
}
|
||||
|
||||
hipDeviceProp_t device_properties;
|
||||
HIPCHECK(hipGetDeviceProperties(&device_properties, device_id));
|
||||
if (device_properties.cooperativeMultiDeviceUnmatchedSharedMem == 0) {
|
||||
return 0;
|
||||
}
|
||||
return 1;
|
||||
}
|
||||
|
||||
__global__ void test_kernel(long long *array) {
|
||||
unsigned int rank = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
array[rank] += clock64();
|
||||
}
|
||||
|
||||
__global__ void second_test_kernel(long long *array) {
|
||||
unsigned int rank = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
array[rank] += clock64();
|
||||
}
|
||||
|
||||
int main(int argc, char** argv) {
|
||||
hipError_t err;
|
||||
/*************************************************************************/
|
||||
/* Parse the command line parameters *************************************/
|
||||
// Arguments to pull out of the command line.
|
||||
int device_num, FailFlag = 0;
|
||||
HIPCHECK(hipGetDeviceCount(&device_num));
|
||||
if (device_num < 2) {
|
||||
std::cout << "This test requires atleast two gpus but the system has ";
|
||||
std::cout << " only "<< device_num <<std::endl;
|
||||
std::cout << "The test is skipping with Pass result" << std::endl;
|
||||
passed();
|
||||
}
|
||||
for (int dev = 0; dev < (device_num-1); ++dev) {
|
||||
std::cout << "First device number: " << dev << std::endl;
|
||||
std::cout << "Second device number: " << (dev + 1) << std::endl;
|
||||
|
||||
/*************************************************************************/
|
||||
/* Test whether target devices support cooperative groups ****************/
|
||||
for (int i = 0; i < 2; i++) {
|
||||
if (!cooperative_groups_support((dev + i))) {
|
||||
std::cout << "Skipping the test with Pass result.\n";
|
||||
passed();
|
||||
}
|
||||
}
|
||||
|
||||
/*************************************************************************/
|
||||
/* We will try to launch more waves than the GPUs can fit. ***************/
|
||||
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 i = 0; i < 2; i++) {
|
||||
HIPCHECK(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];
|
||||
}
|
||||
std::cout << "Device " << (dev + i);
|
||||
std::cout << " name: " << device_properties[i].name << std::endl;
|
||||
}
|
||||
std::cout << std::endl;
|
||||
|
||||
// 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++) {
|
||||
HIPCHECK(hipSetDevice((dev + i)));
|
||||
HIPCHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(
|
||||
&max_blocks_per_sm_arr[i], test_kernel, 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 = max_blocks_per_sm * num_sm;
|
||||
|
||||
/*************************************************************************/
|
||||
/* Create the streams we will use in this test. **************************/
|
||||
hipStream_t streams[2];
|
||||
for (int i = 0; i < 2; i++) {
|
||||
HIPCHECK(hipSetDevice((dev + i)));
|
||||
HIPCHECK(hipStreamCreate(&streams[i]));
|
||||
}
|
||||
|
||||
/*************************************************************************/
|
||||
/* Set up data to pass into the kernel ***********************************/
|
||||
|
||||
// Alocate the host input buffer, and two device-focused buffers per GPU
|
||||
// that we will use for our test.
|
||||
unsigned int *good_dev_array[2];
|
||||
unsigned int *bad_dev_array[2];
|
||||
for (int i = 0; i < 2; i++) {
|
||||
int good_size = desired_blocks * warp_size * sizeof(long long);
|
||||
int bad_size = 2 * desired_blocks * warp_size * sizeof(long long);
|
||||
|
||||
HIPCHECK(hipSetDevice((dev + i)));
|
||||
HIPCHECK(hipMalloc(reinterpret_cast<void**>(&good_dev_array[i]),
|
||||
good_size));
|
||||
HIPCHECK(hipMemsetAsync(good_dev_array[i], 0, good_size, streams[i]));
|
||||
HIPCHECK(hipMalloc(reinterpret_cast<void**>(&bad_dev_array[i]),
|
||||
bad_size));
|
||||
HIPCHECK(hipMemsetAsync(bad_dev_array[i], 0, bad_size, streams[i]));
|
||||
}
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
/*************************************************************************/
|
||||
/* Launch the kernels ****************************************************/
|
||||
std::cout << "Launching a multi-GPU cooperative kernel with too many ";
|
||||
std::cout << "warps..." << std::endl;
|
||||
|
||||
void *dev_params[2][1];
|
||||
hipLaunchParams md_params[2];
|
||||
for (int i = 0; i < 2; i++) {
|
||||
dev_params[i][0] = reinterpret_cast<void*>(&bad_dev_array[i]);
|
||||
|
||||
md_params[i].func = reinterpret_cast<void*>(test_kernel);
|
||||
md_params[i].gridDim = 2 * 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];
|
||||
}
|
||||
|
||||
err = hipLaunchCooperativeKernelMultiDevice(md_params, 2, 0);
|
||||
if (err != hipErrorCooperativeLaunchTooLarge) {
|
||||
std::cerr << "ERROR! Tried to launch a multi-GPU cooperative kernel ";
|
||||
std::cerr << "with too many warps." << std::endl;
|
||||
std::cerr << "This SHOULD have failed with the error ";
|
||||
std::cerr << "hipErrorCooperativeLaunchTooLarge (";
|
||||
std::cerr << hipErrorCooperativeLaunchTooLarge << ")." << std::endl;
|
||||
std::cerr << "Instead, the launch returned " << hipGetErrorName(err);
|
||||
std::cerr << " (" << err << ")" << std::endl;
|
||||
FailFlag = 1;
|
||||
} else {
|
||||
std::cout << "\tProperly saw this return ";
|
||||
std::cout << "hipErrorCooperativeLaunchTooLarge" << std::endl;
|
||||
}
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
std::cout << "Launching a multi-GPU cooperative kernel to the same ";
|
||||
std::cout << "device twice..." << std::endl;
|
||||
for (int i = 0; i < 2; i++) {
|
||||
dev_params[i][0] = reinterpret_cast<void*>(&good_dev_array[i]);
|
||||
md_params[i].gridDim = desired_blocks;
|
||||
md_params[i].stream = streams[0];
|
||||
}
|
||||
err = hipLaunchCooperativeKernelMultiDevice(md_params, 2, 0);
|
||||
if (err != hipErrorInvalidDevice) {
|
||||
std::cerr << "ERROR! Tried to launch a multi-GPU cooperative kernel ";
|
||||
std::cerr << "to the same device twice." << std::endl;
|
||||
std::cerr << "This SHOULD have failed with the error ";
|
||||
std::cerr << "hipErrorInvalidDevice (";
|
||||
std::cerr << hipErrorInvalidDevice << ")." << std::endl;
|
||||
std::cerr << "Instead, the launch returned " << hipGetErrorName(err);
|
||||
std::cerr << " (" << err << ")" << std::endl;
|
||||
FailFlag = 1;
|
||||
} else {
|
||||
std::cout << "\tProperly saw this return ";
|
||||
std::cout << "hipErrorInvalidDevice" << std::endl;
|
||||
}
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
std::cout << "Launching a multi-GPU cooperative kernel to the NULL ";
|
||||
std::cout << "stream" << std::endl;
|
||||
for (int i = 0; i < 2; i++) {
|
||||
md_params[i].stream = NULL;
|
||||
}
|
||||
err = hipLaunchCooperativeKernelMultiDevice(md_params, 2, 0);
|
||||
if (err != hipErrorInvalidResourceHandle) {
|
||||
std::cerr << "ERROR! Tried to launch a multi-GPU cooperative kernel ";
|
||||
std::cerr << "to the NULL stream." << std::endl;
|
||||
std::cerr << "This SHOULD have failed with the error ";
|
||||
std::cerr << "hipErrorInvalidResourceHandle (";
|
||||
std::cerr << hipErrorInvalidResourceHandle << ")." << std::endl;
|
||||
std::cerr << "Instead, the launch returned " << hipGetErrorName(err);
|
||||
std::cerr << " (" << err << ")" << std::endl;
|
||||
FailFlag = 1;
|
||||
} else {
|
||||
std::cout << "\tProperly saw this return ";
|
||||
std::cout << "hipErrorInvalidResourceHandle" << std::endl;
|
||||
}
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
std::cout << "Launching a multi-GPU cooperative kernel with two ";
|
||||
std::cout << "different kernels." << std::endl;
|
||||
bool supports_sep_kernels = true;
|
||||
for (int i = 0; i < 2; i++) {
|
||||
md_params[i].stream = streams[i];
|
||||
if (!support_for_separate_kernels((dev + i))) {
|
||||
supports_sep_kernels = false;
|
||||
}
|
||||
}
|
||||
md_params[1].func = reinterpret_cast<void*>(second_test_kernel);
|
||||
err = hipLaunchCooperativeKernelMultiDevice(md_params, 2, 0);
|
||||
if ((supports_sep_kernels && err != hipSuccess) ||
|
||||
(!supports_sep_kernels && err != hipErrorInvalidValue)) {
|
||||
if (supports_sep_kernels) {
|
||||
std::cerr << "ERROR! Tried to launch a multi-GPU cooperative kernel ";
|
||||
std::cerr << "with two different kernels." << std::endl;
|
||||
std::cerr << "This SHOULD have succeeded with hipSuccess (";
|
||||
std::cerr << hipSuccess << ")." << std::endl;
|
||||
std::cerr << "Instead, the launch returned " << hipGetErrorName(err);
|
||||
std::cerr << " (" << err << ")" << std::endl;
|
||||
} else {
|
||||
std::cerr << "ERROR! Tried to launch a multi-GPU cooperative kernel ";
|
||||
std::cerr << "with two different kernels." << std::endl;
|
||||
std::cerr << "This SHOULD have failed with the error ";
|
||||
std::cerr << "hipErrorInvalidValue (";
|
||||
std::cerr << hipErrorInvalidValue << ")." << std::endl;
|
||||
std::cerr << "Instead, the launch returned " << hipGetErrorName(err);
|
||||
std::cerr << " (" << err << ")" << std::endl;
|
||||
}
|
||||
FailFlag = 1;
|
||||
} else {
|
||||
std::cout << "\tProperly saw this return ";
|
||||
if (supports_sep_kernels) {
|
||||
std::cout << "hipSuccess" << std::endl;
|
||||
} else {
|
||||
std::cout << "hipErrorInvalidValue" << std::endl;
|
||||
}
|
||||
}
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
std::cout << "Launching a multi-GPU cooperative kernel with two ";
|
||||
std::cout << "different grid sizes." << std::endl;
|
||||
bool supports_sep_sizes = true;
|
||||
for (int i = 0; i < 2; i++) {
|
||||
md_params[i].func = reinterpret_cast<void*>(test_kernel);
|
||||
md_params[i].gridDim = i+1;
|
||||
if (!support_for_separate_grid_sizes((dev + i))) {
|
||||
supports_sep_sizes = false;
|
||||
}
|
||||
}
|
||||
err = hipLaunchCooperativeKernelMultiDevice(md_params, 2, 0);
|
||||
if ((supports_sep_sizes && err != hipSuccess) ||
|
||||
(!supports_sep_sizes && err == hipErrorInvalidValue)) {
|
||||
if (supports_sep_sizes) {
|
||||
std::cerr << "ERROR! Tried to launch a multi-GPU cooperative kernel ";
|
||||
std::cerr << "with two different grid sizes." << std::endl;
|
||||
std::cerr << "This SHOULD have succeeded with hipSuccess (";
|
||||
std::cerr << hipSuccess << ")." << std::endl;
|
||||
std::cerr << "Instead, the launch returned " << hipGetErrorName(err);
|
||||
std::cerr << " (" << err << ")" << std::endl;
|
||||
} else {
|
||||
std::cerr << "ERROR! Tried to launch a multi-GPU cooperative kernel ";
|
||||
std::cerr << "with two different grid sizes." << std::endl;
|
||||
std::cerr << "This SHOULD have failed with the error ";
|
||||
std::cerr << "hipErrorInvalidValue (";
|
||||
std::cerr << hipErrorInvalidValue << ")." << std::endl;
|
||||
std::cerr << "Instead, the launch returned " << hipGetErrorName(err);
|
||||
std::cerr << " (" << err << ")" << std::endl;
|
||||
FailFlag = 1;
|
||||
}
|
||||
} else {
|
||||
std::cout << "\tProperly saw this return ";
|
||||
if (supports_sep_kernels) {
|
||||
std::cout << "hipSuccess" << std::endl;
|
||||
} else {
|
||||
std::cout << "hipErrorInvalidValue" << std::endl;
|
||||
}
|
||||
}
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
std::cout << "Launching a multi-GPU cooperative kernel with two ";
|
||||
std::cout << "different block dimensions." << std::endl;
|
||||
supports_sep_sizes = true;
|
||||
for (int i = 0; i < 2; i++) {
|
||||
md_params[i].gridDim = desired_blocks;
|
||||
md_params[i].blockDim = i+1;
|
||||
if (!support_for_separate_block_dims((dev + i))) {
|
||||
supports_sep_sizes = false;
|
||||
}
|
||||
}
|
||||
err = hipLaunchCooperativeKernelMultiDevice(md_params, 2, 0);
|
||||
if ((supports_sep_sizes && err != hipSuccess) ||
|
||||
(!supports_sep_sizes && err == hipErrorInvalidValue)) {
|
||||
if (supports_sep_sizes) {
|
||||
std::cerr << "ERROR! Tried to launch a multi-GPU cooperative kernel ";
|
||||
std::cerr << "with two different block dimensions." << std::endl;
|
||||
std::cerr << "This SHOULD have succeeded with hipSuccess (";
|
||||
std::cerr << hipSuccess << ")." << std::endl;
|
||||
std::cerr << "Instead, the launch returned " << hipGetErrorName(err);
|
||||
std::cerr << " (" << err << ")" << std::endl;
|
||||
} else {
|
||||
std::cerr << "ERROR! Tried to launch a multi-GPU cooperative kernel ";
|
||||
std::cerr << "with two different block dimensions." << std::endl;
|
||||
std::cerr << "This SHOULD have failed with the error ";
|
||||
std::cerr << "hipErrorInvalidValue (";
|
||||
std::cerr << hipErrorInvalidValue << ")." << std::endl;
|
||||
std::cerr << "Instead, the launch returned " << hipGetErrorName(err);
|
||||
std::cerr << " (" << err << ")" << std::endl;
|
||||
FailFlag = 1;
|
||||
}
|
||||
} else {
|
||||
std::cout << "\tProperly saw this return ";
|
||||
if (supports_sep_kernels) {
|
||||
std::cout << "hipSuccess" << std::endl;
|
||||
} else {
|
||||
std::cout << "hipErrorInvalidValue" << std::endl;
|
||||
}
|
||||
}
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
std::cout << "Launching a multi-GPU cooperative kernel with block ";
|
||||
std::cout << "dimensions of zero." << std::endl;
|
||||
for (int i = 0; i < 2; i++) {
|
||||
md_params[i].blockDim = 0;
|
||||
}
|
||||
err = hipLaunchCooperativeKernelMultiDevice(md_params, 2, 0);
|
||||
if (err != hipErrorInvalidConfiguration) {
|
||||
std::cerr << "ERROR! Tried to launch a multi-GPU cooperative kernel ";
|
||||
std::cerr << "with block dimensions of zero." << std::endl;
|
||||
std::cerr << "This SHOULD have failed with the error ";
|
||||
std::cerr << "hipErrorInvalidConfiguration (";
|
||||
std::cerr << hipErrorInvalidConfiguration << ")." << std::endl;
|
||||
std::cerr << "Instead, the launch returned " << hipGetErrorName(err);
|
||||
std::cerr << " (" << err << ")" << std::endl;
|
||||
FailFlag = 1;
|
||||
} else {
|
||||
std::cout << "\tProperly saw this return ";
|
||||
std::cout << "hipErrorInvalidConfiguration" << std::endl;
|
||||
}
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
std::cout << "Launching a multi-GPU cooperative kernel with two ";
|
||||
std::cout << "different shared memory sizes." << std::endl;
|
||||
supports_sep_sizes = true;
|
||||
for (int i = 0; i < 2; i++) {
|
||||
md_params[i].blockDim = warp_size;
|
||||
md_params[i].sharedMem = i;
|
||||
if (!support_for_separate_shared_sizes((dev + i))) {
|
||||
supports_sep_sizes = false;
|
||||
}
|
||||
}
|
||||
err = hipLaunchCooperativeKernelMultiDevice(md_params, 2, 0);
|
||||
if ((supports_sep_sizes && err != hipSuccess) ||
|
||||
(!supports_sep_sizes && err == hipErrorInvalidValue)) {
|
||||
if (supports_sep_sizes) {
|
||||
std::cerr << "ERROR! Tried to launch a multi-GPU cooperative kernel ";
|
||||
std::cerr << "with two different shared memory sizes." << std::endl;
|
||||
std::cerr << "This SHOULD have succeeded with hipSuccess (";
|
||||
std::cerr << hipSuccess << ")." << std::endl;
|
||||
std::cerr << "Instead, the launch returned " << hipGetErrorName(err);
|
||||
std::cerr << " (" << err << ")" << std::endl;
|
||||
} else {
|
||||
std::cerr << "ERROR! Tried to launch a multi-GPU cooperative kernel ";
|
||||
std::cerr << "with two different shared memory sizes." << std::endl;
|
||||
std::cerr << "This SHOULD have failed with the error ";
|
||||
std::cerr << "hipErrorInvalidValue (";
|
||||
std::cerr << hipErrorInvalidValue << ")." << std::endl;
|
||||
std::cerr << "Instead, the launch returned " << hipGetErrorName(err);
|
||||
std::cerr << " (" << err << ")" << std::endl;
|
||||
FailFlag = 1;
|
||||
}
|
||||
} else {
|
||||
std::cout << "\tProperly saw this return ";
|
||||
if (supports_sep_kernels) {
|
||||
std::cout << "hipSuccess" << std::endl;
|
||||
} else {
|
||||
std::cout << "hipErrorInvalidValue" << std::endl;
|
||||
}
|
||||
}
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
std::cout << "Launching a multi-GPU cooperative kernel with maximum ";
|
||||
std::cout << "number of warps..." << std::endl;
|
||||
for (int i = 0; i < 2; i++) {
|
||||
md_params[i].sharedMem = 0;
|
||||
}
|
||||
HIPCHECK(hipLaunchCooperativeKernelMultiDevice(md_params, 2, 0));
|
||||
std::cout << "\tProperly launched." << std::endl;
|
||||
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
for (int m = 0; m < 2; ++m) {
|
||||
HIPCHECK(hipFree(good_dev_array[m]));
|
||||
HIPCHECK(hipFree(bad_dev_array[m]));
|
||||
HIPCHECK(hipStreamDestroy(streams[m]));
|
||||
}
|
||||
if (FailFlag == 1) {
|
||||
break;
|
||||
}
|
||||
}
|
||||
if (FailFlag == 1) {
|
||||
failed("");
|
||||
} else {
|
||||
passed();
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,581 @@
|
||||
/*
|
||||
Copyright (c) 2020-present 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.
|
||||
*/
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include <chrono>
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <hip/hip_cooperative_groups.h>
|
||||
#include "test_common.h"
|
||||
|
||||
static inline void hipCheckAndFail(hipError_t errval,
|
||||
const char *file, int line) {
|
||||
hipError_t last_err = hipGetLastError();
|
||||
if (errval != hipSuccess) {
|
||||
std::cerr << "hip error: " << hipGetErrorString(errval);
|
||||
std::cerr << std::endl;
|
||||
std::cerr << " Location: " << file << ":" << line << std::endl;
|
||||
failed("");
|
||||
}
|
||||
if (last_err != errval) {
|
||||
std::cerr << "Error: the return value of a function was not the same ";
|
||||
std::cerr << "as the value returned by hipGetLastError()" << std::endl;
|
||||
std::cerr << " Location: " << file << ":" << line << std::endl;
|
||||
std::cerr << " Function returned: " << hipGetErrorString(errval);
|
||||
std::cerr << " (" << errval << ")" << std::endl;
|
||||
std::cerr << "hipGetLastError() returned: " << hipGetErrorString(last_err);
|
||||
std::cerr << " (" << last_err << ")" << std::endl;
|
||||
failed("");
|
||||
}
|
||||
}
|
||||
#define hipCheckErr(errval) \
|
||||
do { hipCheckAndFail((errval), __FILE__, __LINE__); } while (0)
|
||||
|
||||
static int cooperative_groups_support(int device_id) {
|
||||
hipError_t err;
|
||||
int cooperative_attribute;
|
||||
HIPCHECK(hipDeviceGetAttribute(&cooperative_attribute,
|
||||
hipDeviceAttributeCooperativeLaunch, device_id));
|
||||
if (!cooperative_attribute) {
|
||||
std::cerr << "Cooperative launch support not available in ";
|
||||
std::cerr << "the device attribute for device " << device_id;
|
||||
std::cerr << std::endl;
|
||||
return 0;
|
||||
}
|
||||
|
||||
int multi_gpu_cooperative_attribute;
|
||||
HIPCHECK(hipDeviceGetAttribute(&multi_gpu_cooperative_attribute,
|
||||
hipDeviceAttributeCooperativeMultiDeviceLaunch, device_id));
|
||||
if (!multi_gpu_cooperative_attribute) {
|
||||
std::cerr << "Multi-GPU cooperative launch support not available in ";
|
||||
std::cerr << "the device attribute for device " << device_id;
|
||||
std::cerr << std::endl;
|
||||
return 0;
|
||||
}
|
||||
|
||||
hipDeviceProp_t device_properties;
|
||||
HIPCHECK(hipGetDeviceProperties(&device_properties, device_id));
|
||||
if (device_properties.cooperativeLaunch == 0) {
|
||||
std::cerr << "Cooperative group support not available in ";
|
||||
std::cerr << "device properties." << std::endl;
|
||||
return 0;
|
||||
}
|
||||
if (device_properties.cooperativeMultiDeviceLaunch == 0) {
|
||||
std::cerr << "Multi-GPU cooperative group support not available in ";
|
||||
std::cerr << "device properties." << std::endl;
|
||||
return 0;
|
||||
}
|
||||
return 1;
|
||||
}
|
||||
|
||||
__global__ void test_coop_kernel(unsigned int loops, long long *array,
|
||||
int fast_gpu) {
|
||||
cooperative_groups::multi_grid_group mgrid =
|
||||
cooperative_groups::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 start_clock = clock64();
|
||||
while (clock64() < (start_clock+1000000)) {}
|
||||
array[rank] += clock64();
|
||||
}
|
||||
}
|
||||
|
||||
__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 start_clock = clock64();
|
||||
while (clock64() < (start_clock+1000000)) {}
|
||||
array[rank] += clock64();
|
||||
}
|
||||
}
|
||||
|
||||
int main(int argc, char** argv) {
|
||||
hipError_t err;
|
||||
int device_num, FailFlag = 0;
|
||||
uint32_t loops = 2000;
|
||||
uint32_t fast_loops = 1;
|
||||
int32_t fast_gpu = -1;
|
||||
HIPCHECK(hipGetDeviceCount(&device_num));
|
||||
if (device_num < 2) {
|
||||
std::cout << "This test requires atleast two gpus but the system has ";
|
||||
std::cout << " only "<< device_num <<std::endl;
|
||||
std::cout << "The test is skipping with Pass result" << std::endl;
|
||||
passed();
|
||||
}
|
||||
for (int dev = 0; dev < (device_num-1); ++dev) {
|
||||
std::cout << "First device number: " << dev << std::endl;
|
||||
std::cout << "Second device number: " << (dev + 1) << std::endl;
|
||||
std::cout << "Loops: " << loops << std::endl;
|
||||
|
||||
/*************************************************************************/
|
||||
/* Test whether target devices support cooperative groups ****************/
|
||||
for (int i = 0; i < 2; i++) {
|
||||
if (!cooperative_groups_support(dev + i)) {
|
||||
std::cout << "Skipping the test with Pass result.\n";
|
||||
passed();
|
||||
}
|
||||
}
|
||||
|
||||
/*************************************************************************/
|
||||
/* 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 i = 0; i < 2; i++) {
|
||||
HIPCHECK(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];
|
||||
}
|
||||
std::cout << "Device " << (i + 1);
|
||||
std::cout << " name: " << device_properties[i].name << std::endl;
|
||||
}
|
||||
std::cout << std::endl;
|
||||
|
||||
// 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++) {
|
||||
HIPCHECK(hipSetDevice(dev + i));
|
||||
HIPCHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(
|
||||
&max_blocks_per_sm_arr[i], test_kernel, 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) {
|
||||
std::cerr << "The requested number of blocks will not fit on the GPU";
|
||||
std::cerr << std::endl;
|
||||
std::cerr << "You requested " << desired_blocks << " but we can only ";
|
||||
std::cerr << "fit " << (max_blocks_per_sm * num_sm) << std::endl;
|
||||
failed("");
|
||||
}
|
||||
|
||||
/*************************************************************************/
|
||||
/* Create the streams we will use in this test. **************************/
|
||||
hipStream_t streams[2];
|
||||
for (int i = 0; i < 2; i++) {
|
||||
HIPCHECK(hipSetDevice(dev + i));
|
||||
HIPCHECK(hipStreamCreate(&streams[i]));
|
||||
}
|
||||
|
||||
/*************************************************************************/
|
||||
/* Set up data to pass into the kernelx **********************************/
|
||||
|
||||
// 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);
|
||||
HIPCHECK(hipSetDevice(dev + i));
|
||||
HIPCHECK(hipMalloc(reinterpret_cast<void**>(&dev_array[i]), good_size));
|
||||
HIPCHECK(hipMemsetAsync(dev_array[i], 0, good_size, streams[i]));
|
||||
}
|
||||
for (int i = 0; i < 2; i++) {
|
||||
HIPCHECK(hipSetDevice(dev + i));
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
}
|
||||
|
||||
/*************************************************************************/
|
||||
/* Launch the kernels ****************************************************/
|
||||
void *dev_params[2][3];
|
||||
hipLaunchParams md_params[2];
|
||||
std::chrono::time_point<std::chrono::system_clock> start_time[6];
|
||||
std::chrono::time_point<std::chrono::system_clock> end_time[6];
|
||||
|
||||
std::cout << "Test 0: Launching a multi-GPU cooperative kernel...\n";
|
||||
std::cout << "This should result in the following pattern:" << std::endl;
|
||||
std::cout << "GPU " << dev << ": Long Coop Kernel" << std::endl;
|
||||
std::cout << "GPU " << (dev + 1) << ": Long Coop Kernel" << std::endl;
|
||||
|
||||
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);
|
||||
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();
|
||||
HIPCHECK(hipLaunchCooperativeKernelMultiDevice(md_params, 2, 0));
|
||||
for (int i = 0; i < 2; i++) {
|
||||
HIPCHECK(hipSetDevice(dev + i));
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
}
|
||||
end_time[0] = std::chrono::system_clock::now();
|
||||
|
||||
std::cout << std::endl;
|
||||
std::cout << "Test 1: Launching a multi-GPU cooperative kernel with the ";
|
||||
std::cout << "following pattern:" << std::endl;
|
||||
std::cout << "GPU " << dev << ": Standard Kernel --> Long Coop Kernel\n";
|
||||
std::cout << "GPU " << (dev + 1) << ": --> Coop ";
|
||||
std::cout << "--> Standard Kernel\n";
|
||||
fast_gpu = 1;
|
||||
start_time[1] = std::chrono::system_clock::now();
|
||||
HIPCHECK(hipSetDevice(dev));
|
||||
hipLaunchKernelGGL(test_kernel, dim3(desired_blocks), dim3(warp_size), 0,
|
||||
streams[0], loops, dev_array[0]);
|
||||
HIPCHECK(hipGetLastError());
|
||||
HIPCHECK(hipLaunchCooperativeKernelMultiDevice(md_params, 2, 0));
|
||||
HIPCHECK(hipSetDevice(dev + 1));
|
||||
hipLaunchKernelGGL(test_kernel, dim3(desired_blocks), dim3(warp_size), 0,
|
||||
streams[1], loops, dev_array[1]);
|
||||
HIPCHECK(hipGetLastError());
|
||||
for (int i = 0; i < 2; i++) {
|
||||
HIPCHECK(hipSetDevice(dev + i));
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
}
|
||||
end_time[1] = std::chrono::system_clock::now();
|
||||
fast_gpu = -1;
|
||||
|
||||
std::cout << std::endl;
|
||||
std::cout << "Test 2: Launching a multi-GPU cooperative kernel with the ";
|
||||
std::cout << "following pattern:" << std::endl;
|
||||
std::cout << "GPU " << dev << ": Standard Kernel --> Coop" << std::endl;
|
||||
std::cout << "GPU " << (dev + 1) << ": --> Long Coop";
|
||||
std::cout << " Kernel --> ";
|
||||
std::cout << "Standard Kernel\n";
|
||||
fast_gpu = 0;
|
||||
start_time[2] = std::chrono::system_clock::now();
|
||||
HIPCHECK(hipSetDevice(dev));
|
||||
hipLaunchKernelGGL(test_kernel, dim3(desired_blocks), dim3(warp_size), 0,
|
||||
streams[0], loops, dev_array[0]);
|
||||
HIPCHECK(hipGetLastError());
|
||||
HIPCHECK(hipLaunchCooperativeKernelMultiDevice(md_params, 2, 0));
|
||||
HIPCHECK(hipSetDevice(dev + 1));
|
||||
hipLaunchKernelGGL(test_kernel, dim3(desired_blocks), dim3(warp_size), 0,
|
||||
streams[1], loops, dev_array[1]);
|
||||
HIPCHECK(hipGetLastError());
|
||||
for (int i = 0; i < 2; i++) {
|
||||
HIPCHECK(hipSetDevice(dev + i));
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
}
|
||||
end_time[2] = std::chrono::system_clock::now();
|
||||
fast_gpu = -1;
|
||||
|
||||
std::cout << std::endl;
|
||||
std::cout << "Test 3: Launching a multi-GPU cooperative kernel with the ";
|
||||
std::cout << "ability to overlap regular and cooperative kernels ";
|
||||
std::cout << "only at the beginning." << std::endl;
|
||||
std::cout << "This should result in the following pattern:" << std::endl;
|
||||
std::cout << "GPU " << dev << ": Standard Kernel --> Coop" << std::endl;
|
||||
std::cout << "GPU " << (dev + 1) << ": Long Coop Kernel --> Standard";
|
||||
std::cout<< " Kernel\n";
|
||||
fast_gpu = 0;
|
||||
start_time[3] = std::chrono::system_clock::now();
|
||||
HIPCHECK(hipSetDevice(dev));
|
||||
hipLaunchKernelGGL(test_kernel, dim3(desired_blocks), dim3(warp_size), 0,
|
||||
streams[0], loops, dev_array[0]);
|
||||
HIPCHECK(hipGetLastError());
|
||||
HIPCHECK(hipLaunchCooperativeKernelMultiDevice(md_params, 2,
|
||||
hipCooperativeLaunchMultiDeviceNoPreSync));
|
||||
HIPCHECK(hipSetDevice(dev + 1));
|
||||
hipLaunchKernelGGL(test_kernel, dim3(desired_blocks), dim3(warp_size), 0,
|
||||
streams[1], loops, dev_array[1]);
|
||||
HIPCHECK(hipGetLastError());
|
||||
for (int i = 0; i < 2; i++) {
|
||||
HIPCHECK(hipSetDevice(dev + i));
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
}
|
||||
end_time[3] = std::chrono::system_clock::now();
|
||||
fast_gpu = -1;
|
||||
|
||||
std::cout << std::endl;
|
||||
std::cout << "Test 4: Launching a multi-GPU cooperative kernel with the ";
|
||||
std::cout << "ability to overlap regular and cooperative kernels ";
|
||||
std::cout << "only at the end." << std::endl;
|
||||
std::cout << "This should result in the following pattern:" << std::endl;
|
||||
std::cout << "GPU " << dev << ": Standard Kernel --> Long Coop Kernel\n";
|
||||
std::cout << "GPU " << (dev + 1) << ": --> Coop --> ";
|
||||
std::cout << "Standard Kernel\n";
|
||||
fast_gpu = 1;
|
||||
start_time[4] = std::chrono::system_clock::now();
|
||||
HIPCHECK(hipSetDevice(dev));
|
||||
hipLaunchKernelGGL(test_kernel, dim3(desired_blocks), dim3(warp_size), 0,
|
||||
streams[0], loops, dev_array[0]);
|
||||
HIPCHECK(hipGetLastError());
|
||||
HIPCHECK(hipLaunchCooperativeKernelMultiDevice(md_params, 2,
|
||||
hipCooperativeLaunchMultiDeviceNoPostSync));
|
||||
HIPCHECK(hipSetDevice(dev + 1));
|
||||
hipLaunchKernelGGL(test_kernel, dim3(desired_blocks), dim3(warp_size), 0,
|
||||
streams[1], loops, dev_array[1]);
|
||||
for (int i = 0; i < 2; i++) {
|
||||
HIPCHECK(hipSetDevice(dev + i));
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
}
|
||||
end_time[4] = std::chrono::system_clock::now();
|
||||
fast_gpu = -1;
|
||||
|
||||
std::cout << std::endl;
|
||||
std::cout << "Test 5: Launching a multi-GPU cooperative kernel with the ";
|
||||
std::cout << "ability to overlap regular and cooperative kernels";
|
||||
std::cout << std::endl;
|
||||
std::cout << "This should result in the following pattern:" << std::endl;
|
||||
std::cout << "GPU " << dev << ": Standard Kernel --> Long Coop Kernel\n";
|
||||
std::cout << "GPU " << (dev + 1) << ": Long Coop Kernel --> Standard";
|
||||
std::cout << " Kernel\n";
|
||||
start_time[5] = std::chrono::system_clock::now();
|
||||
HIPCHECK(hipSetDevice(dev));
|
||||
hipLaunchKernelGGL(test_kernel, dim3(desired_blocks), dim3(warp_size), 0,
|
||||
streams[0], loops, dev_array[0]);
|
||||
HIPCHECK(hipGetLastError());
|
||||
HIPCHECK(hipLaunchCooperativeKernelMultiDevice(md_params, 2,
|
||||
hipCooperativeLaunchMultiDeviceNoPreSync |
|
||||
hipCooperativeLaunchMultiDeviceNoPostSync));
|
||||
HIPCHECK(hipSetDevice(dev + 1));
|
||||
hipLaunchKernelGGL(test_kernel, dim3(desired_blocks), dim3(warp_size), 0,
|
||||
streams[1], loops, dev_array[1]);
|
||||
HIPCHECK(hipGetLastError());
|
||||
for (int i = 0; i < 2; i++) {
|
||||
HIPCHECK(hipSetDevice(dev + i));
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
}
|
||||
end_time[5] = std::chrono::system_clock::now();
|
||||
|
||||
std::chrono::duration<double> single_kernel_time =
|
||||
(end_time[0] - start_time[0]);
|
||||
std::chrono::duration<double> serialized_gpu0_time =
|
||||
(end_time[1] - start_time[1]);
|
||||
std::chrono::duration<double> serialized_gpu1_time =
|
||||
(end_time[2] - start_time[2]);
|
||||
std::chrono::duration<double> pre_overlapped_time =
|
||||
(end_time[3] - start_time[3]);
|
||||
std::chrono::duration<double> post_overlapped_time =
|
||||
(end_time[4] - start_time[4]);
|
||||
std::chrono::duration<double> overlapped_time =
|
||||
(end_time[5] - start_time[5]);
|
||||
|
||||
std::cout << "Test 0: A single kernel on both GPUs took:" << std::endl;
|
||||
std::cout << " " << single_kernel_time.count();
|
||||
std::cout << " seconds" << std::endl;
|
||||
std::cout << std::endl;
|
||||
std::cout << "Test 1: Serialized set of three kernels with GPU0";
|
||||
std::cout << " being long took:";
|
||||
std::cout << " " << serialized_gpu0_time.count();
|
||||
std::cout << " seconds" << std::endl;
|
||||
std::cerr << "Expect between " << (2.7 * single_kernel_time.count());
|
||||
std::cerr << " and ";
|
||||
std::cerr << (3.3 * single_kernel_time.count()) << " seconds.\n";
|
||||
std::cout << std::endl;
|
||||
std::cout << "Test 2: Serialized set of three kernels with GPU1";
|
||||
std::cout << " being long took:" << std::endl;
|
||||
std::cout << " " << serialized_gpu1_time.count();
|
||||
std::cout << " seconds" << std::endl;
|
||||
std::cerr << "Expect between " << (2.7 * single_kernel_time.count());
|
||||
std::cerr << " and ";
|
||||
std::cerr << (3.3 * single_kernel_time.count()) << " seconds.\n";
|
||||
std::cout << std::endl;
|
||||
std::cout << "Test 3: Multiple kernels with pre-overlap allowed took:\n";
|
||||
std::cout << " " << pre_overlapped_time.count();
|
||||
std::cout << " seconds" << std::endl;
|
||||
std::cerr << "Expect between " << (1.7 * single_kernel_time.count());
|
||||
std::cerr << " and ";
|
||||
std::cerr << (2.3 * single_kernel_time.count()) << " seconds.\n";
|
||||
std::cout << std::endl;
|
||||
std::cout << "Test 4: Multiple kernels with post-overlap allowed took:\n";
|
||||
std::cout << " " << post_overlapped_time.count();
|
||||
std::cout << " seconds" << std::endl;
|
||||
std::cerr << "Expect between " << (1.7 * single_kernel_time.count());
|
||||
std::cerr << " and ";
|
||||
std::cerr << (2.3 * single_kernel_time.count()) << " seconds.";
|
||||
std::cout << std::endl;
|
||||
std::cout << "Test 5: Multiple kernels with overlap allowed took:\n";
|
||||
std::cout << " " << overlapped_time.count();
|
||||
std::cout << " seconds" << std::endl;
|
||||
std::cerr << "Expect between " << (1.8 * single_kernel_time.count());
|
||||
std::cerr << " and ";
|
||||
std::cerr << (2.2 * single_kernel_time.count()) << " seconds.\n";
|
||||
|
||||
// Test that fully not-overlapped kernels take roughly 3x as long as one
|
||||
// cooperative kernel.
|
||||
if (serialized_gpu0_time > 3.3 * single_kernel_time ||
|
||||
serialized_gpu0_time < 2.7 * single_kernel_time) {
|
||||
std::cerr << "ERROR!" << std::endl;
|
||||
std::cerr << "Test 1, the first case where all kernels should be ";
|
||||
std::cerr << "serialized, had a runtime that was very different ";
|
||||
std::cerr << "than what was expected." << std::endl;
|
||||
std::cerr << "Was " << serialized_gpu0_time.count() << " seconds.\n";
|
||||
std::cerr << "Expected between ";
|
||||
std::cerr << (2.7 * single_kernel_time.count()) << " and ";
|
||||
std::cerr << (3.3 * single_kernel_time.count()) << " seconds.\n";
|
||||
std::cerr << "Were they truly serialized?" << std::endl;
|
||||
FailFlag = 1;
|
||||
}
|
||||
|
||||
// Test that fully not-overlapped kernels take roughly 3x as long as one
|
||||
// cooperative kernel.
|
||||
if (serialized_gpu1_time > 3.3 * single_kernel_time ||
|
||||
serialized_gpu1_time < 2.7 * single_kernel_time) {
|
||||
std::cerr << "ERROR!" << std::endl;
|
||||
std::cerr << "Test 2, the second case where all kernels should be ";
|
||||
std::cerr << "serialized, had a runtime that was very different ";
|
||||
std::cerr << "than what was expected." << std::endl;
|
||||
std::cerr << "Was " << serialized_gpu1_time.count();
|
||||
std::cerr << " seconds." << std::endl;
|
||||
std::cerr << "Expected between ";
|
||||
std::cerr << (2.7 * single_kernel_time.count()) << " and ";
|
||||
std::cerr << (3.3 * single_kernel_time.count()) << " seconds.\n";
|
||||
std::cerr << "Were they truly serialized?" << std::endl;
|
||||
FailFlag = 1;
|
||||
}
|
||||
|
||||
// Test that kernels that can overlap only before the cooperative kernel
|
||||
// launches kernels take roughly the same time (in this case)
|
||||
if (pre_overlapped_time > 2.3 * single_kernel_time ||
|
||||
pre_overlapped_time < 1.7 * single_kernel_time) {
|
||||
std::cerr << "ERROR!" << std::endl;
|
||||
std::cerr << "Test 3, the case where the last kernel is serialized, had ";
|
||||
std::cerr << "a runtime that was very different than what was ";
|
||||
std::cerr << "expected." << std::endl;
|
||||
std::cerr << "Was " << pre_overlapped_time.count() << " seconds.\n";
|
||||
std::cerr << "Expected between ";
|
||||
std::cerr << (1.7 * single_kernel_time.count()) << " and ";
|
||||
std::cerr << (2.3 * single_kernel_time.count()) << " seconds.\n";
|
||||
FailFlag = 1;
|
||||
}
|
||||
|
||||
// Test that kernels that can overlap only after the cooperative kernel
|
||||
// launches kernels take roughly the same time (in this case)
|
||||
if (post_overlapped_time > 2.3 * single_kernel_time ||
|
||||
post_overlapped_time < 1.7 * single_kernel_time) {
|
||||
std::cerr << "ERROR!" << std::endl;
|
||||
std::cerr << "Teste 4, the case where the first kernel is ";
|
||||
std::cerr << "serialized, had a runtime that was very different ";
|
||||
std::cerr << "than what was expected." << std::endl;
|
||||
std::cerr << "Was " << post_overlapped_time.count() << " seconds.\n";
|
||||
std::cerr << "Expected between ";
|
||||
std::cerr << (1.7 * single_kernel_time.count()) << " and ";
|
||||
std::cerr << (2.3 * single_kernel_time.count()) << " seconds.\n";
|
||||
FailFlag = 1;
|
||||
}
|
||||
|
||||
// Test that, with the right flags on the kernel launch, that we prevent
|
||||
// incomplete launches from serializing the cooperative launch streams.
|
||||
if (overlapped_time > 2.2 * single_kernel_time ||
|
||||
overlapped_time < 1.8 * single_kernel_time) {
|
||||
std::cerr << "ERROR!" << std::endl;
|
||||
std::cerr << "Test 5, the case where normal and cooperative kernel ";
|
||||
std::cerr << "launches should overlap, does not appear to have done so.";
|
||||
std::cerr << std::endl;
|
||||
std::cerr << "Was " << overlapped_time.count() << " seconds.\n";
|
||||
std::cerr << "Expected between ";
|
||||
std::cerr << (1.8 * single_kernel_time.count()) << " and ";
|
||||
std::cerr << (2.2 * single_kernel_time.count()) << " seconds.\n";
|
||||
std::cerr << "Is the normal kernel being serialized with the ";
|
||||
std::cerr << "cooperative kernels on different streams?" << std::endl;
|
||||
FailFlag = 1;
|
||||
}
|
||||
for (int k = 0; k < 2; ++k) {
|
||||
HIPCHECK(hipFree(dev_array[k]));
|
||||
HIPCHECK(hipStreamDestroy(streams[k]));
|
||||
}
|
||||
if (FailFlag == 1) {
|
||||
break;
|
||||
}
|
||||
}
|
||||
if (FailFlag == 1) {
|
||||
failed("");
|
||||
} else {
|
||||
passed();
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,374 @@
|
||||
/*
|
||||
Copyright (c) 2020-present 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 launch N warps to all GPUs detected
|
||||
in the HIP system. N is a command-line parameter, but the user should set N
|
||||
small enough that all warps can be on each of the GPUs at the same time.
|
||||
|
||||
All of the warps do a "work loop". Within the work loop, every warp
|
||||
atomically increments a global variable that is shared between both fo the
|
||||
target GPUs. The value returned from this atomic increment entriely depends
|
||||
on the order the warps from the GPUs arrive at the atomic instruction. Each
|
||||
warp then stores the result into a global array based on its warp ID.
|
||||
|
||||
We also add a sleep/wait loop into the code so that the last warp runs much
|
||||
slower than everyone else. As such, it should store much larger values than
|
||||
all the other warps.
|
||||
|
||||
If there are no barrier within the loop, then warp 0 will likely ge to the
|
||||
global variable the first time while all the other warps have each
|
||||
incremented it many times. If the barrier properly works, then each warp
|
||||
will increment the variable once per time through the loop, and all threads
|
||||
will sleep on the barrier waiting for the last warp to finally catch up.
|
||||
*/
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <hip/hip_cooperative_groups.h>
|
||||
#include "test_common.h"
|
||||
|
||||
static int cooperative_groups_support(int device_id) {
|
||||
hipError_t err;
|
||||
int cooperative_attribute;
|
||||
HIPCHECK(hipDeviceGetAttribute(&cooperative_attribute,
|
||||
hipDeviceAttributeCooperativeLaunch, device_id));
|
||||
if (!cooperative_attribute) {
|
||||
std::cerr << "Cooperative launch support not available in ";
|
||||
std::cerr << "the device attribute for device " << device_id;
|
||||
std::cerr << std::endl;
|
||||
return 0;
|
||||
}
|
||||
|
||||
int multi_gpu_cooperative_attribute;
|
||||
HIPCHECK(hipDeviceGetAttribute(&multi_gpu_cooperative_attribute,
|
||||
hipDeviceAttributeCooperativeMultiDeviceLaunch, device_id));
|
||||
if (!multi_gpu_cooperative_attribute) {
|
||||
std::cerr << "Multi-GPU cooperative launch support not available in ";
|
||||
std::cerr << "the device attribute for device " << device_id;
|
||||
std::cerr << std::endl;
|
||||
return 0;
|
||||
}
|
||||
|
||||
hipDeviceProp_t device_properties;
|
||||
HIPCHECK(hipGetDeviceProperties(&device_properties, device_id));
|
||||
if (device_properties.cooperativeLaunch == 0) {
|
||||
std::cerr << "Cooperative group support not available in ";
|
||||
std::cerr << "device properties." << std::endl;
|
||||
return 0;
|
||||
}
|
||||
if (device_properties.cooperativeMultiDeviceLaunch == 0) {
|
||||
std::cerr << "Multi-GPU cooperative group support not available in ";
|
||||
std::cerr << "device properties." << std::endl;
|
||||
return 0;
|
||||
}
|
||||
return 1;
|
||||
}
|
||||
|
||||
static int 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++) {
|
||||
if (host_buffer[i*warps+j] > max_in_this_loop) {
|
||||
std::cerr << "Barrier failure!" << std::endl;
|
||||
std::cerr << " Buffer entry " << i*warps+j;
|
||||
std::cerr << " contains the value " << host_buffer[i*warps+j];
|
||||
std::cerr << " but it should not be more than ";
|
||||
std::cerr << max_in_this_loop << std::endl;
|
||||
return -1;
|
||||
}
|
||||
}
|
||||
}
|
||||
std::cout << "\tBarriers work properly!" << std::endl;
|
||||
return 0;
|
||||
}
|
||||
|
||||
static int 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;
|
||||
}
|
||||
}
|
||||
std::cout << "Desired value is " << desired_val << std::endl;
|
||||
if (array_val != desired_val) {
|
||||
std::cerr << "ERROR! Multi-grid barrier does not appear to work.";
|
||||
std::cerr << std::endl;
|
||||
std::cerr << "Expected the multi-GPUs to work together to produce ";
|
||||
std::cerr << "the value " << desired_val << std::endl;
|
||||
std::cerr << "However, the entry returned from the multi-GPU ";
|
||||
std::cerr << "kernel was " << array_val << std::endl;
|
||||
return -1;
|
||||
}
|
||||
std::cout << "\tMulti-GPU barriers appear to work here." << std::endl;
|
||||
return 0;
|
||||
}
|
||||
|
||||
__global__ void
|
||||
test_kernel(unsigned int *atomic_val, unsigned int *global_array,
|
||||
unsigned int *array, uint32_t loops) {
|
||||
cooperative_groups::grid_group grid = cooperative_groups::this_grid();
|
||||
cooperative_groups::multi_grid_group mgrid =
|
||||
cooperative_groups::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 start_clock = clock64();
|
||||
while (clock64() < (start_clock+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 start_clock = clock64();
|
||||
while (clock64() < (start_clock+100000000)) {}
|
||||
}
|
||||
// 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;
|
||||
}
|
||||
}
|
||||
|
||||
int main(int argc, char** argv) {
|
||||
hipError_t err;
|
||||
int num_devices = 0;
|
||||
uint32_t loops = 2;
|
||||
uint32_t warps = 10;
|
||||
uint32_t block_size = 1;
|
||||
|
||||
std::cout << "Loops: " << loops << std::endl;
|
||||
std::cout << "Warps: " << warps << std::endl;
|
||||
std::cout << "Block size: " << block_size << std::endl;
|
||||
|
||||
HIPCHECK(hipGetDeviceCount(&num_devices));
|
||||
if (num_devices < 2) {
|
||||
std::cout << "Not enough GPUs to run test." << std::endl;
|
||||
std::cout << "We require at least 2 GPUs, but only found ";
|
||||
std::cout << num_devices << std::endl;
|
||||
std::cout << "Skipping the test with PASSED result\n";
|
||||
passed();
|
||||
}
|
||||
|
||||
uint32_t device_num[num_devices];
|
||||
|
||||
/*************************************************************************/
|
||||
/* Test whether target device supports cooperative groups ****************/
|
||||
for (int i = 0; i < num_devices; i++) {
|
||||
device_num[i] = i;
|
||||
if (!cooperative_groups_support(device_num[i])) {
|
||||
std::cout << "Skipping the test with Pass result.\n";
|
||||
passed();
|
||||
}
|
||||
}
|
||||
|
||||
/*************************************************************************/
|
||||
/* Test whether the requested size will fit on the GPU *******************/
|
||||
int warp_sizes[num_devices];
|
||||
int num_sms[num_devices];
|
||||
hipDeviceProp_t device_properties[num_devices];
|
||||
int warp_size = INT_MAX;
|
||||
int num_sm = INT_MAX;
|
||||
for (int i = 0; i < num_devices; i++) {
|
||||
HIPCHECK(hipGetDeviceProperties(&device_properties[i], device_num[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];
|
||||
}
|
||||
std::cout << "Device " << (i + 1);
|
||||
std::cout << " name: " << device_properties[i].name << std::endl;
|
||||
}
|
||||
std::cout << std::endl;
|
||||
|
||||
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++) {
|
||||
HIPCHECK(hipSetDevice(device_num[i]));
|
||||
HIPCHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(
|
||||
&max_blocks_per_sm_arr[i], test_kernel, 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;
|
||||
if (requested_blocks > max_blocks_per_sm * num_sm) {
|
||||
std::cerr << "Requesting to run " << requested_blocks << " blocks, ";
|
||||
std::cerr << "but we can only guarantee to simultaneously run ";
|
||||
std::cerr << (max_blocks_per_sm * num_sm) << std::endl;
|
||||
failed("");
|
||||
}
|
||||
|
||||
/*************************************************************************/
|
||||
/* Set up data to pass into the kernel ***********************************/
|
||||
// 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] = (unsigned int*)calloc(total_buffer_len,
|
||||
sizeof(unsigned int));
|
||||
HIPCHECK(hipSetDevice(device_num[i]));
|
||||
HIPCHECK(hipMalloc(reinterpret_cast<void**>(&kernel_buffer[i]),
|
||||
total_buffer_len * sizeof(unsigned int)));
|
||||
HIPCHECK(hipMemcpy(kernel_buffer[i], host_buffer[i],
|
||||
total_buffer_len * sizeof(unsigned int),
|
||||
hipMemcpyHostToDevice));
|
||||
HIPCHECK(hipMalloc(reinterpret_cast<void**>(&kernel_atomic[i]),
|
||||
sizeof(unsigned int)));
|
||||
HIPCHECK(hipMemset(kernel_atomic[i], 0, sizeof(unsigned int)));
|
||||
HIPCHECK(hipStreamCreate(&streams[i]));
|
||||
}
|
||||
|
||||
// Single kernel atomic shared between both devices; put it on the host
|
||||
unsigned int* global_array;
|
||||
HIPCHECK(hipHostMalloc(reinterpret_cast<void**>(&global_array),
|
||||
num_devices * sizeof(unsigned int), 0));
|
||||
HIPCHECK(hipMemset(global_array, 0, num_devices * sizeof(unsigned int)));
|
||||
|
||||
/*************************************************************************/
|
||||
/* Launch the kernels ****************************************************/
|
||||
std::cout << "Launching a kernel with " << warps << " warps ";
|
||||
std::cout << "in " << requested_blocks << " thread blocks.";
|
||||
std::cout << std::endl;
|
||||
|
||||
void *dev_params[num_devices][4];
|
||||
hipLaunchParams md_params[num_devices];
|
||||
for (int i = 0; i < num_devices; i++) {
|
||||
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);
|
||||
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];
|
||||
}
|
||||
|
||||
HIPCHECK(hipLaunchCooperativeKernelMultiDevice(md_params, num_devices, 0));
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
/*************************************************************************/
|
||||
/* Read back the buffers and print out its data **************************/
|
||||
for (int dev = 0; dev < num_devices; dev++) {
|
||||
HIPCHECK(hipMemcpy(host_buffer[dev], kernel_buffer[dev],
|
||||
total_buffer_len * sizeof(unsigned int),
|
||||
hipMemcpyDeviceToHost));
|
||||
}
|
||||
|
||||
for (unsigned int i = 0; i < loops; i++) {
|
||||
for (int dev = 0; dev < num_devices; dev++) {
|
||||
std::cout << "+++++++++++++++++ Device " << dev;
|
||||
std::cout << "+++++++++++++++++" << std::endl;
|
||||
for (unsigned int j = 0; j < requested_blocks; j++) {
|
||||
std::cout << "Buffer entry " << (i*warps+j);
|
||||
std::cout << " (written by warp " << j << ")";
|
||||
std::cout << " is " << host_buffer[dev][i*requested_blocks+j];
|
||||
std::cout << std::endl;
|
||||
}
|
||||
}
|
||||
std::cout << "==========================\n";
|
||||
}
|
||||
for (unsigned int dev = 0; dev < num_devices; dev++) {
|
||||
std::cout << "Testing output from device " << dev << std::endl;
|
||||
int local_ret_val = verify_barrier_buffer(loops, requested_blocks,
|
||||
host_buffer[dev], num_devices);
|
||||
if (local_ret_val) {
|
||||
failed("");
|
||||
}
|
||||
}
|
||||
|
||||
std::cout << std::endl << "The multi-GPU shared updates contain:\n";
|
||||
for (int i = 0; i < num_devices; i++) {
|
||||
std::cout << "Entry " << i << ": ";
|
||||
std::cout << global_array[i] << std::endl;
|
||||
}
|
||||
int flag = 0;
|
||||
for (int dev = 0; dev < num_devices; dev++) {
|
||||
std::cout << "Testing multi-GPU output for entry " << dev << std::endl;
|
||||
int local_ret_val = verify_multi_gpu_buffer(loops, global_array[dev]);
|
||||
if (local_ret_val) {
|
||||
flag = 1;
|
||||
}
|
||||
}
|
||||
for (int k = 0; k < num_devices; ++k) {
|
||||
HIPCHECK(hipFree(kernel_buffer[k]));
|
||||
HIPCHECK(hipFree(kernel_atomic[k]));
|
||||
HIPCHECK(hipStreamDestroy(streams[k]));
|
||||
free(host_buffer[k]);
|
||||
}
|
||||
if (flag == 1) {
|
||||
failed("");
|
||||
} else {
|
||||
passed();
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,233 @@
|
||||
/*
|
||||
Copyright (c) 2020-present 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 launch N warps. N is a command-line
|
||||
parameter, but the user should set N small enough that all warps can be on
|
||||
the GPU at the same time.
|
||||
|
||||
All of the warps do a "work loop". Within the work loop, every warp
|
||||
atomically increments a global variable. The value returned from this atomic
|
||||
increment entriely depends on the order the threads arrive at the atomic
|
||||
instruction. Each warp then stores the result into a global array based on its
|
||||
warp ID.
|
||||
|
||||
We also add a sleep/wait loop into the code so that the last warp runs much
|
||||
slower than everyone else. As such, it should store much larger values than
|
||||
all the other warps.
|
||||
|
||||
If there are no barrier within the loop, then the last warp will likely get to
|
||||
the global variable the first time after all the other warps have each
|
||||
incremented it many times. If the barrier properly works, then each warp
|
||||
will increment the variable once per time through the loop, and all threads
|
||||
will sleep on the barrier waiting for the last warp to finally catch up.
|
||||
*/
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <hip/hip_cooperative_groups.h>
|
||||
#include "test_common.h"
|
||||
|
||||
static int cooperative_groups_support(int device_id) {
|
||||
hipError_t err;
|
||||
int cooperative_attribute;
|
||||
HIPCHECK(hipDeviceGetAttribute(&cooperative_attribute,
|
||||
hipDeviceAttributeCooperativeLaunch, device_id));
|
||||
if (!cooperative_attribute) {
|
||||
std::cerr << "Cooperative launch support not available in ";
|
||||
std::cerr << "the device attribute for device " << device_id;
|
||||
std::cerr << std::endl;
|
||||
return 0;
|
||||
}
|
||||
|
||||
hipDeviceProp_t device_properties;
|
||||
HIPCHECK(hipGetDeviceProperties(&device_properties, device_id));
|
||||
if (device_properties.cooperativeLaunch == 0) {
|
||||
std::cerr << "Cooperative group support not available in ";
|
||||
std::cerr << "device properties." << std::endl;
|
||||
return 0;
|
||||
}
|
||||
return 1;
|
||||
}
|
||||
|
||||
static int 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++) {
|
||||
if (host_buffer[i*warps+j] > max_in_this_loop) {
|
||||
std::cerr << "Barrier failure!" << std::endl;
|
||||
std::cerr << " Buffer entry " << i*warps+j;
|
||||
std::cerr << " contains the value " << host_buffer[i*warps+j];
|
||||
std::cerr << " but it should not be more than ";
|
||||
std::cerr << max_in_this_loop << std::endl;
|
||||
return -1;
|
||||
}
|
||||
}
|
||||
}
|
||||
std::cout << "Barriers work properly!" << std::endl;
|
||||
return 0;
|
||||
}
|
||||
|
||||
__global__ void
|
||||
test_kernel(unsigned int *atomic_val, unsigned int *array,
|
||||
unsigned int loops) {
|
||||
cooperative_groups::grid_group grid = cooperative_groups::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 start_clock = clock64();
|
||||
while (clock64() < (start_clock+1000000)) {}
|
||||
}
|
||||
|
||||
if (threadIdx.x == 0) {
|
||||
array[offset] = atomicInc(&atomic_val[0], UINT_MAX);
|
||||
}
|
||||
grid.sync();
|
||||
offset += gridDim.x;
|
||||
}
|
||||
}
|
||||
|
||||
int main(int argc, char** argv) {
|
||||
hipError_t err;
|
||||
int device_num;
|
||||
uint32_t loops = 2;
|
||||
uint32_t warps = 10;
|
||||
uint32_t block_size = 1;
|
||||
HIPCHECK(hipGetDeviceCount(&device_num));
|
||||
for (int dev = 0; dev < device_num; ++dev) {
|
||||
std::cout << "Device number: " << dev << std::endl;
|
||||
std::cout << "Loops: " << loops << std::endl;
|
||||
std::cout << "Warps: " << warps << std::endl;
|
||||
std::cout << "Block size: " << block_size << std::endl;
|
||||
|
||||
/*************************************************************************/
|
||||
/* Test whether target device supports cooperative groups ****************/
|
||||
HIPCHECK(hipSetDevice(dev));
|
||||
if (!cooperative_groups_support(dev)) {
|
||||
std::cout << "Skipping the test with Pass result.\n";
|
||||
passed();
|
||||
}
|
||||
|
||||
/*************************************************************************/
|
||||
/* Test whether the requested size will fit on the GPU *******************/
|
||||
int warp_size;
|
||||
int num_sms;
|
||||
int max_blocks_per_sm;
|
||||
hipDeviceProp_t device_properties;
|
||||
HIPCHECK(hipGetDeviceProperties(&device_properties, dev));
|
||||
warp_size = device_properties.warpSize;
|
||||
num_sms = device_properties.multiProcessorCount;
|
||||
|
||||
std::cout << "Device name: " << device_properties.name << std::endl;
|
||||
std::cout << std::endl;
|
||||
|
||||
int num_threads_in_block = block_size * warp_size;
|
||||
|
||||
// Calculate the device occupancy to know how many blocks can be run.
|
||||
HIPCHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&max_blocks_per_sm,
|
||||
test_kernel, num_threads_in_block, 0));
|
||||
|
||||
int requested_blocks = warps / block_size;
|
||||
if (requested_blocks > max_blocks_per_sm * num_sms) {
|
||||
std::cerr << "Requesting to run " << requested_blocks << " blocks, ";
|
||||
std::cerr << "but we can only guarantee to simultaneously run ";
|
||||
std::cerr << (max_blocks_per_sm * num_sms) << std::endl;
|
||||
failed("");
|
||||
}
|
||||
|
||||
/*************************************************************************/
|
||||
/* Set up data to pass into the kernel ***********************************/
|
||||
// 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 = (unsigned int*)calloc(total_buffer_len,
|
||||
sizeof(unsigned int));
|
||||
|
||||
unsigned int *kernel_buffer;
|
||||
HIPCHECK(hipMalloc(reinterpret_cast<void**>(&kernel_buffer),
|
||||
total_buffer_len * sizeof(unsigned int)));
|
||||
HIPCHECK(hipMemcpy(kernel_buffer, host_buffer,
|
||||
total_buffer_len * sizeof(unsigned int),
|
||||
hipMemcpyHostToDevice));
|
||||
|
||||
unsigned int *kernel_atomic;
|
||||
HIPCHECK(hipMalloc(reinterpret_cast<void**>(&kernel_atomic),
|
||||
sizeof(unsigned int)));
|
||||
HIPCHECK(hipMemset(kernel_atomic, 0, sizeof(unsigned int)));
|
||||
|
||||
/*************************************************************************/
|
||||
/* Launch the kernel *****************************************************/
|
||||
std::cout << "Launching a kernel with " << warps << " warps ";
|
||||
std::cout << "in " << requested_blocks << " thread blocks.";
|
||||
std::cout << std::endl;
|
||||
|
||||
void *params[3];
|
||||
params[0] = reinterpret_cast<void*>(&kernel_atomic);
|
||||
params[1] = reinterpret_cast<void*>(&kernel_buffer);
|
||||
params[2] = reinterpret_cast<void*>(&loops);
|
||||
HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast<void*>(test_kernel),
|
||||
requested_blocks,
|
||||
num_threads_in_block, params, 0, NULL));
|
||||
|
||||
/*************************************************************************/
|
||||
/* Read back the buffer and print out its data****************************/
|
||||
HIPCHECK(hipMemcpy(host_buffer, kernel_buffer,
|
||||
total_buffer_len * sizeof(unsigned int),
|
||||
hipMemcpyDeviceToHost));
|
||||
|
||||
for (unsigned int i = 0; i < loops; i++) {
|
||||
for (unsigned int j = 0; j < requested_blocks; j++) {
|
||||
std::cout << "Buffer entry " << (i*warps+j);
|
||||
std::cout << " (written by warp " << j << ")";
|
||||
std::cout << " is " << host_buffer[i * requested_blocks + j];
|
||||
std::cout << std::endl;
|
||||
}
|
||||
std::cout << "==========================\n";
|
||||
}
|
||||
int ret_val = verify_barrier_buffer(loops, requested_blocks, host_buffer);
|
||||
HIPCHECK(hipFree(kernel_buffer));
|
||||
HIPCHECK(hipFree(kernel_atomic));
|
||||
if (ret_val == -1) {
|
||||
failed("");
|
||||
} else {
|
||||
passed();
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,374 @@
|
||||
/*
|
||||
Copyright (c) 2020-present 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 launch N warps to each of two GPUs.
|
||||
N is a command-line parameter, but the user should set N small enough that all
|
||||
warps can be on each of the GPUs at the same time.
|
||||
|
||||
All of the warps do a "work loop". Within the work loop, every warp
|
||||
atomically increments a global variable that is shared between both fo the
|
||||
target GPUs. The value returned from this atomic increment entriely depends
|
||||
on the order the warps from the GPUs arrive at the atomic instruction. Each
|
||||
warp then stores the result into a global array based on its warp ID.
|
||||
|
||||
We also add a sleep/wait loop into the code so that the last warp runs much
|
||||
slower than everyone else. As such, it should store much larger values than
|
||||
all the other warps.
|
||||
|
||||
If there are no barrier within the loop, then warp 0 will likely ge to the
|
||||
global variable the first time while all the other warps have each
|
||||
incremented it many times. If the barrier properly works, then each warp
|
||||
will increment the variable once per time through the loop, and all threads
|
||||
will sleep on the barrier waiting for the last warp to finally catch up.
|
||||
*/
|
||||
|
||||
/* HIT_START
|
||||
* BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc
|
||||
* TEST: %t
|
||||
* HIT_END
|
||||
*/
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
#include <hip/hip_cooperative_groups.h>
|
||||
#include "test_common.h"
|
||||
|
||||
static int cooperative_groups_support(int device_id) {
|
||||
hipError_t err;
|
||||
int cooperative_attribute;
|
||||
HIPCHECK(hipDeviceGetAttribute(&cooperative_attribute,
|
||||
hipDeviceAttributeCooperativeLaunch, device_id));
|
||||
if (!cooperative_attribute) {
|
||||
std::cerr << "Cooperative launch support not available in ";
|
||||
std::cerr << "the device attribute for device " << device_id;
|
||||
std::cerr << std::endl;
|
||||
return 0;
|
||||
}
|
||||
|
||||
int multi_gpu_cooperative_attribute;
|
||||
HIPCHECK(hipDeviceGetAttribute(&multi_gpu_cooperative_attribute,
|
||||
hipDeviceAttributeCooperativeMultiDeviceLaunch, device_id));
|
||||
if (!multi_gpu_cooperative_attribute) {
|
||||
std::cerr << "Multi-GPU cooperative launch support not available in ";
|
||||
std::cerr << "the device attribute for device " << device_id;
|
||||
std::cerr << std::endl;
|
||||
return 0;
|
||||
}
|
||||
|
||||
hipDeviceProp_t device_properties;
|
||||
HIPCHECK(hipGetDeviceProperties(&device_properties, device_id));
|
||||
if (device_properties.cooperativeLaunch == 0) {
|
||||
std::cerr << "Cooperative group support not available in ";
|
||||
std::cerr << "device properties." << std::endl;
|
||||
return 0;
|
||||
}
|
||||
if (device_properties.cooperativeMultiDeviceLaunch == 0) {
|
||||
std::cerr << "Multi-GPU cooperative group support not available in ";
|
||||
std::cerr << "device properties." << std::endl;
|
||||
return 0;
|
||||
}
|
||||
return 1;
|
||||
}
|
||||
|
||||
static int 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++) {
|
||||
if (host_buffer[i*warps+j] > max_in_this_loop) {
|
||||
std::cerr << "Barrier failure!" << std::endl;
|
||||
std::cerr << " Buffer entry " << i*warps+j;
|
||||
std::cerr << " contains the value " << host_buffer[i*warps+j];
|
||||
std::cerr << " but it should not be more than ";
|
||||
std::cerr << max_in_this_loop << std::endl;
|
||||
return -1;
|
||||
}
|
||||
}
|
||||
}
|
||||
std::cout << "\tBarriers work properly!" << std::endl;
|
||||
return 0;
|
||||
}
|
||||
|
||||
static int 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;
|
||||
}
|
||||
}
|
||||
std::cout << "Desired value is " << desired_val << std::endl;
|
||||
if (array_val != desired_val) {
|
||||
std::cerr << "ERROR! Multi-grid barrier does not appear to work.";
|
||||
std::cerr << std::endl;
|
||||
std::cerr << "Expected the multi-GPUs to work together to produce ";
|
||||
std::cerr << "the value " << desired_val << std::endl;
|
||||
std::cerr << "However, the entry returned from the multi-GPU ";
|
||||
std::cerr << "kernel was " << array_val << std::endl;
|
||||
return -1;
|
||||
}
|
||||
std::cout << "\tMulti-GPU barriers appear to work here." << std::endl;
|
||||
return 0;
|
||||
}
|
||||
|
||||
__global__ void
|
||||
test_kernel(unsigned int *atomic_val, unsigned int *global_array,
|
||||
unsigned int *array, uint32_t loops) {
|
||||
cooperative_groups::grid_group grid = cooperative_groups::this_grid();
|
||||
cooperative_groups::multi_grid_group mgrid =
|
||||
cooperative_groups::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 start_clock = clock64();
|
||||
while (clock64() < (start_clock + 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 start_clock = clock64();
|
||||
while (clock64() < (start_clock + 100000000)) {}
|
||||
}
|
||||
// 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;
|
||||
}
|
||||
}
|
||||
|
||||
int main(int argc, char** argv) {
|
||||
hipError_t err;
|
||||
int device_num = 0, flag = 0;
|
||||
uint32_t loops = 2;
|
||||
uint32_t warps = 10;
|
||||
uint32_t block_size = 1;
|
||||
HIPCHECK(hipGetDeviceCount(&device_num));
|
||||
if (device_num < 2) {
|
||||
std::cout << "This test needs atleast two gpus but found only";
|
||||
std::cout << device_num << std::endl;
|
||||
std::cout << "Hence skipping the test with pass result\n";
|
||||
passed();
|
||||
}
|
||||
|
||||
for (int d = 0; d < (device_num - 1); ++d) {
|
||||
std::cout << "First device number: " << d << std::endl;
|
||||
std::cout << "Second device number: " << (d + 1) << std::endl;
|
||||
std::cout << "Loops: " << loops << std::endl;
|
||||
std::cout << "Warps: " << warps << std::endl;
|
||||
std::cout << "Block size: " << block_size << std::endl;
|
||||
|
||||
/*************************************************************************/
|
||||
/* Test whether target device supports cooperative groups ****************/
|
||||
for (int i = 0; i < 2; i++) {
|
||||
if (!cooperative_groups_support((d + i))) {
|
||||
std::cout << "Skipping the test with Pass result.\n";
|
||||
passed();
|
||||
}
|
||||
}
|
||||
|
||||
/*************************************************************************/
|
||||
/* Test whether the requested size will fit on 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 i = 0; i < 2; i++) {
|
||||
HIPCHECK(hipGetDeviceProperties(&device_properties[i], (d + 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];
|
||||
}
|
||||
std::cout << "Device " << (d + i);
|
||||
std::cout << " name: " << device_properties[i].name << std::endl;
|
||||
}
|
||||
std::cout << std::endl;
|
||||
|
||||
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[2];
|
||||
int max_blocks_per_sm = INT_MAX;
|
||||
for (int i = 0; i < 2; i++) {
|
||||
HIPCHECK(hipSetDevice((d + i)));
|
||||
HIPCHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(
|
||||
&max_blocks_per_sm_arr[i], test_kernel, 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;
|
||||
if (requested_blocks > max_blocks_per_sm * num_sm) {
|
||||
std::cerr << "Requesting to run " << requested_blocks << " blocks, ";
|
||||
std::cerr << "but we can only guarantee to simultaneously run ";
|
||||
std::cerr << (max_blocks_per_sm * num_sm) << std::endl;
|
||||
failed("");
|
||||
}
|
||||
|
||||
/*************************************************************************/
|
||||
/* Set up data to pass into the kernel ***********************************/
|
||||
// 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[2];
|
||||
unsigned int *kernel_buffer[2];
|
||||
unsigned int *kernel_atomic[2];
|
||||
hipStream_t streams[2];
|
||||
for (int i = 0; i < 2; i++) {
|
||||
host_buffer[i] = (unsigned int*)calloc(total_buffer_len,
|
||||
sizeof(unsigned int));
|
||||
HIPCHECK(hipSetDevice((d + i)));
|
||||
HIPCHECK(hipMalloc(reinterpret_cast<void**>(&kernel_buffer[i]),
|
||||
total_buffer_len * sizeof(unsigned int)));
|
||||
HIPCHECK(hipMemcpy(kernel_buffer[i], host_buffer[i],
|
||||
total_buffer_len * sizeof(unsigned int), hipMemcpyHostToDevice));
|
||||
HIPCHECK(hipMalloc(reinterpret_cast<void**>(&kernel_atomic[i]),
|
||||
sizeof(unsigned int)));
|
||||
HIPCHECK(hipMemset(kernel_atomic[i], 0, sizeof(unsigned int)));
|
||||
HIPCHECK(hipStreamCreate(&streams[i]));
|
||||
}
|
||||
|
||||
// Single kernel atomic shared between both devices; put it on the host
|
||||
unsigned int* global_array;
|
||||
HIPCHECK(hipHostMalloc(reinterpret_cast<void**>(&global_array),
|
||||
2 * sizeof(unsigned int), 0));
|
||||
HIPCHECK(hipMemset(global_array, 0, 2 * sizeof(unsigned int)));
|
||||
|
||||
/*************************************************************************/
|
||||
/* Launch the kernels ****************************************************/
|
||||
std::cout << "Launching a kernel with " << warps << " warps ";
|
||||
std::cout << "in " << requested_blocks << " thread blocks.";
|
||||
std::cout << std::endl;
|
||||
|
||||
void *dev_params[2][4];
|
||||
hipLaunchParams md_params[2];
|
||||
for (int i = 0; i < 2; i++) {
|
||||
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);
|
||||
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];
|
||||
}
|
||||
|
||||
HIPCHECK(hipLaunchCooperativeKernelMultiDevice(md_params, 2, 0));
|
||||
HIPCHECK(hipDeviceSynchronize());
|
||||
|
||||
/*************************************************************************/
|
||||
/* Read back the buffers and print out its data **************************/
|
||||
for (int dev = 0; dev < 2; dev++) {
|
||||
HIPCHECK(hipMemcpy(host_buffer[d + dev], kernel_buffer[d + dev],
|
||||
total_buffer_len * sizeof(unsigned int),
|
||||
hipMemcpyDeviceToHost));
|
||||
}
|
||||
|
||||
for (unsigned int i = 0; i < loops; i++) {
|
||||
for (int dev = 0; dev < 2; dev++) {
|
||||
std::cout << "+++++++++++++++++ Device " << (d + dev);
|
||||
std::cout << "+++++++++++++++++" << std::endl;
|
||||
for (unsigned int j = 0; j < requested_blocks; j++) {
|
||||
std::cout << "Buffer entry " << (i * warps + j);
|
||||
std::cout << " (written by warp " << j << ")";
|
||||
std::cout << " is " << host_buffer[dev][i * requested_blocks + j];
|
||||
std::cout << std::endl;
|
||||
}
|
||||
}
|
||||
std::cout << "==========================\n";
|
||||
}
|
||||
for (unsigned int dev = 0; dev < 2; dev++) {
|
||||
std::cout << "Testing output from device " << (d + dev) << std::endl;
|
||||
int local_ret_val = verify_barrier_buffer(loops, requested_blocks,
|
||||
host_buffer[dev], 2);
|
||||
if (local_ret_val == -1) {
|
||||
flag = 1;
|
||||
}
|
||||
}
|
||||
|
||||
std::cout << std::endl << "The multi-GPU shared updates contain:";
|
||||
std::cout << std::endl;
|
||||
for (int i = 0; i < 2; i++) {
|
||||
std::cout << "Entry " << i << ": ";
|
||||
std::cout << global_array[i] << std::endl;
|
||||
}
|
||||
for (int dev = 0; dev < 2; dev++) {
|
||||
std::cout << "Testing multi-GPU output for entry " << (d + dev);
|
||||
std::cout << std::endl;
|
||||
int local_ret_val = verify_multi_gpu_buffer(loops, global_array[dev]);
|
||||
if (local_ret_val) {
|
||||
flag = 1;
|
||||
}
|
||||
}
|
||||
for (int k = 0; k < 2; ++k) {
|
||||
HIPCHECK(hipFree(kernel_buffer[k]));
|
||||
HIPCHECK(hipFree(kernel_atomic[k]));
|
||||
HIPCHECK(hipStreamDestroy(streams[k]));
|
||||
free(host_buffer[k]);
|
||||
}
|
||||
}
|
||||
if (flag == 1) {
|
||||
failed("");
|
||||
} else {
|
||||
passed();
|
||||
}
|
||||
}
|
||||
Reference in New Issue
Block a user