From ba6ef165ecf6fa046efc75824fbaf8be6c8ce793 Mon Sep 17 00:00:00 2001 From: jiabaxie Date: Fri, 8 Mar 2024 16:20:28 -0500 Subject: [PATCH] SWDEV-450204 - test case for hipGetProcAddress Change-Id: I1b14211a1c11087f34ab526dbddd8ef1e02e46ce [ROCm/hip-tests commit: cf045b879a0d131de91726153e6d1eb7af708c30] --- .../unit/dynamicLoading/hipApiDynamicLoad.cc | 227 +++++++++++------- 1 file changed, 139 insertions(+), 88 deletions(-) diff --git a/projects/hip-tests/catch/unit/dynamicLoading/hipApiDynamicLoad.cc b/projects/hip-tests/catch/unit/dynamicLoading/hipApiDynamicLoad.cc index d12c97c4c7..9f81bcbcbc 100644 --- a/projects/hip-tests/catch/unit/dynamicLoading/hipApiDynamicLoad.cc +++ b/projects/hip-tests/catch/unit/dynamicLoading/hipApiDynamicLoad.cc @@ -63,12 +63,138 @@ THE SOFTWARE. * - HIP_VERSION >= 5.6 */ -TEST_CASE("Unit_hipApiDynamicLoad") { - uint32_t *A_d, *C_d; - uint32_t *A_h, *C_h; - size_t N = 1000000; - size_t Nbytes = N * sizeof(uint32_t); +void test_dynamicLoading(void* sym_hipGetDevice, + void* sym_hipMalloc, + void* sym_hipMemcpyHtoD, + void* sym_hipMemcpyDtoH, + void* sym_hipModuleLoad, + void* sym_hipGetDeviceProperties, + void* sym_hipModuleGetFunction, + void* sym_hipModuleLaunchKernel){ + uint32_t *A_d, *C_d; + uint32_t *A_h, *C_h; + size_t N = 1000000; + size_t Nbytes = N * sizeof(uint32_t); + hipError_t (*dyn_hipGetDevice)(hipDevice_t*, int) = reinterpret_cast + (sym_hipGetDevice); + hipError_t (*dyn_hipMalloc)(void**, uint32_t) = reinterpret_cast + (sym_hipMalloc); + + hipError_t (*dyn_hipMemcpyHtoD)(hipDeviceptr_t, void*, size_t) = + reinterpret_cast(sym_hipMemcpyHtoD); + + hipError_t (*dyn_hipMemcpyDtoH)(void*, hipDeviceptr_t, size_t) = + reinterpret_cast(sym_hipMemcpyDtoH); + + hipError_t (*dyn_hipModuleLoad)(hipModule_t*, const char*) = reinterpret_cast + (sym_hipModuleLoad); + + hipError_t (*dyn_hipGetDeviceProperties)(hipDeviceProp_t*, int) = + reinterpret_cast + (sym_hipGetDeviceProperties); + + hipError_t (*dyn_hipModuleGetFunction)(hipFunction_t*, hipModule_t, + const char*) = reinterpret_cast < hipError_t (*)(hipFunction_t*, + hipModule_t, const char*)>(sym_hipModuleGetFunction); + + hipError_t (*dyn_hipModuleLaunchKernel)(hipFunction_t, unsigned int, + unsigned int, unsigned int, unsigned int, unsigned int, + unsigned int, unsigned int, hipStream_t, void**, void**) + = reinterpret_cast(sym_hipModuleLaunchKernel); + + hipDevice_t device; + HIPCHECK(dyn_hipGetDevice(&device, 0)); + + hipDeviceProp_t props; + HIPCHECK(dyn_hipGetDeviceProperties(&props, device)); + A_h = reinterpret_cast(malloc(Nbytes)); + REQUIRE(A_h != NULL); + C_h = reinterpret_cast(malloc(Nbytes)); + REQUIRE(C_h != NULL); + + for (size_t i = 0; i < N; i++) { + A_h[i] = i; + } + + HIPCHECK(dyn_hipMalloc(reinterpret_cast(&A_d), Nbytes)); + HIPCHECK(dyn_hipMalloc(reinterpret_cast(&C_d), Nbytes)); + + HIPCHECK(dyn_hipMemcpyHtoD((hipDeviceptr_t)(A_d), A_h, Nbytes)); + + struct { + void* _Cd; + void* _Ad; + size_t _N; + } args; + args._Cd = reinterpret_cast (C_d); + args._Ad = reinterpret_cast (A_d); + args._N = static_cast (N); + size_t size = sizeof(args); + + void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, + HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, HIP_LAUNCH_PARAM_END}; + + hipModule_t Module; + HIPCHECK(dyn_hipModuleLoad(&Module, fileName)); + + hipFunction_t Function; + HIPCHECK(dyn_hipModuleGetFunction(&Function, Module, "bit_extract_kernel")); + + HIPCHECK(dyn_hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, + reinterpret_cast(&config))); + + HIPCHECK(dyn_hipMemcpyDtoH(C_h, (hipDeviceptr_t)(C_d), Nbytes)); + + for (size_t i = 0; i < N; i++) { + unsigned Agold = ((A_h[i] & 0xf00) >> 8); + REQUIRE(C_h[i] == Agold); + } + HIPCHECK(hipFree(A_d)); + HIPCHECK(hipFree(C_d)); + free(A_h); + free(C_h); + +} +TEST_CASE("Unit_hipApiDynamicLoad_hipGetProcAddress") { + void* sym_hipGetDevice; + void* sym_hipMalloc; + void* sym_hipMemcpyHtoD; + void* sym_hipMemcpyDtoH; + void* sym_hipModuleLoad; + void* sym_hipGetDeviceProperties; + void* sym_hipModuleGetFunction; + void* sym_hipModuleLaunchKernel; + + int currentHipVersion = 0; + HIPCHECK(hipRuntimeGetVersion(¤tHipVersion)); + + HIPCHECK(hipGetProcAddress("hipGetDevice", &sym_hipGetDevice, currentHipVersion, 0, nullptr)); + HIPCHECK(hipGetProcAddress("hipMalloc", &sym_hipMalloc, currentHipVersion, 0, nullptr)); + HIPCHECK(hipGetProcAddress("hipMemcpyHtoD", &sym_hipMemcpyHtoD, currentHipVersion, 0, nullptr)); + HIPCHECK(hipGetProcAddress("hipMemcpyDtoH", &sym_hipMemcpyDtoH, currentHipVersion, 0, nullptr)); + HIPCHECK(hipGetProcAddress("hipModuleLoad", &sym_hipModuleLoad, currentHipVersion, 0, nullptr)); + HIPCHECK(hipGetProcAddress("hipGetDeviceProperties", &sym_hipGetDeviceProperties, currentHipVersion, 0, nullptr)); + HIPCHECK(hipGetProcAddress("hipModuleGetFunction", &sym_hipModuleGetFunction, currentHipVersion, 0, nullptr)); + HIPCHECK(hipGetProcAddress("hipModuleLaunchKernel", &sym_hipModuleLaunchKernel, currentHipVersion, 0, nullptr)); + + test_dynamicLoading(sym_hipGetDevice, + sym_hipMalloc, + sym_hipMemcpyHtoD, + sym_hipMemcpyDtoH, + sym_hipModuleLoad, + sym_hipGetDeviceProperties, + sym_hipModuleGetFunction, + sym_hipModuleLaunchKernel); +} + + +TEST_CASE("Unit_hipApiDynamicLoad") { void* handle = dlopen("libamdhip64.so", RTLD_LAZY); REQUIRE(handle != NULL); @@ -82,90 +208,15 @@ TEST_CASE("Unit_hipApiDynamicLoad") { void* sym_hipModuleLaunchKernel = dlsym(handle, "hipModuleLaunchKernel"); dlclose(handle); - hipError_t (*dyn_hipGetDevice)(hipDevice_t*, int) = reinterpret_cast - (sym_hipGetDevice); - hipError_t (*dyn_hipMalloc)(void**, uint32_t) = reinterpret_cast - (sym_hipMalloc); - - hipError_t (*dyn_hipMemcpyHtoD)(hipDeviceptr_t, void*, size_t) = - reinterpret_cast(sym_hipMemcpyHtoD); - - hipError_t (*dyn_hipMemcpyDtoH)(void*, hipDeviceptr_t, size_t) = - reinterpret_cast(sym_hipMemcpyDtoH); - - hipError_t (*dyn_hipModuleLoad)(hipModule_t*, const char*) = reinterpret_cast - (sym_hipModuleLoad); - - hipError_t (*dyn_hipGetDeviceProperties)(hipDeviceProp_t*, int) = - reinterpret_cast - (sym_hipGetDeviceProperties); - - hipError_t (*dyn_hipModuleGetFunction)(hipFunction_t*, hipModule_t, - const char*) = reinterpret_cast < hipError_t (*)(hipFunction_t*, - hipModule_t, const char*)>(sym_hipModuleGetFunction); - - hipError_t (*dyn_hipModuleLaunchKernel)(hipFunction_t, unsigned int, - unsigned int, unsigned int, unsigned int, unsigned int, - unsigned int, unsigned int, hipStream_t, void**, void**) - = reinterpret_cast(sym_hipModuleLaunchKernel); - - hipDevice_t device; - HIPCHECK(dyn_hipGetDevice(&device, 0)); - - hipDeviceProp_t props; - HIPCHECK(dyn_hipGetDeviceProperties(&props, device)); - A_h = reinterpret_cast(malloc(Nbytes)); - REQUIRE(A_h != NULL); - C_h = reinterpret_cast(malloc(Nbytes)); - REQUIRE(C_h != NULL); - - for (size_t i = 0; i < N; i++) { - A_h[i] = i; - } - - HIPCHECK(dyn_hipMalloc(reinterpret_cast(&A_d), Nbytes)); - HIPCHECK(dyn_hipMalloc(reinterpret_cast(&C_d), Nbytes)); - - HIPCHECK(dyn_hipMemcpyHtoD((hipDeviceptr_t)(A_d), A_h, Nbytes)); - - struct { - void* _Cd; - void* _Ad; - size_t _N; - } args; - args._Cd = reinterpret_cast (C_d); - args._Ad = reinterpret_cast (A_d); - args._N = static_cast (N); - size_t size = sizeof(args); - - void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, - HIP_LAUNCH_PARAM_BUFFER_SIZE, &size, HIP_LAUNCH_PARAM_END}; - - hipModule_t Module; - HIPCHECK(dyn_hipModuleLoad(&Module, fileName)); - - hipFunction_t Function; - HIPCHECK(dyn_hipModuleGetFunction(&Function, Module, "bit_extract_kernel")); - - HIPCHECK(dyn_hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, NULL, - reinterpret_cast(&config))); - - HIPCHECK(dyn_hipMemcpyDtoH(C_h, (hipDeviceptr_t)(C_d), Nbytes)); - - for (size_t i = 0; i < N; i++) { - unsigned Agold = ((A_h[i] & 0xf00) >> 8); - REQUIRE(C_h[i] == Agold); - } - HIPCHECK(hipFree(A_d)); - HIPCHECK(hipFree(C_d)); - free(A_h); - free(C_h); + test_dynamicLoading(sym_hipGetDevice, + sym_hipMalloc, + sym_hipMemcpyHtoD, + sym_hipMemcpyDtoH, + sym_hipModuleLoad, + sym_hipGetDeviceProperties, + sym_hipModuleGetFunction, + sym_hipModuleLaunchKernel); } /**