SWDEV-523137 - Enable and fix failing tests on NV (#602)

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

@@ -331,7 +331,6 @@
         "Unit_hipMallocMipmappedArray_MultiThread",
         "Unit_hipMallocMipmappedArray_Negative_InvalidFlags",
         "Unit_hipGetMipmappedArrayLevel_Negative",
-        "Unit_hipFreeMipmappedArray_Negative_DoubleFree",
         "Unit_hipFreeMipmappedArrayMultiTArray - int",
         "Unit_Thread_Block_Tile_Dynamic_Getters_Positive_Basic",
         "Performance_hipMemcpy2D_HostToHost",

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

@@ -29,7 +29,6 @@
         "Unit_hipMemcpyParam2D_multiDevice-D2D - long double",
         "Unit_hipMemcpyParam2DAsync_multiDevice-StreamOnDiffDevice - char",
         "Unit_hipMemsetFunctional_ZeroValue_hipMemsetD16",
-        "Unit_hipStreamAttachMemAsync_Negative_Parameters",
         "hipStreamPerThread_CoopLaunch",
         "hipCGMultiGridGroupType",
         "Grid_Group_Getters_Positive_Basic",
@@ -52,20 +51,12 @@
         "Unit_Device_Complex_hipCfma_Negative_Parameters_RTC",
         "Unit_Device_make_Complex_Negative_Parameters_RTC",
         "Unit_Device_Complex_Cast_Negative_Parameters_RTC",
-        "=== Below 2 tests are disabled due to defect EXSWHTEC-342 ===",
-        "Unit_hipDeviceSetLimit_Negative_Parameters",
-        "Unit_hipDeviceGetLimit_Negative_Parameters",
         "=== Below tests are failing PSDB ===",
-        "Unit_hipGraphMem_Alloc_Free_NodeGetParams_Functional_3",
-        "Unit_hipMemPoolSetAccess_Negative_Parameters",
-        "Unit_hipMallocMipmappedArray_Negative_NumLevels",
-        "Unit_hipFreeMipmappedArray_Negative_Nullptr",
         "Unit_hipFreeMipmappedArrayMultiTArray - int",
         "Unit_hipFreeMipmappedArray_Negative_Parameters",
         "Unit_hipCreateSurfaceObject_Negative_Parameters",
         "Unit_hipDestroySurfaceObject_Negative_Parameters",
         "Unit_hipMemcpy2D_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",
@@ -223,11 +214,7 @@
         "Unit_hipDeviceSetLimit_Negative_MallocHeapSize",
         "=== Disabling tests which no longer behave the same on nvidia platform ===",
         "Unit_hipGraphInstantiateWithParams_Negative",
-        "Unit_hipGraphAddChildGraphNode_OrgGraphAsChildGraph",
-        "Unit_hipDeviceSynchronize_Positive_Nullstream",
-        "Unit_hipDeviceSynchronize_Functional",
         "Unit_hipDeviceReset_Positive_Basic",
         "Unit_hipDeviceReset_Positive_Threaded",
-        "Unit_hipModuleGetTexRef_Positive_Basic"
     ]
 }

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

@@ -26,7 +26,6 @@
         "Unit_hipMemcpy3D_Positive_Synchronization_Behavior",
         "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",

projects/hip-tests/catch/unit/device/hipDeviceSetGetLimit.cc

