diff --git a/projects/hip-tests/catch/unit/module/CMakeLists.txt b/projects/hip-tests/catch/unit/module/CMakeLists.txt index 5fd06d0893..62e382784f 100644 --- a/projects/hip-tests/catch/unit/module/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/module/CMakeLists.txt @@ -34,6 +34,8 @@ set(TEST_SRC hipFuncGetAttribute.cc hipGetFuncBySymbol.cc hipDrvLaunchKernelEx.cc + hipModuleGetFunctionCount.cc + hipModuleLoadFatBinary.cc ) add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/get_function_module.code @@ -296,31 +298,24 @@ set(AMD_TEST_SRC if(HIP_PLATFORM MATCHES "amd") set(TEST_SRC ${TEST_SRC} ${AMD_TEST_SRC}) endif() - hip_add_exe_to_target(NAME module TEST_SRC ${TEST_SRC} TEST_TARGET_NAME build_tests) - add_custom_target(managed_kernel.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/managed_kernel.cpp -o ${CMAKE_CURRENT_BINARY_DIR}/../module/managed_kernel.code -I${HIP_PATH}/include -I${CMAKE_CURRENT_SOURCE_DIR}/../../include --hip-path=${HIP_PATH}) -hip_add_exe_to_target(NAME managedKernel - TEST_SRC ${LINUX_TEST_SRC} - TEST_TARGET_NAME build_tests - LINKER_LIBS ${CMAKE_DL_LIBS}) - add_custom_target(vcpy_kernel.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/vcpy_kernel.cpp -o ${CMAKE_CURRENT_BINARY_DIR}/../module/vcpy_kernel.code -I${HIP_PATH}/include -I${CMAKE_CURRENT_SOURCE_DIR}/../../include --hip-path=${HIP_PATH}) -hip_add_exe_to_target(NAME VcpyKernel - TEST_SRC ${LINUX_TEST_SRC} - TEST_TARGET_NAME build_tests - LINKER_LIBS ${CMAKE_DL_LIBS}) - add_custom_target(matmul.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/matmul.cpp -o ${CMAKE_CURRENT_BINARY_DIR}/../module/matmul.code -I${HIP_PATH}/include -I${CMAKE_CURRENT_SOURCE_DIR}/../../include --hip-path=${HIP_PATH}) -hip_add_exe_to_target(NAME matmul - TEST_SRC ${LINUX_TEST_SRC} - TEST_TARGET_NAME build_tests - LINKER_LIBS ${CMAKE_DL_LIBS}) +add_custom_target(kernel_count.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/kernel_count.cpp -o ${CMAKE_CURRENT_BINARY_DIR}/../module/kernel_count.code -I${HIP_PATH}/include -I${CMAKE_CURRENT_SOURCE_DIR}/../../include --hip-path=${HIP_PATH}) + +add_custom_target(emptyModuleCount.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/emptyModuleCount.cpp -o ${CMAKE_CURRENT_BINARY_DIR}/../module/emptyModuleCount.code -I${HIP_PATH}/include -I${CMAKE_CURRENT_SOURCE_DIR}/../../include --hip-path=${HIP_PATH}) + +add_dependencies(ModuleTest managed_kernel.code) +add_dependencies(ModuleTest vcpy_kernel.code) +add_dependencies(ModuleTest matmul.code) +add_dependencies(ModuleTest kernel_count.code) +add_dependencies(ModuleTest emptyModuleCount.code) add_custom_target(kernel_composite_test.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/kernel_composite_test.cpp -o ${CMAKE_CURRENT_BINARY_DIR}/../module/kernel_composite_test.code -I${HIP_PATH}/include -I${CMAKE_CURRENT_SOURCE_DIR}/../../include --hip-path=${HIP_PATH}) set_property(GLOBAL APPEND PROPERTY G_INSTALL_CUSTOM_TARGETS @@ -328,12 +323,11 @@ set_property(GLOBAL APPEND PROPERTY G_INSTALL_CUSTOM_TARGETS ${CMAKE_CURRENT_BINARY_DIR}/matmul.code ${CMAKE_CURRENT_BINARY_DIR}/vcpy_kernel.code ${CMAKE_CURRENT_BINARY_DIR}/managed_kernel.code + ${CMAKE_CURRENT_BINARY_DIR}/kernel_count.code + ${CMAKE_CURRENT_BINARY_DIR}/emptyModuleCount.code ) add_executable(testhipModuleLoadUnloadFunc_exe EXCLUDE_FROM_ALL testhipModuleLoadUnloadFunc_exe.cc) set_property(GLOBAL APPEND PROPERTY G_INSTALL_EXE_TARGETS testhipModuleLoadUnloadFunc_exe) -hip_add_exe_to_target(NAME compositeKernel - TEST_SRC ${LINUX_TEST_SRC} - TEST_TARGET_NAME build_tests - LINKER_LIBS ${CMAKE_DL_LIBS}) -add_dependencies(module managed_kernel.code vcpy_kernel.code matmul.code kernel_composite_test.code testhipModuleLoadUnloadFunc_exe) +add_dependencies(module managed_kernel.code vcpy_kernel.code matmul.code kernel_composite_test.code +testhipModuleLoadUnloadFunc_exe) diff --git a/projects/hip-tests/catch/unit/module/emptyModuleCount.cpp b/projects/hip-tests/catch/unit/module/emptyModuleCount.cpp new file mode 100644 index 0000000000..e69de29bb2 diff --git a/projects/hip-tests/catch/unit/module/hipModuleGetFunctionCount.cc b/projects/hip-tests/catch/unit/module/hipModuleGetFunctionCount.cc new file mode 100644 index 0000000000..2a8270d468 --- /dev/null +++ b/projects/hip-tests/catch/unit/module/hipModuleGetFunctionCount.cc @@ -0,0 +1,114 @@ +/* +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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +#include "hip_module_common.hh" +#include +#include +/** + * @addtogroup hipModuleGetFunctionCount hipModuleGetFunctionCount + * @{ + * @ingroup ModuleTest + * `hipError_t hipModuleGetFunctionCount (unsigned int* count, hipModule_t mod)` + * - Returns the number of functions within a module + */ + +/** + * Test Description + * ------------------------ + * - Test case verifies the below positive cases of hipModuleGetFunctionCount + * API. + * - Get the device count from different kinds of modules. + * 1. Module is built for Single architecture which contain Single global + * function + * 2. Module is built for multi architecture which contain Single global + * function + * 3. Empty Module which doesn't contain any function. + * 4. Module which contain both __global__, __device__ functions in it. + * 5. RTC module. + * Test source + * ------------------------ + * - catch/unit/module/hipModuleGetFunctionCount.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 7.1 + */ +TEST_CASE("Unit_hipModuleGetFunctionCount_Functional") { + hipModule_t moduleSingleArch, moduleEmpty, doubleKernelModule, rtcModule; + unsigned int count = 0; + SECTION("Single arch, Single global function") { + HIP_CHECK(hipModuleLoad(&moduleSingleArch, "vcpy_kernel.code")); + HIP_CHECK(hipModuleGetFunctionCount(&count, moduleSingleArch)); + REQUIRE(count == 1); + HIP_CHECK(hipModuleUnload(moduleSingleArch)); + } +#if HT_AMD + SECTION("Multi arch, Single global function") { + hipModule_t moduleMultiArch; + const auto loaded_module = + LoadModuleIntoBuffer("copyKernelCompressed.code"); + HIP_CHECK(hipModuleLoadData(&moduleMultiArch, loaded_module.data())); + HIP_CHECK(hipModuleGetFunctionCount(&count, moduleMultiArch)); + REQUIRE(count == 1); + HIP_CHECK(hipModuleUnload(moduleMultiArch)); + } +#endif + SECTION("Empty Module Count") { + HIP_CHECK(hipModuleLoad(&moduleEmpty, "emptyModuleCount.code")); + HIP_CHECK(hipModuleGetFunctionCount(&count, moduleEmpty)); + REQUIRE(count == 0); + HIP_CHECK(hipModuleUnload(moduleEmpty)); + } + SECTION("__global__, __device__ functions module") { + HIP_CHECK(hipModuleLoad(&doubleKernelModule, "kernel_count.code")); + HIP_CHECK(hipModuleGetFunctionCount(&count, doubleKernelModule)); + REQUIRE(count == 1); + HIP_CHECK(hipModuleUnload(doubleKernelModule)); + } + + SECTION("Load RTCd module") { + const auto rtc = + CreateRTCCharArray(R"(extern "C" __global__ void kernel() {})"); + HIP_CHECK(hipModuleLoadData(&rtcModule, rtc.data())); + REQUIRE(rtcModule != nullptr); + HIP_CHECK(hipModuleGetFunctionCount(&count, rtcModule)); + REQUIRE(count == 1); + HIP_CHECK(hipModuleUnload(rtcModule)); + } +} +/** + * Test Description + * ------------------------ + * - Test case verifies the negative case of hipModuleGetFunctionCount API. + * Test source + * ------------------------ + * - catch/unit/module/hipModuleGetFunctionCount.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 7.1 + */ +TEST_CASE("Unit_hipModuleGetFunctionCount_NegativeTsts") { + unsigned int count = 0; + SECTION("Input module as nullptr") { + HIP_CHECK_ERROR(hipModuleGetFunctionCount(&count, nullptr), + hipErrorInvalidHandle); + } +} +/** + * End doxygen group ModuleTest. + * @} + */ diff --git a/projects/hip-tests/catch/unit/module/hipModuleLoadFatBinary.cc b/projects/hip-tests/catch/unit/module/hipModuleLoadFatBinary.cc new file mode 100644 index 0000000000..d72b97a8df --- /dev/null +++ b/projects/hip-tests/catch/unit/module/hipModuleLoadFatBinary.cc @@ -0,0 +1,162 @@ +/* +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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +#include "hip_module_common.hh" +#include +#include +#include +/** + * @addtogroup hipModuleLoadFatBinary hipModuleLoadFatBinary + * @{ + * @ingroup ModuleTest + * `hipError_t hipModuleLoadFatBinary(hipModule_t* module, const void* fatbin)` + * - Loads fatbin object + */ + +/** + * Test Description + * ------------------------ + * - Test case verifies the negative cases of hipModuleLoadFatBinary API. + * Test source + * ------------------------ + * - catch/unit/module/hipModuleLoadFatBinary.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 7.1 + */ +TEST_CASE("Unit_hipModuleLoadFatBinary_NegativeTsts") { + hipModule_t Module; + SECTION("fatCubin as nullptr") { + HIP_CHECK_ERROR(hipModuleLoadFatBinary(&Module, nullptr), + hipErrorInvalidValue); + } + SECTION("Load Module with No Kernel function") { + const auto loaded_module = LoadModuleIntoBuffer("emptyModuleCount.code"); + HIP_CHECK(hipModuleLoadFatBinary(&Module, loaded_module.data())); + REQUIRE(Module != nullptr); + HIP_CHECK(hipModuleUnload(Module)); + } +} +#if HT_AMD +void loadKernelData(hipFunction_t kernel) { + constexpr int LEN = 64; + constexpr int SIZE = LEN * sizeof(float); + float *Ad, *Bd; + + std::array A; + std::array B; + + for (uint32_t i = 0; i < LEN; i++) { + A.fill(i * 1.0f); + B.fill(0.0f); + } + HIP_CHECK(hipMalloc(&Ad, SIZE)); + HIP_CHECK(hipMalloc(&Bd, SIZE)); + + HIP_CHECK(hipMemcpy(Ad, A.data(), SIZE, hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(Bd, B.data(), SIZE, hipMemcpyHostToDevice)); + + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + + struct { + void *_Ad; + void *_Bd; + size_t size; + } args; + args._Ad = reinterpret_cast(Ad); + args._Bd = reinterpret_cast(Bd); + args.size = LEN; + size_t size = sizeof(args); + + void *config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, + HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, HIP_LAUNCH_PARAM_END}; + HIP_CHECK(hipModuleLaunchKernel(kernel, 1, 1, 1, LEN, 1, 1, 0, stream, NULL, + reinterpret_cast(&config))); + + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipMemcpy(B.data(), Bd, SIZE, hipMemcpyDeviceToHost)); + // Validation + for (size_t i = 0; i < A.size(); i++) { + INFO("Array Failed at index: " << i + << "\nA value at failed index: " << A[i] + << "\nB value at failed index: " << B[i]); + REQUIRE(A[i] == B[i]); + } + + HIP_CHECK(hipFree(Ad)); + HIP_CHECK(hipFree(Bd)); +} +/** + * Test Description + * ------------------------ + * - Test case verifies the below positive cases of hipModuleLoadFatBinary API. + * case-1 : Loads Compiled module with regular target in compressed fatbin + * case-2 : Loads Compiled module with Generic target in regular fatbin + * case-3 : Loads Compiled module with Generic target in compressed fatbin + * Test source + * ------------------------ + * - catch/unit/module/hipModuleLoadFatBinary.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 7.1 + */ +TEST_CASE("Unit_hipModuleLoadFatBinary_PosiiveTsts") { + hipModule_t Module; + SECTION("Compiled module with regular target in compressed fatbin") { + const auto loaded_module = + LoadModuleIntoBuffer("copyKernelCompressed.code"); + HIP_CHECK(hipModuleLoadFatBinary(&Module, loaded_module.data())); + REQUIRE(Module != nullptr); + hipFunction_t kernel = nullptr; + HIP_CHECK(hipModuleGetFunction(&kernel, Module, "copy_ker")); + loadKernelData(kernel); + REQUIRE(kernel != nullptr); + HIP_CHECK(hipModuleUnload(Module)); + } + if (isGenericTargetSupported()) { + SECTION("Compiled module with Generic target in regular fatbin") { + const auto loaded_module = + LoadModuleIntoBuffer("copyKernelGenericTarget.code"); + HIP_CHECK(hipModuleLoadFatBinary(&Module, loaded_module.data())); + REQUIRE(Module != nullptr); + hipFunction_t kernel = nullptr; + HIP_CHECK(hipModuleGetFunction(&kernel, Module, "copy_ker")); + REQUIRE(kernel != nullptr); + loadKernelData(kernel); + HIP_CHECK(hipModuleUnload(Module)); + } + + SECTION("Compiled module with Generic target in compressed fatbin") { + const auto loaded_module = + LoadModuleIntoBuffer("copyKernelGenericTargetCompressed.code"); + HIP_CHECK(hipModuleLoadFatBinary(&Module, loaded_module.data())); + REQUIRE(Module != nullptr); + hipFunction_t kernel = nullptr; + HIP_CHECK(hipModuleGetFunction(&kernel, Module, "copy_ker")); + REQUIRE(kernel != nullptr); + loadKernelData(kernel); + HIP_CHECK(hipModuleUnload(Module)); + } + } +} +#endif +/** + * End doxygen group ModuleTest. + * @} + */ diff --git a/projects/hip-tests/catch/unit/module/kernel_count.cpp b/projects/hip-tests/catch/unit/module/kernel_count.cpp new file mode 100644 index 0000000000..55a4ad593c --- /dev/null +++ b/projects/hip-tests/catch/unit/module/kernel_count.cpp @@ -0,0 +1,29 @@ +/* +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 WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS 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 INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +#include "hip/hip_runtime.h" + +extern "C" __device__ void hello_world_2(float* a, float* b) { + int tx = threadIdx.x; + b[tx] = a[tx]; +} + +extern "C" __global__ void hello_world(float* a, float* b) { + int tx = threadIdx.x; + b[tx] = a[tx]; +}