diff --git a/projects/hip-tests/catch/unit/module/CMakeLists.txt b/projects/hip-tests/catch/unit/module/CMakeLists.txt index 905cd1bf52..5fd06d0893 100644 --- a/projects/hip-tests/catch/unit/module/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/module/CMakeLists.txt @@ -85,7 +85,8 @@ if(HIP_PLATFORM MATCHES "amd") set(TEST_SRC ${TEST_SRC} hipExtModuleLaunchKernel.cc - hipHccModuleLaunchKernel.cc) + hipHccModuleLaunchKernel.cc + hipLinkCreate.cc) if(BUILD_SHARED_LIBS) set(TEST_SRC @@ -120,6 +121,21 @@ add_custom_target(addKernel.code -I${HIP_PATH}/include/ --hip-path=${HIP_PATH} -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) +add_custom_target(addKernel.spv + COMMAND ${CMAKE_CXX_COMPILER} --genco --offload-arch=amdgcnspirv + --no-gpu-bundle-output + ${CMAKE_CURRENT_SOURCE_DIR}/addKernel.cc + -o ${CMAKE_CURRENT_BINARY_DIR}/addKernel.spv + -I${HIP_PATH}/include/ --hip-path=${HIP_PATH} + -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) + +add_custom_target(addKernel-bundle.spv + COMMAND ${CMAKE_CXX_COMPILER} --genco --offload-arch=amdgcnspirv + ${CMAKE_CURRENT_SOURCE_DIR}/addKernel.cc + -o ${CMAKE_CURRENT_BINARY_DIR}/addKernel-bundle.spv + -I${HIP_PATH}/include/ --hip-path=${HIP_PATH} + -I${CMAKE_CURRENT_SOURCE_DIR}/../../include) + add_custom_target(copyKernelCompressed.code COMMAND ${CMAKE_CXX_COMPILER} -mcode-object-version=5 --offload-compress --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/copyKernel.cc @@ -146,6 +162,8 @@ set_property(GLOBAL APPEND PROPERTY G_INSTALL_CUSTOM_TARGETS ${CMAKE_CURRENT_BINARY_DIR}/copyKernel.code ${CMAKE_CURRENT_BINARY_DIR}/copyKernel.s ${CMAKE_CURRENT_BINARY_DIR}/addKernel.code + ${CMAKE_CURRENT_BINARY_DIR}/addKernel.spv + ${CMAKE_CURRENT_BINARY_DIR}/addKernel-bundle.spv ${CMAKE_CURRENT_BINARY_DIR}/copyKernelCompressed.code ${CMAKE_CURRENT_BINARY_DIR}/copyKernelGenericTarget.code ${CMAKE_CURRENT_BINARY_DIR}/copyKernelGenericTargetCompressed.code @@ -245,6 +263,8 @@ if(HIP_PLATFORM MATCHES "amd") add_dependencies(build_tests empty_module.code) add_dependencies(build_tests copyKernel.code copyKernel.s) add_dependencies(build_tests addKernel.code) +add_dependencies(build_tests addKernel.spv) +add_dependencies(build_tests addKernel-bundle.spv) add_dependencies(build_tests copyKernelCompressed.code) add_dependencies(build_tests copyKernelGenericTarget.code) add_dependencies(build_tests copyKernelGenericTargetCompressed.code) diff --git a/projects/hip-tests/catch/unit/module/hipLinkCreate.cc b/projects/hip-tests/catch/unit/module/hipLinkCreate.cc new file mode 100644 index 0000000000..7c7f998958 --- /dev/null +++ b/projects/hip-tests/catch/unit/module/hipLinkCreate.cc @@ -0,0 +1,281 @@ +/* +Copyright (c) 2025 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 + +#include +#include +#include + +static constexpr const char* SPIRV_FILE = "addKernel.spv"; +static constexpr const char* SPIRV_BUNDLED_FILE = "addKernel-bundle.spv"; +static constexpr int ARRAY_SIZE = 1; +static constexpr int REF_VALUE = 5; + +#if HT_AMD +static inline bool load_co_from_file(const char *filename, std::vector &co_source) { + std::ifstream file_stream{filename, std::ios_base::in | std::ios_base::binary}; + if (!file_stream.good()) { + return false; + } + + file_stream.seekg(0, std::ios::end); + std::streampos file_size = file_stream.tellg(); + file_stream.seekg(0, std::ios::beg); + + // Read the file contents + co_source.resize(file_size); + file_stream.read(co_source.data(), file_size); + + file_stream.close(); + + return true; +} + +static inline void JitLink(hipModule_t *Module, hipFunction_t *Kernel, hipLinkState_t *LinkState, + hipJitInputType input_type, bool from_file, const char *filename) { + + const char* isaopts[] = {"-mllvm", "-inline-threshold=1", "-mllvm", "-inlinehint-threshold=1"}; + std::vector jit_options = {hipJitOptionIRtoISAOptExt, + hipJitOptionIRtoISAOptCountExt}; + size_t isaoptssize = 4; + const void* lopts[] = {(void*)isaopts, (void*)(isaoptssize)}; + + + HIP_CHECK(hipLinkCreate(jit_options.size(), jit_options.data(), (void **)lopts, LinkState)); + + if (!from_file) { + std::vector co_source; + REQUIRE(load_co_from_file(filename, co_source) == true); + HIP_CHECK(hipLinkAddData(*LinkState, input_type, (void*)co_source.data(), co_source.size(), + "LinkSPIRV1", 0, nullptr, nullptr)); + } else { + HIP_CHECK(hipLinkAddFile(*LinkState,input_type, filename, 0, nullptr, nullptr)); + } + + void *linkOut; + size_t linkSize = 0; + // Complete Linker Step + HIP_CHECK(hipLinkComplete(*LinkState, &linkOut, &linkSize)); + + // Load codeobject into module + HIP_CHECK(hipModuleLoadData(Module, linkOut)); + + // Locate Kernel Entry Point + HIP_CHECK(hipModuleGetFunction(Kernel, *Module, "addKernel")); + + // Destroy Linker invocation + HIP_CHECK(hipLinkDestroy(*LinkState)); + +} + +/** + * Test Description + * ------------------------ + * - Validates SPIR-V linking and kernel execution using HIP's JIT linker with both bundled and + * unbundled SPIR-V code objects. Tests both hipLinkAddData and hipLinkAddFile APIs + * + * Test source + * ------------------------ + * - unit/module/hipLinkCreate.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.4 + */ +TEST_CASE("Unit_hip_linker_spirv_input") { + size_t N = ARRAY_SIZE; + size_t sizeBytes = N * sizeof(int); + int *A_h = new int[sizeBytes]; + REQUIRE(A_h != nullptr); + int *A_d; HIP_CHECK(hipMalloc(&A_d, sizeBytes)); + REQUIRE(A_d != nullptr); + + for ( int i = 0; i < N; i++) { + A_h[i] = REF_VALUE + i; + } + HIP_CHECK(hipMemcpy(A_d, A_h, sizeBytes, hipMemcpyHostToDevice)); + + hipModule_t Module; + hipFunction_t Kernel; + hipLinkState_t Linkstate; + + bool from_file = false; + hipJitInputType input_type= hipJitInputSpirv; + const char *filename; + SECTION("Link Add Data with Bundled Spirv") { + from_file = false; + input_type = hipJitInputSpirv; + filename = SPIRV_BUNDLED_FILE; + } + SECTION("Link Add Data with UnBundled Spirv") { + from_file = false; + input_type = hipJitInputSpirv; + filename = SPIRV_FILE; + } + SECTION("Link Add File with Bundled Spirv") { + from_file = true; + input_type = hipJitInputSpirv; + filename = SPIRV_BUNDLED_FILE; + } + SECTION("Link Add File with UnBundled Spirv") { + from_file = true; + input_type = hipJitInputSpirv; + filename = SPIRV_FILE; + } + + JitLink(&Module, &Kernel, &Linkstate, input_type, from_file,filename); + void *args[2] = {&A_d, &N}; + HIP_CHECK(hipModuleLaunchKernel(Kernel, 1, 1, 1, 1, 1, 1, 0, nullptr, args, nullptr)); + HIP_CHECK(hipModuleUnload(Module)); + HIP_CHECK(hipMemcpy(A_h, A_d, sizeBytes, hipMemcpyDeviceToHost)); + REQUIRE(A_h[0] == REF_VALUE + 2); + HIP_CHECK(hipFree(A_d)); + delete[] A_h; +} + +/** + * Test Description + * ------------------------ + * Negative test cases for hipLinkCreate to verify it fails gracefully with invalid arguments like + * null pointers and mismatched option arrays. + * + * Test source + * ------------------------ + * - unit/module/hipLinkCreate.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.4 + */ +TEST_CASE("Unit_hipLinkCreate_Negative") { + + hipLinkState_t linkstate; + hipJitOption options[4]; + void *optionVals[4]; + options[0] = hipJitOptionInfoLogBuffer; + options[1] = hipJitOptionInfoLogBufferSizeBytes; + options[2] = hipJitOptionErrorLogBuffer; + options[3] = hipJitOptionErrorLogBufferSizeBytes; + + SECTION("linkstate == nullptr") { + HIP_CHECK_ERROR(hipLinkCreate(0, options, optionVals, nullptr), hipErrorInvalidValue); + } + SECTION("num_options != 0 & option val pptr == nullptr") { + HIP_CHECK_ERROR(hipLinkCreate(4, options, nullptr, &linkstate), hipErrorInvalidValue); + } + SECTION("num_options != 0 & option val ptr == nullptr") { + HIP_CHECK_ERROR(hipLinkCreate(4, options, optionVals, &linkstate), hipErrorInvalidValue); + } + +} +/** + * Test Description + * ------------------------ + * - Validates that unsupported CUDA only options don't crash the linker and result in the correct + * error. + * + * Test source + * ------------------------ + * - unit/module/hipLinkCreate.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.4 + */ +TEST_CASE("Unit_hipLinkCreate_AddLinker_CUDA_only_options") { + hipLinkState_t linkstate; + // Random options so that it is not null + const char* isaopts[] = {"-mllvm", "-inline-threshold=1", "-mllvm", "-inlinehint-threshold=1"}; + size_t isaoptssize = 4; + const void* lopts[] = {(void*)isaopts, (void*)(isaoptssize)}; + + std::vector options = {hipJitOptionMaxRegisters, + hipJitOptionThreadsPerBlock, + hipJitOptionWallTime, + hipJitOptionInfoLogBuffer, + hipJitOptionInfoLogBufferSizeBytes, + hipJitOptionErrorLogBuffer, + hipJitOptionErrorLogBufferSizeBytes, + hipJitOptionOptimizationLevel, + hipJitOptionTargetFromContext, + hipJitOptionTarget, + hipJitOptionFallbackStrategy, + hipJitOptionGenerateDebugInfo, + hipJitOptionLogVerbose, + hipJitOptionGenerateLineInfo, + hipJitOptionCacheMode, + hipJitOptionSm3xOpt, + hipJitOptionFastCompile, + hipJitOptionGlobalSymbolNames, + hipJitOptionGlobalSymbolAddresses, + hipJitOptionGlobalSymbolCount, + hipJitOptionLto, + hipJitOptionFtz, + hipJitOptionPrecDiv, + hipJitOptionPrecSqrt, + hipJitOptionFma, + hipJitOptionPositionIndependentCode, + hipJitOptionMinCTAPerSM, + hipJitOptionMaxThreadsPerBlock, + hipJitOptionOverrideDirectiveValues, + hipJitOptionNumOptions}; + + HIP_CHECK_ERROR(hipLinkCreate(options.size(), options.data(), (void**)lopts, &linkstate), + hipErrorInvalidValue); +} + +/** + * Test Description + * ------------------------ + * Verifies error handling of hipLinkAddFile when given invalid parameters, input types, or + * nonexistent files. + * + * Test source + * ------------------------ + * - unit/module/hipLinkCreate.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.4 + */ +TEST_CASE("Unit_hipLinkAddFile_Negative") { + hipLinkState_t linkstate; + HIP_CHECK(hipLinkCreate(0, nullptr, nullptr, &linkstate)); + + SECTION("linkstate == nullptr") { + HIP_CHECK_ERROR(hipLinkAddFile(nullptr, hipJitInputSpirv, SPIRV_FILE, 0, nullptr, nullptr), + hipErrorInvalidHandle); + } + SECTION("Jit input Cubin") { + HIP_CHECK_ERROR( + hipLinkAddFile(linkstate, hipJitInputCubin, SPIRV_FILE, 0, nullptr, nullptr), + hipErrorInvalidValue); + } + SECTION("Jit Input PTX") { + HIP_CHECK_ERROR(hipLinkAddFile(linkstate, hipJitInputPtx, SPIRV_FILE, 0, nullptr, nullptr), + hipErrorInvalidValue); + } + SECTION("Input File not valid") { + HIP_CHECK_ERROR( + hipLinkAddFile(linkstate, hipJitInputSpirv, "unknown_file", 0, nullptr, nullptr), + hipErrorInvalidConfiguration); + } + HIP_CHECK(hipLinkDestroy(linkstate)); +} +#endif \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/rtc/RtcFunctions.cpp b/projects/hip-tests/catch/unit/rtc/RtcFunctions.cpp index 403662b332..cf012e1d22 100644 --- a/projects/hip-tests/catch/unit/rtc/RtcFunctions.cpp +++ b/projects/hip-tests/catch/unit/rtc/RtcFunctions.cpp @@ -22,8 +22,8 @@ This file contains functions for idividual HIPRTC supported compiler options validation. For PASS senario the function returns 1 or 0 otherwise. */ -#include #include +#include #include #include #include diff --git a/projects/hip-tests/catch/unit/rtc/RtcUtility.cpp b/projects/hip-tests/catch/unit/rtc/RtcUtility.cpp index 5773df29de..bc31565fbc 100644 --- a/projects/hip-tests/catch/unit/rtc/RtcUtility.cpp +++ b/projects/hip-tests/catch/unit/rtc/RtcUtility.cpp @@ -48,8 +48,8 @@ parameters of the respective block name. */ -#include #include +#include #include #include #include diff --git a/projects/hip-tests/catch/unit/rtc/hipRtcComplexHeader.cc b/projects/hip-tests/catch/unit/rtc/hipRtcComplexHeader.cc index 550cffebab..5c237ff0db 100644 --- a/projects/hip-tests/catch/unit/rtc/hipRtcComplexHeader.cc +++ b/projects/hip-tests/catch/unit/rtc/hipRtcComplexHeader.cc @@ -28,9 +28,8 @@ THE SOFTWARE. * string and compile using the api mentioned above. */ -#include -#include #include +#include #include static constexpr auto hip_complex_basic_string{ diff --git a/projects/hip-tests/catch/unit/rtc/hipStreamCaptureRtc.cc b/projects/hip-tests/catch/unit/rtc/hipStreamCaptureRtc.cc index e8c1a94330..f25d27b7f6 100644 --- a/projects/hip-tests/catch/unit/rtc/hipStreamCaptureRtc.cc +++ b/projects/hip-tests/catch/unit/rtc/hipStreamCaptureRtc.cc @@ -17,10 +17,10 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ +#include #include #include #include -#include static constexpr auto kernel_src{ R"_KERN_EMBED_( diff --git a/projects/hip-tests/catch/unit/rtc/hiprtcComplrOptnTesting.cc b/projects/hip-tests/catch/unit/rtc/hiprtcComplrOptnTesting.cc index 53809539cd..2d50477813 100644 --- a/projects/hip-tests/catch/unit/rtc/hiprtcComplrOptnTesting.cc +++ b/projects/hip-tests/catch/unit/rtc/hiprtcComplrOptnTesting.cc @@ -17,15 +17,14 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ +#include #include -#include #include #include #include #include #include #include -#include #include "headers/RtcUtility.h" #include "headers/RtcFunctions.h" #include "headers/RtcKernels.h" diff --git a/projects/hip-tests/catch/unit/rtc/hiprtcGetLoweredName.cc b/projects/hip-tests/catch/unit/rtc/hiprtcGetLoweredName.cc index 0e0a52a7e1..93d65b04d7 100644 --- a/projects/hip-tests/catch/unit/rtc/hiprtcGetLoweredName.cc +++ b/projects/hip-tests/catch/unit/rtc/hiprtcGetLoweredName.cc @@ -18,10 +18,10 @@ THE SOFTWARE. */ #include +#include #include #include #include -#include /** diff --git a/projects/hip-tests/catch/unit/rtc/hiprtc_MathConstants_HeaderTst.cc b/projects/hip-tests/catch/unit/rtc/hiprtc_MathConstants_HeaderTst.cc index b5a536d7e3..5ba6c3d763 100644 --- a/projects/hip-tests/catch/unit/rtc/hiprtc_MathConstants_HeaderTst.cc +++ b/projects/hip-tests/catch/unit/rtc/hiprtc_MathConstants_HeaderTst.cc @@ -28,9 +28,8 @@ THE SOFTWARE. * string and compile using the api mentioned above. */ -#include -#include #include +#include static constexpr auto mathConstants_string{ R"( diff --git a/projects/hip-tests/catch/unit/rtc/hiprtc_MathFunctions_HeaderTst.cc b/projects/hip-tests/catch/unit/rtc/hiprtc_MathFunctions_HeaderTst.cc index f9cb922c25..fb876e5bbf 100644 --- a/projects/hip-tests/catch/unit/rtc/hiprtc_MathFunctions_HeaderTst.cc +++ b/projects/hip-tests/catch/unit/rtc/hiprtc_MathFunctions_HeaderTst.cc @@ -28,11 +28,10 @@ THE SOFTWARE. * string and compile using the api mentioned above. */ -#include -#include -#include #include #include "hip_test_rtc.hh" +#include +#include static constexpr auto mathFuntn_string{ R"( diff --git a/projects/hip-tests/catch/unit/rtc/hiprtc_TextureTypes_HeaderTst.cc b/projects/hip-tests/catch/unit/rtc/hiprtc_TextureTypes_HeaderTst.cc index a7b096cd0a..bd8aea1c05 100644 --- a/projects/hip-tests/catch/unit/rtc/hiprtc_TextureTypes_HeaderTst.cc +++ b/projects/hip-tests/catch/unit/rtc/hiprtc_TextureTypes_HeaderTst.cc @@ -27,9 +27,8 @@ THE SOFTWARE. * string and compile using the api mentioned above. */ -#include -#include #include +#include static constexpr auto TextureTypes_string{ R"( diff --git a/projects/hip-tests/catch/unit/rtc/hiprtc_VectorTypes_HeaderTst.cc b/projects/hip-tests/catch/unit/rtc/hiprtc_VectorTypes_HeaderTst.cc index e343acc25a..c4aa8767d6 100644 --- a/projects/hip-tests/catch/unit/rtc/hiprtc_VectorTypes_HeaderTst.cc +++ b/projects/hip-tests/catch/unit/rtc/hiprtc_VectorTypes_HeaderTst.cc @@ -28,9 +28,8 @@ THE SOFTWARE. * string and compile using the api mentioned above. */ -#include -#include #include +#include static constexpr auto vectorTypes_string{ R"( diff --git a/projects/hip-tests/catch/unit/rtc/hiprtc_bfloat16_HeaderTst.cc b/projects/hip-tests/catch/unit/rtc/hiprtc_bfloat16_HeaderTst.cc index 707598a173..61ef7e4ee5 100644 --- a/projects/hip-tests/catch/unit/rtc/hiprtc_bfloat16_HeaderTst.cc +++ b/projects/hip-tests/catch/unit/rtc/hiprtc_bfloat16_HeaderTst.cc @@ -28,9 +28,8 @@ THE SOFTWARE. * string and compile using the api mentioned above. */ -#include -#include #include +#include static constexpr auto bfloat16_string{ R"( extern "C" diff --git a/projects/hip-tests/catch/unit/rtc/hiprtc_fp16_HeaderTst.cc b/projects/hip-tests/catch/unit/rtc/hiprtc_fp16_HeaderTst.cc index b4c241d84f..5b2e0e82f8 100644 --- a/projects/hip-tests/catch/unit/rtc/hiprtc_fp16_HeaderTst.cc +++ b/projects/hip-tests/catch/unit/rtc/hiprtc_fp16_HeaderTst.cc @@ -27,9 +27,8 @@ THE SOFTWARE. * string and compile using the api mentioned above. */ -#include -#include #include +#include static constexpr auto fp16_string{ R"(