SWDEV-519340 - Enable and fix hipModuleLoad test

catch/hipTestMain/config/config_amd_linux

@@ -168,10 +168,6 @@
         "Unit_hipGraphicsResourceGetMappedPointer_Negative_Parameters",
         "Unit_hipGraphicsUnmapResources_Negative_Parameters",
         "Unit_hipGraphicsUnregisterResource_Negative_Parameters",
-        "SWDEV-443760: This test fails when device memory is used for kernel args",
-        "=== Below tests fail in external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/356 ===",
-        "Note: Test disabled due to defect - EXSWHTEC-151",
-        "Unit_hipModuleLoad_Negative_Load_From_A_File_That_Is_Not_A_Module",
         "Note: Following two tests disabled due to defect - EXSWHTEC-153",
         "Unit_hipModuleLoadData_Negative_Image_Is_An_Empty_String",
         "Unit_hipModuleLoadDataEx_Negative_Image_Is_An_Empty_String",

catch/hipTestMain/config/config_amd_windows

@@ -644,8 +644,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",

catch/hipTestMain/config/config_nvidia_linux.json

@@ -70,12 +70,6 @@
         "Unit_hipMemcpy2D_Positive_Synchronization_Behavior",
         "Unit_hipDrvMemcpy3D_Positive_Synchronization_Behavior",
         "Unit_hipFreeMipmappedArray_Negative_DoubleFree",
-        "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",
         "Performance_hipMemsetD16",
         "Performance_hipMemsetD16Async",
         "Performance_hipMemsetD32",

catch/unit/module/hipExtModuleLaunchKernel.cc

