diff --git a/catch/unit/CMakeLists.txt b/catch/unit/CMakeLists.txt index 7d44709d1..7b4eb0f75 100644 --- a/catch/unit/CMakeLists.txt +++ b/catch/unit/CMakeLists.txt @@ -35,6 +35,7 @@ add_subdirectory(multiThread) add_subdirectory(compiler) add_subdirectory(errorHandling) add_subdirectory(cooperativeGrps) +add_subdirectory(module) add_subdirectory(context) add_subdirectory(warp) add_subdirectory(dynamicLoading) diff --git a/catch/unit/module/CMakeLists.txt b/catch/unit/module/CMakeLists.txt index afdf8c4f0..36a9e27dd 100644 --- a/catch/unit/module/CMakeLists.txt +++ b/catch/unit/module/CMakeLists.txt @@ -19,11 +19,21 @@ # SOFTWARE. # Common Tests - Test independent of all platforms -if(HIP_PLATFORM MATCHES "amd") set(TEST_SRC - hipExtModuleLaunchKernel.cc + hip_module_common.cc + hipModuleLaunchKernel.cc ) +add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/launch_kernel_module.code + COMMAND ${CMAKE_CXX_COMPILER} --genco --std=c++17 ${CMAKE_CURRENT_SOURCE_DIR}/launch_kernel_module.cc -o launch_kernel_module.code + DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/launch_kernel_module.cc) +add_custom_target(launch_kernel_module ALL DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/launch_kernel_module.code) + +if(HIP_PLATFORM MATCHES "amd") +set(TEST_SRC + ${TEST_SRC} + hipExtModuleLaunchKernel.cc) + # Note to pass arch use format like -DOFFLOAD_ARCH_STR="--offload-arch=gfx900 --offload-arch=gfx906" # having space at the start/end of OFFLOAD_ARCH_STR can cause build failures @@ -100,14 +110,25 @@ add_custom_target(copiousArgKernel17.code -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include --rocm-path=${ROCM_PATH}) endif() +endif() +if(HIP_PLATFORM MATCHES "amd") + set(RTCLIB "hiprtc") +else() + set(RTCLIB "nvrtc") +endif() hip_add_exe_to_target(NAME ModuleTest TEST_SRC ${TEST_SRC} - TEST_TARGET_NAME build_tests COMMON_SHARED_SRC ${COMMON_SHARED_SRC}) + TEST_TARGET_NAME build_tests + LINKER_LIBS ${RTCLIB} + COMMON_SHARED_SRC ${COMMON_SHARED_SRC}) +add_dependencies(ModuleTest launch_kernel_module) + +if(HIP_PLATFORM MATCHES "amd") add_dependencies(build_tests copyKernel.code copyKernel.s) if(UNIX) add_dependencies(build_tests copiousArgKernel.code copiousArgKernel0.code copiousArgKernel1.code copiousArgKernel2.code copiousArgKernel3.code copiousArgKernel16.code copiousArgKernel17.code) endif() -endif() \ No newline at end of file +endif() diff --git a/catch/unit/module/hipExtModuleLaunchKernel.cc b/catch/unit/module/hipExtModuleLaunchKernel.cc index b8bceb26e..8c77b796d 100644 --- a/catch/unit/module/hipExtModuleLaunchKernel.cc +++ b/catch/unit/module/hipExtModuleLaunchKernel.cc @@ -50,6 +50,8 @@ THE SOFTWARE. #include "hip/hip_ext.h" #include // NOLINT +#include "hip_module_launch_kernel_common.hh" + static constexpr auto totalWorkGroups{1024}; static constexpr auto localWorkSize{512}; static constexpr auto lastWorkSizeEven{256}; @@ -69,7 +71,7 @@ static bool searchRegExpr(const std::regex& expr, const char* filename) { assemblyfile.seekg(0, assemblyfile.end); int len = assemblyfile.tellg(); assemblyfile.seekg(0, assemblyfile.beg); - char *fbuf = new char[len + 1]; + char* fbuf = new char[len + 1]; assemblyfile.read(fbuf, len); fbuf[len] = '\0'; @@ -124,8 +126,7 @@ TEST_CASE("Unit_hipExtModuleLaunchKernel_NonUniformWorkGroup") { auto isEven = GENERATE(0, 1); // Calculate size auto lastWorkSize = isEven ? lastWorkSizeEven : lastWorkSizeOdd; - size_t arraylength = - (totalWorkGroups - 1)*localWorkSize + lastWorkSize; + size_t arraylength = (totalWorkGroups - 1) * localWorkSize + lastWorkSize; size_t sizeBytes{arraylength * sizeof(int)}; // Get module and function from module hipModule_t Module; @@ -133,9 +134,9 @@ TEST_CASE("Unit_hipExtModuleLaunchKernel_NonUniformWorkGroup") { HIP_CHECK(hipModuleLoad(&Module, fileName)); HIP_CHECK(hipModuleGetFunction(&Function, Module, kernel_name)); // Allocate resources - int *A = new int[arraylength]; + int* A = new int[arraylength]; REQUIRE(A != nullptr); - int *B = new int[arraylength]; + int* B = new int[arraylength]; REQUIRE(B != nullptr); // Inititialize data for (size_t i = 0; i < arraylength; i++) { @@ -155,14 +156,13 @@ TEST_CASE("Unit_hipExtModuleLaunchKernel_NonUniformWorkGroup") { args.buffersize = arraylength; size_t size = sizeof(args); - void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, - HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, + void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, HIP_LAUNCH_PARAM_END}; // Memcpy from A to Ad HIP_CHECK(hipMemcpy(Ad, A, sizeBytes, hipMemcpyDefault)); - REQUIRE(hipErrorInvalidValue == hipExtModuleLaunchKernel(Function, - arraylength, 1, 1, localWorkSize, 1, 1, 0, 0, NULL, - reinterpret_cast(&config), 0)); + REQUIRE(hipErrorInvalidValue == + hipExtModuleLaunchKernel(Function, arraylength, 1, 1, localWorkSize, 1, 1, 0, 0, NULL, + reinterpret_cast(&config), 0)); HIP_CHECK(hipDeviceSynchronize()); HIP_CHECK(hipFree(Ad)); HIP_CHECK(hipFree(Bd)); @@ -194,9 +194,9 @@ TEST_CASE("Unit_hipExtModuleLaunchKernel_UniformWorkGroup") { HIP_CHECK(hipModuleLoad(&Module, fileName)); HIP_CHECK(hipModuleGetFunction(&Function, Module, kernel_name)); // Allocate resources - int *A = new int[arraylength]; + int* A = new int[arraylength]; REQUIRE(A != nullptr); - int *B = new int[arraylength]; + int* B = new int[arraylength]; REQUIRE(B != nullptr); // Inititialize data for (size_t i = 0; i < arraylength; i++) { @@ -216,14 +216,12 @@ TEST_CASE("Unit_hipExtModuleLaunchKernel_UniformWorkGroup") { args.buffersize = arraylength; size_t size = sizeof(args); - void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, - HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, + void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, HIP_LAUNCH_PARAM_END}; // Memcpy from A to Ad HIP_CHECK(hipMemcpy(Ad, A, sizeBytes, hipMemcpyDefault)); - HIP_CHECK(hipExtModuleLaunchKernel(Function, arraylength, 1, 1, - localWorkSize, 1, 1, 0, 0, NULL, - reinterpret_cast(&config), 0)); + HIP_CHECK(hipExtModuleLaunchKernel(Function, arraylength, 1, 1, localWorkSize, 1, 1, 0, 0, NULL, + reinterpret_cast(&config), 0)); // Memcpy results back to host HIP_CHECK(hipMemcpy(B, Bd, sizeBytes, hipMemcpyDefault)); HIP_CHECK(hipDeviceSynchronize()); @@ -237,3 +235,53 @@ TEST_CASE("Unit_hipExtModuleLaunchKernel_UniformWorkGroup") { delete[] B; HIP_CHECK(hipModuleUnload(Module)); } + +TEST_CASE("Unit_hipExtModuleLaunchKernel_Positive_Basic") { + ModuleLaunchKernelPositiveBasic(); + + SECTION("Timed kernel launch with events") { + hipEvent_t start_event = nullptr, stop_event = nullptr; + HIP_CHECK(hipEventCreate(&start_event)); + HIP_CHECK(hipEventCreate(&stop_event)); + const auto kernel = GetKernel(mg.module(), "Delay"); + int clock_rate = 0; + HIP_CHECK(hipDeviceGetAttribute(&clock_rate, hipDeviceAttributeClockRate, 0)); + uint32_t interval = 100; + uint32_t ticks_per_second = clock_rate; + void* kernel_params[2] = {&interval, &ticks_per_second}; + HIP_CHECK(hipExtModuleLaunchKernel(kernel, 1, 1, 1, 1, 1, 1, 0, nullptr, kernel_params, nullptr, + start_event, stop_event)); + HIP_CHECK(hipDeviceSynchronize()); + auto elapsed = 0.0f; + HIP_CHECK(hipEventElapsedTime(&elapsed, start_event, stop_event)); + REQUIRE(static_cast(elapsed) >= interval); + } +} + +TEST_CASE("Unit_hipExtModuleLaunchKernel_Positive_Parameters") { + ModuleLaunchKernelPositiveParameters(); + + SECTION("Pass only start event") { + hipEvent_t start_event = nullptr; + HIP_CHECK(hipEventCreate(&start_event)); + const auto kernel = GetKernel(mg.module(), "NOPKernel"); + HIP_CHECK(hipExtModuleLaunchKernel(kernel, 1, 1, 1, 1, 1, 1, 0, nullptr, nullptr, nullptr, + start_event, nullptr)); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipEventQuery(start_event)); + } + + SECTION("Pass only stop event") { + hipEvent_t stop_event = nullptr; + HIP_CHECK(hipEventCreate(&stop_event)); + const auto kernel = GetKernel(mg.module(), "NOPKernel"); + HIP_CHECK(hipExtModuleLaunchKernel(kernel, 1, 1, 1, 1, 1, 1, 0, nullptr, nullptr, nullptr, + nullptr, stop_event)); + HIP_CHECK(hipDeviceSynchronize()); + HIP_CHECK(hipEventQuery(stop_event)); + } +} + +TEST_CASE("Unit_hipExtModuleLaunchKernel_Negative_Parameters") { + ModuleLaunchKernelNegativeParameters(); +} diff --git a/catch/unit/module/hipModuleLaunchKernel.cc b/catch/unit/module/hipModuleLaunchKernel.cc new file mode 100644 index 000000000..f440e8c01 --- /dev/null +++ b/catch/unit/module/hipModuleLaunchKernel.cc @@ -0,0 +1,49 @@ +/* +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 "hip_module_launch_kernel_common.hh" + +#include +#include + +static hipError_t hipModuleLaunchKernelWrapper(hipFunction_t f, uint32_t gridX, uint32_t gridY, + uint32_t gridZ, uint32_t blockX, uint32_t blockY, + uint32_t blockZ, size_t sharedMemBytes, + hipStream_t hStream, void** kernelParams, + void** extra, hipEvent_t, hipEvent_t, uint32_t) { + return hipModuleLaunchKernel(f, gridX, gridY, gridZ, blockX, blockY, blockZ, sharedMemBytes, + hStream, kernelParams, extra); +} + +TEST_CASE("Unit_hipModuleLaunchKernel_Positive_Basic") { + HIP_CHECK(hipFree(nullptr)); + ModuleLaunchKernelPositiveBasic(); +} + +TEST_CASE("Unit_hipModuleLaunchKernel_Positive_Parameters") { + HIP_CHECK(hipFree(nullptr)); + ModuleLaunchKernelPositiveParameters(); +} + +TEST_CASE("Unit_hipModuleLaunchKernel_Negative_Parameters") { + HIP_CHECK(hipFree(nullptr)); + ModuleLaunchKernelNegativeParameters(); +} \ No newline at end of file diff --git a/catch/unit/module/hip_module_common.cc b/catch/unit/module/hip_module_common.cc new file mode 100644 index 000000000..4e5e7de58 --- /dev/null +++ b/catch/unit/module/hip_module_common.cc @@ -0,0 +1,71 @@ +/* +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 "hip_module_common.hh" + +#include +#include + +#include +#include + +ModuleGuard ModuleGuard::LoadModule(const char* fname) { + hipModule_t module = nullptr; + HIP_CHECK(hipModuleLoad(&module, fname)); + return ModuleGuard{module}; +} + +ModuleGuard ModuleGuard::LoadModuleDataFile(const char* fname) { + const auto loaded_module = LoadModuleIntoBuffer(fname); + hipModule_t module = nullptr; + HIP_CHECK(hipModuleLoadData(&module, loaded_module.data())); + return ModuleGuard{module}; +} + +ModuleGuard ModuleGuard::LoadModuleDataRTC(const char* code) { + const auto rtc = CreateRTCCharArray(code); + hipModule_t module = nullptr; + HIP_CHECK(hipModuleLoadData(&module, rtc.data())); + return ModuleGuard{module}; +} + +// Load module into buffer instead of mapping file to avoid platform specific mechanisms +std::vector LoadModuleIntoBuffer(const char* path_string) { + std::experimental::filesystem::path p(path_string); + const auto file_size = std::experimental::filesystem::file_size(p); + std::ifstream f(p, std::ios::binary | std::ios::in); + REQUIRE(f); + std::vector empty_module(file_size); + REQUIRE(f.read(empty_module.data(), file_size)); + return empty_module; +} + +std::vector CreateRTCCharArray(const char* src) { + hiprtcProgram prog; + HIPRTC_CHECK(hiprtcCreateProgram(&prog, src, "prog", 0, nullptr, nullptr)); + HIPRTC_CHECK(hiprtcCompileProgram(prog, 0, nullptr)); + size_t code_size = 0; + HIPRTC_CHECK(hiprtcGetCodeSize(prog, &code_size)); + std::vector code(code_size, '\0'); + HIPRTC_CHECK(hiprtcGetCode(prog, code.data())); + HIPRTC_CHECK(hiprtcDestroyProgram(&prog)); + return code; +} \ No newline at end of file diff --git a/catch/unit/module/hip_module_common.hh b/catch/unit/module/hip_module_common.hh new file mode 100644 index 000000000..41b153b3a --- /dev/null +++ b/catch/unit/module/hip_module_common.hh @@ -0,0 +1,57 @@ +/* +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. +*/ + +#pragma once + +#include + +#include + +class ModuleGuard { + public: + ~ModuleGuard() { static_cast(hipModuleUnload(module_)); } + + ModuleGuard(const ModuleGuard&) = delete; + ModuleGuard(ModuleGuard&&) = delete; + + static ModuleGuard LoadModule(const char* fname); + + static ModuleGuard LoadModuleDataFile(const char* fname); + + static ModuleGuard LoadModuleDataRTC(const char* code); + + hipModule_t module() const { return module_; } + + private: + ModuleGuard(const hipModule_t module) : module_{module} {} + hipModule_t module_ = nullptr; +}; + +// Load module into buffer instead of mapping file to avoid platform specific mechanisms +std::vector LoadModuleIntoBuffer(const char* path_string); + +std::vector CreateRTCCharArray(const char* src); + +inline hipFunction_t GetKernel(const hipModule_t module, const char* kname) { + hipFunction_t kernel = nullptr; + HIP_CHECK(hipModuleGetFunction(&kernel, module, kname)); + return kernel; +} \ No newline at end of file diff --git a/catch/unit/module/hip_module_launch_kernel_common.hh b/catch/unit/module/hip_module_launch_kernel_common.hh new file mode 100644 index 000000000..91ce3a9fd --- /dev/null +++ b/catch/unit/module/hip_module_launch_kernel_common.hh @@ -0,0 +1,269 @@ +/* +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. +*/ + +#pragma once + +#include "hip_module_common.hh" + +#include +#include +#include + +inline int GetDeviceAttribute(const int device, + const hipDeviceAttribute_t attr) { + int value = 0; + HIP_CHECK(hipDeviceGetAttribute(&value, attr, device)); + return value; +} + +inline ModuleGuard InitModule() { + HIP_CHECK(hipFree(nullptr)); + return ModuleGuard::LoadModule("launch_kernel_module.code"); +} + +inline ModuleGuard mg{InitModule()}; + +using ExtModuleLaunchKernelSig = hipError_t(hipFunction_t, uint32_t, uint32_t, uint32_t, uint32_t, + uint32_t, uint32_t, size_t, hipStream_t, void**, void**, + hipEvent_t, hipEvent_t, uint32_t); + +template void ModuleLaunchKernelPositiveBasic() { + SECTION("Kernel with no arguments") { + hipFunction_t f = GetKernel(mg.module(), "NOPKernel"); + HIP_CHECK(func(f, 1, 1, 1, 1, 1, 1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u)); + HIP_CHECK(hipDeviceSynchronize()); + } + + SECTION("Kernel with arguments using kernelParams") { + hipFunction_t f = GetKernel(mg.module(), "Kernel42"); + 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(func(f, 1, 1, 1, 1, 1, 1, 0, nullptr, kernel_args, nullptr, nullptr, nullptr, 0u)); + int result = 0; + HIP_CHECK(hipMemcpy(&result, result_dev.ptr(), sizeof(result), hipMemcpyDefault)); + REQUIRE(result == 42); + } + + SECTION("Kernel with arguments using extra") { + hipFunction_t f = GetKernel(mg.module(), "Kernel42"); + LinearAllocGuard result_dev(LinearAllocs::hipMalloc, sizeof(int)); + HIP_CHECK(hipMemset(result_dev.ptr(), 0, sizeof(*result_dev.ptr()))); + int* result_ptr = result_dev.ptr(); + size_t size = sizeof(result_ptr); + // clang-format off + void *extra[] = { + HIP_LAUNCH_PARAM_BUFFER_POINTER, &result_ptr, + HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, + HIP_LAUNCH_PARAM_END + }; + // clang-format on + HIP_CHECK(func(f, 1, 1, 1, 1, 1, 1, 0, nullptr, nullptr, extra, nullptr, nullptr, 0u)); + int result = 0; + HIP_CHECK(hipMemcpy(&result, result_dev.ptr(), sizeof(result), hipMemcpyDefault)); + REQUIRE(result == 42); + } +} + +template void ModuleLaunchKernelPositiveParameters() { + const auto LaunchNOPKernel = [=](unsigned int gridDimX, unsigned int gridDimY, + unsigned int gridDimZ, unsigned int blockDimX, + unsigned int blockDimY, unsigned int blockDimZ) { + hipFunction_t f = GetKernel(mg.module(), "NOPKernel"); + HIP_CHECK(func(f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, 0, nullptr, + nullptr, nullptr, nullptr, nullptr, 0u)); + HIP_CHECK(hipDeviceSynchronize()); + }; + + SECTION("gridDimX == maxGridDimX") { + const unsigned int x = GetDeviceAttribute(0, hipDeviceAttributeMaxGridDimX); + LaunchNOPKernel(x, 1, 1, 1, 1, 1); + } + + SECTION("gridDimY == maxGridDimY") { + const unsigned int y = GetDeviceAttribute(0, hipDeviceAttributeMaxGridDimY); + LaunchNOPKernel(1, y, 1, 1, 1, 1); + } + + SECTION("gridDimZ == maxGridDimZ") { + const unsigned int z = GetDeviceAttribute(0, hipDeviceAttributeMaxGridDimZ); + LaunchNOPKernel(1, 1, z, 1, 1, 1); + } + + SECTION("blockDimX == maxBlockDimX") { + const unsigned int x = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimX); + LaunchNOPKernel(1, 1, 1, x, 1, 1); + } + + SECTION("blockDimY == maxBlockDimY") { + const unsigned int y = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimY); + LaunchNOPKernel(1, 1, 1, 1, y, 1); + } + + SECTION("blockDimZ == maxBlockDimZ") { + const unsigned int z = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimZ); + LaunchNOPKernel(1, 1, 1, 1, 1, z); + } +} + +template void ModuleLaunchKernelNegativeParameters() { + hipFunction_t f = GetKernel(mg.module(), "NOPKernel"); + +// Disabled on AMD due to defect - EXSWHTEC-157 +#if HT_NVIDIA + SECTION("f == nullptr") { + HIP_CHECK_ERROR( + func(nullptr, 1, 1, 1, 1, 1, 1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u), + hipErrorInvalidResourceHandle); + } +#endif + + SECTION("gridDimX == 0") { + HIP_CHECK_ERROR(func(f, 0, 1, 1, 1, 1, 1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u), + hipErrorInvalidValue); + } + + SECTION("gridDimY == 0") { + HIP_CHECK_ERROR(func(f, 1, 0, 1, 1, 1, 1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u), + hipErrorInvalidValue); + } + + SECTION("gridDimZ == 0") { + HIP_CHECK_ERROR(func(f, 1, 1, 0, 1, 1, 1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u), + hipErrorInvalidValue); + } + + SECTION("blockDimX == 0") { + HIP_CHECK_ERROR(func(f, 1, 1, 1, 0, 1, 1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u), + hipErrorInvalidValue); + } + + SECTION("blockDimY == 0") { + HIP_CHECK_ERROR(func(f, 1, 1, 1, 1, 0, 1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u), + hipErrorInvalidValue); + } + + SECTION("blockDimZ == 0") { + HIP_CHECK_ERROR(func(f, 1, 1, 1, 1, 1, 0, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u), + hipErrorInvalidValue); + } + +// Disabled on AMD due to defect - EXSWHTEC-158 +#if HT_NVIDIA + SECTION("gridDimX > maxGridDimX") { + const unsigned int x = GetDeviceAttribute(0, hipDeviceAttributeMaxGridDimX) + 1u; + HIP_CHECK_ERROR(func(f, x, 1, 1, 1, 1, 1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u), + hipErrorInvalidValue); + } + + SECTION("gridDimY > maxGridDimY") { + const unsigned int y = GetDeviceAttribute(0, hipDeviceAttributeMaxGridDimY) + 1u; + HIP_CHECK_ERROR(func(f, 1, y, 1, 1, 1, 1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u), + hipErrorInvalidValue); + } + + SECTION("gridDimZ > maxGridDimZ") { + const unsigned int z = GetDeviceAttribute(0, hipDeviceAttributeMaxGridDimZ) + 1u; + HIP_CHECK_ERROR(func(f, 1, 1, z, 1, 1, 1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u), + hipErrorInvalidValue); + } +#endif + +// Disabled on AMD due to defect - EXSWHTEC-156 +#if HT_NVIDIA + SECTION("blockDimX > maxBlockDimX") { + const unsigned int x = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimX) + 1u; + HIP_CHECK_ERROR(func(f, 1, 1, 1, x, 1, 1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u), + hipErrorInvalidValue); + } + + SECTION("blockDimY > maxBlockDimY") { + const unsigned int y = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimY) + 1u; + HIP_CHECK_ERROR(func(f, 1, 1, 1, 1, y, 1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u), + hipErrorInvalidValue); + } + + SECTION("blockDimZ > maxBlockDimZ") { + const unsigned int z = GetDeviceAttribute(0, hipDeviceAttributeMaxBlockDimZ) + 1u; + HIP_CHECK_ERROR(func(f, 1, 1, 1, 1, 1, z, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u), + hipErrorInvalidValue); + } +#endif + +// Disabled on AMD due to defect - EXSWHTEC-162 +#if HT_NVIDIA + SECTION("blockDimX * blockDimY * blockDimZ > MaxThreadsPerBlock") { + const unsigned int max = GetDeviceAttribute(0, hipDeviceAttributeMaxThreadsPerBlock); + const unsigned int dim = std::ceil(std::cbrt(max)) + 1; + HIP_CHECK_ERROR( + func(f, 1, 1, 1, dim, dim, dim, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u), + hipErrorInvalidValue); + } +#endif + +// Disabled on AMD due to defect - EXSWHTEC-159 +#if HT_NVIDIA + SECTION("sharedMemBytes > max shared memory per block") { + const unsigned int max = GetDeviceAttribute(0, hipDeviceAttributeMaxSharedMemoryPerBlock) + 1u; + HIP_CHECK_ERROR(func(f, 1, 1, 1, 1, 1, 1, max, nullptr, nullptr, nullptr, nullptr, nullptr, 0u), + hipErrorInvalidValue); + } +#endif + +// Disabled on AMD due to defect - EXSWHTEC-160 +#if HT_NVIDIA + SECTION("Invalid stream") { + hipStream_t stream = nullptr; + HIP_CHECK(hipStreamCreate(&stream)); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK_ERROR(func(f, 1, 1, 1, 1, 1, 0, 0, stream, nullptr, nullptr, nullptr, nullptr, 0u), + hipErrorContextIsDestroyed); + } +#endif + + SECTION("Passing kernel_args and extra simultaneously") { + hipFunction_t f = GetKernel(mg.module(), "Kernel42"); + LinearAllocGuard result_dev(LinearAllocs::hipMalloc, sizeof(int)); + int* result_ptr = result_dev.ptr(); + size_t size = sizeof(result_ptr); + void* kernel_args[1] = {&result_ptr}; + // clang-format off + void *extra[] = { + HIP_LAUNCH_PARAM_BUFFER_POINTER, &result_ptr, + HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, + HIP_LAUNCH_PARAM_END + }; + // clang-format on + HIP_CHECK_ERROR(func(f, 1, 1, 1, 1, 1, 1, 0, nullptr, kernel_args, extra, nullptr, nullptr, 0u), + hipErrorInvalidValue); + } + +// Disabled on AMD due to defect - EXSWHTEC-161 +#if HT_NVIDIA + SECTION("Invalid extra") { + hipFunction_t f = GetKernel(mg.module(), "Kernel42"); + void* extra[0] = {}; + HIP_CHECK_ERROR(func(f, 1, 1, 1, 1, 1, 1, 0, nullptr, nullptr, extra, nullptr, nullptr, 0u), + hipErrorInvalidValue); + } +#endif +} \ No newline at end of file diff --git a/catch/unit/module/launch_kernel_module.cc b/catch/unit/module/launch_kernel_module.cc new file mode 100644 index 000000000..01c04b45d --- /dev/null +++ b/catch/unit/module/launch_kernel_module.cc @@ -0,0 +1,37 @@ +/* +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 + +extern "C" { +__global__ void NOPKernel() {} + +__global__ void Kernel42(int* out) { *out = 42; } + +// Interval is in millisecond +__global__ void Delay(uint32_t interval, const uint32_t ticks_per_ms) { + while (interval--) { + uint64_t start = clock(); + while (clock() - start < ticks_per_ms) { + } + } +} +} \ No newline at end of file