SWDEV-519340 - Enable and fix hipModuleLoad test (#607)

2025-11-12 09:28:49 +01:00
@@ -77,8 +77,6 @@
        "Unit_hipGraphicsResourceGetMappedPointer_Negative_Parameters",
        "Unit_hipGraphicsUnmapResources_Negative_Parameters",
        "Unit_hipGraphicsUnregisterResource_Negative_Parameters",
-        "Note: Test disabled due to defect - EXSWHTEC-151",
-        "Unit_hipModuleLoad_Negative_Load_From_A_File_That_Is_Not_A_Module",
        "SWDEV-442583: Below tests failing in stress test on 12/01/24 ===",
        "Unit_hipLaunchCooperativeKernelMultiDevice_Negative_Parameters",
        "Unit_hipLaunchCooperativeKernelMultiDevice_Negative_MultiKernelSameDevice",
@@ -366,8 +366,6 @@
        "Unit___syncthreads_count_Positive_Basic",
        "Unit___syncthreads_and_Positive_Basic",
        "Unit___syncthreads_or_Positive_Basic",
-        "Note: Test disabled due to defect - EXSWHTEC-151",
-        "Unit_hipModuleLoad_Negative_Load_From_A_File_That_Is_Not_A_Module",
        "Note: Test disabled due to defect - EXSWHTEC-152",
        "Unit_hipModuleUnload_Negative_Module_Is_Nullptr",
        "Note: Following two tests disabled due to defect - EXSWHTEC-153",
@@ -10,12 +10,6 @@
        "Unit_atomicExch_system_Positive_Host_And_GPU - float",
        "Unit_hipModuleUnload_Negative_Double_Unload",
        "=== Below tests are failing PSDB ===",
-        "Unit_hipModuleLoad_Positive_Basic",
-        "Unit_hipModuleLoad_Negative_Load_From_A_File_That_Is_Not_A_Module",
-        "Unit_hipModuleLoadData_Positive_Basic",
-        "Unit_hipModuleLoadData_Negative_Parameters",
-        "Unit_hipModuleLoadDataEx_Positive_Basic",
-        "Unit_hipModuleLoadDataEx_Negative_Parameters",
        "Unit_Assert_Positive_Basic_KernelFail",
        "Unit_hipMemMapArrayAsync_Positive_Basic",
    ]
@@ -73,12 +73,35 @@ add_custom_target(coopKernel.code
                  -I${HIP_PATH}/include/ --hip-path=${HIP_PATH}
                  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include)

+add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/not_a_module.txt
+                   COMMAND ${CMAKE_COMMAND} -E copy
+                   ${CMAKE_CURRENT_SOURCE_DIR}/not_a_module.txt
+                   ${CMAKE_CURRENT_BINARY_DIR}/not_a_module.txt
+                   DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/not_a_module.txt)
+add_custom_target(not_a_module ALL DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/not_a_module.txt)
+add_custom_command(OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/empty_file.txt
+                   COMMAND ${CMAKE_COMMAND} -E copy
+                   ${CMAKE_CURRENT_SOURCE_DIR}/empty_file.txt
+                   ${CMAKE_CURRENT_BINARY_DIR}/empty_file.txt
+                   DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/empty_file.txt)
+add_custom_target(empty_file ALL DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/empty_file.txt)
+
+add_custom_target(empty_module
+                  COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR}
+                  ${CMAKE_CURRENT_SOURCE_DIR}/empty_module.cc
+                  -o ${CMAKE_CURRENT_BINARY_DIR}/empty_module.code
+                  -I${HIP_PATH}/include/ --hip-path=${HIP_PATH}
+                  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include)
+
 set_property(GLOBAL APPEND PROPERTY G_INSTALL_CUSTOM_TARGETS
  ${CMAKE_CURRENT_BINARY_DIR}/get_function_module.code
  ${CMAKE_CURRENT_BINARY_DIR}/launch_kernel_module.code
  ${CMAKE_CURRENT_BINARY_DIR}/get_global_test_module.code
  ${CMAKE_CURRENT_BINARY_DIR}/get_tex_ref_module.code
  ${CMAKE_CURRENT_BINARY_DIR}/coopKernel.code
+  ${CMAKE_CURRENT_BINARY_DIR}/not_a_module.txt
+  ${CMAKE_CURRENT_BINARY_DIR}/empty_file.txt
+  ${CMAKE_CURRENT_BINARY_DIR}/empty_module.code
 )
 # Note to pass arch use format like -DOFFLOAD_ARCH_STR="--offload-arch=gfx900  --offload-arch=gfx906"
 # having space at the start/end of OFFLOAD_ARCH_STR can cause build failures