@@ -265,7 +265,7 @@ TEST_CASE("Unit_hipExtModuleLaunchKernel_Positive_Parameters") {
   SECTION("Pass only start event") {
     hipEvent_t start_event = nullptr;
     HIP_CHECK(hipEventCreate(&start_event));
-    const auto kernel = GetKernel(mg.module(), "NOPKernel");
+    const auto kernel = GetKernel(GetModule(), "NOPKernel");
     HIP_CHECK(hipExtModuleLaunchKernel(kernel, 1, 1, 1, 1, 1, 1, 0, nullptr,
                                        nullptr, nullptr,
                                        start_event, nullptr));
@@ -277,7 +277,7 @@ TEST_CASE("Unit_hipExtModuleLaunchKernel_Positive_Parameters") {
   SECTION("Pass only stop event") {
     hipEvent_t stop_event = nullptr;
     HIP_CHECK(hipEventCreate(&stop_event));
-    const auto kernel = GetKernel(mg.module(), "NOPKernel");
+    const auto kernel = GetKernel(GetModule(), "NOPKernel");
     HIP_CHECK(hipExtModuleLaunchKernel(kernel, 1, 1, 1, 1, 1, 1, 0, nullptr,
                                        nullptr, nullptr,
                                        nullptr, stop_event));

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

catch/unit/module/hipModuleGetTexRef.cc

@@ -59,7 +59,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") {
@@ -68,4 +70,4 @@ TEST_CASE("Unit_hipModuleGetTexRef_Negative_Name_Is_Empty_String") {
   hipTexRef tex_ref = nullptr;
 
   HIP_CHECK_ERROR(hipModuleGetTexRef(&tex_ref, module, ""), hipErrorInvalidValue);
-}
+}

catch/unit/module/hipModuleLaunchCooperativeKernel.cc

@@ -56,13 +56,13 @@ TEST_CASE("Unit_hipModuleLaunchCooperativeKernel_Positive_Basic") {
   }
 
   SECTION("Cooperative kernel with no arguments") {
-    hipFunction_t f = GetKernel(mg.module(), "CoopKernel");
+    hipFunction_t f = GetKernel(GetModule(), "CoopKernel");
     HIP_CHECK(hipModuleLaunchCooperativeKernel(f, 2, 2, 1, 1, 1, 1, 0, nullptr, nullptr));
     HIP_CHECK(hipDeviceSynchronize());
   }
 
   SECTION("Kernel with arguments using kernelParams") {
-    hipFunction_t f = GetKernel(mg.module(), "Kernel42");
+    hipFunction_t f = GetKernel(GetModule(), "Kernel42");
 
     LinearAllocGuard<int> result_dev(LinearAllocs::hipMalloc, sizeof(int));
     HIP_CHECK(hipMemset(result_dev.ptr(), 0, sizeof(*result_dev.ptr())));
@@ -94,7 +94,7 @@ TEST_CASE("Unit_hipModuleLaunchCooperativeKernel_Positive_Parameters") {
     return;
   }
 
-  hipFunction_t f = GetKernel(mg.module(), "NOPKernel");
+  hipFunction_t f = GetKernel(GetModule(), "NOPKernel");
 
   SECTION("blockDim.x == maxBlockDimX") {
     const unsigned int x = GetDeviceAttribute(hipDeviceAttributeMaxBlockDimX, 0);
@@ -129,7 +129,7 @@ TEST_CASE("Unit_hipModuleLaunchCooperativeKernel_Negative_Parameters") {
     return;
   }
 
-  hipFunction_t f = GetKernel(mg.module(), "NOPKernel");
+  hipFunction_t f = GetKernel(GetModule(), "NOPKernel");
 
   SECTION("f == nullptr") {
     HIP_CHECK_ERROR(

catch/unit/module/hipModuleLaunchCooperativeKernelMultiDevice.cc

@@ -163,7 +163,7 @@ TEST_CASE("Unit_hipModuleLaunchCooperativeKernelMultiDevice_Negative_Parameters"
 
   if (device_count > 1) {
     SECTION("launchParamsList.func doesn't match across all devices") {
-      params_list[1].function = GetKernel(mg.module(), "NOPKernel");
+      params_list[1].function = GetKernel(GetModule(), "NOPKernel");
 #if HT_AMD
       HIP_CHECK_ERROR(
           hipModuleLaunchCooperativeKernelMultiDevice(params_list.data(), device_count, 0u),

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

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") {
-    constexpr auto fileName = "vcpy_kernel.code";
-    hipModule_t module;
-    HIP_CHECK(hipModuleLoad(&module, fileName));
-    REQUIRE(module != nullptr);
-    HIP_CHECK(hipModuleUnload(module));
+  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();
 }

catch/unit/module/hip_module_common.cc

@@ -49,12 +49,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;
 }
 
@@ -68,4 +68,4 @@ std::vector<char> CreateRTCCharArray(const char* src) {
   HIPRTC_CHECK(hiprtcGetCode(prog, code.data()));
   HIPRTC_CHECK(hiprtcDestroyProgram(&prog));
   return code;
-}
+}

catch/unit/module/hip_module_launch_kernel_common.hh

@@ -27,26 +27,25 @@ THE SOFTWARE.
 #include <resource_guards.hh>
 #include <utils.hh>
 
-inline ModuleGuard InitModule() {
+static hipModule_t GetModule() {
   HIP_CHECK(hipFree(nullptr));
-  return ModuleGuard::LoadModule("launch_kernel_module.code");
+  static const auto mg = ModuleGuard::LoadModule("launch_kernel_module.code");
+  return mg.module();
 }
 
-inline ModuleGuard mg{InitModule()};
-
 using ExtModuleLaunchKernelSig = hipError_t(hipFunction_t, uint32_t, uint32_t, uint32_t, uint32_t,
                                             uint32_t, uint32_t, size_t, hipStream_t, void**, void**,
                                             hipEvent_t, hipEvent_t, uint32_t);
 
 template <ExtModuleLaunchKernelSig* func> void ModuleLaunchKernelPositiveBasic() {
   SECTION("Kernel with no arguments") {
-    hipFunction_t f = GetKernel(mg.module(), "NOPKernel");
+    hipFunction_t f = GetKernel(GetModule(), "NOPKernel");
     HIP_CHECK(func(f, 1, 1, 1, 1, 1, 1, 0, nullptr, nullptr, nullptr, nullptr, nullptr, 0u));
     HIP_CHECK(hipDeviceSynchronize());
   }
 
   SECTION("Kernel with arguments using kernelParams") {
-    hipFunction_t f = GetKernel(mg.module(), "Kernel42");
+    hipFunction_t f = GetKernel(GetModule(), "Kernel42");
     LinearAllocGuard<int> result_dev(LinearAllocs::hipMalloc, sizeof(int));
     HIP_CHECK(hipMemset(result_dev.ptr(), 0, sizeof(*result_dev.ptr())));
     int* result_ptr = result_dev.ptr();
@@ -58,7 +57,7 @@ template <ExtModuleLaunchKernelSig* func> void ModuleLaunchKernelPositiveBasic()
   }
 
   SECTION("Kernel with arguments using extra") {
-    hipFunction_t f = GetKernel(mg.module(), "Kernel42");
+    hipFunction_t f = GetKernel(GetModule(), "Kernel42");
     LinearAllocGuard<int> result_dev(LinearAllocs::hipMalloc, sizeof(int));
     HIP_CHECK(hipMemset(result_dev.ptr(), 0, sizeof(*result_dev.ptr())));
     int* result_ptr = result_dev.ptr();
@@ -81,7 +80,7 @@ template <ExtModuleLaunchKernelSig* func> void ModuleLaunchKernelPositiveParamet
   const auto LaunchNOPKernel = [=](unsigned int gridDimX, unsigned int gridDimY,
                                    unsigned int gridDimZ, unsigned int blockDimX,
                                    unsigned int blockDimY, unsigned int blockDimZ) {
-    hipFunction_t f = GetKernel(mg.module(), "NOPKernel");
+    hipFunction_t f = GetKernel(GetModule(), "NOPKernel");
     HIP_CHECK(func(f, gridDimX, gridDimY, gridDimZ, blockDimX, blockDimY, blockDimZ, 0, nullptr,
                    nullptr, nullptr, nullptr, nullptr, 0u));
     HIP_CHECK(hipDeviceSynchronize());
@@ -120,7 +119,7 @@ template <ExtModuleLaunchKernelSig* func> void ModuleLaunchKernelPositiveParamet
 
 template <ExtModuleLaunchKernelSig* func> void ModuleLaunchKernelNegativeParameters(
                                                            bool extLaunch = false) {
-  hipFunction_t f = GetKernel(mg.module(), "NOPKernel");
+  hipFunction_t f = GetKernel(GetModule(), "NOPKernel");
   hipError_t expectedErrorLaunchParam = (extLaunch == true) ? hipErrorInvalidConfiguration
                                                              : hipErrorInvalidValue;
   hipError_t expectedErrorOverCapacityGridDim = (extLaunch == true) ? hipSuccess
@@ -213,7 +212,7 @@ template <ExtModuleLaunchKernelSig* func> void ModuleLaunchKernelNegativeParamet
   }
 
   SECTION("Passing kernel_args and extra simultaneously") {
-    hipFunction_t f = GetKernel(mg.module(), "Kernel42");
+    hipFunction_t f = GetKernel(GetModule(), "Kernel42");
     LinearAllocGuard<int> result_dev(LinearAllocs::hipMalloc, sizeof(int));
     int* result_ptr = result_dev.ptr();
     size_t size = sizeof(result_ptr);
@@ -230,7 +229,7 @@ template <ExtModuleLaunchKernelSig* func> void ModuleLaunchKernelNegativeParamet
   }
 
   SECTION("Invalid extra") {
-    hipFunction_t f = GetKernel(mg.module(), "Kernel42");
+    hipFunction_t f = GetKernel(GetModule(), "Kernel42");
     void* extra[0] = {};
     HIP_CHECK_ERROR(func(f, 1, 1, 1, 1, 1, 1, 0, nullptr, nullptr, extra, nullptr, nullptr, 0u),
                     hipErrorInvalidValue);

catch/unit/module/not_a_module.txt

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