SWDEV-416603 - [catch2][dtest] Adding test for linking unused libraries
Change-Id: Ib3109c404ad7d04a8416eb785b5d1b0885da8ce3
[ROCm/hip-tests commit: efc7afbf22]
This commit is contained in:
Rupam Chetia
Rakesh Roy
@@ -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()
|
||||
|
||||
@@ -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 <hip_test_common.hh>
|
||||
#include <hip_test_defgroups.hh>
|
||||
#include <hip_test_process.hh>
|
||||
#include <map>
|
||||
/**
|
||||
* @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<std::string, std::string> 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);
|
||||
}
|
||||
@@ -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 <iostream>
|
||||
#include <vector>
|
||||
#include <random>
|
||||
|
||||
#include <hip/hip_runtime.h>
|
||||
|
||||
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<<<grdsize, blksize>>>(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;
|
||||
}
|
||||
@@ -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 <hip/hip_runtime.h>
|
||||
|
||||
__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];
|
||||
}
|
||||
Reference in New Issue
Block a user