diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux index 61f3d68b04..398ced33b3 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_linux +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_linux @@ -132,8 +132,6 @@ "Unit_hipEventIpc", "=== SWDEV-427101:Below test fails randomly in PSDB ===", "Unit_deviceAllocation_InOneThread_AccessInAllThreads", - "=== Below test failed in stress test on 11/10/23 ===", - "Unit_hipApiLinkUnusedLibs", "=== StreamCreate Performance tests may potentially fail on all configs ===", "Unit_hipStreamCreate_Performance", "Unit_hipStreamCreate_WithFlagsPerformance_Nonblocking", diff --git a/projects/hip-tests/catch/unit/dynamicLoading/CMakeLists.txt b/projects/hip-tests/catch/unit/dynamicLoading/CMakeLists.txt index a9c3390d74..430b4cbd6d 100644 --- a/projects/hip-tests/catch/unit/dynamicLoading/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/dynamicLoading/CMakeLists.txt @@ -21,8 +21,7 @@ if(UNIX) set(TEST_SRC ${TEST_SRC} complex_loading_behavior.cc) set(AMD_TEST_SRC - hipApiDynamicLoad.cc - hipApiLinkUnusedLib.cc) + hipApiDynamicLoad.cc) if(HIP_PLATFORM MATCHES "amd") set(TEST_SRC ${TEST_SRC} ${AMD_TEST_SRC}) @@ -41,7 +40,6 @@ 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 @@ -49,5 +47,5 @@ hip_add_exe_to_target(NAME Dynamic TEST_TARGET_NAME build_tests LINKER_LIBS ${CMAKE_DL_LIBS}) endif() -add_dependencies(build_tests bit_extract_kernel.code libLazyLoad.so vecadd.cc hipApiLinkUnusedLibAppExe.cc) +add_dependencies(build_tests bit_extract_kernel.code libLazyLoad.so vecadd.cc) endif() diff --git a/projects/hip-tests/catch/unit/dynamicLoading/hipApiLinkUnusedLib.cc b/projects/hip-tests/catch/unit/dynamicLoading/hipApiLinkUnusedLib.cc deleted file mode 100644 index 126929ad13..0000000000 --- a/projects/hip-tests/catch/unit/dynamicLoading/hipApiLinkUnusedLib.cc +++ /dev/null @@ -1,86 +0,0 @@ -/* -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 deleted file mode 100644 index a8f377532d..0000000000 --- a/projects/hip-tests/catch/unit/dynamicLoading/hipApiLinkUnusedLibAppExe.cc +++ /dev/null @@ -1,65 +0,0 @@ -/* -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; -}