SWDEV-546347-Added tests for Module Management APIs (#539)

* SWDEV-546347-Added tests for hipModuleGetFunctionCount

* SWDEV-546347-Added tests for hipModuleLoadFatBinary

* SWDEV-546347-Added 3 more tests for hipModuleLoadFatBinary

* SWDEV-546347-Updated hipModuleLoadFatBinary.cc

* SWDEV-546347-Added test for hipModuleGetLoadingMode API

* SWDEV-546347-Clang-format issues

* SWDEV-546347-Removed tests of hipModuleGetLoadingMode API

* SWDEV-546347 - Updated CMake file

* SWDEV-546347-Addressed review comments

* SWDEV-546347-Updated CmakeLists

* SWDEV-546347-Addressed review comments

* SWDEV-546347-Addressed Review comments

---------

Co-authored-by: SrinivasRao Gollamandala <Srinivasarao.Gollamandala@amd.com>
Co-authored-by: systems-assistant[bot] <systems-assistant[bot]@users.noreply.github.com>
Co-authored-by: jainprad <92369414+jainprad@users.noreply.github.com>
Co-authored-by: Rambabu Swargam <rambabu.swargam@amd.com>
이 커밋은 다음에 포함됨:
systems-assistant[bot]
2025-10-10 15:29:20 +05:30
커밋한 사람 GitHub
부모 446fbd3191
커밋 9021b43cd1
5개의 변경된 파일320개의 추가작업 그리고 21개의 파일을 삭제
+15 -21
파일 보기
@@ -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)
파일 보기
+114
파일 보기
@@ -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 <hip_test_common.hh>
#include <hip_test_defgroups.hh>
/**
* @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.
* @}
*/
+162
파일 보기
@@ -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 <array>
#include <hip_test_common.hh>
#include <hip_test_defgroups.hh>
/**
* @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<float, LEN> A;
std::array<float, LEN> 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<void *>(Ad);
args._Bd = reinterpret_cast<void *>(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<void **>(&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.
* @}
*/
+29
파일 보기
@@ -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];
}