diff --git a/catch/hipTestMain/config/config_amd_linux b/catch/hipTestMain/config/config_amd_linux index 73ec536866..d19f61b90b 100644 --- a/catch/hipTestMain/config/config_amd_linux +++ b/catch/hipTestMain/config/config_amd_linux @@ -863,8 +863,6 @@ "Unit_safeAtomicMin_Positive_SameAddress - float", "=== SWDEV-454220 : Below test hanged in stress test on 22/03/24 ===", "Unit_hipExtLaunchKernel_Positive_Basic", - "=== Temporarily disable the test that failed in mi300 ===", - "Unit_test_generic_target_only_codeobject", #endif #if defined gfx1030 "=== SWDEV-445961: These tests hang in PSDB stress test on 09/02/2024 ===", diff --git a/catch/unit/compiler/CMakeLists.txt b/catch/unit/compiler/CMakeLists.txt index 4837ebfc73..ab630f7adf 100644 --- a/catch/unit/compiler/CMakeLists.txt +++ b/catch/unit/compiler/CMakeLists.txt @@ -18,15 +18,18 @@ if(HIP_PLATFORM MATCHES "amd") TEST_SRC ${TEST_SRC} TEST_TARGET_NAME build_tests) - 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") + set(OFFLOAD_ARCH_GENERIC_STR "--offload-arch=gfx9-generic --offload-arch=gfx9-4-generic:sramecc+:xnack- --offload-arch=gfx9-4-generic:sramecc-:xnack- --offload-arch=gfx9-4-generic:xnack+ --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(GENERIC_TARGET_ONLY_COMPRESSED_EXE hipSquareGenericTargetOnlyCompressed) + set(LIBFS) if(WIN32) set(GENERIC_TARGET_ONLY_EXE ${GENERIC_TARGET_ONLY_EXE}.exe) + set(GENERIC_TARGET_ONLY_COMPRESSED_EXE ${GENERIC_TARGET_ONLY_COMPRESSED_EXE}.exe) else() set(LIBFS -lstdc++fs) endif() @@ -42,7 +45,19 @@ if(HIP_PLATFORM MATCHES "amd") -I${CMAKE_CURRENT_SOURCE_DIR}/../../include -I${CMAKE_CURRENT_SOURCE_DIR}/../../external/Catch2 -I${CMAKE_CURRENT_SOURCE_DIR}/../../external/picojson ${LIBFS}) + add_custom_target(hipSquareGenericTargetOnlyCompressed ALL + COMMAND ${CMAKE_CXX_COMPILER} -DNO_GENERIC_TARGET_ONLY_TEST -DGENERIC_COMPRESSED --std=c++17 -mcode-object-version=6 --offload-compress -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_COMPRESSED_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}) + set_property(GLOBAL APPEND PROPERTY G_INSTALL_CUSTOM_TARGETS ${CMAKE_CURRENT_BINARY_DIR}/${GENERIC_TARGET_ONLY_COMPRESSED_EXE}) # Build hipSquareGenericTarget to cover generic targets and the specific target set(TEST_SRC @@ -51,7 +66,14 @@ if(HIP_PLATFORM MATCHES "amd") 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}") + set_target_properties(hipSquareGenericTarget PROPERTIES COMPILE_FLAGS "-mcode-object-version=6 -w ${OFFLOAD_ARCH_GENERIC_STR}") + + hip_add_exe_to_target(NAME hipSquareGenericTargetCompressed + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME build_tests) + set_target_properties(hipSquareGenericTargetCompressed PROPERTIES COMPILE_FLAGS " -DGENERIC_COMPRESSED -mcode-object-version=6 --offload-compress -w ${OFFLOAD_ARCH_GENERIC_STR}") + add_dependencies(hipSquareGenericTarget hipSquareGenericTargetOnly) + add_dependencies(hipSquareGenericTarget hipSquareGenericTargetCompressed) + add_dependencies(hipSquareGenericTarget hipSquareGenericTargetOnlyCompressed) endif() diff --git a/catch/unit/compiler/hipSquareGenericTarget.cc b/catch/unit/compiler/hipSquareGenericTarget.cc index 01971f63f6..fe937b619f 100644 --- a/catch/unit/compiler/hipSquareGenericTarget.cc +++ b/catch/unit/compiler/hipSquareGenericTarget.cc @@ -34,8 +34,11 @@ static __global__ void vector_square_generic(T* C_d, const T* A_d, size_t N) { } } - -TEST_CASE("Unit_test_generic_target_codeobject") { +#ifdef GENERIC_COMPRESSED +TEST_CASE("Unit_test_generic_target_in_compressed_fatbin") { +#else +TEST_CASE("Unit_test_generic_target_in_regular_fatbin ") { +#endif if (!isGenericTargetSupported()) { fprintf(stderr, "Generic target test is skipped\n"); return; @@ -95,14 +98,24 @@ TEST_CASE("Unit_test_generic_target_codeobject") { } #ifndef NO_GENERIC_TARGET_ONLY_TEST -TEST_CASE("Unit_test_generic_target_only_codeobject") { +#ifdef GENERIC_COMPRESSED +TEST_CASE("Unit_test_generic_target_only_in_compressed_fatbin") { +#ifdef __linux__ + char *cmd = "chmod u+x ./hipSquareGenericTargetOnlyCompressed && ./hipSquareGenericTargetOnlyCompressed"; +#else + char *cmd = "hipSquareGenericTargetOnlyCompressed.exe"; +#endif +#else // else GENERIC_COMPRESSED +TEST_CASE("Unit_test_generic_target_only_in_regular_fatbin ") { #ifdef __linux__ char *cmd = "chmod u+x ./hipSquareGenericTargetOnly && ./hipSquareGenericTargetOnly"; #else char *cmd = "hipSquareGenericTargetOnly.exe"; #endif +#endif // GENERIC_COMPRESSED + printf("Run %s\n", cmd); REQUIRE(std::system(cmd) == 0); printf("PASSED!\n"); } -#endif +#endif // NO_GENERIC_TARGET_ONLY_TEST diff --git a/catch/unit/module/CMakeLists.txt b/catch/unit/module/CMakeLists.txt index 778a89aaf7..3139f221c1 100644 --- a/catch/unit/module/CMakeLists.txt +++ b/catch/unit/module/CMakeLists.txt @@ -113,7 +113,7 @@ 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") +set(OFFLOAD_ARCH_GENERIC_STR "--offload-arch=gfx9-generic --offload-arch=gfx9-4-generic:sramecc+:xnack- --offload-arch=gfx9-4-generic:sramecc-:xnack- --offload-arch=gfx9-4-generic:xnack+ --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 @@ -121,6 +121,12 @@ add_custom_target(copyKernelGenericTarget.code -I${HIP_PATH}/include/ --rocm-path=${ROCM_PATH} -I${CMAKE_CURRENT_SOURCE_DIR}/../../include --rocm-path=${ROCM_PATH}) +add_custom_target(copyKernelGenericTargetCompressed.code + COMMAND ${CMAKE_CXX_COMPILER} -mcode-object-version=6 --offload-compress --genco ${OFFLOAD_ARCH_GENERIC_STR} + ${CMAKE_CURRENT_SOURCE_DIR}/copyKernel.cc + -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/copyKernelGenericTargetCompressed.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 @@ -128,6 +134,7 @@ set_property(GLOBAL APPEND PROPERTY G_INSTALL_CUSTOM_TARGETS ${CMAKE_CURRENT_BINARY_DIR}/addKernel.code ${CMAKE_CURRENT_BINARY_DIR}/copyKernelCompressed.code ${CMAKE_CURRENT_BINARY_DIR}/copyKernelGenericTarget.code + ${CMAKE_CURRENT_BINARY_DIR}/copyKernelGenericTargetCompressed.code ) if(UNIX) @@ -224,6 +231,7 @@ 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) +add_dependencies(build_tests copyKernelGenericTargetCompressed.code) if(UNIX) add_dependencies(build_tests copiousArgKernel.code copiousArgKernel0.code copiousArgKernel1.code copiousArgKernel2.code diff --git a/catch/unit/module/hipExtModuleLaunchKernel.cc b/catch/unit/module/hipExtModuleLaunchKernel.cc index faf4fa3708..3fe64b8bb7 100644 --- a/catch/unit/module/hipExtModuleLaunchKernel.cc +++ b/catch/unit/module/hipExtModuleLaunchKernel.cc @@ -53,6 +53,7 @@ constexpr auto fileName = "copyKernel.code"; constexpr auto kernel_name = "copy_ker"; constexpr auto fileNameCompressed = "copyKernelCompressed.code"; constexpr auto fileNameGenericTarget = "copyKernelGenericTarget.code"; +constexpr auto fileNameGenericTargetCompressed = "copyKernelGenericTargetCompressed.code"; static constexpr auto totalWorkGroups{1024}; static constexpr auto localWorkSize{512}; @@ -192,19 +193,26 @@ TEST_CASE("Unit_hipExtModuleLaunchKernel_UniformWorkGroup") { // Get module and function from module hipModule_t Module; hipFunction_t Function; - SECTION("uncompressed codeobjects") { + SECTION("regular fatbin") { HIP_CHECK(hipModuleLoad(&Module, fileName)); } - SECTION("compressed codeobjects") { + SECTION("compressed fatbin") { HIP_CHECK(hipModuleLoad(&Module, fileNameCompressed)); } - SECTION("generic target codeobjects") { + SECTION("generic target in regular fatbin") { if (!isGenericTargetSupported()) { fprintf(stderr, "Generic target test is skipped\n"); return; } HIP_CHECK(hipModuleLoad(&Module, fileNameGenericTarget)); } + SECTION("generic target in compressed fatbin") { + if (!isGenericTargetSupported()) { + fprintf(stderr, "Generic target test is skipped\n"); + return; + } + HIP_CHECK(hipModuleLoad(&Module, fileNameGenericTargetCompressed)); + } HIP_CHECK(hipModuleGetFunction(&Function, Module, kernel_name)); // Allocate resources int* A = new int[arraylength]; diff --git a/catch/unit/module/hipModuleLoadData.cc b/catch/unit/module/hipModuleLoadData.cc index 3210dd90c4..2eb3af63a5 100644 --- a/catch/unit/module/hipModuleLoadData.cc +++ b/catch/unit/module/hipModuleLoadData.cc @@ -37,7 +37,7 @@ TEST_CASE("Unit_hipModuleLoadData_Positive_Basic") { } #if HT_AMD - SECTION("Load compiled module from file with compressed code objects") { + SECTION("Load compiled module from file with regular target in compressed fatbin") { const auto loaded_module = LoadModuleIntoBuffer("copyKernelCompressed.code"); HIP_CHECK(hipModuleLoadData(&module, loaded_module.data())); REQUIRE(module != nullptr); @@ -47,7 +47,7 @@ TEST_CASE("Unit_hipModuleLoadData_Positive_Basic") { HIP_CHECK(hipModuleUnload(module)); } - SECTION("Load compiled module from file with generic target code objects") { + SECTION("Load compiled module from file with generic target in regular fatbin") { if (!isGenericTargetSupported()) { fprintf(stderr, "Generic target test is skipped\n"); return; @@ -60,6 +60,20 @@ TEST_CASE("Unit_hipModuleLoadData_Positive_Basic") { REQUIRE(kernel != nullptr); HIP_CHECK(hipModuleUnload(module)); } + + SECTION("Load compiled module from file with generic target in compressed fatbin") { + if (!isGenericTargetSupported()) { + fprintf(stderr, "Generic target test is skipped\n"); + return; + } + const auto loaded_module = LoadModuleIntoBuffer("copyKernelGenericTargetCompressed.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") { diff --git a/samples/0_Intro/generic_target/README.md b/samples/0_Intro/generic_target/README.md index 5c0688ea0f..10158ee41b 100644 --- a/samples/0_Intro/generic_target/README.md +++ b/samples/0_Intro/generic_target/README.md @@ -55,7 +55,7 @@ info: copy Device2Host info: check result PASSED: generic targets! -$./saxyGenericTarget +$./saxpyGenericTarget Find generic target gfx11-generic SAXPY test passed ``` \ No newline at end of file