SWDEV-447973 - Add generic target codeobj test
Add simple tests to verify generic target code
objects.
Change-Id: Iae148c3c938b18247624938512918dbb3cbc462e
[ROCm/hip-tests commit: f581518103]
Este cometimento está contido em:
taosang2
cometido por
Tao Sang
Tao Sang
ascendente
86432880c6
cometimento
53ddaca565
@@ -17,5 +17,41 @@ if(HIP_PLATFORM MATCHES "amd")
|
||||
hip_add_exe_to_target(NAME SimpleCompressedCodeObjectTest
|
||||
TEST_SRC ${TEST_SRC}
|
||||
TEST_TARGET_NAME build_tests)
|
||||
endif()
|
||||
|
||||
set(OFFLOAD_ARCH_GENERIC_STR "--offload-arch=gfx9-generic --offload-arch=gfx10-1-generic --offload-arch=gfx10-3-generic --offload-arch=gfx11-generic --offload-arch=gfx12-generic")
|
||||
|
||||
# Build hipSquareGenericTargetOnly to cover generic targets only
|
||||
# Because default catch2 build will reference CMAKE_CXX_FLAGS that contains specific targets which will hijack generic
|
||||
# target in hip-rt, we have to use custom build to contain generic targets only.
|
||||
set(GENERIC_TARGET_ONLY_EXE hipSquareGenericTargetOnly)
|
||||
set(LIBFS)
|
||||
if(WIN32)
|
||||
set(GENERIC_TARGET_ONLY_EXE ${GENERIC_TARGET_ONLY_EXE}.exe)
|
||||
else()
|
||||
set(LIBFS -lstdc++fs)
|
||||
endif()
|
||||
|
||||
add_custom_target(hipSquareGenericTargetOnly ALL
|
||||
COMMAND ${CMAKE_CXX_COMPILER} -DNO_GENERIC_TARGET_ONLY_TEST --std=c++17 -mcode-object-version=6 -w "${OFFLOAD_ARCH_GENERIC_STR}"
|
||||
${CMAKE_CURRENT_SOURCE_DIR}/hipSquareGenericTarget.cc
|
||||
${CMAKE_CURRENT_SOURCE_DIR}/../../hipTestMain/hip_test_context.cc
|
||||
${CMAKE_CURRENT_SOURCE_DIR}/../../hipTestMain/hip_test_features.cc
|
||||
${CMAKE_CURRENT_SOURCE_DIR}/../../hipTestMain/main.cc
|
||||
-o ${CMAKE_CURRENT_BINARY_DIR}/${GENERIC_TARGET_ONLY_EXE}
|
||||
-I${HIP_PATH}/include/ --rocm-path=${ROCM_PATH}
|
||||
-I${CMAKE_CURRENT_SOURCE_DIR}/../../include
|
||||
-I${CMAKE_CURRENT_SOURCE_DIR}/../../external/Catch2
|
||||
-I${CMAKE_CURRENT_SOURCE_DIR}/../../external/picojson ${LIBFS})
|
||||
set_property(GLOBAL APPEND PROPERTY G_INSTALL_CUSTOM_TARGETS ${CMAKE_CURRENT_BINARY_DIR}/${GENERIC_TARGET_ONLY_EXE})
|
||||
|
||||
# Build hipSquareGenericTarget to cover generic targets and the specific target
|
||||
set(TEST_SRC
|
||||
hipSquareGenericTarget.cc
|
||||
)
|
||||
hip_add_exe_to_target(NAME hipSquareGenericTarget
|
||||
TEST_SRC ${TEST_SRC}
|
||||
TEST_TARGET_NAME build_tests)
|
||||
set_source_files_properties(hipSquareGenericTarget.cc
|
||||
PROPERTIES COMPILE_FLAGS "-mcode-object-version=6 -w ${OFFLOAD_ARCH_GENERIC_STR}")
|
||||
add_dependencies(hipSquareGenericTarget hipSquareGenericTargetOnly)
|
||||
endif()
|
||||
|
||||
@@ -0,0 +1,108 @@
|
||||
/*
|
||||
Copyright (c) 2024 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 WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS 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 IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
/*
|
||||
* Square each element in the array A and write to array C.
|
||||
*/
|
||||
template <typename T>
|
||||
static __global__ void vector_square_generic(T* C_d, const T* A_d, size_t N) {
|
||||
size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
|
||||
size_t stride = blockDim.x * gridDim.x;
|
||||
|
||||
for (size_t i = offset; i < N; i += stride) {
|
||||
C_d[i] = A_d[i] * A_d[i];
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
TEST_CASE("Unit_test_generic_target_codeobject") {
|
||||
if (!isGenericTargetSupported()) {
|
||||
fprintf(stderr, "Generic target test is skipped\n");
|
||||
return;
|
||||
}
|
||||
float *A_d, *C_d;
|
||||
float *A_h, *C_h;
|
||||
size_t N = 1000000;
|
||||
size_t Nbytes = N * sizeof(float);
|
||||
static int device = 0;
|
||||
HIP_CHECK(hipSetDevice(device));
|
||||
hipDeviceProp_t props;
|
||||
HIP_CHECK(hipGetDeviceProperties(&props, device /*deviceID*/));
|
||||
printf("info: running on device %s\n", props.name);
|
||||
#ifdef __HIP_PLATFORM_AMD__
|
||||
printf("info: architecture on AMD GPU device is: %s\n", props.gcnArchName);
|
||||
// check the scope of supportted types
|
||||
#endif
|
||||
printf("info: allocate host mem (%6.2f MB)\n", 2 * Nbytes / 1024.0 / 1024.0);
|
||||
A_h = (float*)malloc(Nbytes);
|
||||
HIP_CHECK(A_h == 0 ? hipErrorOutOfMemory : hipSuccess);
|
||||
C_h = (float*)malloc(Nbytes);
|
||||
HIP_CHECK(C_h == 0 ? hipErrorOutOfMemory : hipSuccess);
|
||||
// Fill with Phi + i
|
||||
for (size_t i = 0; i < N; i++) {
|
||||
A_h[i] = 1.618f + i;
|
||||
}
|
||||
printf("info: allocate device mem (%6.2f MB)\n", 2 * Nbytes / 1024.0 / 1024.0);
|
||||
HIP_CHECK(hipMalloc(&A_d, Nbytes));
|
||||
HIP_CHECK(hipMalloc(&C_d, Nbytes));
|
||||
|
||||
printf("info: copy Host2Device\n");
|
||||
HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
|
||||
|
||||
const unsigned blocks = 512;
|
||||
const unsigned threadsPerBlock = 256;
|
||||
|
||||
printf("info: launch 'vector_square' kernel\n");
|
||||
hipLaunchKernelGGL(vector_square_generic, dim3(blocks), dim3(threadsPerBlock),
|
||||
0, 0, C_d, A_d, N);
|
||||
|
||||
printf("info: copy Device2Host\n");
|
||||
HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
|
||||
HIP_CHECK(hipDeviceSynchronize());
|
||||
|
||||
printf("info: check result\n");
|
||||
for (size_t i = 0; i < N; i++) {
|
||||
if (C_h[i] != A_h[i] * A_h[i]) {
|
||||
HIP_CHECK(hipErrorUnknown);
|
||||
}
|
||||
}
|
||||
HIP_CHECK(hipFree(A_d));
|
||||
HIP_CHECK(hipFree(C_d));
|
||||
free(A_h);
|
||||
free(C_h);
|
||||
printf("PASSED!\n");
|
||||
REQUIRE(true);
|
||||
}
|
||||
|
||||
#ifndef NO_GENERIC_TARGET_ONLY_TEST
|
||||
TEST_CASE("Unit_test_generic_target_only_codeobject") {
|
||||
#ifdef __linux__
|
||||
char *cmd = "chmod u+x ./hipSquareGenericTargetOnly && ./hipSquareGenericTargetOnly";
|
||||
#else
|
||||
char *cmd = "hipSquareGenericTargetOnly.exe";
|
||||
#endif
|
||||
printf("Run %s\n", cmd);
|
||||
REQUIRE(std::system(cmd) == 0);
|
||||
printf("PASSED!\n");
|
||||
}
|
||||
#endif
|
||||
@@ -113,12 +113,21 @@ add_custom_target(copyKernelCompressed.code
|
||||
-I${HIP_PATH}/include/ --rocm-path=${ROCM_PATH}
|
||||
-I${CMAKE_CURRENT_SOURCE_DIR}/../../include --rocm-path=${ROCM_PATH})
|
||||
|
||||
set(OFFLOAD_ARCH_GENERIC_STR "--offload-arch=gfx9-generic --offload-arch=gfx9-4-generic --offload-arch=gfx10-1-generic --offload-arch=gfx10-3-generic --offload-arch=gfx11-generic --offload-arch=gfx12-generic")
|
||||
add_custom_target(copyKernelGenericTarget.code
|
||||
COMMAND ${CMAKE_CXX_COMPILER} -mcode-object-version=6 --genco ${OFFLOAD_ARCH_GENERIC_STR}
|
||||
${CMAKE_CURRENT_SOURCE_DIR}/copyKernel.cc
|
||||
-o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/copyKernelGenericTarget.code
|
||||
-I${HIP_PATH}/include/ --rocm-path=${ROCM_PATH}
|
||||
-I${CMAKE_CURRENT_SOURCE_DIR}/../../include --rocm-path=${ROCM_PATH})
|
||||
|
||||
set_property(GLOBAL APPEND PROPERTY G_INSTALL_CUSTOM_TARGETS
|
||||
${CMAKE_CURRENT_BINARY_DIR}/empty_module.code
|
||||
${CMAKE_CURRENT_BINARY_DIR}/copyKernel.code
|
||||
${CMAKE_CURRENT_BINARY_DIR}/copyKernel.s
|
||||
${CMAKE_CURRENT_BINARY_DIR}/addKernel.code
|
||||
${CMAKE_CURRENT_BINARY_DIR}/copyKernelCompressed.code
|
||||
${CMAKE_CURRENT_BINARY_DIR}/copyKernelGenericTarget.code
|
||||
)
|
||||
|
||||
if(UNIX)
|
||||
@@ -214,6 +223,8 @@ add_dependencies(build_tests empty_module.code)
|
||||
add_dependencies(build_tests copyKernel.code copyKernel.s)
|
||||
add_dependencies(build_tests addKernel.code)
|
||||
add_dependencies(build_tests copyKernelCompressed.code)
|
||||
add_dependencies(build_tests copyKernelGenericTarget.code)
|
||||
|
||||
if(UNIX)
|
||||
add_dependencies(build_tests copiousArgKernel.code copiousArgKernel0.code copiousArgKernel1.code copiousArgKernel2.code
|
||||
copiousArgKernel3.code copiousArgKernel16.code copiousArgKernel17.code)
|
||||
|
||||
@@ -52,6 +52,8 @@ THE SOFTWARE.
|
||||
constexpr auto fileName = "copyKernel.code";
|
||||
constexpr auto kernel_name = "copy_ker";
|
||||
constexpr auto fileNameCompressed = "copyKernelCompressed.code";
|
||||
constexpr auto fileNameGenericTarget = "copyKernelGenericTarget.code";
|
||||
|
||||
static constexpr auto totalWorkGroups{1024};
|
||||
static constexpr auto localWorkSize{512};
|
||||
static constexpr auto lastWorkSizeEven{256};
|
||||
@@ -196,6 +198,13 @@ TEST_CASE("Unit_hipExtModuleLaunchKernel_UniformWorkGroup") {
|
||||
SECTION("compressed codeobjects") {
|
||||
HIP_CHECK(hipModuleLoad(&Module, fileNameCompressed));
|
||||
}
|
||||
SECTION("generic target codeobjects") {
|
||||
if (!isGenericTargetSupported()) {
|
||||
fprintf(stderr, "Generic target test is skipped\n");
|
||||
return;
|
||||
}
|
||||
HIP_CHECK(hipModuleLoad(&Module, fileNameGenericTarget));
|
||||
}
|
||||
HIP_CHECK(hipModuleGetFunction(&Function, Module, kernel_name));
|
||||
// Allocate resources
|
||||
int* A = new int[arraylength];
|
||||
|
||||
@@ -36,7 +36,7 @@ TEST_CASE("Unit_hipModuleLoadData_Positive_Basic") {
|
||||
HIP_CHECK(hipModuleUnload(module));
|
||||
}
|
||||
|
||||
#if defined(__HIP_PLATFORM_AMD__)
|
||||
#if HT_AMD
|
||||
SECTION("Load compiled module from file with compressed code objects") {
|
||||
const auto loaded_module = LoadModuleIntoBuffer("copyKernelCompressed.code");
|
||||
HIP_CHECK(hipModuleLoadData(&module, loaded_module.data()));
|
||||
@@ -46,6 +46,20 @@ TEST_CASE("Unit_hipModuleLoadData_Positive_Basic") {
|
||||
REQUIRE(kernel != nullptr);
|
||||
HIP_CHECK(hipModuleUnload(module));
|
||||
}
|
||||
|
||||
SECTION("Load compiled module from file with generic target code objects") {
|
||||
if (!isGenericTargetSupported()) {
|
||||
fprintf(stderr, "Generic target test is skipped\n");
|
||||
return;
|
||||
}
|
||||
const auto loaded_module = LoadModuleIntoBuffer("copyKernelGenericTarget.code");
|
||||
HIP_CHECK(hipModuleLoadData(&module, loaded_module.data()));
|
||||
REQUIRE(module != nullptr);
|
||||
hipFunction_t kernel = nullptr;
|
||||
HIP_CHECK(hipModuleGetFunction(&kernel, module, "copy_ker"));
|
||||
REQUIRE(kernel != nullptr);
|
||||
HIP_CHECK(hipModuleUnload(module));
|
||||
}
|
||||
#endif
|
||||
|
||||
SECTION("Load RTCd module") {
|
||||
|
||||
Criar uma nova questão referindo esta
Bloquear um utilizador