From 40ca4a5ea8b45f5e44e6e17cb67944b6a432b16a Mon Sep 17 00:00:00 2001 From: Lakhan singh Thakur Date: Thu, 3 Sep 2020 17:37:43 +0530 Subject: [PATCH] [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 --- .../cooperativeGrps/api_failure_tests.cpp | 280 +++++++++ .../cooperativeGrps/cooperative_streams.cpp | 283 +++++++++ .../grid_group_data_sharing.cpp | 303 +++++++++ .../multi_gpu_api_failure_tests.cpp | 568 +++++++++++++++++ .../cooperativeGrps/multi_gpu_streams.cpp | 581 ++++++++++++++++++ .../multi_grid_group_all_gpus.cpp | 374 +++++++++++ .../simple_grid_group_barrier.cpp | 233 +++++++ .../simple_multi_grid_group_barrier.cpp | 374 +++++++++++ 8 files changed, 2996 insertions(+) create mode 100644 hipamd/tests/src/runtimeApi/cooperativeGrps/api_failure_tests.cpp create mode 100644 hipamd/tests/src/runtimeApi/cooperativeGrps/cooperative_streams.cpp create mode 100644 hipamd/tests/src/runtimeApi/cooperativeGrps/grid_group_data_sharing.cpp create mode 100644 hipamd/tests/src/runtimeApi/cooperativeGrps/multi_gpu_api_failure_tests.cpp create mode 100644 hipamd/tests/src/runtimeApi/cooperativeGrps/multi_gpu_streams.cpp create mode 100644 hipamd/tests/src/runtimeApi/cooperativeGrps/multi_grid_group_all_gpus.cpp create mode 100644 hipamd/tests/src/runtimeApi/cooperativeGrps/simple_grid_group_barrier.cpp create mode 100644 hipamd/tests/src/runtimeApi/cooperativeGrps/simple_multi_grid_group_barrier.cpp diff --git a/hipamd/tests/src/runtimeApi/cooperativeGrps/api_failure_tests.cpp b/hipamd/tests/src/runtimeApi/cooperativeGrps/api_failure_tests.cpp new file mode 100644 index 0000000000..81a733eb07 --- /dev/null +++ b/hipamd/tests/src/runtimeApi/cooperativeGrps/api_failure_tests.cpp @@ -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 +#include +#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(&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(&dev_array[i]); + } + + err = hipLaunchCooperativeKernel(reinterpret_cast(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(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(); +} diff --git a/hipamd/tests/src/runtimeApi/cooperativeGrps/cooperative_streams.cpp b/hipamd/tests/src/runtimeApi/cooperativeGrps/cooperative_streams.cpp new file mode 100644 index 0000000000..580bd1aa78 --- /dev/null +++ b/hipamd/tests/src/runtimeApi/cooperativeGrps/cooperative_streams.cpp @@ -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 +#include +#include +#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(&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(&loops); + coop_params[i][1] = reinterpret_cast(&dev_array[i]); + } + + std::cout << "Launching a single cooperative kernel..." << std::endl; + auto single_start = std::chrono::system_clock::now(); + HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast(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(test_kernel), + desired_blocks, warp_size, + coop_params[0], 0, streams[0])); + HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast(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(test_kernel), + desired_blocks, warp_size, + coop_params[0], 0, streams[0])); + HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast(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 single_kernel_time = + (single_end - single_start); + std::chrono::duration double_kernel_time = + (double_end - double_start); + std::chrono::duration 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(); +} diff --git a/hipamd/tests/src/runtimeApi/cooperativeGrps/grid_group_data_sharing.cpp b/hipamd/tests/src/runtimeApi/cooperativeGrps/grid_group_data_sharing.cpp new file mode 100644 index 0000000000..8fd0ac4776 --- /dev/null +++ b/hipamd/tests/src/runtimeApi/cooperativeGrps/grid_group_data_sharing.cpp @@ -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 +#include +#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(&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(&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(&first_dev_array); + coop_params[1] = reinterpret_cast(&second_dev_array); + coop_params[2] = reinterpret_cast(&loops); + coop_params[3] = reinterpret_cast(&width); + HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast(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(""); + } +} diff --git a/hipamd/tests/src/runtimeApi/cooperativeGrps/multi_gpu_api_failure_tests.cpp b/hipamd/tests/src/runtimeApi/cooperativeGrps/multi_gpu_api_failure_tests.cpp new file mode 100644 index 0000000000..002afaafdd --- /dev/null +++ b/hipamd/tests/src/runtimeApi/cooperativeGrps/multi_gpu_api_failure_tests.cpp @@ -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 +#include +#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 <(&good_dev_array[i]), + good_size)); + HIPCHECK(hipMemsetAsync(good_dev_array[i], 0, good_size, streams[i])); + HIPCHECK(hipMalloc(reinterpret_cast(&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(&bad_dev_array[i]); + + md_params[i].func = reinterpret_cast(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(&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(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(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(); + } +} diff --git a/hipamd/tests/src/runtimeApi/cooperativeGrps/multi_gpu_streams.cpp b/hipamd/tests/src/runtimeApi/cooperativeGrps/multi_gpu_streams.cpp new file mode 100644 index 0000000000..155cc1924f --- /dev/null +++ b/hipamd/tests/src/runtimeApi/cooperativeGrps/multi_gpu_streams.cpp @@ -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 +#include +#include +#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 < 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(&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 start_time[6]; + std::chrono::time_point 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(&loops); + dev_params[i][1] = reinterpret_cast(&dev_array[i]); + dev_params[i][2] = reinterpret_cast(&fast_gpu); + md_params[i].func = reinterpret_cast(test_coop_kernel); + 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 single_kernel_time = + (end_time[0] - start_time[0]); + std::chrono::duration serialized_gpu0_time = + (end_time[1] - start_time[1]); + std::chrono::duration serialized_gpu1_time = + (end_time[2] - start_time[2]); + std::chrono::duration pre_overlapped_time = + (end_time[3] - start_time[3]); + std::chrono::duration post_overlapped_time = + (end_time[4] - start_time[4]); + std::chrono::duration 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(); + } +} diff --git a/hipamd/tests/src/runtimeApi/cooperativeGrps/multi_grid_group_all_gpus.cpp b/hipamd/tests/src/runtimeApi/cooperativeGrps/multi_grid_group_all_gpus.cpp new file mode 100644 index 0000000000..b68c909180 --- /dev/null +++ b/hipamd/tests/src/runtimeApi/cooperativeGrps/multi_grid_group_all_gpus.cpp @@ -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 +#include +#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(&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(&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(&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(&kernel_atomic[i]); + dev_params[i][1] = reinterpret_cast(&global_array); + dev_params[i][2] = reinterpret_cast(&kernel_buffer[i]); + dev_params[i][3] = reinterpret_cast(&loops); + md_params[i].func = reinterpret_cast(test_kernel); + 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(); + } +} diff --git a/hipamd/tests/src/runtimeApi/cooperativeGrps/simple_grid_group_barrier.cpp b/hipamd/tests/src/runtimeApi/cooperativeGrps/simple_grid_group_barrier.cpp new file mode 100644 index 0000000000..174a3c7f85 --- /dev/null +++ b/hipamd/tests/src/runtimeApi/cooperativeGrps/simple_grid_group_barrier.cpp @@ -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 +#include +#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(&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(&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(&kernel_atomic); + params[1] = reinterpret_cast(&kernel_buffer); + params[2] = reinterpret_cast(&loops); + HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast(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(); + } + } +} diff --git a/hipamd/tests/src/runtimeApi/cooperativeGrps/simple_multi_grid_group_barrier.cpp b/hipamd/tests/src/runtimeApi/cooperativeGrps/simple_multi_grid_group_barrier.cpp new file mode 100644 index 0000000000..736dac1a2c --- /dev/null +++ b/hipamd/tests/src/runtimeApi/cooperativeGrps/simple_multi_grid_group_barrier.cpp @@ -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 +#include +#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(&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(&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(&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(&kernel_atomic[i]); + dev_params[i][1] = reinterpret_cast(&global_array); + dev_params[i][2] = reinterpret_cast(&kernel_buffer[i]); + dev_params[i][3] = reinterpret_cast(&loops); + md_params[i].func = reinterpret_cast(test_kernel); + 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(); + } +}