SWDEV-461794 Fix Unit_hipMemcpyAsync_Positive_Synchronization_Behavior

- Fixed wrong assumptions in the test:
   - Pinned Host memory to Pinned Host memory is async (before:sync)
   - Pinned Host memory to Device is async (before:sync)

Change-Id: Ib826e177854cbcdad04181e245382cf0bec02c46


[ROCm/hip-tests commit: 12a5d4c24c]

projects/hip-tests/catch/include/memcpy1d_tests_common.hh

@@ -247,12 +247,19 @@ void MemcpySyncBehaviorCheck(F memcpy_func, const bool should_sync,
 }
 
 template <typename F>
-void MemcpyHtoDSyncBehavior(F memcpy_func, const bool should_sync,
-                            const hipStream_t kernel_stream = nullptr) {
-  using LA = LinearAllocs;
-  const auto host_alloc_type = GENERATE(LA::malloc, LA::hipHostMalloc);
-  LinearAllocGuard<int> host_alloc(host_alloc_type, kPageSize);
-  LinearAllocGuard<int> device_alloc(LA::hipMalloc, kPageSize);
+void MemcpyHPageabletoDSyncBehavior(F memcpy_func, const bool should_sync,
+                                    const hipStream_t kernel_stream = nullptr) {
+  LinearAllocGuard<int> host_alloc(LinearAllocs::malloc, kPageSize);
+  LinearAllocGuard<int> device_alloc(LinearAllocs::hipMalloc, kPageSize);
+  MemcpySyncBehaviorCheck(std::bind(memcpy_func, device_alloc.ptr(), host_alloc.ptr(), kPageSize),
+                          should_sync, kernel_stream);
+}
+
+template <typename F>
+void MemcpyHPinnedtoDSyncBehavior(F memcpy_func, const bool should_sync,
+                                  const hipStream_t kernel_stream = nullptr) {
+  LinearAllocGuard<int> host_alloc(LinearAllocs::hipHostMalloc, kPageSize);
+  LinearAllocGuard<int> device_alloc(LinearAllocs::hipMalloc, kPageSize);
   MemcpySyncBehaviorCheck(std::bind(memcpy_func, device_alloc.ptr(), host_alloc.ptr(), kPageSize),
                           should_sync, kernel_stream);
 }
@@ -288,8 +295,9 @@ template <typename F>
 void MemcpyHtoHSyncBehavior(F memcpy_func, const bool should_sync,
                             const hipStream_t kernel_stream = nullptr) {
   using LA = LinearAllocs;
-  const auto src_alloc_type = GENERATE(LA::malloc, LA::hipHostMalloc);
-  const auto dst_alloc_type = GENERATE(LA::malloc, LA::hipHostMalloc);
+  const auto [src_alloc_type, dst_alloc_type] = GENERATE(
+      std::make_tuple(LA::malloc, LA::hipHostMalloc),
+      std::make_tuple(LA::hipHostMalloc, LA::malloc), std::make_tuple(LA::malloc, LA::malloc));
 
   LinearAllocGuard<int> src_alloc(src_alloc_type, kPageSize);
   LinearAllocGuard<int> dst_alloc(dst_alloc_type, kPageSize);
@@ -297,6 +305,15 @@ void MemcpyHtoHSyncBehavior(F memcpy_func, const bool should_sync,
                           should_sync, kernel_stream);
 }
 
+template <typename F>
+void MemcpyHPinnedtoHPinnedSyncBehavior(F memcpy_func, const bool should_sync,
+                            const hipStream_t kernel_stream = nullptr) {
+  LinearAllocGuard<int> src_alloc(LinearAllocs::hipHostMalloc, kPageSize);
+  LinearAllocGuard<int> dst_alloc(LinearAllocs::hipHostMalloc, kPageSize);
+  MemcpySyncBehaviorCheck(std::bind(memcpy_func, dst_alloc.ptr(), src_alloc.ptr(), kPageSize),
+                          should_sync, kernel_stream);
+}
+
 // Common negative tests
 template <typename F> void MemcpyCommonNegativeTests(F f, void* dst, void* src, size_t count) {
   SECTION("dst == nullptr") { HIP_CHECK_ERROR(f(nullptr, src, count), hipErrorInvalidValue); }
@@ -316,4 +333,4 @@ void MemcpyWithDirectionCommonNegativeTests(F f, void* dst, void* src, size_t co
                     hipErrorInvalidMemcpyDirection);
   }
 #endif
-}
+}

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

