From 86677752e9d3790a71cb59e31d2f312aef1b95fd Mon Sep 17 00:00:00 2001 From: Rupam Chetia Date: Tue, 22 Aug 2023 17:03:19 +0530 Subject: [PATCH] SWDEV-416603 - [catch2][dtest] Adding test for linking unused libraries Change-Id: Ib3109c404ad7d04a8416eb785b5d1b0885da8ce3 [ROCm/hip-tests commit: efc7afbf22383b63f3a7b64c8039e4916734150c] --- .../catch/unit/dynamicLoading/CMakeLists.txt | 8 +- .../dynamicLoading/hipApiLinkUnusedLib.cc | 86 +++++++++++++++++++ .../hipApiLinkUnusedLibAppExe.cc | 65 ++++++++++++++ .../catch/unit/dynamicLoading/vecadd.cc | 27 ++++++ 4 files changed, 184 insertions(+), 2 deletions(-) create mode 100644 projects/hip-tests/catch/unit/dynamicLoading/hipApiLinkUnusedLib.cc create mode 100644 projects/hip-tests/catch/unit/dynamicLoading/hipApiLinkUnusedLibAppExe.cc create mode 100644 projects/hip-tests/catch/unit/dynamicLoading/vecadd.cc diff --git a/projects/hip-tests/catch/unit/dynamicLoading/CMakeLists.txt b/projects/hip-tests/catch/unit/dynamicLoading/CMakeLists.txt index b187d7dabd..a9c3390d74 100644 --- a/projects/hip-tests/catch/unit/dynamicLoading/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/dynamicLoading/CMakeLists.txt @@ -21,7 +21,8 @@ if(UNIX) set(TEST_SRC ${TEST_SRC} complex_loading_behavior.cc) set(AMD_TEST_SRC - hipApiDynamicLoad.cc) + hipApiDynamicLoad.cc + hipApiLinkUnusedLib.cc) if(HIP_PLATFORM MATCHES "amd") set(TEST_SRC ${TEST_SRC} ${AMD_TEST_SRC}) @@ -39,11 +40,14 @@ endif() add_custom_target(bit_extract_kernel.code COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR} ${CMAKE_CURRENT_SOURCE_DIR}/bit_extract_kernel.cpp -o ${CMAKE_CURRENT_BINARY_DIR}/../dynamicLoading/bit_extract_kernel.code -I${CMAKE_CURRENT_SOURCE_DIR}/../../../../include/ -I${CMAKE_CURRENT_SOURCE_DIR}/../../include -L${HIP_PATH}/${CMAKE_INSTALL_LIBDIR}) +add_custom_target(vecadd.cc COMMAND cp ${CMAKE_CURRENT_SOURCE_DIR}/vecadd.cc ${CMAKE_CURRENT_BINARY_DIR}/../dynamicLoading/) +add_custom_target(hipApiLinkUnusedLibAppExe.cc COMMAND cp ${CMAKE_CURRENT_SOURCE_DIR}/hipApiLinkUnusedLibAppExe.cc ${CMAKE_CURRENT_BINARY_DIR}/../dynamicLoading/) + if(HIP_PLATFORM MATCHES "amd") hip_add_exe_to_target(NAME Dynamic TEST_SRC ${LINUX_TEST_SRC} TEST_TARGET_NAME build_tests LINKER_LIBS ${CMAKE_DL_LIBS}) endif() -add_dependencies(build_tests bit_extract_kernel.code libLazyLoad.so) +add_dependencies(build_tests bit_extract_kernel.code libLazyLoad.so vecadd.cc hipApiLinkUnusedLibAppExe.cc) endif() diff --git a/projects/hip-tests/catch/unit/dynamicLoading/hipApiLinkUnusedLib.cc b/projects/hip-tests/catch/unit/dynamicLoading/hipApiLinkUnusedLib.cc new file mode 100644 index 0000000000..126929ad13 --- /dev/null +++ b/projects/hip-tests/catch/unit/dynamicLoading/hipApiLinkUnusedLib.cc @@ -0,0 +1,86 @@ +/* +Copyright (c) 2023 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 +#include +#include +#include +/** +* @addtogroup hipLaunchKernelGGL hipLaunchKernelGGL +* @{ +* @ingroup DynamicLoading +* Launches Kernel with launch parameters and shared memory on stream with arguments passed. +*/ + +/** + * Test Description + * ------------------------ + * - This test links a HIP Application with an unused shared object built + * using a different architecture. Launch a kernel built with current GPU + * architecture. Launching the kernel should return error. + * Test source + * ------------------------ + * - catch/unit/dynamicLoading/hipApiLinkUnusedLib.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.0 + */ + +TEST_CASE("Unit_hipApiLinkUnusedLibs") { + hipDeviceProp_t prop; + HIP_CHECK(hipGetDeviceProperties(&prop, 0)); + std::string curArch = prop.gcnArchName; + std::string difArch = ""; + // Build the application for current GPU architecture + // and build the unrelated library for a different GPU. + std::map arch = {{"gfx9", "gfx1010"}, + {"gfx10", "gfx1100"}, + {"gfx11", "gfx900"}}; + + for (auto& elem : arch) { + if (std::string::npos != curArch.find(elem.first)) { + difArch = elem.second; + } + } + if (difArch == "") { + HipTest::HIP_SKIP_TEST("offload-arch not Found. Skipping Test ..."); + return; + } + std::string cmd1 = std::string("hipcc vecadd.cc --offload-arch=") + + difArch; + std::string cmd2 = " -fPIC -c -o vecadd.o"; + cmd1 = cmd1 + cmd2; + std::system(cmd1.data()); + cmd1 = std::string("hipcc vecadd.o -shared --offload-arch=") + difArch; + cmd2 = " -o vecadd.so"; + cmd1 = cmd1 + cmd2; + std::system(cmd1.data()); + cmd1 = std::string( + "hipcc hipApiLinkUnusedLibAppExe.cc -c -o hipApiLinkUnusedLibAppExe.o"); + std::system(cmd1.data()); + cmd1 = std::string( + "hipcc hipApiLinkUnusedLibAppExe.o -o hipApiLinkUnusedLibAppExe"); + cmd2 = std::string(" -Wl,-rpath,. vecadd.so"); + cmd1 = cmd1 + cmd2; + std::system(cmd1.data()); + hip::SpawnProc proc("hipApiLinkUnusedLibAppExe", true); + REQUIRE(proc.run() == 0); + int result = std::stof(proc.getOutput()); + INFO("result: " << result); + REQUIRE(result == 0); +} diff --git a/projects/hip-tests/catch/unit/dynamicLoading/hipApiLinkUnusedLibAppExe.cc b/projects/hip-tests/catch/unit/dynamicLoading/hipApiLinkUnusedLibAppExe.cc new file mode 100644 index 0000000000..a8f377532d --- /dev/null +++ b/projects/hip-tests/catch/unit/dynamicLoading/hipApiLinkUnusedLibAppExe.cc @@ -0,0 +1,65 @@ +/* +Copyright (c) 2023 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 +#include +#include + +#include + +constexpr int arSize = 1024; +constexpr int arBytes = arSize * sizeof(int); + +__global__ void kerSquare(int *inout_a) { + int id = blockIdx.x * blockDim.x + threadIdx.x; + inout_a[id] = inout_a[id] * inout_a[id]; +} + +int main() { + int testPassed = 0; // 0 = passed, -1 = failed + int *a_h = new int[arSize]; + if (a_h == nullptr) { + testPassed = -1; + } + // Inititialize data + for (size_t i = 0; i < arSize; i++) { + a_h[i] = i; + } + int *a_d; + if (hipSuccess != hipMalloc(&a_d, arBytes)) { + testPassed = -1; + } + constexpr int blksize = 128; + constexpr int grdsize = arSize / blksize; + if (hipSuccess != hipMemcpy(a_d, a_h, arSize, hipMemcpyHostToDevice)) { + testPassed = -1; + } + kerSquare<<>>(a_d); + // kernel invocation should return error hipErrorSharedObjectInitFailed + if (hipErrorSharedObjectInitFailed != hipGetLastError()) { + testPassed = -1; + } + if (hipSuccess != hipFree(a_d)) { + testPassed = -1; + } + delete[] a_h; + // std::cout prints the results back to output stream for parent + // process to read. + std::cout << testPassed << std::endl; + return testPassed; +} diff --git a/projects/hip-tests/catch/unit/dynamicLoading/vecadd.cc b/projects/hip-tests/catch/unit/dynamicLoading/vecadd.cc new file mode 100644 index 0000000000..b870f1176e --- /dev/null +++ b/projects/hip-tests/catch/unit/dynamicLoading/vecadd.cc @@ -0,0 +1,27 @@ +/* +Copyright (c) 2023 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 + +__global__ void kerAdd(int *in_a, int *in_b, int *out_c, int nelem) { + int id = blockIdx.x * blockDim.x + threadIdx.x; + if (id >= nelem) { + return; + } + out_c[id] = in_a[id] + in_b[id]; +}