diff --git a/catch/hipTestMain/config/config_amd_linux b/catch/hipTestMain/config/config_amd_linux index 5c4cce0510..ddb82cfb77 100644 --- a/catch/hipTestMain/config/config_amd_linux +++ b/catch/hipTestMain/config/config_amd_linux @@ -257,6 +257,11 @@ "Unit_Device_Complex_hipCfma_Negative_Parameters_RTC", "Unit_Device_make_Complex_Negative_Parameters_RTC", "Unit_Device_Complex_Cast_Negative_Parameters_RTC", + "Note: Test disabled due to defect - EXSWHTEC-151", + "Unit_hipModuleLoad_Negative_Load_From_A_File_That_Is_Not_A_Module", + "Note: Following two tests disabled due to defect - EXSWHTEC-153", + "Unit_hipModuleLoadData_Negative_Image_Is_An_Empty_String", + "Unit_hipModuleLoadDataEx_Negative_Image_Is_An_Empty_String", #endif #if defined VEGA20 "=== SWDEV-419112 Below tests fail in stress test on 29/08/23 ===", diff --git a/catch/hipTestMain/config/config_amd_windows b/catch/hipTestMain/config/config_amd_windows index 9b2639734a..25436e508f 100644 --- a/catch/hipTestMain/config/config_amd_windows +++ b/catch/hipTestMain/config/config_amd_windows @@ -356,6 +356,13 @@ "Unit_hipGetMipmappedArrayLevel_Negative", "Unit_hipFreeMipmappedArray_Negative_DoubleFree", "Unit_hipFreeMipmappedArrayMultiTArray - int", + "Note: Test disabled due to defect - EXSWHTEC-151", + "Unit_hipModuleLoad_Negative_Load_From_A_File_That_Is_Not_A_Module", + "Note: Test disabled due to defect - EXSWHTEC-152", + "Unit_hipModuleUnload_Negative_Module_Is_Nullptr", + "Note: Following two tests disabled due to defect - EXSWHTEC-153", + "Unit_hipModuleLoadData_Negative_Image_Is_An_Empty_String", + "Unit_hipModuleLoadDataEx_Negative_Image_Is_An_Empty_String", #endif "End of json" ] diff --git a/catch/hipTestMain/config/config_nvidia_linux.json b/catch/hipTestMain/config/config_nvidia_linux.json index f5d75df1b1..79d9d577ee 100644 --- a/catch/hipTestMain/config/config_nvidia_linux.json +++ b/catch/hipTestMain/config/config_nvidia_linux.json @@ -53,6 +53,7 @@ "Unit_atomicExch_system_Positive_Host_And_GPU - unsigned int", "Unit_atomicExch_system_Positive_Host_And_GPU - unsigned long long", "Unit_atomicExch_system_Positive_Host_And_GPU - float", + "Unit_hipModuleUnload_Negative_Double_Unload", "=== Below tests fail in external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/356 ===", "Unit_Device_Complex_Unary_Negative_Parameters_RTC", "Unit_Device_Complex_Binary_Negative_Parameters_RTC", diff --git a/catch/hipTestMain/config/config_nvidia_windows.json b/catch/hipTestMain/config/config_nvidia_windows.json index 4bd8c6dde7..3ae7e43a08 100644 --- a/catch/hipTestMain/config/config_nvidia_windows.json +++ b/catch/hipTestMain/config/config_nvidia_windows.json @@ -16,6 +16,7 @@ "Unit_ChannelDescriptor_Positive_Basic_3D - long3", "Unit_ChannelDescriptor_Positive_Basic_4D - ulong4", "Unit_ChannelDescriptor_Positive_Basic_4D - long4", + "Unit_hipModuleUnload_Negative_Double_Unload", "=== Below tests fail in external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/356 ===", "Unit_Device_Complex_Unary_Negative_Parameters_RTC", "Unit_Device_Complex_Binary_Negative_Parameters_RTC", diff --git a/catch/unit/CMakeLists.txt b/catch/unit/CMakeLists.txt index 8fac51f5c6..ab39835390 100644 --- a/catch/unit/CMakeLists.txt +++ b/catch/unit/CMakeLists.txt @@ -35,12 +35,11 @@ add_subdirectory(multiThread) add_subdirectory(compiler) add_subdirectory(errorHandling) add_subdirectory(cooperativeGrps) -add_subdirectory(warp) add_subdirectory(context) -add_subdirectory(device_memory) +add_subdirectory(module) +add_subdirectory(warp) add_subdirectory(dynamicLoading) add_subdirectory(g++) -add_subdirectory(module) add_subdirectory(channelDescriptor) add_subdirectory(executionControl) add_subdirectory(vector_types) diff --git a/catch/unit/module/CMakeLists.txt b/catch/unit/module/CMakeLists.txt index afdf8c4f0f..262f8ff7dd 100644 --- a/catch/unit/module/CMakeLists.txt +++ b/catch/unit/module/CMakeLists.txt @@ -19,14 +19,29 @@ # SOFTWARE. # Common Tests - Test independent of all platforms -if(HIP_PLATFORM MATCHES "amd") set(TEST_SRC - hipExtModuleLaunchKernel.cc + hip_module_common.cc + hipModuleLoad.cc + hipModuleLoadData.cc + hipModuleLoadDataEx.cc + hipModuleUnload.cc ) +add_custom_target(empty_module.code + COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} + ${CMAKE_CURRENT_SOURCE_DIR}/empty_module.cc + -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/empty_module.code + -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ + -I${CMAKE_CURRENT_SOURCE_DIR}/../../include --rocm-path=${ROCM_PATH}) + # 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 +if(HIP_PLATFORM MATCHES "amd") +set(TEST_SRC + ${TEST_SRC} + hipExtModuleLaunchKernel.cc) + add_custom_target(copyKernel.code COMMAND ${CMAKE_CXX_COMPILER} -mcode-object-version=5 --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/copyKernel.cc @@ -100,11 +115,22 @@ 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(build_tests empty_module.code) + +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 diff --git a/catch/unit/module/empty_module.cc b/catch/unit/module/empty_module.cc new file mode 100644 index 0000000000..d77f3e5733 --- /dev/null +++ b/catch/unit/module/empty_module.cc @@ -0,0 +1,20 @@ +/* +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. +*/ \ No newline at end of file diff --git a/catch/unit/module/hipModuleLoad.cc b/catch/unit/module/hipModuleLoad.cc new file mode 100644 index 0000000000..5812c6a698 --- /dev/null +++ b/catch/unit/module/hipModuleLoad.cc @@ -0,0 +1,59 @@ +/* +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 +#include + +TEST_CASE("Unit_hipModuleLoad_Positive_Basic") { + HIP_CHECK(hipFree(nullptr)); + hipModule_t module = nullptr; + HIP_CHECK(hipModuleLoad(&module, "empty_module.code")); + REQUIRE(module != nullptr); + HIP_CHECK(hipModuleUnload(module)); +} + +TEST_CASE("Unit_hipModuleLoad_Negative_Parameters") { + HIP_CHECK(hipFree(nullptr)); + hipModule_t module; + + SECTION("module == nullptr") { + HIP_CHECK_ERROR(hipModuleLoad(nullptr, "empty_module.code"), hipErrorInvalidValue); + } + + SECTION("fname == nullptr") { + HIP_CHECK_ERROR(hipModuleLoad(&module, nullptr), hipErrorInvalidValue); + } + + SECTION("fname == empty string") { + HIP_CHECK_ERROR(hipModuleLoad(&module, ""), hipErrorInvalidValue); + } + + SECTION("fname == non existent file") { + HIP_CHECK_ERROR(hipModuleLoad(&module, "non existent file"), hipErrorFileNotFound); + } +} + +TEST_CASE("Unit_hipModuleLoad_Negative_Load_From_A_File_That_Is_Not_A_Module") { + HIP_CHECK(hipFree(nullptr)); + hipModule_t module; + + HIP_CHECK_ERROR(hipModuleLoad(&module, "not_a_module.txt"), hipErrorInvalidImage); +} \ No newline at end of file diff --git a/catch/unit/module/hipModuleLoadData.cc b/catch/unit/module/hipModuleLoadData.cc new file mode 100644 index 0000000000..4d364f1d33 --- /dev/null +++ b/catch/unit/module/hipModuleLoadData.cc @@ -0,0 +1,67 @@ +/* +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 + + +TEST_CASE("Unit_hipModuleLoadData_Positive_Basic") { + HIP_CHECK(hipFree(nullptr)); + hipModule_t module = nullptr; + + SECTION("Load compiled module from file") { + const auto loaded_module = LoadModuleIntoBuffer("empty_module.code"); + HIP_CHECK(hipModuleLoadData(&module, loaded_module.data())); + REQUIRE(module != nullptr); + HIP_CHECK(hipModuleUnload(module)); + } + + SECTION("Load RTCd module") { + const auto rtc = CreateRTCCharArray(R"(extern "C" __global__ void kernel() {})"); + HIP_CHECK(hipModuleLoadData(&module, rtc.data())); + REQUIRE(module != nullptr); + HIP_CHECK(hipModuleUnload(module)); + } +} + +TEST_CASE("Unit_hipModuleLoadData_Negative_Parameters") { + HIP_CHECK(hipFree(nullptr)); + hipModule_t module; + + SECTION("module == nullptr") { + const auto loaded_module = LoadModuleIntoBuffer("empty_module.code"); + HIP_CHECK_ERROR(hipModuleLoadData(nullptr, loaded_module.data()), hipErrorInvalidValue); + LoadModuleIntoBuffer("empty_module.code"); + } + + SECTION("image == nullptr") { + HIP_CHECK_ERROR(hipModuleLoadData(&module, nullptr), hipErrorInvalidValue); + } +} + +TEST_CASE("Unit_hipModuleLoadData_Negative_Image_Is_An_Empty_String") { + HIP_CHECK(hipFree(nullptr)); + hipModule_t module; + + HIP_CHECK_ERROR(hipModuleLoadData(&module, ""), hipErrorInvalidImage); +} \ No newline at end of file diff --git a/catch/unit/module/hipModuleLoadDataEx.cc b/catch/unit/module/hipModuleLoadDataEx.cc new file mode 100644 index 0000000000..2ee3f833fe --- /dev/null +++ b/catch/unit/module/hipModuleLoadDataEx.cc @@ -0,0 +1,69 @@ +/* +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 + + +TEST_CASE("Unit_hipModuleLoadDataEx_Positive_Basic") { + HIP_CHECK(hipFree(nullptr)); + hipModule_t module = nullptr; + + SECTION("Load compiled module from file") { + const auto loaded_module = LoadModuleIntoBuffer("empty_module.code"); + HIP_CHECK(hipModuleLoadDataEx(&module, loaded_module.data(), 0, nullptr, nullptr)); + REQUIRE(module != nullptr); + HIP_CHECK(hipModuleUnload(module)); + } + + SECTION("Load RTCd module") { + const auto rtc = CreateRTCCharArray(R"(extern "C" __global__ void kernel() {})"); + HIP_CHECK(hipModuleLoadDataEx(&module, rtc.data(), 0, nullptr, nullptr)); + REQUIRE(module != nullptr); + HIP_CHECK(hipModuleUnload(module)); + } +} + +TEST_CASE("Unit_hipModuleLoadDataEx_Negative_Parameters") { + HIP_CHECK(hipFree(nullptr)); + hipModule_t module = nullptr; + + SECTION("module == nullptr") { + const auto loaded_module = LoadModuleIntoBuffer("empty_module.code"); + HIP_CHECK_ERROR(hipModuleLoadDataEx(nullptr, loaded_module.data(), 0, nullptr, nullptr), + hipErrorInvalidValue); + LoadModuleIntoBuffer("empty_module.code"); + } + + SECTION("image == nullptr") { + HIP_CHECK_ERROR(hipModuleLoadDataEx(&module, nullptr, 0, nullptr, nullptr), + hipErrorInvalidValue); + } +} + +TEST_CASE("Unit_hipModuleLoadDataEx_Negative_Image_Is_An_Empty_String") { + HIP_CHECK(hipFree(nullptr)); + hipModule_t module; + + HIP_CHECK_ERROR(hipModuleLoadDataEx(&module, "", 0, nullptr, nullptr), hipErrorInvalidImage); +} \ No newline at end of file diff --git a/catch/unit/module/hipModuleUnload.cc b/catch/unit/module/hipModuleUnload.cc new file mode 100644 index 0000000000..914b66c89f --- /dev/null +++ b/catch/unit/module/hipModuleUnload.cc @@ -0,0 +1,38 @@ +/* +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 +#include + +TEST_CASE("Unit_hipModuleUnload_Negative_Module_Is_Nullptr") { + HIP_CHECK(hipFree(nullptr)); + + HIP_CHECK_ERROR(hipModuleUnload(nullptr), hipErrorInvalidResourceHandle); +} + +TEST_CASE("Unit_hipModuleUnload_Negative_Double_Unload") { + HIP_CHECK(hipFree(nullptr)); + + hipModule_t module = nullptr; + HIP_CHECK(hipModuleLoad(&module, "empty_module.code")); + HIP_CHECK(hipModuleUnload(module)); + HIP_CHECK_ERROR(hipModuleUnload(module), hipErrorNotFound); +} diff --git a/catch/unit/module/hip_module_common.cc b/catch/unit/module/hip_module_common.cc new file mode 100644 index 0000000000..4e5e7de581 --- /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 0000000000..41b153b3a3 --- /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/not_a_module.txt b/catch/unit/module/not_a_module.txt new file mode 100644 index 0000000000..e69de29bb2