SWDEV-508863 - Add generic target in compressed fatbin test (#28)

Change-Id: Iae148c3c998b18247624938512918dbb31bc462d

[ROCm/hip-tests commit: 5a33b33ec8]

projects/hip-tests/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 ===",

projects/hip-tests/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()

projects/hip-tests/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

projects/hip-tests/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

projects/hip-tests/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];

projects/hip-tests/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") {

projects/hip-tests/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
 ```