From f3a4512a78fa29582a8393bf9641325b9b2b8c9a Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Mirza=20Halil=C4=8Devi=C4=87?= <109971222+mirza-halilcevic@users.noreply.github.com> Date: Thu, 28 Dec 2023 19:31:47 +0100 Subject: [PATCH] EXSWHTEC-115 - Implement tests for hipLaunchCooperativeKernel APIs #59 Change-Id: I5bda5ee3787a73aeeab5c25c05278e7aa2f8bfa2 [ROCm/hip-tests commit: 380ace735f1a373a90690a410de8164f6d962e5d] --- .../catch/hipTestMain/config/config_amd_linux | 5 - .../hipTestMain/config/config_amd_windows | 5 - .../unit/executionControl/CMakeLists.txt | 5 +- .../execution_control_common.cc | 12 +- .../execution_control_common.hh | 8 +- .../hipLaunchCooperativeKernel.cc | 188 ++++++++++++++++++ .../hipLaunchCooperativeKernelMultiDevice.cc | 159 +++++++++++++++ 7 files changed, 369 insertions(+), 13 deletions(-) create mode 100644 projects/hip-tests/catch/unit/executionControl/hipLaunchCooperativeKernel.cc create mode 100644 projects/hip-tests/catch/unit/executionControl/hipLaunchCooperativeKernelMultiDevice.cc diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux index 5cf8241d8d..f479b39e15 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux @@ -48,11 +48,6 @@ "Unit_hipFuncSetAttribute_Positive_PreferredSharedMemoryCarveout", "Unit_hipFuncSetAttribute_Positive_Parameters", "Unit_hipFuncSetAttribute_Negative_Parameters", - "NOTE: The following 4 tests are disabled due to defect - EXSWHTEC-240", - "Unit_hipFuncSetCacheConfig_Negative_Not_Supported", - "Unit_hipFuncSetSharedMemConfig_Negative_Not_Supported", - "Unit_hipFuncSetAttribute_Positive_MaxDynamicSharedMemorySize_Not_Supported", - "Unit_hipFuncSetAttribute_Positive_PreferredSharedMemoryCarveout_Not_Supported", "NOTE: The following test is disabled due to defect - EXSWHTEC-243", "Unit_hipExtLaunchKernel_Negative_Parameters", "NOTE: The following test is disabled due to defect - EXSWHTEC-244", diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows index 44a6dd0e0c..097f165983 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows @@ -115,11 +115,6 @@ "Unit_hipEventCreateWithFlags_DefaultFlg_NonCohHstMem", "Unit_hipEventCreateWithFlags_DisableSystemFence_CohHstMem", "Unit_hipEventCreateWithFlags_DefaultFlg_CohHstMem", - "NOTE: The following 4 tests are disabled due to defect - EXSWHTEC-240", - "Unit_hipFuncSetCacheConfig_Negative_Not_Supported", - "Unit_hipFuncSetSharedMemConfig_Negative_Not_Supported", - "Unit_hipFuncSetAttribute_Positive_MaxDynamicSharedMemorySize_Not_Supported", - "Unit_hipFuncSetAttribute_Positive_PreferredSharedMemoryCarveout_Not_Supported", "Unit_hipOccupancyMaxActiveBlocksPerMultiprocessor_Negative_Parameters", "Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_Negative_Parameters", "Unit_hipModuleOccupancyMaxPotentialBlockSizeWithFlags_Negative_Parameters", diff --git a/projects/hip-tests/catch/unit/executionControl/CMakeLists.txt b/projects/hip-tests/catch/unit/executionControl/CMakeLists.txt index 293da7cb99..a27f9dc4f1 100644 --- a/projects/hip-tests/catch/unit/executionControl/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/executionControl/CMakeLists.txt @@ -3,6 +3,9 @@ set(TEST_SRC hipFuncSetCacheConfig.cc hipFuncSetSharedMemConfig.cc hipFuncSetAttribute.cc + hipFuncGetAttributes.cc + hipLaunchCooperativeKernel.cc + hipLaunchCooperativeKernelMultiDevice.cc ) if(HIP_PLATFORM MATCHES "amd") @@ -15,4 +18,4 @@ endif() hip_add_exe_to_target(NAME ExecutionControlTest TEST_SRC ${TEST_SRC} TEST_TARGET_NAME build_tests - COMPILE_OPTIONS -std=c++17) \ No newline at end of file + COMPILE_OPTIONS -std=c++17) diff --git a/projects/hip-tests/catch/unit/executionControl/execution_control_common.cc b/projects/hip-tests/catch/unit/executionControl/execution_control_common.cc index f9a2cc1675..e4c04de1f3 100644 --- a/projects/hip-tests/catch/unit/executionControl/execution_control_common.cc +++ b/projects/hip-tests/catch/unit/executionControl/execution_control_common.cc @@ -23,5 +23,15 @@ THE SOFTWARE. #include "execution_control_common.hh" #include +#include -__global__ void kernel() {} \ No newline at end of file +__global__ void kernel() {} + +__global__ void kernel2() {} + +__global__ void kernel_42(int* val) { *val = 42; } + +__global__ void coop_kernel() { + cooperative_groups::grid_group grid = cooperative_groups::this_grid(); + grid.sync(); +} \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/executionControl/execution_control_common.hh b/projects/hip-tests/catch/unit/executionControl/execution_control_common.hh index e66fcf28f3..425c5309ed 100644 --- a/projects/hip-tests/catch/unit/executionControl/execution_control_common.hh +++ b/projects/hip-tests/catch/unit/executionControl/execution_control_common.hh @@ -22,4 +22,10 @@ THE SOFTWARE. #pragma once -__global__ void kernel(); \ No newline at end of file +__global__ void kernel(); + +__global__ void kernel2(); + +__global__ void kernel_42(int* val); + +__global__ void coop_kernel(); \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/executionControl/hipLaunchCooperativeKernel.cc b/projects/hip-tests/catch/unit/executionControl/hipLaunchCooperativeKernel.cc new file mode 100644 index 0000000000..eb7eb2293f --- /dev/null +++ b/projects/hip-tests/catch/unit/executionControl/hipLaunchCooperativeKernel.cc @@ -0,0 +1,188 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "execution_control_common.hh" + +#include +#include +#include +#include + +TEST_CASE("Unit_hipLaunchCooperativeKernel_Positive_Basic") { + if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) { + HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported"); + return; + } + + SECTION("Cooperative kernel with no arguments") { + HIP_CHECK(hipLaunchCooperativeKernel(reinterpret_cast(coop_kernel), dim3{2, 2, 1}, + dim3{1, 1, 1}, nullptr, 0, nullptr)); + HIP_CHECK(hipDeviceSynchronize()); + } + + SECTION("Kernel with arguments using kernelParams") { + LinearAllocGuard result_dev(LinearAllocs::hipMalloc, sizeof(int)); + HIP_CHECK(hipMemset(result_dev.ptr(), 0, sizeof(*result_dev.ptr()))); + + int* result_ptr = result_dev.ptr(); + void* kernel_args[1] = {&result_ptr}; + HIP_CHECK(hipLaunchCooperativeKernel(reinterpret_cast(kernel_42), dim3{1, 1, 1}, + dim3{1, 1, 1}, kernel_args, 0, nullptr)); + + int result = 0; + HIP_CHECK(hipMemcpy(&result, result_dev.ptr(), sizeof(result), hipMemcpyDefault)); + REQUIRE(result == 42); + } +} + +TEST_CASE("Unit_hipLaunchCooperativeKernel_Positive_Parameters") { + if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) { + HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported"); + return; + } + + SECTION("blockDim.x == maxBlockDimX") { + const unsigned int x = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimX); + HIP_CHECK(hipLaunchCooperativeKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, + dim3{x, 1, 1}, nullptr, 0, nullptr)); + } + + SECTION("blockDim.y == maxBlockDimY") { + const unsigned int y = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimY); + HIP_CHECK(hipLaunchCooperativeKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, + dim3{y, 1, 1}, nullptr, 0, nullptr)); + } + + SECTION("blockDim.z == maxBlockDimZ") { + const unsigned int z = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimZ); + HIP_CHECK(hipLaunchCooperativeKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, + dim3{z, 1, 1}, nullptr, 0, nullptr)); + } +} + +TEST_CASE("Unit_hipLaunchCooperativeKernel_Negative_Parameters") { + if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) { + HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported"); + return; + } + + SECTION("f == nullptr") { + HIP_CHECK_ERROR(hipLaunchCooperativeKernel(static_cast(nullptr), dim3{1, 1, 1}, + dim3{1, 1, 1}, nullptr, 0, nullptr), + hipErrorInvalidDeviceFunction); + } + + SECTION("gridDim.x == 0") { + HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast(kernel), dim3{0, 1, 1}, + dim3{1, 1, 1}, nullptr, 0, nullptr), + hipErrorInvalidConfiguration); + } + + SECTION("gridDim.y == 0") { + HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast(kernel), dim3{1, 0, 1}, + dim3{1, 1, 1}, nullptr, 0, nullptr), + hipErrorInvalidConfiguration); + } + + SECTION("gridDim.z == 0") { + HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast(kernel), dim3{1, 1, 0}, + dim3{1, 1, 1}, nullptr, 0, nullptr), + hipErrorInvalidConfiguration); + } + + SECTION("blockDim.x == 0") { + HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, + dim3{0, 1, 1}, nullptr, 0, nullptr), + hipErrorInvalidConfiguration); + } + + SECTION("blockDim.y == 0") { + HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, + dim3{1, 0, 1}, nullptr, 0, nullptr), + hipErrorInvalidConfiguration); + } + + SECTION("blockDim.z == 0") { + HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, + dim3{1, 1, 0}, nullptr, 0, nullptr), + hipErrorInvalidConfiguration); + } + + SECTION("blockDim.x > maxBlockDimX") { + const unsigned int x = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimX) + 1u; + HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, + dim3{x, 1, 1}, nullptr, 0, nullptr), + hipErrorInvalidConfiguration); + } + + SECTION("blockDim.y > maxBlockDimY") { + const unsigned int y = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimY) + 1u; + HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, + dim3{1, y, 1}, nullptr, 0, nullptr), + hipErrorInvalidConfiguration); + } + + SECTION("blockDim.z > maxBlockDimZ") { + const unsigned int z = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimZ) + 1u; + HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, + dim3{1, 1, z}, nullptr, 0, nullptr), + hipErrorInvalidConfiguration); + } + + SECTION("blockDim.x * blockDim.y * blockDim.z > maxThreadsPerBlock") { + const unsigned int max = GetDeviceAttribute(0, hipDeviceAttributeMaxThreadsPerBlock); + const unsigned int dim = std::ceil(std::cbrt(max)); + HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, + dim3{dim, dim, dim}, nullptr, 0, nullptr), + hipErrorInvalidConfiguration); + } + + SECTION( + "gridDim.x * gridDim.y * gridDim.z > maxActiveBlocksPerMultiprocessor * " + "multiProcessorCount") { + int max_blocks; + HIP_CHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&max_blocks, + reinterpret_cast(kernel), 1, 0)); + const unsigned int multiproc_count = + GetDeviceAttribute(0, hipDeviceAttributeMultiprocessorCount); + const unsigned int dim = std::ceil(std::cbrt(max_blocks * multiproc_count)); + HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast(kernel), dim3{dim, dim, dim}, + dim3{1, 1, 1}, nullptr, 0, nullptr), + hipErrorCooperativeLaunchTooLarge); + } + + SECTION("sharedMemBytes > maxSharedMemoryPerBlock") { + const unsigned int max = GetDeviceAttribute(0, hipDeviceAttributeMaxSharedMemoryPerBlock) + 1u; + HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, + dim3{1, 1, 1}, nullptr, max, nullptr), + hipErrorCooperativeLaunchTooLarge); + } + + SECTION("Invalid stream") { + hipStream_t stream = nullptr; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK_ERROR(hipLaunchCooperativeKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, + dim3{1, 1, 1}, nullptr, 0, stream), + hipErrorContextIsDestroyed); + } +} \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/executionControl/hipLaunchCooperativeKernelMultiDevice.cc b/projects/hip-tests/catch/unit/executionControl/hipLaunchCooperativeKernelMultiDevice.cc new file mode 100644 index 0000000000..c6b8503203 --- /dev/null +++ b/projects/hip-tests/catch/unit/executionControl/hipLaunchCooperativeKernelMultiDevice.cc @@ -0,0 +1,159 @@ +/* +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +#include "execution_control_common.hh" + +#include +#include +#include +#include + +TEST_CASE("Unit_hipLaunchCooperativeKernelMultiDevice_Positive_Basic") { + if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) { + HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported"); + return; + } + + const auto device_count = HipTest::getDeviceCount(); + + std::vector params_list(device_count); + + int device = 0; + for (auto& params : params_list) { + params.func = reinterpret_cast(coop_kernel); + params.gridDim = dim3{1, 1, 1}; + params.blockDim = dim3{1, 1, 1}; + params.args = nullptr; + params.sharedMem = 0; + HIP_CHECK(hipSetDevice(device++)); + HIP_CHECK(hipStreamCreate(¶ms.stream)); + } + + HIP_CHECK(hipLaunchCooperativeKernelMultiDevice(params_list.data(), device_count, 0u)); + + for (const auto params : params_list) { + HIP_CHECK(hipStreamSynchronize(params.stream)); + } + + for (const auto params : params_list) { + HIP_CHECK(hipStreamDestroy(params.stream)); + } +} + +TEST_CASE("Unit_hipLaunchCooperativeKernelMultiDevice_Negative_Parameters") { + if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) { + HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported"); + return; + } + + const auto device_count = HipTest::getDeviceCount(); + + std::vector params_list(device_count); + + int device = 0; + for (auto& params : params_list) { + params.func = reinterpret_cast(coop_kernel); + params.gridDim = dim3{1, 1, 1}; + params.blockDim = dim3{1, 1, 1}; + params.args = nullptr; + params.sharedMem = 0; + HIP_CHECK(hipSetDevice(device++)); + HIP_CHECK(hipStreamCreate(¶ms.stream)); + } + + SECTION("launchParamsList == nullptr") { + HIP_CHECK_ERROR(hipLaunchCooperativeKernelMultiDevice(nullptr, device_count, 0u), + hipErrorInvalidValue); + } + + SECTION("numDevices == 0") { + HIP_CHECK_ERROR(hipLaunchCooperativeKernelMultiDevice(params_list.data(), 0, 0u), + hipErrorInvalidValue); + } + + SECTION("numDevices > device count") { + HIP_CHECK_ERROR(hipLaunchCooperativeKernelMultiDevice(params_list.data(), device_count + 1, 0u), + hipErrorInvalidValue); + } + + SECTION("invalid flags") { + HIP_CHECK_ERROR(hipLaunchCooperativeKernelMultiDevice(params_list.data(), device_count, 999), + hipErrorInvalidValue); + } + + if (device_count > 1) { + SECTION("launchParamsList.func doesn't match across all devices") { + params_list[1].func = reinterpret_cast(kernel); + HIP_CHECK_ERROR(hipLaunchCooperativeKernelMultiDevice(params_list.data(), device_count, 0u), + hipErrorInvalidValue); + } + + SECTION("launchParamsList.gridDim doesn't match across all kernels") { + params_list[1].gridDim = dim3{2, 2, 2}; + HIP_CHECK_ERROR(hipLaunchCooperativeKernelMultiDevice(params_list.data(), device_count, 0u), + hipErrorInvalidValue); + } + + SECTION("launchParamsList.blockDim doesn't match across all kernels") { + params_list[1].blockDim = dim3{2, 2, 2}; + HIP_CHECK_ERROR(hipLaunchCooperativeKernelMultiDevice(params_list.data(), device_count, 0u), + hipErrorInvalidValue); + } + + SECTION("launchParamsList.sharedMem doesn't match across all kernels") { + params_list[1].sharedMem = 1024; + HIP_CHECK_ERROR(hipLaunchCooperativeKernelMultiDevice(params_list.data(), device_count, 0u), + hipErrorInvalidValue); + } + } + + for (const auto params : params_list) { + HIP_CHECK(hipStreamDestroy(params.stream)); + } +} + +TEST_CASE("Unit_hipLaunchCooperativeKernelMultiDevice_Negative_MultiKernelSameDevice") { + if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) { + HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported"); + return; + } + + HIP_CHECK(hipSetDevice(0)); + + std::vector params_list(2); + + for (auto& params : params_list) { + params.func = reinterpret_cast(coop_kernel); + params.gridDim = dim3{1, 1, 1}; + params.blockDim = dim3{1, 1, 1}; + params.args = nullptr; + params.sharedMem = 0; + HIP_CHECK(hipStreamCreate(¶ms.stream)); + } + + HIP_CHECK_ERROR(hipLaunchCooperativeKernelMultiDevice(params_list.data(), 2, 0u), + hipErrorInvalidValue); + + for (const auto params : params_list) { + HIP_CHECK(hipStreamDestroy(params.stream)); + } +} \ No newline at end of file