diff --git a/INSTALL.md b/INSTALL.md index 7327a7611f..c939c34c50 100644 --- a/INSTALL.md +++ b/INSTALL.md @@ -35,7 +35,7 @@ HIP-Clang is the compiler for compiling HIP programs on AMD platform. HIP-Clang can be built manually: ``` -git clone -b rocm-4.3.x https://github.com/RadeonOpenCompute/llvm-project.git +git clone -b amd-stg-open https://github.com/RadeonOpenCompute/llvm-project.git cd llvm-project mkdir -p build && cd build cmake -DCMAKE_INSTALL_PREFIX=/opt/rocm/llvm -DCMAKE_BUILD_TYPE=Release -DLLVM_ENABLE_ASSERTIONS=1 -DLLVM_TARGETS_TO_BUILD="AMDGPU;X86" -DLLVM_ENABLE_PROJECTS="clang;lld;compiler-rt" ../llvm @@ -45,7 +45,7 @@ sudo make install Rocm device library can be manually built as following, ``` export PATH=/opt/rocm/llvm/bin:$PATH -git clone -b rocm-4.3.x https://github.com/RadeonOpenCompute/ROCm-Device-Libs.git +git clone -b amd-stg-open https://github.com/RadeonOpenCompute/ROCm-Device-Libs.git cd ROCm-Device-Libs mkdir -p build && cd build CC=clang CXX=clang++ cmake -DLLVM_DIR=/opt/rocm/llvm -DCMAKE_BUILD_TYPE=Release -DLLVM_ENABLE_WERROR=1 -DLLVM_ENABLE_ASSERTIONS=1 -DCMAKE_INSTALL_PREFIX=/opt/rocm .. @@ -72,10 +72,10 @@ apt-get install hip-runtime-nvidia hip-devel ## Get HIP source code ``` -git clone -b rocm-4.4.x https://github.com/ROCm-Developer-Tools/hipamd.git -git clone -b rocm-4.4.x https://github.com/ROCm-Developer-Tools/hip.git -git clone -b rocm-4.4.x https://github.com/ROCm-Developer-Tools/ROCclr.git -git clone -b rocm-4.4.x https://github.com/RadeonOpenCompute/ROCm-OpenCL-Runtime.git +git clone -b develop https://github.com/ROCm-Developer-Tools/hipamd.git +git clone -b develop https://github.com/ROCm-Developer-Tools/hip.git +git clone -b develop https://github.com/ROCm-Developer-Tools/ROCclr.git +git clone -b develop https://github.com/RadeonOpenCompute/ROCm-OpenCL-Runtime.git ``` ## Set the environment variables diff --git a/README.md b/README.md index ff292dee17..d4e16e70c7 100644 --- a/README.md +++ b/README.md @@ -22,8 +22,9 @@ The information contained herein is for informational purposes only, and is subj The HIP repository maintains several branches. The branches that are of importance are: -* Main branch: This is the stable branch. It is up to date with the latest release branch, for example, if the latest HIP release is rocm-3.9, main branch will be the repository based on this release. -* Release branches. These are branches corresponding to each ROCM release, listed with release tags, such as rocm-3.8, rocm-3.9, etc. +* develop branch: This is the default branch, on which the new features are still under development and visible. While this maybe of interest to many, it should be noted that this branch and the features under development might not be stable. +* Main branch: This is the stable branch. It is up to date with the latest release branch, for example, if the latest HIP release is rocm-4.3, main branch will be the repository based on this release. +* Release branches. These are branches corresponding to each ROCM release, listed with release tags, such as rocm-4.2, rocm-4.3, etc. ## Release tagging: diff --git a/docs/markdown/hip_debugging.md b/docs/markdown/hip_debugging.md old mode 100755 new mode 100644 diff --git a/tests/performance/module/hipPerfModuleLoad.cpp b/tests/performance/module/hipPerfModuleLoad.cpp old mode 100755 new mode 100644 diff --git a/tests/src/hiprtc/saxpy.cpp b/tests/src/hiprtc/saxpy.cpp old mode 100755 new mode 100644 diff --git a/tests/src/hiprtc/saxpy_pch.cpp b/tests/src/hiprtc/saxpy_pch.cpp old mode 100755 new mode 100644 diff --git a/tests/src/runtimeApi/cooperativeGrps/cooperative_streams.cpp b/tests/src/runtimeApi/cooperativeGrps/cooperative_streams.cpp index 8a25f6f5d7..2300528397 100644 --- a/tests/src/runtimeApi/cooperativeGrps/cooperative_streams.cpp +++ b/tests/src/runtimeApi/cooperativeGrps/cooperative_streams.cpp @@ -50,8 +50,10 @@ regular kernels can properly run in parallel with cooperative kernels. */ /* HIT_START - * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 EXCLUDE_HIP_PLATFORM nvidia - * TEST: %t + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 + * TEST_NAMED: %t cooperative_streams_least_capacity --tests 0x0 + * TEST_NAMED: %t cooperative_streams_half_capacity --tests 0x1 + * TEST_NAMED: %t cooperative_streams_full_capacity --tests 0x2 * HIT_END */ @@ -67,7 +69,7 @@ static inline void hipCheckAndFail(hipError_t errval, std::cerr << "hip error: " << hipGetErrorString(errval); std::cerr << std::endl; std::cerr << "Location: " << file << ":" << line << std::endl; - failed(""); + failed(" "); } if (last_err != errval) { std::cerr << "Error: the return value of a function was not the same "; @@ -77,35 +79,13 @@ static inline void hipCheckAndFail(hipError_t errval, std::cerr << " (" << errval << ")" << std::endl; std::cerr << "hipGetLastError() returned: " << hipGetErrorString(last_err); std::cerr << " (" << last_err << ")" << std::endl; - failed(""); + 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) { +__global__ void test_kernel(uint32_t loops, unsigned long long *array, long long totalTicks) { unsigned int rank = blockIdx.x * blockDim.x + threadIdx.x; for (int i = 0; i < loops; i++) { @@ -119,15 +99,137 @@ __global__ void test_kernel(uint32_t loops, unsigned long long *array) { // If it rolls over, we don't know how much to add to catch up. // So just ignore those slipped cycles. last_clock = cur_clock; - } while(time_diff < 1000000); + } while(time_diff < totalTicks); array[rank] += clock64(); } } +template +bool verifyLeastCapacity(T& single_kernel_time, T& double_kernel_time, T& triple_kernel_time) +{ +#ifdef __HIP_PLATFORM_AMD__ + // hipLaunchCooperativeKernel() follows serialization policy on AMD devices + // Test that the two cooperative kernels took roughly twice as long as the one + if (double_kernel_time < 1.8 * single_kernel_time || + double_kernel_time > 2.2 * 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; + return false; + } +#else + // hipLaunchCooperativeKernel() doesn't follow serialization policy on NV devices + // Test that the two cooperative kernels took roughly as long as the one + if (double_kernel_time < 0.8 * single_kernel_time || + double_kernel_time > 1.2 * single_kernel_time ) { + std::cerr << "ERROR!" << std::endl; + std::cerr << "Two cooperative kernels launched at the same "; + std::cerr << "time did not take roughly as long as a single "; + std::cerr << "cooperative kernel." << std::endl; + return false; + } +#endif + + // Test that the three kernels together took roughly as long as the 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; + return false; + } + return true; +} + +template +bool verifyHalfCapacity(T& single_kernel_time, T& double_kernel_time, T& triple_kernel_time) +{ + // Test that the two cooperative kernels took roughly twice as long as the one + if (double_kernel_time < 1.8 * single_kernel_time || + double_kernel_time > 2.2 * 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; + return false; + } + + // Test that the three kernels together took roughly as long as the 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; + return false; + } + return true; +} + +template +bool verifyFullCapacity(T& single_kernel_time, T& double_kernel_time, T& triple_kernel_time) +{ + // Test that the two cooperative kernels took roughly twice as long as the one + if (double_kernel_time < 1.8 * single_kernel_time || + double_kernel_time > 2.2 * 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; + return false; + } + + // Test that the three kernels together took roughly 1.6 times as long as the two + // cooperative kernels. If the first 2 kernels run very fast, the third + // won't share much time with the second kernel. + if (triple_kernel_time > 1.7 * 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 70% longer than the two cooperative kernels "; + std::cerr << "alone." << std::endl; + return false; + } + return true; +} + +template +bool verify(int tests, T &single_kernel_time, T &double_kernel_time, + T &triple_kernel_time) { + switch (tests) { + case 0: + return verifyLeastCapacity(single_kernel_time, double_kernel_time, + triple_kernel_time); + case 1: + return verifyHalfCapacity(single_kernel_time, double_kernel_time, + triple_kernel_time); + case 2: + return verifyFullCapacity(single_kernel_time, double_kernel_time, + triple_kernel_time); + default: + return false; + } +} + int main(int argc, char** argv) { - hipError_t err; + p_tests = 1; // Default for half capacity + HipTest::parseStandardArguments(argc, argv, true); + + if (p_tests < 0 || p_tests > 2) { + printf("--tests 0: test least capacity\n"); + printf(" 1: test half capacity\n"); + printf(" 2: test full capacity\n"); + failed("Wrong p_tests %d\n", p_tests); + } + hipError_t err = hipSuccess; /*************************************************************************/ - int device_num = 0, loops = 1000, FailFlag = 0; + int device_num = 0, loops = 1000; + bool FailFlag = false; /* 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 @@ -136,36 +238,62 @@ int main(int argc, char** argv) { 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)); + + /* Test whether target device supports cooperative groups ****************/ + if (device_properties.cooperativeLaunch == 0) { + std::cout << "Cooperative group support not available in device " << dev << std::endl; + continue; + } + + /* We will launch enough waves to fill up all of the GPU *****************/ 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; + long long totalTicks = device_properties.clockRate ; + int max_blocks_per_sm = 0; // Calculate the device occupancy to know how many blocks can be run. HIPCHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&max_blocks_per_sm, test_kernel, warp_size, 0)); + int max_active_blocks = max_blocks_per_sm * num_sms; + int coop_blocks = 0; + int reg_blocks = 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(""); + switch (p_tests) { + case 0: + // 1 block + coop_blocks = 1; + reg_blocks = 1; + break; + case 1: + // Half capacity + // To make sure the second kernel launched by hipLaunchCooperativeKernel + // is invoked after the first kernel finished + coop_blocks = max_active_blocks / 2 + 1; + // To make sure the third kernel launched by hipLaunchKernelGGL is invoked + // concurrently with the second kernel + reg_blocks = max_active_blocks - coop_blocks; + break; + case 2: + // Full capacity + coop_blocks = max_active_blocks; + reg_blocks = max_active_blocks; + break; + default: + failed("wrong p_tests %d", p_tests); } + std::cout << "p_tests: " << p_tests << std::endl; + std::cout << "Device: " << dev << std::endl; + std::cout << "Device name: " << device_properties.name << std::endl; + std::cout << "clockRate: " << device_properties.clockRate << " khz" <(&loops); coop_params[i][1] = reinterpret_cast(&dev_array[i]); + coop_params[i][2] = reinterpret_cast(&totalTicks); } + // Verify over capacity + HIPCHECK_API(hipLaunchCooperativeKernel(reinterpret_cast(test_kernel), + max_active_blocks + 1, warp_size, + coop_params[0], 0, streams[0]), hipErrorCooperativeLaunchTooLarge); + + std::cout << "Launching an initial single cooperative kernel..." << std::endl; + // We need exclude the the initial launching as it will need time to load code obj. + auto single_start0 = std::chrono::system_clock::now(); + HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast(test_kernel), + max_active_blocks, warp_size, + coop_params[0], 0, streams[0])); + HIPCHECK(hipDeviceSynchronize()); + auto single_end0 = std::chrono::system_clock::now(); + 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_blocks, warp_size, coop_params[0], 0, streams[0])); HIPCHECK(hipDeviceSynchronize()); @@ -205,10 +348,10 @@ int main(int argc, char** argv) { auto double_start = std::chrono::system_clock::now(); HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast(test_kernel), - desired_blocks, warp_size, + coop_blocks, warp_size, coop_params[0], 0, streams[0])); HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast(test_kernel), - desired_blocks, warp_size, + coop_blocks, warp_size, coop_params[1], 0, streams[1])); HIPCHECK(hipDeviceSynchronize()); @@ -218,18 +361,18 @@ int main(int argc, char** argv) { auto triple_start = std::chrono::system_clock::now(); HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast(test_kernel), - desired_blocks, warp_size, + coop_blocks, warp_size, coop_params[0], 0, streams[0])); HIPCHECK(hipLaunchCooperativeKernel(reinterpret_cast(test_kernel), - desired_blocks, warp_size, + coop_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); + hipLaunchKernelGGL(test_kernel, dim3(reg_blocks), dim3(warp_size), + 0, streams[2], loops, dev_array[2], totalTicks); HIPCHECK(hipDeviceSynchronize()); auto triple_end = std::chrono::system_clock::now(); + std::chrono::duration single_kernel_time0 = + (single_end0 - single_start0); std::chrono::duration single_kernel_time = (single_end - single_start); std::chrono::duration double_kernel_time = @@ -237,6 +380,10 @@ int main(int argc, char** argv) { std::chrono::duration triple_kernel_time = (triple_end - triple_start); + std::cout << "Initial single kernel took:" << std::endl; + std::cout << " " << single_kernel_time0.count(); + std::cout << " seconds" << std::endl; + std::cout << std::endl; std::cout << "A single kernel took:" << std::endl; std::cout << " " << single_kernel_time.count(); std::cout << " seconds" << std::endl; @@ -252,41 +399,17 @@ int main(int argc, char** argv) { 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; - } + FailFlag = !verify(p_tests, single_kernel_time, + double_kernel_time, triple_kernel_time); + 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])); + if (FailFlag) { + failed("Failed to exit "); } - failed(""); } passed(); } diff --git a/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupType.cpp b/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupType.cpp index f3af976d35..0b170de1f7 100755 --- a/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupType.cpp +++ b/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupType.cpp @@ -22,8 +22,8 @@ THE SOFTWARE. /* HIT_START - * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 -rdc=true -gencode arch=compute_60,code=sm_60 - * TEST: %t EXCLUDE_HIP_PLATFORM nvidia + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 -rdc=true -gencode arch=compute_60,code=sm_60 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_80,code=sm_80 + * TEST: %t * HIT_END */ @@ -37,6 +37,9 @@ THE SOFTWARE. #define ASSERT_LE(lhs, rhs) assert(lhs <= rhs) #define ASSERT_GE(lhs, rhs) assert(lhs >= rhs) +constexpr int MaxGPUs = 8; +int nGpu = 0; + using namespace cooperative_groups; static __global__ @@ -82,7 +85,7 @@ void kernel_cg_multi_grid_group_type(int* numGridsTestD, // multi-grid level sync mg.sync(); // grid (gpu) 0 does final reduction across all grids (gpus) - if (mg.grid_rank() == 0) { + if (mg.grid_rank() == 0 && blockIdx.x == 0 && threadIdx.x == 0) { syncResultD[0] = 0; for (int i = 1; i <= mg.num_grids(); ++i) { syncResultD[0] += syncResultD[i]; @@ -92,29 +95,12 @@ void kernel_cg_multi_grid_group_type(int* numGridsTestD, static void test_cg_multi_grid_group_type(int blockSize) { - // Get device count - constexpr int MaxGPUs = 8; - int nGpu = 0; - ASSERT_EQUAL(hipGetDeviceCount(&nGpu), hipSuccess); - - // Check if device suppurts multi gpu cooperative group support - hipDeviceProp_t deviceProp[MaxGPUs]; - for (int i = 0; i < nGpu; i++) { - ASSERT_EQUAL(hipSetDevice(i), hipSuccess); - hipGetDeviceProperties(&deviceProp[i], 0); - if (!deviceProp[i].cooperativeMultiDeviceLaunch) { - printf("Device doesn't support multi gpu cooperative launch"); - return; - } - hipDeviceSynchronize(); - } - // Create a stream each device hipStream_t stream[MaxGPUs]; for (int i = 0; i < nGpu; i++) { ASSERT_EQUAL(hipSetDevice(i), hipSuccess); + hipDeviceSynchronize(); // Make sure work is done on this device ASSERT_EQUAL(hipStreamCreate(&stream[i]), hipSuccess); - hipDeviceSynchronize(); } // Allocate host and device memory @@ -146,8 +132,6 @@ static void test_cg_multi_grid_group_type(int blockSize) hipHostMalloc(&syncResultD, sizeof(int) * (nGpu + 1), hipHostMallocCoherent), hipSuccess); } - - hipDeviceSynchronize(); } // Launch Kernel @@ -171,15 +155,12 @@ static void test_cg_multi_grid_group_type(int blockSize) launchParamsList[i].sharedMem = 0; launchParamsList[i].stream = stream[i]; launchParamsList[i].args = &args[i * NumKernelArgs]; - - hipDeviceSynchronize(); } - hipLaunchCooperativeKernelMultiDevice(launchParamsList, nGpu, 0); + HIPCHECK(hipLaunchCooperativeKernelMultiDevice(launchParamsList, nGpu, 0)); // Copy result from device to host for (int i = 0; i < nGpu; i++) { ASSERT_EQUAL(hipSetDevice(i), hipSuccess); - ASSERT_EQUAL(hipMemcpy(numGridsTestH[i], numGridsTestD[i], nBytes, hipMemcpyDeviceToHost), hipSuccess); ASSERT_EQUAL(hipMemcpy(gridRankTestH[i], gridRankTestD[i], nBytes, hipMemcpyDeviceToHost), @@ -191,7 +172,6 @@ static void test_cg_multi_grid_group_type(int blockSize) ASSERT_EQUAL(hipMemcpy(isValidTestH[i], isValidTestD[i], nBytes, hipMemcpyDeviceToHost), hipSuccess); - hipDeviceSynchronize(); } // Validate results @@ -231,31 +211,34 @@ static void test_cg_multi_grid_group_type(int blockSize) ASSERT_EQUAL(hipFree(isValidTestD[i]), hipSuccess); ASSERT_EQUAL(hipFree(syncTestD[i]), hipSuccess); - if (i == 0) - ASSERT_EQUAL(hipFree(syncResultD), hipSuccess); - + if (i == 0) { + ASSERT_EQUAL(hipHostFree(syncResultD), hipSuccess); + } ASSERT_EQUAL(hipHostFree(numGridsTestH[i]), hipSuccess); ASSERT_EQUAL(hipHostFree(gridRankTestH[i]), hipSuccess); ASSERT_EQUAL(hipHostFree(sizeTestH[i]), hipSuccess); ASSERT_EQUAL(hipHostFree(thdRankTestH[i]), hipSuccess); ASSERT_EQUAL(hipHostFree(isValidTestH[i]), hipSuccess); - - hipDeviceSynchronize(); } } int main() { // Set `maxThreadsPerBlock` by taking minimum among all available devices - int nGpu = 0; ASSERT_EQUAL(hipGetDeviceCount(&nGpu), hipSuccess); + if (nGpu > MaxGPUs) { + nGpu = MaxGPUs; + } int maxThreadsPerBlock = INT_MAX; for (int i = 0; i < nGpu; i++) { hipDeviceProp_t deviceProperties; ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, i), hipSuccess); + if (!deviceProperties.cooperativeMultiDeviceLaunch) { + printf("Device doesn't support cooperative launch!"); + passed(); + } int curDeviceMaxThreadsPerBlock = deviceProperties.maxThreadsPerBlock; maxThreadsPerBlock = min(maxThreadsPerBlock, curDeviceMaxThreadsPerBlock); - hipDeviceSynchronize(); } // Test block sizes which are powers of 2 diff --git a/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupTypeViaBaseType.cpp b/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupTypeViaBaseType.cpp index 4431a8a18f..0c5dc7b104 100644 --- a/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupTypeViaBaseType.cpp +++ b/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupTypeViaBaseType.cpp @@ -22,8 +22,8 @@ THE SOFTWARE. /* HIT_START - * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 -rdc=true -gencode arch=compute_60,code=sm_60 - * TEST: %t EXCLUDE_HIP_PLATFORM nvidia + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 -D_CG_ABI_EXPERIMENTAL -rdc=true -gencode arch=compute_60,code=sm_60 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_80,code=sm_80 + * TEST: %t * HIT_END */ @@ -37,6 +37,9 @@ THE SOFTWARE. #define ASSERT_LE(lhs, rhs) assert(lhs <= rhs) #define ASSERT_GE(lhs, rhs) assert(lhs >= rhs) +constexpr int MaxGPUs = 8; +int nGpu = 0; + using namespace cooperative_groups; static __global__ @@ -47,7 +50,8 @@ void kernel_cg_multi_grid_group_type_via_base_type(int *sizeTestD, int *syncTestD, int *syncResultD) { - multi_grid_group tg = this_multi_grid(); + thread_group tg = this_multi_grid(); // This can work if _CG_ABI_EXPERIMENTAL defined on Cuda + int gIdx = (blockIdx.x * blockDim.x) + threadIdx.x; // Test size @@ -58,8 +62,12 @@ void kernel_cg_multi_grid_group_type_via_base_type(int *sizeTestD, thdRankTestD[gIdx] = tg.thread_rank(); // Test is_valid +#ifdef __HIP_PLATFORM_AMD__ isValidTestD[gIdx] = tg.is_valid(); - +#else + // Cuda has no thread_group.is_valid() + isValidTestD[gIdx] = true; +#endif // Test sync // // Eech thread assign 1 to their respective location @@ -76,7 +84,7 @@ void kernel_cg_multi_grid_group_type_via_base_type(int *sizeTestD, // multi-grid level sync tg.sync(); // grid (gpu) 0 does final reduction across all grids (gpus) - if (this_multi_grid().grid_rank() == 0) { + if (this_multi_grid().grid_rank() == 0 && blockIdx.x == 0 && threadIdx.x == 0) { syncResultD[0] = 0; for (int i = 1; i <= this_multi_grid().num_grids(); ++i) { syncResultD[0] += syncResultD[i]; @@ -86,29 +94,12 @@ void kernel_cg_multi_grid_group_type_via_base_type(int *sizeTestD, static void test_cg_multi_grid_group_type_via_base_type(int blockSize) { - // Get device count - constexpr int MaxGPUs = 8; - int nGpu = 0; - ASSERT_EQUAL(hipGetDeviceCount(&nGpu), hipSuccess); - - // Check if device suppurts multi gpu cooperative group support - hipDeviceProp_t deviceProp[MaxGPUs]; - for (int i = 0; i < nGpu; i++) { - ASSERT_EQUAL(hipSetDevice(i), hipSuccess); - hipGetDeviceProperties(&deviceProp[i], 0); - if (!deviceProp[i].cooperativeMultiDeviceLaunch) { - printf("Device doesn't support multi gpu cooperative launch"); - return; - } - hipDeviceSynchronize(); - } - // Create a stream each device hipStream_t stream[MaxGPUs]; for (int i = 0; i < nGpu; i++) { ASSERT_EQUAL(hipSetDevice(i), hipSuccess); + hipDeviceSynchronize(); // Make sure work is done on this device ASSERT_EQUAL(hipStreamCreate(&stream[i]), hipSuccess); - hipDeviceSynchronize(); } // Allocate host and device memory @@ -137,8 +128,6 @@ static void test_cg_multi_grid_group_type_via_base_type(int blockSize) hipHostMalloc(&syncResultD, sizeof(int) * (nGpu + 1), hipHostMallocCoherent), hipSuccess); } - - hipDeviceSynchronize(); } // Launch Kernel @@ -161,10 +150,8 @@ static void test_cg_multi_grid_group_type_via_base_type(int blockSize) launchParamsList[i].sharedMem = 0; launchParamsList[i].stream = stream[i]; launchParamsList[i].args = &args[i * NumKernelArgs]; - - hipDeviceSynchronize(); } - hipLaunchCooperativeKernelMultiDevice(launchParamsList, nGpu, 0); + HIPCHECK(hipLaunchCooperativeKernelMultiDevice(launchParamsList, nGpu, 0)); // Copy result from device to host for (int i = 0; i < nGpu; i++) { @@ -178,8 +165,6 @@ static void test_cg_multi_grid_group_type_via_base_type(int blockSize) hipSuccess); ASSERT_EQUAL(hipMemcpy(isValidTestH[i], isValidTestD[i], nBytes, hipMemcpyDeviceToHost), hipSuccess); - - hipDeviceSynchronize(); } // Validate results @@ -218,29 +203,32 @@ static void test_cg_multi_grid_group_type_via_base_type(int blockSize) ASSERT_EQUAL(hipFree(syncTestD[i]), hipSuccess); if (i == 0) - ASSERT_EQUAL(hipFree(syncResultD), hipSuccess); + ASSERT_EQUAL(hipHostFree(syncResultD), hipSuccess); ASSERT_EQUAL(hipHostFree(sizeTestH[i]), hipSuccess); ASSERT_EQUAL(hipHostFree(gridRankTestH[i]), hipSuccess); ASSERT_EQUAL(hipHostFree(thdRankTestH[i]), hipSuccess); ASSERT_EQUAL(hipHostFree(isValidTestH[i]), hipSuccess); - - hipDeviceSynchronize(); } } int main() { // Set `maxThreadsPerBlock` by taking minimum among all available devices - int nGpu = 0; ASSERT_EQUAL(hipGetDeviceCount(&nGpu), hipSuccess); + if (nGpu > MaxGPUs) { + nGpu = MaxGPUs; + } int maxThreadsPerBlock = INT_MAX; for (int i = 0; i < nGpu; i++) { hipDeviceProp_t deviceProperties; ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, i), hipSuccess); + if (!deviceProperties.cooperativeMultiDeviceLaunch) { + printf("Device doesn't support cooperative launch!"); + passed(); + } int curDeviceMaxThreadsPerBlock = deviceProperties.maxThreadsPerBlock; maxThreadsPerBlock = min(maxThreadsPerBlock, curDeviceMaxThreadsPerBlock); - hipDeviceSynchronize(); } // Test block sizes which are powers of 2 diff --git a/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupTypeViaPublicApi.cpp b/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupTypeViaPublicApi.cpp index 60014e4959..a00a40345c 100644 --- a/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupTypeViaPublicApi.cpp +++ b/tests/src/runtimeApi/cooperativeGrps/hipCGMultiGridGroupTypeViaPublicApi.cpp @@ -22,8 +22,8 @@ THE SOFTWARE. /* HIT_START - * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 -rdc=true -gencode arch=compute_60,code=sm_60 - * TEST: %t EXCLUDE_HIP_PLATFORM all + * BUILD: %t %s ../../test_common.cpp NVCC_OPTIONS --std=c++11 -rdc=true -gencode arch=compute_60,code=sm_60 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_80,code=sm_80 + * TEST: %t * HIT_END */ @@ -37,6 +37,9 @@ THE SOFTWARE. #define ASSERT_LE(lhs, rhs) assert(lhs <= rhs) #define ASSERT_GE(lhs, rhs) assert(lhs >= rhs) +constexpr int MaxGPUs = 8; +int nGpu = 0; + using namespace cooperative_groups; static __global__ @@ -76,7 +79,7 @@ void kernel_cg_multi_grid_group_type_via_public_api(int *sizeTestD, // multi-grid level sync via public api sync(mg); // grid (gpu) 0 does final reduction across all grids (gpus) - if (this_multi_grid().grid_rank() == 0) { + if (this_multi_grid().grid_rank() == 0 && blockIdx.x == 0 && threadIdx.x == 0) { syncResultD[0] = 0; for (int i = 1; i <= this_multi_grid().num_grids(); ++i) { syncResultD[0] += syncResultD[i]; @@ -86,29 +89,12 @@ void kernel_cg_multi_grid_group_type_via_public_api(int *sizeTestD, static void test_cg_multi_grid_group_type_via_public_api(int blockSize) { - // Get device count - constexpr int MaxGPUs = 8; - int nGpu = 0; - ASSERT_EQUAL(hipGetDeviceCount(&nGpu), hipSuccess); - - // Check if device suppurts multi gpu cooperative group support - hipDeviceProp_t deviceProp[MaxGPUs]; - for (int i = 0; i < nGpu; i++) { - ASSERT_EQUAL(hipSetDevice(i), hipSuccess); - hipGetDeviceProperties(&deviceProp[i], 0); - if (!deviceProp[i].cooperativeMultiDeviceLaunch) { - printf("Device doesn't support multi gpu cooperative launch"); - return; - } - hipDeviceSynchronize(); - } - // Create a stream each device hipStream_t stream[MaxGPUs]; for (int i = 0; i < nGpu; i++) { ASSERT_EQUAL(hipSetDevice(i), hipSuccess); + hipDeviceSynchronize(); // Make sure work is done on this device ASSERT_EQUAL(hipStreamCreate(&stream[i]), hipSuccess); - hipDeviceSynchronize(); } // Allocate host and device memory @@ -137,8 +123,6 @@ static void test_cg_multi_grid_group_type_via_public_api(int blockSize) hipHostMalloc(&syncResultD, sizeof(int) * (nGpu + 1), hipHostMallocCoherent), hipSuccess); } - - hipDeviceSynchronize(); } // Launch Kernel @@ -161,10 +145,8 @@ static void test_cg_multi_grid_group_type_via_public_api(int blockSize) launchParamsList[i].sharedMem = 0; launchParamsList[i].stream = stream[i]; launchParamsList[i].args = &args[i * NumKernelArgs]; - - hipDeviceSynchronize(); } - hipLaunchCooperativeKernelMultiDevice(launchParamsList, nGpu, 0); + HIPCHECK(hipLaunchCooperativeKernelMultiDevice(launchParamsList, nGpu, 0)); // Copy result from device to host for (int i = 0; i < nGpu; i++) { @@ -178,8 +160,6 @@ static void test_cg_multi_grid_group_type_via_public_api(int blockSize) hipSuccess); ASSERT_EQUAL(hipMemcpy(isValidTestH[i], isValidTestD[i], nBytes, hipMemcpyDeviceToHost), hipSuccess); - - hipDeviceSynchronize(); } // Validate results @@ -218,29 +198,32 @@ static void test_cg_multi_grid_group_type_via_public_api(int blockSize) ASSERT_EQUAL(hipFree(syncTestD[i]), hipSuccess); if (i == 0) - ASSERT_EQUAL(hipFree(syncResultD), hipSuccess); + ASSERT_EQUAL(hipHostFree(syncResultD), hipSuccess); ASSERT_EQUAL(hipHostFree(sizeTestH[i]), hipSuccess); ASSERT_EQUAL(hipHostFree(gridRankTestH[i]), hipSuccess); ASSERT_EQUAL(hipHostFree(thdRankTestH[i]), hipSuccess); ASSERT_EQUAL(hipHostFree(isValidTestH[i]), hipSuccess); - - hipDeviceSynchronize(); } } int main() { // Set `maxThreadsPerBlock` by taking minimum among all available devices - int nGpu = 0; ASSERT_EQUAL(hipGetDeviceCount(&nGpu), hipSuccess); + if (nGpu > MaxGPUs) { + nGpu = MaxGPUs; + } int maxThreadsPerBlock = INT_MAX; for (int i = 0; i < nGpu; i++) { hipDeviceProp_t deviceProperties; ASSERT_EQUAL(hipGetDeviceProperties(&deviceProperties, i), hipSuccess); + if (!deviceProperties.cooperativeMultiDeviceLaunch) { + printf("Device doesn't support cooperative launch!"); + passed(); + } int curDeviceMaxThreadsPerBlock = deviceProperties.maxThreadsPerBlock; maxThreadsPerBlock = min(maxThreadsPerBlock, curDeviceMaxThreadsPerBlock); - hipDeviceSynchronize(); } // Test block sizes which are powers of 2 diff --git a/tests/src/test_common.h b/tests/src/test_common.h old mode 100755 new mode 100644 diff --git a/tests/unit/test_common.h b/tests/unit/test_common.h old mode 100755 new mode 100644