@@ -38,9 +38,14 @@ TEST_CASE("Unit_hipMemcpyAsync_Positive_Synchronization_Behavior") {
   using namespace std::placeholders;
   HIP_CHECK(hipDeviceSynchronize());
 
-  SECTION("Host memory to device memory") {
-    MemcpyHtoDSyncBehavior(std::bind(hipMemcpyAsync, _1, _2, _3, hipMemcpyHostToDevice, nullptr),
-                           true);
+  SECTION("Host pageable memory to device memory") {
+    MemcpyHPageabletoDSyncBehavior(
+        std::bind(hipMemcpyAsync, _1, _2, _3, hipMemcpyHostToDevice, nullptr), true);
+  }
+
+  SECTION("Host pinned memory to device memory") {
+    MemcpyHPinnedtoDSyncBehavior(
+        std::bind(hipMemcpyAsync, _1, _2, _3, hipMemcpyHostToDevice, nullptr), false);
   }
 
   SECTION("Device memory to pageable host memory") {
@@ -66,6 +71,9 @@ TEST_CASE("Unit_hipMemcpyAsync_Positive_Synchronization_Behavior") {
   SECTION("Host memory to host memory") {
     MemcpyHtoHSyncBehavior(std::bind(hipMemcpyAsync, _1, _2, _3, hipMemcpyHostToHost, nullptr),
                            true);
+
+    MemcpyHPinnedtoHPinnedSyncBehavior(
+        std::bind(hipMemcpyAsync, _1, _2, _3, hipMemcpyHostToHost, nullptr), false);
   }
 }
 

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

@@ -98,7 +98,7 @@ TEST_CASE("Unit_hipMemcpyHtoDAsync_Positive_Synchronization_Behavior") {
       "Nvidia");
   return;
 #endif
-  MemcpyHtoDSyncBehavior(
+  MemcpyHPinnedtoDSyncBehavior(
       [](void* dst, void* src, size_t count) {
         return hipMemcpyHtoDAsync(reinterpret_cast<hipDeviceptr_t>(dst), src, count, nullptr);
       },

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

@@ -38,7 +38,7 @@ TEST_CASE("Unit_hipMemcpy_Positive_Synchronization_Behavior") {
   // For transfers from pinned host memory to device memory, the function is synchronous with
   // respect to the host
   SECTION("Host memory to device memory") {
-    MemcpyHtoDSyncBehavior(std::bind(hipMemcpy, _1, _2, _3, hipMemcpyHostToDevice), true);
+    MemcpyHPageabletoDSyncBehavior(std::bind(hipMemcpy, _1, _2, _3, hipMemcpyHostToDevice), true);
   }
 
   // For transfers from device to either pageable or pinned host memory, the function returns only
@@ -65,6 +65,8 @@ TEST_CASE("Unit_hipMemcpy_Positive_Synchronization_Behavior") {
   // respect to the host
   SECTION("Host memory to host memory") {
     MemcpyHtoHSyncBehavior(std::bind(hipMemcpy, _1, _2, _3, hipMemcpyHostToHost), true);
+    MemcpyHPinnedtoHPinnedSyncBehavior(
+        std::bind(hipMemcpyAsync, _1, _2, _3, hipMemcpyHostToHost, nullptr), true);
   }
 }
 

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

@@ -60,7 +60,7 @@ TEST_CASE("Unit_hipMemcpyHtoD_Positive_Basic") {
 }
 
 TEST_CASE("Unit_hipMemcpyHtoD_Positive_Synchronization_Behavior") {
-  MemcpyHtoDSyncBehavior(
+  MemcpyHPageabletoDSyncBehavior(
       [](void* dst, void* src, size_t count) {
         return hipMemcpyHtoD(reinterpret_cast<hipDeviceptr_t>(dst), src, count);
       },
@@ -116,4 +116,4 @@ TEST_CASE("Unit_hipMemcpyDtoD_Negative_Parameters") {
                              reinterpret_cast<hipDeviceptr_t>(src), count);
       },
       dst_alloc.ptr(), src_alloc.ptr(), kPageSize);
-}
+}