SWDEV-519340 - Enable and fix hipModuleLoad test

projects/hip-tests/catch/hipTestMain/config/config_amd_linux

@@ -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",

projects/hip-tests/catch/hipTestMain/config/config_amd_windows

@@ -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",

projects/hip-tests/catch/hipTestMain/config/config_nvidia_linux.json

@@ -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",
     ]

projects/hip-tests/catch/unit/module/CMakeLists.txt

@@ -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)

projects/hip-tests/catch/unit/module/hipDrvLaunchKernelEx.cc

@@ -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.

projects/hip-tests/catch/unit/module/hipModuleGetFunction.cc

@@ -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);
 }

projects/hip-tests/catch/unit/module/hipModuleGetFunctionCount.cc

@@ -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

projects/hip-tests/catch/unit/module/hipModuleGetGlobal.cc

@@ -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") {

projects/hip-tests/catch/unit/module/hipModuleGetTexRef.cc

@@ -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") {

projects/hip-tests/catch/unit/module/hipModuleLoad.cc

@@ -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);
+}

projects/hip-tests/catch/unit/module/hipModuleLoadFatBinary.cc

@@ -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) {

projects/hip-tests/catch/unit/module/hipModuleUnload.cc

@@ -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();
 }

projects/hip-tests/catch/unit/module/hip_module_common.cc

@@ -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;
-}
+}

projects/hip-tests/catch/unit/module/not_a_module.txt

@@ -0,0 +1 @@
+This is not a module!