@@ -96,13 +119,6 @@ if(BUILD_SHARED_LIBS)
      hipGetProcAddressModuleApis.cc)
 endif()

-add_custom_target(empty_module.code
-                  COMMAND ${CMAKE_CXX_COMPILER} --genco ${OFFLOAD_ARCH_STR}
-                  ${CMAKE_CURRENT_SOURCE_DIR}/empty_module.cc
-                  -o ${CMAKE_CURRENT_BINARY_DIR}/../../unit/module/empty_module.code
-                  -I${HIP_PATH}/include/ --hip-path=${HIP_PATH}
-                  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include)
-
 add_custom_target(copyKernel.code
                  COMMAND ${CMAKE_CXX_COMPILER} -mcode-object-version=5 --genco ${OFFLOAD_ARCH_STR}
                  ${CMAKE_CURRENT_SOURCE_DIR}/copyKernel.cc
@@ -160,7 +176,6 @@ add_custom_target(copyKernelGenericTargetCompressed.code
                  -I${HIP_PATH}/include/ --hip-path=${HIP_PATH}
                  -I${CMAKE_CURRENT_SOURCE_DIR}/../../include)
 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
@@ -260,9 +275,11 @@ add_dependencies(ModuleTest get_function_module)
 add_dependencies(ModuleTest launch_kernel_module)
 add_dependencies(ModuleTest get_global_test_module)
 add_dependencies(ModuleTest get_tex_ref_module)
+add_dependencies(ModuleTest not_a_module)
+add_dependencies(ModuleTest empty_file)
+add_dependencies(ModuleTest empty_module)

 if(HIP_PLATFORM MATCHES "amd")
-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 addKernel.spv)
@@ -239,6 +239,7 @@ TEST_CASE("Unit_hipDrvLaunchKernelEx_Functional") {
 *  - HIP_VERSION >= 6.5
 */
 TEST_CASE("Unit_hipDrvLaunchKernelEx_With_Different_Kernels") {
+  CTX_CREATE();
  if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) {
    HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported");
    return;
@@ -300,6 +301,7 @@ TEST_CASE("Unit_hipDrvLaunchKernelEx_With_Different_Kernels") {
  }

  HIP_CHECK(hipModuleUnload(module));
+  CTX_DESTROY();
 }

 /**
@@ -317,6 +319,7 @@ TEST_CASE("Unit_hipDrvLaunchKernelEx_With_Different_Kernels") {
 *  - HIP_VERSION >= 6.5
 */
 TEST_CASE("Unit_hipDrvLaunchKernelEx_With_CooperativeKernelWithArgs") {
+  CTX_CREATE();
  if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) {
    HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported");
    return;
@@ -371,6 +374,7 @@ TEST_CASE("Unit_hipDrvLaunchKernelEx_With_CooperativeKernelWithArgs") {
  HIP_CHECK(hipFree(devMem1));
  HIP_CHECK(hipFree(devMem2));
  HIP_CHECK(hipModuleUnload(module));
+  CTX_DESTROY();
 }

 /**
@@ -391,6 +395,7 @@ TEST_CASE("Unit_hipDrvLaunchKernelEx_With_CooperativeKernelWithArgs") {
 *  - HIP_VERSION >= 6.5
 */
 TEST_CASE("Unit_hipDrvLaunchKernelEx_With_MaxBlockDims") {
+  CTX_CREATE();
  if (!DeviceAttributesSupport(0, hipDeviceAttributeCooperativeLaunch)) {
    HipTest::HIP_SKIP_TEST("CooperativeLaunch not supported");
    return;
@@ -444,6 +449,7 @@ TEST_CASE("Unit_hipDrvLaunchKernelEx_With_MaxBlockDims") {
  }

  HIP_CHECK(hipModuleUnload(module));
+  CTX_DESTROY();
 }
 /**
 * End doxygen group ModuleTest.
@@ -24,23 +24,19 @@ THE SOFTWARE.
 #include <hip_test_common.hh>
 #include <hip/hip_runtime_api.h>

-static hipModule_t GetModule() {
-  HIP_CHECK(hipFree(nullptr));
-  static const auto mg = ModuleGuard::LoadModule("get_function_module.code");
-  return mg.module();
-}
-
 TEST_CASE("Unit_hipModuleGetFunction_Positive_Basic") {
+  auto mg = ModuleGuard::InitModule("get_function_module.code");
  hipFunction_t kernel = nullptr;
-  HIP_CHECK(hipModuleGetFunction(&kernel, GetModule(), "GlobalKernel"));
+  HIP_CHECK(hipModuleGetFunction(&kernel, mg.module(), "GlobalKernel"));
  REQUIRE(kernel != nullptr);
 }

 TEST_CASE("Unit_hipModuleGetFunction_Negative_Parameters") {
+  auto mg = ModuleGuard::InitModule("get_function_module.code");
  hipFunction_t kernel = nullptr;

  SECTION("function == nullptr") {
-    HIP_CHECK_ERROR(hipModuleGetFunction(nullptr, GetModule(), "GlobalKernel"),
+    HIP_CHECK_ERROR(hipModuleGetFunction(nullptr, mg.module(), "GlobalKernel"),
                    hipErrorInvalidValue);
  }

@@ -53,23 +49,23 @@ TEST_CASE("Unit_hipModuleGetFunction_Negative_Parameters") {
 #endif

  SECTION("kname == nullptr") {
-    HIP_CHECK_ERROR(hipModuleGetFunction(&kernel, GetModule(), nullptr), hipErrorInvalidValue);
+    HIP_CHECK_ERROR(hipModuleGetFunction(&kernel, mg.module(), nullptr), hipErrorInvalidValue);
  }

 // Disabled on AMD due to defect - EXSWHTEC-155
 #if HT_NVIDIA
  SECTION("kname == empty string") {
-    HIP_CHECK_ERROR(hipModuleGetFunction(&kernel, GetModule(), ""), hipErrorInvalidValue);
+    HIP_CHECK_ERROR(hipModuleGetFunction(&kernel, mg.module(), ""), hipErrorInvalidValue);
  }
 #endif

  SECTION("kname == non existent kernel") {
-    HIP_CHECK_ERROR(hipModuleGetFunction(&kernel, GetModule(), "NonExistentKernel"),
+    HIP_CHECK_ERROR(hipModuleGetFunction(&kernel, mg.module(), "NonExistentKernel"),
                    hipErrorNotFound);
  }

  SECTION("kname == __device__ kernel") {
-    HIP_CHECK_ERROR(hipModuleGetFunction(&kernel, GetModule(), "DeviceKernel"), hipErrorNotFound);
+    HIP_CHECK_ERROR(hipModuleGetFunction(&kernel, mg.module(), "DeviceKernel"), hipErrorNotFound);
  }
 }

@@ -83,9 +79,9 @@ TEST_CASE("Unit_hipModuleGetFunction_DiffDevice") {
    return;
  }

+  auto mg = ModuleGuard::InitModule("get_function_module.code");
  hipFunction_t kernel = nullptr;
-  auto module = GetModule();
  HIP_CHECK(hipSetDevice(1));
-  HIP_CHECK(hipModuleGetFunction(&kernel, module, "GlobalKernel"));
+  HIP_CHECK(hipModuleGetFunction(&kernel, mg.module(), "GlobalKernel"));
  REQUIRE(kernel != nullptr);
 }
@@ -48,6 +48,7 @@ THE SOFTWARE.
 * - HIP_VERSION >= 7.1
 */
 TEST_CASE("Unit_hipModuleGetFunctionCount_Functional") {
+  CTX_CREATE();
  hipModule_t moduleSingleArch, moduleEmpty, doubleKernelModule, rtcModule;
  unsigned int count = 0;
  SECTION("Single arch, Single global function") {
@@ -89,6 +90,7 @@ TEST_CASE("Unit_hipModuleGetFunctionCount_Functional") {
    REQUIRE(count == 1);
    HIP_CHECK(hipModuleUnload(rtcModule));
  }
+  CTX_DESTROY();
 }
 /**
 * Test Description
@@ -127,8 +127,10 @@ TEST_CASE("Unit_hipModuleGetGlobal_Negative_Hmod_Is_Nullptr") {
  hipDeviceptr_t global = 0;
  size_t global_size = 0;

+  CTX_CREATE();
  HIP_CHECK_ERROR(hipModuleGetGlobal(&global, &global_size, nullptr, "int_var"),
                  hipErrorInvalidResourceHandle);
+  CTX_DESTROY();
 }

 TEST_CASE("Unit_hipModuleGetGlobal_Negative_Name_Is_Empty_String") {
@@ -61,7 +61,9 @@ TEST_CASE("Unit_hipModuleGetTexRef_Negative_Hmod_Is_Nullptr") {
  CHECK_IMAGE_SUPPORT
  hipTexRef tex_ref = nullptr;

+  CTX_CREATE();
  HIP_CHECK_ERROR(hipModuleGetTexRef(&tex_ref, nullptr, "tex"), hipErrorInvalidResourceHandle);
+  CTX_DESTROY();
 }

 TEST_CASE("Unit_hipModuleGetTexRef_Negative_Name_Is_Empty_String") {
@@ -56,4 +56,5 @@ TEST_CASE("Unit_hipModuleLoad_Negative_Load_From_A_File_That_Is_Not_A_Module") {
  hipModule_t module;

  HIP_CHECK_ERROR(hipModuleLoad(&module, "not_a_module.txt"), hipErrorInvalidImage);
-}
+  HIP_CHECK_ERROR(hipModuleLoad(&module, "empty_file.txt"), hipErrorInvalidImage);
+}
@@ -40,6 +40,7 @@ THE SOFTWARE.
 * - HIP_VERSION >= 7.1
 */
 TEST_CASE("Unit_hipModuleLoadFatBinary_NegativeTsts") {
+  CTX_CREATE();
  hipModule_t Module;
  SECTION("fatCubin as nullptr") {
    HIP_CHECK_ERROR(hipModuleLoadFatBinary(&Module, nullptr),
@@ -51,6 +52,7 @@ TEST_CASE("Unit_hipModuleLoadFatBinary_NegativeTsts") {
    REQUIRE(Module != nullptr);
    HIP_CHECK(hipModuleUnload(Module));
  }
+  CTX_DESTROY();
 }
 #if HT_AMD
 void loadKernelData(hipFunction_t kernel) {
@@ -32,7 +32,11 @@ TEST_CASE("Unit_hipModuleUnload_Negative_Double_Unload") {
  hipModule_t module = nullptr;
  HIP_CHECK(hipModuleLoad(&module, "empty_module.code"));
  HIP_CHECK(hipModuleUnload(module));
+#if HT_AMD
  HIP_CHECK_ERROR(hipModuleUnload(module), hipErrorNotFound);
+#else
+  HIP_CHECK_ERROR(hipModuleUnload(module), hipErrorInvalidResourceHandle);
+#endif
 }
 /**
 * @addtogroup hipModuleUnload
@@ -54,9 +58,11 @@ TEST_CASE("Unit_hipModuleUnload_Negative_Double_Unload") {
 * - HIP_VERSION >= 5.6
 */
 TEST_CASE("Unit_hipModuleLoad_basic") {
+  CTX_CREATE();
  constexpr auto fileName = "vcpy_kernel.code";
  hipModule_t module;
  HIP_CHECK(hipModuleLoad(&module, fileName));
  REQUIRE(module != nullptr);
  HIP_CHECK(hipModuleUnload(module));
+  CTX_DESTROY();
 }
@@ -54,12 +54,12 @@ ModuleGuard ModuleGuard::LoadModuleDataRTC(const char* code) {

 // Load module into buffer instead of mapping file to avoid platform specific mechanisms
 std::vector<char> LoadModuleIntoBuffer(const char* path_string) {
-  fs::path p(path_string);
-  const auto file_size = fs::file_size(p);
-  std::ifstream f(p, std::ios::binary | std::ios::in);
-  REQUIRE(f);
-  std::vector<char> empty_module(file_size);
-  REQUIRE(f.read(empty_module.data(), file_size));
+  std::ifstream file_stream(path_string, std::ios::binary | std::ios::in);
+  REQUIRE(file_stream);
+  std::vector<char> empty_module((std::istreambuf_iterator<char>(file_stream)),
+                                 std::istreambuf_iterator<char>());
+  file_stream.close();
+  empty_module.push_back('\0');
  return empty_module;
 }

@@ -73,4 +73,4 @@ std::vector<char> CreateRTCCharArray(const char* src) {
  HIPRTC_CHECK(hiprtcGetCode(prog, code.data()));
  HIPRTC_CHECK(hiprtcDestroyProgram(&prog));
  return code;
-}
+}
@@ -0,0 +1 @@
+This is not a module!