@@ -152,7 +152,11 @@ TEST_CASE("Unit_hipDeviceSetLimit_Negative_MallocHeapSize") {
  *  - HIP_VERSION >= 5.3
  */
 TEST_CASE("Unit_hipDeviceSetLimit_Negative_Parameters") {
+#if HT_AMD
   HIP_CHECK_ERROR(hipDeviceSetLimit(static_cast<hipLimit_t>(-1), 1024), hipErrorUnsupportedLimit);
+#else
+  HIP_CHECK_ERROR(hipDeviceSetLimit(static_cast<hipLimit_t>(-1), 1024), hipErrorInvalidValue);
+#endif
 }
 
 /**
@@ -186,7 +190,11 @@ TEST_CASE("Unit_hipDeviceGetLimit_Negative_Parameters") {
 
   SECTION("unsupported limit") {
     size_t val;
+#if HT_AMD
     HIP_CHECK_ERROR(hipDeviceGetLimit(&val, static_cast<hipLimit_t>(-1)), hipErrorUnsupportedLimit);
+#else
+    HIP_CHECK_ERROR(hipDeviceGetLimit(&val, static_cast<hipLimit_t>(-1)), hipErrorInvalidValue);
+#endif
   }
 }
 

projects/hip-tests/catch/unit/device/hipDeviceSynchronize.cc

@@ -89,19 +89,15 @@ TEST_CASE("Unit_hipDeviceSynchronize_Positive_Nullstream") {
   INFO("Current device: " << device);
 
   int *A_h = nullptr, *A_d = nullptr;
-  HipTest::BlockingContext b_context{nullptr};
   HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), _SIZE, hipHostMallocDefault));
   A_h[0] = 1;
   HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&A_d), _SIZE));
 
   HIP_CHECK(hipMemcpyAsync(A_d, A_h, _SIZE, hipMemcpyHostToDevice, NULL));
-  b_context.block_stream();
-  REQUIRE(b_context.is_blocked());
   hipLaunchKernelGGL(HIP_KERNEL_NAME(Iter), dim3(1), dim3(1), 0, NULL, A_d, 1 << 30);
   HIP_CHECK(hipMemcpyAsync(A_h, A_d, _SIZE, hipMemcpyDeviceToHost, NULL));
 
   REQUIRE(1 << 30 != A_h[0] - 1);
-  b_context.unblock_stream();
   HIP_CHECK(hipDeviceSynchronize());
   REQUIRE(1 << 30 == A_h[0] - 1);
   HIP_CHECK(hipHostFree(A_h));
@@ -124,22 +120,17 @@ TEST_CASE("Unit_hipDeviceSynchronize_Functional") {
   int* A[NUM_STREAMS];
   int* Ad[NUM_STREAMS];
   hipStream_t stream[NUM_STREAMS];
-  std::vector<HipTest::BlockingContext> b_context;
-  b_context.reserve(NUM_STREAMS);
 
   for (int i = 0; i < NUM_STREAMS; i++) {
     HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A[i]), _SIZE, hipHostMallocDefault));
     A[i][0] = 1;
     HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&Ad[i]), _SIZE));
     HIP_CHECK(hipStreamCreate(&stream[i]));
-    b_context.emplace_back(HipTest::BlockingContext(stream[i]));
   }
   for (int i = 0; i < NUM_STREAMS; i++) {
     HIP_CHECK(hipMemcpyAsync(Ad[i], A[i], _SIZE, hipMemcpyHostToDevice, stream[i]));
   }
   for (int i = 0; i < NUM_STREAMS; i++) {
-    b_context[i].block_stream();
-    REQUIRE(b_context[i].is_blocked());
     hipLaunchKernelGGL(HIP_KERNEL_NAME(Iter), dim3(1), dim3(1), 0, stream[i], Ad[i], NUM_ITERS);
   }
   for (int i = 0; i < NUM_STREAMS; i++) {
@@ -154,9 +145,6 @@ TEST_CASE("Unit_hipDeviceSynchronize_Functional") {
   // fail, ie if HIP_LAUNCH_BLOCKING=true.
 
   REQUIRE(NUM_ITERS != A[NUM_STREAMS - 1][0] - 1);
-  for (int i = 0; i < NUM_STREAMS; i++) {
-    b_context[i].unblock_stream();
-  }
   HIP_CHECK(hipDeviceSynchronize());
   REQUIRE(NUM_ITERS == A[NUM_STREAMS - 1][0] - 1);
   for (int i = 0; i < NUM_STREAMS; i++) {

projects/hip-tests/catch/unit/device/hipGetSetDeviceFlags.cc

@@ -233,15 +233,11 @@ TEST_CASE("Unit_hipGetDeviceFlags_Positive_Context") {
  *  - HIP_VERSION >= 5.2
  */
 TEST_CASE("Unit_hipGetSetDeviceFlags_InvalidFlag") {
-#if HT_AMD
-  HipTest::HIP_SKIP_TEST("EXSWCPHIPT-115");
-  return;
-#endif
-  const unsigned int invalidFlag = GENERATE(0b011,  // schedule flags should not overlap
-                                            0b101,  // schedule flags should not overlap
-                                            0b110,  // schedule flags should not overlap
-                                            0b111,  // schedule flags should not overlap
-                                            // 0b100000,  // out of bounds is no longer invalid
+  const unsigned int invalidFlag = GENERATE(0xb011,     // schedule flags should not overlap
+                                            0xb101,     // schedule flags should not overlap
+                                            0xb110,     // schedule flags should not overlap
+                                            0xb111,     // schedule flags should not overlap
+                                            0xb100000,  // out of bounds
                                             0xFFFF);
   CAPTURE(invalidFlag);
   HIP_CHECK_ERROR(hipSetDeviceFlags(invalidFlag), hipErrorInvalidValue);

projects/hip-tests/catch/unit/graph/hipGraphAddChildGraphNode.cc

@@ -119,13 +119,14 @@ and verify the number of the nodes in the original graph
 TEST_CASE("Unit_hipGraphAddChildGraphNode_OrgGraphAsChildGraph") {
   constexpr size_t N = 1024;
   constexpr size_t Nbytes = N * sizeof(int);
-  hipGraph_t graph;
+  hipGraph_t graph, childGraph;
   hipGraphExec_t graphExec;
   int *A_d{nullptr}, *B_d{nullptr};
   int *A_h{nullptr}, *B_h{nullptr};
   HipTest::initArrays<int>(&A_d, &B_d, nullptr, &A_h, &B_h, nullptr, N, false);
 
   HIP_CHECK(hipGraphCreate(&graph, 0));
+  HIP_CHECK(hipGraphCreate(&childGraph, 0));
   hipGraphNode_t memcpyH2D_A, memcpyH2D_B, childGraphNode1;
   size_t numNodes;
   hipStream_t streamForGraph;
@@ -134,7 +135,7 @@ TEST_CASE("Unit_hipGraphAddChildGraphNode_OrgGraphAsChildGraph") {
                                     hipMemcpyHostToDevice));
   HIP_CHECK(hipGraphAddMemcpyNode1D(&memcpyH2D_A, graph, nullptr, 0, A_h, B_d, Nbytes,
                                     hipMemcpyDeviceToHost));
-  HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode1, graph, nullptr, 0, graph));
+  HIP_CHECK(hipGraphAddChildGraphNode(&childGraphNode1, graph, nullptr, 0, childGraph));
 
   HIP_CHECK(hipGraphAddDependencies(graph, &memcpyH2D_B, &memcpyH2D_A, 1));
 

projects/hip-tests/catch/unit/graph/hipGraphInstantiateWithParams.cc

@@ -57,6 +57,7 @@ TEST_CASE("Unit_hipGraphInstantiateWithParams_Negative") {
     hipGraphInstantiateParams params;
     HIP_CHECK(hipGraphCreate(&graph, 0));
     REQUIRE(hipGraphInstantiateWithParams(nullptr, graph, &params) == hipErrorInvalidValue);
+    HIP_CHECK(hipGraphDestroy(graph));
   }
 
   SECTION("Passing nullptr to graph") {
@@ -70,6 +71,7 @@ TEST_CASE("Unit_hipGraphInstantiateWithParams_Negative") {
     HIP_CHECK(hipGraphCreate(&graph, 0));
     hipGraphExec_t graphExec;
     REQUIRE(hipGraphInstantiateWithParams(&graphExec, graph, nullptr) == hipErrorInvalidValue);
+    HIP_CHECK(hipGraphDestroy(graph));
   }
 
   SECTION("Passing invalid flag") {
@@ -80,6 +82,7 @@ TEST_CASE("Unit_hipGraphInstantiateWithParams_Negative") {
     params.flags = 10;
     REQUIRE(hipGraphInstantiateWithParams(&graphExec, graph, &params) == hipErrorInvalidValue);
     REQUIRE(params.result_out == hipGraphInstantiateError);
+    HIP_CHECK(hipGraphDestroy(graph));
   }
 }
 

projects/hip-tests/catch/unit/graph/hipGraphMemAllocNodeGetParams.cc

@@ -51,12 +51,25 @@ THE SOFTWARE.
  *  - HIP_VERSION >= 6.0
  */
 
-static bool validateAllocParam(hipMemAllocNodeParams in, hipMemAllocNodeParams out) {
+static bool validateAllocParam(hipMemAllocNodeParams in, hipMemAllocNodeParams out,
+                               bool accessDesc = false) {
   if (in.bytesize != out.bytesize) return false;
   if (in.poolProps.allocType != out.poolProps.allocType) return false;
   if (in.poolProps.location.id != out.poolProps.location.id) return false;
   if (in.poolProps.location.type != out.poolProps.location.type) return false;
 
+  if (accessDesc) {
+    if (in.accessDescs->location.type != out.accessDescs->location.type) {
+      return false;
+    }
+    if (in.accessDescs->location.id != out.accessDescs->location.id) {
+      return false;
+    }
+    if (in.accessDescs->flags != out.accessDescs->flags) {
+      return false;
+    }
+  }
+
   return true;
 }
 
@@ -299,7 +312,7 @@ TEST_CASE("Unit_hipGraphMem_Alloc_Free_NodeGetParams_Functional_3") {
 
   hipMemAllocNodeParams get_alloc_params;
   HIP_CHECK(hipGraphMemAllocNodeGetParams(alloc_node, &get_alloc_params));
-  REQUIRE(memcmp(&alloc_param, &get_alloc_params, sizeof(hipMemAllocNodeParams)) == 0);
+  REQUIRE(validateAllocParam(alloc_param, get_alloc_params, true) == true);
 
   constexpr int fill_value = 11;
   hipGraphNode_t memset_node;

projects/hip-tests/catch/unit/memory/hipFreeMipmappedArray.cc

@@ -85,31 +85,11 @@ TEMPLATE_TEST_CASE("Unit_hipFreeMipmappedArrayImplicitSyncArray", "", char, floa
 }
 
 TEST_CASE("Unit_hipFreeMipmappedArray_Negative_Nullptr") {
-  HIP_CHECK_ERROR(hipFreeMipmappedArray(nullptr), hipErrorInvalidValue);
-}
-
-TEST_CASE("Unit_hipFreeMipmappedArray_Negative_DoubleFree") {
-  hipMipmappedArray_t arrayPtr{};
-  hipExtent extent{};
-  hipChannelFormatDesc desc = hipCreateChannelDesc<char>();
-
 #if HT_AMD
-  const unsigned int flags = hipArrayDefault;
+  HIP_CHECK_ERROR(hipFreeMipmappedArray(nullptr), hipErrorInvalidValue);
 #else
-  const unsigned int flags = GENERATE(hipArrayDefault, hipArraySurfaceLoadStore);
+  HIP_CHECK(hipFreeMipmappedArray(nullptr));
 #endif
-
-  extent.width = GENERATE(64, 512, 1024);
-  extent.height = GENERATE(64, 512, 1024);
-  extent.depth = GENERATE(0, 64, 512, 1024);
-
-  const unsigned int numLevels = GENERATE(1, 5, 7);
-
-  HIP_CHECK_IGNORED_RETURN(hipMallocMipmappedArray(&arrayPtr, &desc, extent, numLevels, flags),
-                           hipErrorNotSupported);
-
-  HIP_CHECK(hipFreeMipmappedArray(arrayPtr));
-  HIP_CHECK_ERROR(hipFreeMipmappedArray(arrayPtr), hipErrorContextIsDestroyed);
 }
 
 TEMPLATE_TEST_CASE("Unit_hipFreeMipmappedArrayMultiTArray", "", char, int) {

projects/hip-tests/catch/unit/memory/hipMallocMipmappedArray.cc

@@ -390,10 +390,15 @@ TEST_CASE("Unit_hipMallocMipmappedArray_Negative_NumLevels") {
   unsigned int numLevels = floor(log2(size)) + 2;
   hipChannelFormatDesc desc = hipCreateChannelDesc<float>();
 
-  const auto flag = GENERATE(from_range(std::begin(validFlags), std::end(validFlags)));
+  const auto flag = hipArrayDefault;
+#if HT_AMD
   HIP_CHECK_ERRORS(
       hipMallocMipmappedArray(&array, &desc, makeMipmappedExtent(flag, size), numLevels, flag),
       hipErrorInvalidValue, hipErrorNotSupported);
+#else
+  HIP_CHECK(
+      hipMallocMipmappedArray(&array, &desc, makeMipmappedExtent(flag, size), numLevels, flag));
+#endif
 }
 
 TEST_CASE("Unit_hipGetMipmappedArrayLevel_Negative") {

projects/hip-tests/catch/unit/memory/hipMemPoolSetGetAccess.cc

@@ -260,32 +260,41 @@ TEST_CASE("Unit_hipMemPoolSetAccess_Negative_Parameters") {
   SECTION("Mempool is nullptr") {
     HIP_CHECK_ERROR(hipMemPoolSetAccess(nullptr, &desc, 1), hipErrorInvalidValue);
   }
+
+  // Cuda segfaults here!
 #if HT_AMD
   SECTION("Desc is nullptr and count is > 0") {
     HIP_CHECK_ERROR(hipMemPoolSetAccess(mempool.mempool(), nullptr, 1), hipErrorInvalidValue);
   }
 #endif
+
   SECTION("Count > num_device") {
+#if HT_AMD
     HIP_CHECK_ERROR(hipMemPoolSetAccess(mempool.mempool(), &desc, (num_dev + 1)),
                     hipErrorInvalidDevice);
+#else
+    HIP_CHECK_ERROR(hipMemPoolSetAccess(mempool.mempool(), &desc, (num_dev + 1)),
+                    hipErrorNotSupported);
+#endif
   }
 
   SECTION("Passing invalid desc location type") {
     desc.location.type = hipMemLocationTypeInvalid;
+#if HT_AMD
     HIP_CHECK_ERROR(hipMemPoolSetAccess(mempool.mempool(), &desc, 1), hipErrorInvalidValue);
-    desc.location.type = hipMemLocationTypeDevice;
+#else
+    HIP_CHECK_ERROR(hipMemPoolSetAccess(mempool.mempool(), &desc, 1), hipErrorNotSupported);
+#endif
   }
 
   SECTION("Passing invalid desc location id") {
     desc.location.id = num_dev;
     HIP_CHECK_ERROR(hipMemPoolSetAccess(mempool.mempool(), &desc, 1), hipErrorInvalidDevice);
-    desc.location.id = device_id;
   }
 
   SECTION("Revoking access to own memory pool") {
     desc.flags = hipMemAccessFlagsProtNone;
     HIP_CHECK_ERROR(hipMemPoolSetAccess(mempool.mempool(), &desc, 1), hipErrorInvalidDevice);
-    desc.flags = hipMemAccessFlagsProtReadWrite;
   }
 }
 

projects/hip-tests/catch/unit/memory/hipMemcpyWithStream.cc

@@ -31,6 +31,16 @@ TEST_CASE("Unit_hipMemcpy_Positive_Synchronization_Behavior") {
   using namespace std::placeholders;
   HIP_CHECK(hipDeviceSynchronize());
 
+  // For transfers from pageable host memory to device memory, a stream sync is performed before
+  // the copy is initiated. The function will return once the pageable buffer has been copied to
+  // the staging memory for DMA transfer to device memory, but the DMA to final destination may
+  // not have completed.
+  // For transfers from pinned host memory to device memory, the function is synchronous with
+  // respect to the host
+  SECTION("Host memory to device memory") {
+    MemcpyHPageabletoDSyncBehavior(std::bind(hipMemcpy, _1, _2, _3, hipMemcpyHostToDevice), true);
+  }
+
   // For transfers from device to either pageable or pinned host memory, the function returns only
   // once the copy has completed
   SECTION("Device memory to host memory") {
@@ -41,6 +51,13 @@ TEST_CASE("Unit_hipMemcpy_Positive_Synchronization_Behavior") {
 
   // For transfers from device memory to device memory, no host-side synchronization is performed.
   SECTION("Device memory to device memory") {
+    // This behavior differs on NVIDIA and AMD, on AMD the hipMemcpy calls is synchronous with
+    // respect to the host
+#if HT_AMD
+    HipTest::HIP_SKIP_TEST(
+        "EXSWCPHIPT-127 - Memcpy from device to device memory behavior differs on AMD and Nvidia");
+    return;
+#endif
     MemcpyDtoDSyncBehavior(std::bind(hipMemcpy, _1, _2, _3, hipMemcpyDeviceToDevice), false);
   }
 

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

@@ -24,6 +24,8 @@ THE SOFTWARE.
 #include <hip_test_common.hh>
 #include <hip/hip_runtime_api.h>
 
+#if defined(__HIP_PLATFORM_AMD__) || CUDA_VERSION < CUDA_12000
+
 static hipModule_t GetModule() {
   HIP_CHECK(hipFree(nullptr));
   static const auto mg = ModuleGuard::LoadModule("get_tex_ref_module.code");
@@ -68,4 +70,6 @@ TEST_CASE("Unit_hipModuleGetTexRef_Negative_Name_Is_Empty_String") {
   hipTexRef tex_ref = nullptr;
 
   HIP_CHECK_ERROR(hipModuleGetTexRef(&tex_ref, module, ""), hipErrorInvalidValue);
-}
+}
+
+#endif