From 3f65d7a4104d7a181a79d0f37a9c70abf0421fc8 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:32:13 +0100 Subject: [PATCH] EXSWHTEC-114 - Implement tests for hipExtLaunchKernel APIs #60 Change-Id: I9152c31c2c07ddbfb48865cd68a42557e763be28 --- catch/hipTestMain/config/config_amd_linux | 4 + catch/hipTestMain/config/config_amd_windows | 4 + catch/unit/executionControl/CMakeLists.txt | 7 + .../executionControl/hipExtLaunchKernel.cc | 176 ++++++++++++++++++ .../hipExtLaunchMultiKernelMultiDevice.cc | 144 ++++++++++++++ 5 files changed, 335 insertions(+) create mode 100644 catch/unit/executionControl/hipExtLaunchKernel.cc create mode 100644 catch/unit/executionControl/hipExtLaunchMultiKernelMultiDevice.cc diff --git a/catch/hipTestMain/config/config_amd_linux b/catch/hipTestMain/config/config_amd_linux index 8d89762cf3..5cf8241d8d 100644 --- a/catch/hipTestMain/config/config_amd_linux +++ b/catch/hipTestMain/config/config_amd_linux @@ -53,6 +53,10 @@ "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", + "Unit_hipExtLaunchMultiKernelMultiDevice_Negative_Parameters", "Unit_hipOccupancyMaxActiveBlocksPerMultiprocessor_Negative_Parameters", "Unit_hipModuleOccupancyMaxActiveBlocksPerMultiprocessorWithFlags_Negative_Parameters", "Unit_hipModuleOccupancyMaxPotentialBlockSizeWithFlags_Negative_Parameters", diff --git a/catch/hipTestMain/config/config_amd_windows b/catch/hipTestMain/config/config_amd_windows index 623b25dfba..44a6dd0e0c 100644 --- a/catch/hipTestMain/config/config_amd_windows +++ b/catch/hipTestMain/config/config_amd_windows @@ -218,6 +218,10 @@ "Unit_hipVectorTypes_test_on_device", "=== Patch which removes the typetraits implementation from std namespace in hiprtc is reverted ===", "Unit_hiprtc_stdheaders", + "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", + "Unit_hipExtLaunchMultiKernelMultiDevice_Negative_Parameters", "Unit_hipMemAddressReserve_AlignmentTest", "Unit_hipGraphAddMemcpyNode_Negative_Parameters", "Unit_hipMemCreate_ChkWithKerLaunch", diff --git a/catch/unit/executionControl/CMakeLists.txt b/catch/unit/executionControl/CMakeLists.txt index 4c28438009..293da7cb99 100644 --- a/catch/unit/executionControl/CMakeLists.txt +++ b/catch/unit/executionControl/CMakeLists.txt @@ -5,6 +5,13 @@ set(TEST_SRC hipFuncSetAttribute.cc ) +if(HIP_PLATFORM MATCHES "amd") + set(TEST_SRC ${TEST_SRC} + hipExtLaunchKernel.cc + hipExtLaunchMultiKernelMultiDevice.cc + ) +endif() + hip_add_exe_to_target(NAME ExecutionControlTest TEST_SRC ${TEST_SRC} TEST_TARGET_NAME build_tests diff --git a/catch/unit/executionControl/hipExtLaunchKernel.cc b/catch/unit/executionControl/hipExtLaunchKernel.cc new file mode 100644 index 0000000000..8b85507de5 --- /dev/null +++ b/catch/unit/executionControl/hipExtLaunchKernel.cc @@ -0,0 +1,176 @@ +/* +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_hipExtLaunchKernel_Positive_Basic") { + SECTION("Kernel with no arguments") { + HIP_CHECK(hipExtLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{1, 1, 1}, + nullptr, 0, nullptr, nullptr, nullptr, 0u)); + 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(hipExtLaunchKernel(reinterpret_cast(kernel_42), dim3{1, 1, 1}, dim3{1, 1, 1}, + kernel_args, 0, nullptr, nullptr, nullptr, 0u)); + int result = 0; + HIP_CHECK(hipMemcpy(&result, result_dev.ptr(), sizeof(result), hipMemcpyDefault)); + REQUIRE(result == 42); + } +} + +TEST_CASE("Unit_hipExtLaunchKernel_Positive_Parameters") { + SECTION("blockDim.x == maxBlockDimX") { + const unsigned int x = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimX); + HIP_CHECK(hipExtLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{x, 1, 1}, + nullptr, 0, nullptr, nullptr, nullptr, 0u)); + } + + SECTION("blockDim.y == maxBlockDimY") { + const unsigned int y = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimY); + HIP_CHECK(hipExtLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{y, 1, 1}, + nullptr, 0, nullptr, nullptr, nullptr, 0u)); + } + + SECTION("blockDim.z == maxBlockDimZ") { + const unsigned int z = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimZ); + HIP_CHECK(hipExtLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{z, 1, 1}, + nullptr, 0, nullptr, nullptr, nullptr, 0u)); + } +} + +TEST_CASE("Unit_hipExtLaunchKernel_Negative_Parameters") { + SECTION("f == nullptr") { + HIP_CHECK_ERROR(hipExtLaunchKernel(nullptr, dim3{1, 1, 1}, dim3{1, 1, 1}, nullptr, 0, nullptr, + nullptr, nullptr, 0u), + hipErrorInvalidDeviceFunction); + } + + SECTION("gridDim.x == 0") { + HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast(kernel), dim3{0, 1, 1}, + dim3{1, 1, 1}, nullptr, 0, nullptr, nullptr, nullptr, 0u), + hipErrorInvalidValue); + } + + SECTION("gridDim.y == 0") { + HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast(kernel), dim3{1, 0, 1}, + dim3{1, 1, 1}, nullptr, 0, nullptr, nullptr, nullptr, 0u), + hipErrorInvalidValue); + } + + SECTION("gridDim.z == 0") { + HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 0}, + dim3{1, 1, 1}, nullptr, 0, nullptr, nullptr, nullptr, 0u), + hipErrorInvalidValue); + } + + SECTION("blockDim.x == 0") { + HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, + dim3{0, 1, 1}, nullptr, 0, nullptr, nullptr, nullptr, 0u), + hipErrorInvalidValue); + } + + SECTION("blockDim.y == 0") { + HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, + dim3{1, 0, 1}, nullptr, 0, nullptr, nullptr, nullptr, 0u), + hipErrorInvalidValue); + } + + SECTION("blockDim.z == 0") { + HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, + dim3{1, 1, 0}, nullptr, 0, nullptr, nullptr, nullptr, 0u), + hipErrorInvalidValue); + } + + SECTION("blockDim.x > maxBlockDimX") { + const unsigned int x = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimX) + 1u; + HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, + dim3{x, 1, 1}, nullptr, 0, nullptr, nullptr, nullptr, 0u), + hipErrorInvalidConfiguration); + } + + SECTION("blockDim.y > maxBlockDimY") { + const unsigned int y = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimY) + 1u; + HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, + dim3{1, y, 1}, nullptr, 0, nullptr, nullptr, nullptr, 0u), + hipErrorInvalidConfiguration); + } + + SECTION("blockDim.z > maxBlockDimZ") { + const unsigned int z = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimZ) + 1u; + HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, + dim3{1, 1, z}, nullptr, 0, nullptr, nullptr, nullptr, 0u), + 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( + hipExtLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, dim3{dim, dim, dim}, + nullptr, 0, nullptr, nullptr, nullptr, 0u), + hipErrorInvalidConfiguration); + } + + SECTION("sharedMemBytes > maxSharedMemoryPerBlock") { + const unsigned int max = GetDeviceAttribute(0, hipDeviceAttributeMaxSharedMemoryPerBlock) + 1u; + HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, + dim3{1, 1, 1}, nullptr, max, nullptr, nullptr, nullptr, 0u), + hipErrorOutOfMemory); + } + + SECTION("Invalid stream") { + hipStream_t stream = nullptr; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, + dim3{1, 1, 1}, nullptr, 0, stream, nullptr, nullptr, 0u), + hipErrorInvalidValue); + } + + SECTION("Invalid startEvent") { + hipEvent_t event = nullptr; + HIP_CHECK(hipEventCreate(&event)); + HIP_CHECK(hipEventDestroy(event)); + HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, + dim3{1, 1, 1}, nullptr, 0, nullptr, event, nullptr, 0u), + hipErrorInvalidValue); + } + + SECTION("Invalid endEvent") { + hipEvent_t event = nullptr; + HIP_CHECK(hipEventCreate(&event)); + HIP_CHECK(hipEventDestroy(event)); + HIP_CHECK_ERROR(hipExtLaunchKernel(reinterpret_cast(kernel), dim3{1, 1, 1}, + dim3{1, 1, 1}, nullptr, 0, nullptr, nullptr, event, 0u), + hipErrorInvalidValue); + } +} \ No newline at end of file diff --git a/catch/unit/executionControl/hipExtLaunchMultiKernelMultiDevice.cc b/catch/unit/executionControl/hipExtLaunchMultiKernelMultiDevice.cc new file mode 100644 index 0000000000..97b1420b9c --- /dev/null +++ b/catch/unit/executionControl/hipExtLaunchMultiKernelMultiDevice.cc @@ -0,0 +1,144 @@ +/* +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_hipExtLaunchMultiKernelMultiDevice_Positive_Basic") { + const auto device_count = HipTest::getDeviceCount(); + + std::vector params_list(device_count); + + int device = 0; + for (auto& params : params_list) { + params.func = reinterpret_cast(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(hipExtLaunchMultiKernelMultiDevice(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_hipExtLaunchMultiKernelMultiDevice_Negative_Parameters") { + const auto device_count = HipTest::getDeviceCount(); + + std::vector params_list(device_count); + + int device = 0; + for (auto& params : params_list) { + params.func = reinterpret_cast(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(hipExtLaunchMultiKernelMultiDevice(nullptr, device_count, 0u), + hipErrorInvalidValue); + } + + SECTION("numDevices == 0") { + HIP_CHECK_ERROR(hipExtLaunchMultiKernelMultiDevice(params_list.data(), 0, 0u), + hipErrorInvalidValue); + } + + SECTION("numDevices > device count") { + HIP_CHECK_ERROR(hipExtLaunchMultiKernelMultiDevice(params_list.data(), device_count + 1, 0u), + hipErrorInvalidValue); + } + + SECTION("invalid flags") { + HIP_CHECK_ERROR(hipExtLaunchMultiKernelMultiDevice(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(kernel2); + HIP_CHECK_ERROR(hipExtLaunchMultiKernelMultiDevice(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(hipExtLaunchMultiKernelMultiDevice(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(hipExtLaunchMultiKernelMultiDevice(params_list.data(), device_count, 0u), + hipErrorInvalidValue); + } + + SECTION("launchParamsList.sharedMem doesn't match across all kernels") { + params_list[1].sharedMem = 1024; + HIP_CHECK_ERROR(hipExtLaunchMultiKernelMultiDevice(params_list.data(), device_count, 0u), + hipErrorInvalidValue); + } + } + + for (const auto params : params_list) { + HIP_CHECK(hipStreamDestroy(params.stream)); + } +} + +TEST_CASE("Unit_hipExtLaunchMultiKernelMultiDevice_Negative_MultiKernelSameDevice") { + HIP_CHECK(hipSetDevice(0)); + + std::vector params_list(2); + + for (auto& params : params_list) { + params.func = reinterpret_cast(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(hipExtLaunchMultiKernelMultiDevice(params_list.data(), 2, 0u), + hipErrorInvalidValue); + + for (const auto params : params_list) { + HIP_CHECK(hipStreamDestroy(params.stream)); + } +} \ No newline at end of file