EXSWHTEC-217 - Implement new and update existing tests for the hipGra…

…ph*MemcpyNode family of APIs #48 Change-Id: Iae7ac9c855ba6e3288257e99e49f8b16cebb1bac
ROCm · Feb 19, 2024 · cefbaed · cefbaed
1 parent 9a3fd8e
commit cefbaed
Show file tree

Hide file tree

Showing 13 changed files with 2,060 additions and 1,128 deletions.
diff --git a/catch/hipTestMain/config/config_amd_linux b/catch/hipTestMain/config/config_amd_linux
@@ -127,6 +127,7 @@
         "Unit_deviceAllocation_InOneThread_AccessInAllThreads",
         "=== Patch which removes the typetraits implementation from std namespace in hiprtc is reverted ===",
         "Unit_hiprtc_stdheaders",
+        "Unit_hipGraphAddMemcpyNode_Negative_Parameters",
         "Unit_hipMemAddressFree_negative",
         "Unit_hipMemAddressReserve_AlignmentTest",
         "Unit_hipMemAddressReserve_Negative",

diff --git a/catch/hipTestMain/config/config_amd_windows b/catch/hipTestMain/config/config_amd_windows
@@ -219,6 +219,7 @@
         "=== Patch which removes the typetraits implementation from std namespace in hiprtc is reverted ===",
         "Unit_hiprtc_stdheaders",
         "Unit_hipMemAddressReserve_AlignmentTest",
+        "Unit_hipGraphAddMemcpyNode_Negative_Parameters",
         "Unit_hipMemCreate_ChkWithKerLaunch",
         "Unit_hipMemCreate_MapNonContiguousChunks",
         "Unit_hipMemMap_MapPartialPhysicalMem",

diff --git a/catch/include/memcpy1d_tests_common.hh b/catch/include/memcpy1d_tests_common.hh
@@ -141,6 +141,7 @@ void MemcpyDeviceToDeviceShell(F memcpy_func, const hipStream_t kernel_stream =
     int can_access_peer = 0;
     HIP_CHECK(hipDeviceCanAccessPeer(&can_access_peer, src_device, dst_device));
     if (!can_access_peer) {
+      INFO("Peer access cannot be enabled between devices " << src_device << " " << dst_device);
       return;
     }
     HIP_CHECK(hipDeviceEnablePeerAccess(dst_device, 0));
@@ -168,8 +169,8 @@ void MemcpyDeviceToDeviceShell(F memcpy_func, const hipStream_t kernel_stream =
   HIP_CHECK(
       hipMemcpy(result.host_ptr(), dst_allocation.ptr(), allocation_size, hipMemcpyDeviceToHost));
   if constexpr (enable_peer_access) {
-    // If we've gotten this far, EnablePeerAccess must have succeeded, so we only need to check this
-    // condition
+    // If we've gotten this far, EnablePeerAccess must have succeeded, so we
+    // only need to check this condition
     HIP_CHECK(hipDeviceDisablePeerAccess(dst_device));
   }
 
@@ -237,7 +238,6 @@ void MemcpySyncBehaviorCheck(F memcpy_func, const bool should_sync,
   LaunchDelayKernel(std::chrono::milliseconds{100}, kernel_stream);
   HIP_CHECK(memcpy_func());
   if (should_sync) {
-    HIP_CHECK(hipStreamSynchronize(kernel_stream));
     HIP_CHECK(hipStreamQuery(kernel_stream));
   } else {
     HIP_CHECK_ERROR(hipStreamQuery(kernel_stream), hipErrorNotReady);

diff --git a/catch/include/memcpy3d_tests_common.hh b/catch/include/memcpy3d_tests_common.hh
@@ -23,6 +23,7 @@ THE SOFTWARE.
 #pragma once
 #pragma clang diagnostic ignored "-Wmissing-field-initializers"
 #pragma clang diagnostic ignored "-Wunused-lambda-capture"
+
 #include <variant>
 
 #include <hip_test_common.hh>

diff --git a/catch/unit/graph/CMakeLists.txt b/catch/unit/graph/CMakeLists.txt
@@ -32,6 +32,7 @@ set(TEST_SRC
   hipGraph.cc
   hipSimpleGraphWithKernel.cc
   hipGraphAddMemcpyNode.cc
+  hipGraphAddMemcpyNode_old.cc
   hipGraphClone.cc
   hipGraphInstantiateWithFlags.cc
   hipGraphAddHostNode.cc
@@ -72,6 +73,7 @@ set(TEST_SRC
   hipGraphEventRecordNodeSetEvent.cc
   hipGraphEventWaitNodeGetEvent.cc
   hipGraphExecMemcpyNodeSetParams.cc
+  hipGraphExecMemcpyNodeSetParams_old.cc
   hipStreamBeginCapture.cc
   hipGraphAddMemcpyNode1D_old.cc
   hipGraphAddMemcpyNode1D.cc
@@ -101,7 +103,9 @@ set(TEST_SRC
   hipGraphAddMemsetNode.cc
   hipGraphAddKernelNode.cc
   hipGraphMemcpyNodeGetParams.cc
+  hipGraphMemcpyNodeGetParams_old.cc
   hipGraphMemcpyNodeSetParams.cc
+  hipGraphMemcpyNodeSetParams_old.cc
   hipGraphKernelNodeGetParams.cc
   hipGraphKernelNodeSetParams.cc
   hipGraphExecKernelNodeSetParams.cc