SWDEV-486969 - Reuse existing tests to verify behavior of capturing sync Memcpy APis

Change-Id: I7e94743d6957ffaadae9cff297e3f0d93f9ff806


[ROCm/hip-tests commit: 6c5468f802]

projects/hip-tests/catch/unit/memory/array_memcpy_tests_common.hh

@@ -172,7 +172,6 @@ template <bool should_synchronize, typename T, typename F>
 void Memcpy2DHosttoAShell(F memcpy_func, size_t width, size_t height,
                           const hipStream_t kernel_stream = nullptr) {
   const unsigned int flag = hipArrayDefault;
-  ;
 
   size_t allocation_size = width * height * sizeof(T);
 

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

@@ -85,7 +85,7 @@ TEST_CASE("Unit_hipDrvMemcpy3D_Positive_Parameters") {
   CHECK_IMAGE_SUPPORT
 
   constexpr bool async = false;
-  Memcpy3DZeroWidthHeightDepth<async>(DrvMemcpy3DWrapper<async>);
+  Memcpy3DZeroWidthHeightDepth<async>(DrvMemcpy3DWrapper<>);
 }
 
 // Disabled on AMD due to defect - EXSWHTEC-238
@@ -93,8 +93,8 @@ TEST_CASE("Unit_hipDrvMemcpy3D_Positive_Array") {
   CHECK_IMAGE_SUPPORT
 
   constexpr bool async = false;
-  SECTION("Array from/to Host") { DrvMemcpy3DArrayHostShell<async>(DrvMemcpy3DWrapper<async>); }
-  SECTION("Array from/to Device") { DrvMemcpy3DArrayDeviceShell<async>(DrvMemcpy3DWrapper<async>); }
+  SECTION("Array from/to Host") { DrvMemcpy3DArrayHostShell<async>(DrvMemcpy3DWrapper<>); }
+  SECTION("Array from/to Device") { DrvMemcpy3DArrayDeviceShell<async>(DrvMemcpy3DWrapper<>); }
 }
 
 TEST_CASE("Unit_hipDrvMemcpy3D_Negative_Parameters") {
@@ -234,4 +234,24 @@ TEST_CASE("Unit_hipDrvMemcpy3D_Negative_Parameters") {
     NegativeTests(dst_alloc.pitched_ptr(), make_hipPos(0, 0, 0), src_alloc.pitched_ptr(),
                   make_hipPos(0, 0, 0), extent, hipMemcpyDeviceToDevice);
   }
+}
+
+TEST_CASE("Unit_hipDrvMemcpy3D_Capture") {
+  CHECK_IMAGE_SUPPORT
+
+  constexpr hipExtent extent{128 * sizeof(int), 128, 8};
+  LinearAllocGuard3D<int> device_alloc(extent);
+  LinearAllocGuard<int> host_alloc(
+      LinearAllocs::hipHostMalloc,
+      device_alloc.pitch() * device_alloc.height() * device_alloc.depth());
+
+  auto params = GetDrvMemcpy3DParms(device_alloc.pitched_ptr(), make_hipPos(0, 0, 0),
+                                    make_hipPitchedPtr(host_alloc.ptr(), device_alloc.pitch(),
+                                                       device_alloc.width(), device_alloc.height()),
+                                    make_hipPos(0, 0, 0), extent, hipMemcpyHostToDevice);
+
+  hipError_t memcpy_err = hipSuccess;
+  BEGIN_CAPTURE_SYNC(memcpy_err, false);
+  HIP_CHECK_ERROR(hipDrvMemcpy3D(&params), memcpy_err);
+  END_CAPTURE_SYNC(memcpy_err);
 }

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

@@ -151,3 +151,21 @@ TEST_CASE("Unit_hipMemcpy2D_Negative_Parameters") {
                   dst_alloc.width(), dst_alloc.height(), hipMemcpyDeviceToDevice);
   }
 }
+
+TEST_CASE("Unit_hipMemcpy2D_Capture") {
+  CHECK_IMAGE_SUPPORT
+
+  constexpr size_t width = 16;
+  constexpr size_t height = 16;
+
+  LinearAllocGuard2D<int> device_alloc(width, height);
+  LinearAllocGuard<int> host_alloc(LinearAllocs::hipHostMalloc, device_alloc.pitch() * width);
+
+  hipError_t memcpy_err = hipSuccess;
+  BEGIN_CAPTURE_SYNC(memcpy_err, false);
+  HIP_CHECK_ERROR(
+      hipMemcpy2D(device_alloc.ptr(), device_alloc.pitch(), host_alloc.ptr(), device_alloc.pitch(),
+                  device_alloc.width(), device_alloc.height(), hipMemcpyHostToDevice),
+      memcpy_err);
+  END_CAPTURE_SYNC(memcpy_err);
+}

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

@@ -132,9 +132,9 @@ TEST_CASE("Unit_hipMemcpy2DFromArray_Positive_ZeroWidthHeight") {
           width, height);
     }
     SECTION("Width is 0") {
-      Memcpy2DFromArrayZeroWidthHeight<false>(
-          std::bind(hipMemcpy2DFromArray, _1, _2, _3, 0, 0, 0, height, hipMemcpyDeviceToHost),
-          width, height);
+      Memcpy2DFromArrayZeroWidthHeight<false>(std::bind(hipMemcpy2DFromArray, _1, _2, _3,
+                                                        0, 0, 0, height, hipMemcpyDeviceToHost),
+                                              width, height);
     }
   }
   SECTION("Array to device") {
@@ -260,3 +260,25 @@ TEST_CASE("Unit_hipMemcpy2DFromArray_Negative_Parameters") {
 #endif
   }
 }
+
+TEST_CASE("Unit_hipMemcpy2DFromArray_Capture") {
+  CHECK_IMAGE_SUPPORT
+
+  const auto width = 16;
+  const auto height = 16;
+  const auto size = width * height * sizeof(int);
+
+  ArrayAllocGuard<int> A_d(make_hipExtent(width, height, 0), hipArrayDefault);
+  LinearAllocGuard<int> A_h(LinearAllocs::hipHostMalloc, size);
+  LinearAllocGuard<int> B_h(LinearAllocs::hipHostMalloc, size);
+
+  HIP_CHECK(hipMemcpy2DToArray(A_d.ptr(), 0, 0, A_h.ptr(), width * sizeof(int),
+                               width * sizeof(int), height, hipMemcpyHostToDevice));
+
+  hipError_t memcpy_err = hipSuccess;
+  BEGIN_CAPTURE_SYNC(memcpy_err, false);
+  HIP_CHECK_ERROR(hipMemcpy2DFromArray(B_h.host_ptr(), width * sizeof(int), A_d.ptr(), 0, 0,
+                                       width * sizeof(int), height, hipMemcpyDeviceToHost),
+                  memcpy_err);
+  END_CAPTURE_SYNC(memcpy_err);
+}

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

@@ -127,9 +127,9 @@ TEST_CASE("Unit_hipMemcpy2DToArray_Positive_ZeroWidthHeight") {
           width, height);
     }
     SECTION("Width is 0") {
-      Memcpy2DToArrayZeroWidthHeight<false>(
-          std::bind(hipMemcpy2DToArray, _1, 0, 0, _2, _3, 0, height, hipMemcpyHostToDevice), width,
-          height);
+      Memcpy2DToArrayZeroWidthHeight<false>(std::bind(hipMemcpy2DToArray, _1, 0, 0, _2,
+                                                      _3, 0, height, hipMemcpyHostToDevice),
+                                            width, height);
     }
   }
   SECTION("Array to device") {
@@ -140,9 +140,9 @@ TEST_CASE("Unit_hipMemcpy2DToArray_Positive_ZeroWidthHeight") {
           width, height);
     }
     SECTION("Width is 0") {
-      Memcpy2DToArrayZeroWidthHeight<false>(
-          std::bind(hipMemcpy2DToArray, _1, 0, 0, _2, _3, 0, height, hipMemcpyDeviceToDevice),
-          width, height);
+      Memcpy2DToArrayZeroWidthHeight<false>(std::bind(hipMemcpy2DToArray, _1, 0, 0, _2,
+                                                      _3, 0, height, hipMemcpyDeviceToDevice),
+                                            width, height);
     }
   }
 }
@@ -255,3 +255,21 @@ TEST_CASE("Unit_hipMemcpy2DToArray_Negative_Parameters") {
 #endif
   }
 }
+
+TEST_CASE("Unit_hipMemcpy2DToArray_Capture") {
+  CHECK_IMAGE_SUPPORT
+
+  const auto width = 16;
+  const auto height = 16;
+  const auto size = width * height * sizeof(int);
+
+  ArrayAllocGuard<int> array_alloc(make_hipExtent(width, height, 0), hipArrayDefault);
+  LinearAllocGuard<int> host_alloc(LinearAllocs::hipHostMalloc, size);
+
+  hipError_t memcpy_err = hipSuccess;
+  BEGIN_CAPTURE_SYNC(memcpy_err, false);
+  HIP_CHECK_ERROR(hipMemcpy2DToArray(array_alloc.ptr(), 0, 0, host_alloc.ptr(), width * sizeof(int),
+                                     width * sizeof(int), height, hipMemcpyHostToDevice),
+                  memcpy_err);
+  END_CAPTURE_SYNC(memcpy_err);
+}

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

@@ -41,7 +41,9 @@ TEST_CASE("Unit_hipMemcpy3D_Positive_Basic") {
     SECTION("Peer access disabled") {
       Memcpy3DDeviceToDeviceShell<async, false>(Memcpy3DWrapper<>);
     }
-    SECTION("Peer access enabled") { Memcpy3DDeviceToDeviceShell<async, true>(Memcpy3DWrapper<>); }
+    SECTION("Peer access enabled") {
+      Memcpy3DDeviceToDeviceShell<async, true>(Memcpy3DWrapper<>);
+    }
   }
 
   SECTION("Host to Device") { Memcpy3DHostToDeviceShell<async>(Memcpy3DWrapper<>); }
@@ -78,7 +80,12 @@ TEST_CASE("Unit_hipMemcpy3D_Positive_DeviceToDevice_Synchronization_Behavior") {
 
   b_context.block_stream();
   REQUIRE(b_context.is_blocked());
-  HIP_CHECK(hipMemcpy3D(&parms));
+
+  hipError_t memcpy_err = hipSuccess;
+  BEGIN_CAPTURE_SYNC(memcpy_err, false);
+  HIP_CHECK_ERROR(hipMemcpy3D(&parms), memcpy_err);
+  END_CAPTURE_SYNC(memcpy_err);
+
   HIP_CHECK_ERROR(hipStreamQuery(kernel_stream), hipErrorNotReady);
   b_context.unblock_stream();
   HIP_CHECK(hipDeviceSynchronize());
@@ -89,16 +96,16 @@ TEST_CASE("Unit_hipMemcpy3D_Positive_Parameters") {
   CHECK_IMAGE_SUPPORT
 
   constexpr bool async = false;
-  Memcpy3DZeroWidthHeightDepth<async>(Memcpy3DWrapper<async>);
+  Memcpy3DZeroWidthHeightDepth<async>(Memcpy3DWrapper<>);
 }
 
 TEST_CASE("Unit_hipMemcpy3D_Positive_Array") {
   CHECK_IMAGE_SUPPORT
 
   constexpr bool async = false;
-  SECTION("Array from/to Host") { Memcpy3DArrayHostShell<async>(Memcpy3DWrapper<async>); }
+  SECTION("Array from/to Host") { Memcpy3DArrayHostShell<async>(Memcpy3DWrapper<>); }
 #if HT_NVIDIA  // Disabled on AMD due to defect - EXSWHTEC-238
-  SECTION("Array from/to Device") { Memcpy3DArrayDeviceShell<async>(Memcpy3DWrapper<async>); }
+  SECTION("Array from/to Device") { Memcpy3DArrayDeviceShell<async>(Memcpy3DWrapper<>); }
 #endif
 }
 
@@ -250,3 +257,22 @@ TEST_CASE("Unit_hipMemcpy3D_Negative_Parameters") {
                   make_hipPos(0, 0, 0), extent, hipMemcpyDeviceToDevice);
   }
 }
+
+TEST_CASE("Unit_hipMemcpy3D_Capture") {
+  CHECK_IMAGE_SUPPORT
+
+  constexpr hipExtent extent{16 * sizeof(int), 16, 16};
+  LinearAllocGuard3D<int> dev_alloc(extent);
+  LinearAllocGuard<int> host_alloc(LinearAllocs::hipHostMalloc,
+                                   dev_alloc.pitch() * dev_alloc.height() * dev_alloc.depth());
+
+  auto params = GetMemcpy3DParms(dev_alloc.pitched_ptr(), make_hipPos(0, 0, 0),
+                                 make_hipPitchedPtr(host_alloc.ptr(), dev_alloc.pitch(),
+                                                    dev_alloc.width(), dev_alloc.height()),
+                                 make_hipPos(0, 0, 0), extent, hipMemcpyHostToDevice);
+
+  hipError_t memcpy_err = hipSuccess;
+  BEGIN_CAPTURE_SYNC(memcpy_err, false);
+  HIP_CHECK_ERROR(hipMemcpy3D(&params), memcpy_err);
+  END_CAPTURE_SYNC(memcpy_err);
+}

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

@@ -65,12 +65,18 @@ TEST_CASE("Unit_hipMemcpyAtoA_Basic") {
   HIP_CHECK(hipMallocArray(&B_a, &desc, col, row, hipArrayDefault));
   HIP_CHECK(hipMemcpy2DToArray(A_a, 0, 0, A_h, col * sizeof(int),
                                col * sizeof(int), row, hipMemcpyHostToDevice));
-  HIP_CHECK(hipMemcpyAtoA(B_a, 0, A_a, 0, sizeof(int) * row * col));
-  HIP_CHECK(hipMemcpy2DFromArray(B_h, sizeof(int) * col, B_a, 0, 0,
-                                 sizeof(int) * col, row,
-                                 hipMemcpyDeviceToHost));
-  for (int i = 0; i < (row * col); i++) {
-    REQUIRE(A_h[i] == B_h[i]);
+
+  hipError_t memcpy_err = hipSuccess;
+  BEGIN_CAPTURE_SYNC(memcpy_err, false);
+  HIP_CHECK_ERROR(hipMemcpyAtoA(B_a, 0, A_a, 0, sizeof(int) * row * col), memcpy_err);
+  END_CAPTURE_SYNC(memcpy_err);
+
+  if (memcpy_err == hipSuccess) {
+    HIP_CHECK(hipMemcpy2DFromArray(B_h, sizeof(int) * col, B_a, 0, 0, sizeof(int) * col, row,
+                                   hipMemcpyDeviceToHost));
+    for (int i = 0; i < (row * col); i++) {
+      REQUIRE(A_h[i] == B_h[i]);
+    }
   }
   HIP_CHECK(hipFreeArray(A_a));
   HIP_CHECK(hipFreeArray(B_a));

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

@@ -62,12 +62,20 @@ TEST_CASE("Unit_hipMemcpyAtoD_Basic") {
   hipChannelFormatDesc desc = hipCreateChannelDesc<int>();
   HIP_CHECK(hipMallocArray(&A_a, &desc, col, row, hipArrayDefault));
   HIP_CHECK(hipMalloc(&A_d, sizeof(int) * row * col));
-  HIP_CHECK(hipMemcpy2DToArray(A_a, 0, 0, A_h, col * sizeof(int),
-                               col * sizeof(int), row, hipMemcpyHostToDevice));
-  HIP_CHECK(hipMemcpyAtoD(A_d, A_a, 0, sizeof(int) * col * row));
-  HIP_CHECK(hipMemcpyDtoH(B_h, A_d, sizeof(int) * row * col));
-  for (int i = 0; i < (row * col); i++) {
-    REQUIRE(A_h[i] == B_h[i]);
+
+  hipError_t memcpy_err = hipSuccess;
+  BEGIN_CAPTURE_SYNC(memcpy_err, false);
+  HIP_CHECK_ERROR(hipMemcpy2DToArray(A_a, 0, 0, A_h, col * sizeof(int), col * sizeof(int), row,
+                                     hipMemcpyHostToDevice),
+                  memcpy_err);
+  END_CAPTURE_SYNC(memcpy_err);
+
+  if (memcpy_err == hipSuccess) {
+    HIP_CHECK(hipMemcpyAtoD(A_d, A_a, 0, sizeof(int) * col * row));
+    HIP_CHECK(hipMemcpyDtoH(B_h, A_d, sizeof(int) * row * col));
+    for (int i = 0; i < (row * col); i++) {
+      REQUIRE(A_h[i] == B_h[i]);
+    }
   }
   HIP_CHECK(hipFreeArray(A_a));
   HIP_CHECK(hipFree(A_d));

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

@@ -74,7 +74,14 @@ TEST_CASE("Unit_hipMemcpyAtoH_Positive_ZeroCount") {
                                sizeof(int) * width, 1, hipMemcpyHostToDevice));
   fill_value = 41;
   std::fill_n(host_alloc.host_ptr(), width, fill_value);
-  HIP_CHECK(hipMemcpyAtoH(host_alloc.ptr(), array_alloc.ptr(), 0, 0));
+
+  hipError_t memcpy_err = hipSuccess;
+  BEGIN_CAPTURE_SYNC(memcpy_err, false);
+  HIP_CHECK_ERROR(hipMemcpyAtoH(host_alloc.ptr(), array_alloc.ptr(), 0, 0), memcpy_err);
+  END_CAPTURE_SYNC(memcpy_err);
+  if (memcpy_err == hipErrorStreamCaptureImplicit) {
+    return;
+  }
 
   ArrayFindIfNot(host_alloc.host_ptr(), static_cast<uint8_t>(fill_value), width);
 }
@@ -122,3 +129,22 @@ TEST_CASE("Unit_hipMemcpyAtoH_Negative_Parameters") {
                     hipErrorInvalidValue);
   }
 }
+
+TEST_CASE("Unit_hipMemcpyAtoH_Capture") {
+  CHECK_IMAGE_SUPPORT
+
+  const auto width = 1024;
+  const auto height = 0;
+  const auto allocation_size = width * sizeof(int);
+
+  const unsigned int flag = hipArrayDefault;
+
+  ArrayAllocGuard<int> array_alloc(make_hipExtent(width, height, 0), flag);
+  LinearAllocGuard<int> host_alloc(LinearAllocs::hipHostMalloc, allocation_size);
+
+  hipError_t memcpy_err = hipSuccess;
+  BEGIN_CAPTURE_SYNC(memcpy_err, false);
+  HIP_CHECK_ERROR(hipMemcpyAtoH(host_alloc.ptr(), array_alloc.ptr(), 0, allocation_size),
+                  memcpy_err);
+  END_CAPTURE_SYNC(memcpy_err);
+}

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

@@ -77,19 +77,22 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpyDtoD_Basic", "",
     HipTest::checkVectorADD<TestType>(A_h, B_h, C_h, NUM_ELM);
 
     HIP_CHECK(hipSetDevice(1));
-    HIP_CHECK(hipMemcpyDtoD((hipDeviceptr_t)X_d, (hipDeviceptr_t)A_d,
-                            Nbytes));
-    HIP_CHECK(hipMemcpyDtoD((hipDeviceptr_t)Y_d, (hipDeviceptr_t)B_d,
-                            Nbytes));
 
-    hipLaunchKernelGGL(HipTest::vectorADD, dim3(1),
-                        dim3(1), 0, 0,
-                        static_cast<const TestType*>(X_d),
-                        static_cast<const TestType*>(Y_d), Z_d, NUM_ELM);
-    HIP_CHECK(hipGetLastError());
-    HIP_CHECK(hipMemcpyDtoH(C_h, (hipDeviceptr_t)Z_d, Nbytes));
-    HIP_CHECK(hipDeviceSynchronize());
-    HipTest::checkVectorADD<TestType>(A_h, B_h, C_h, NUM_ELM);
+    hipError_t memcpy_err = hipSuccess;
+    BEGIN_CAPTURE_SYNC(memcpy_err, false);
+    HIP_CHECK_ERROR(hipMemcpyDtoD((hipDeviceptr_t)X_d, (hipDeviceptr_t)A_d, Nbytes), memcpy_err);
+    HIP_CHECK_ERROR(hipMemcpyDtoD((hipDeviceptr_t)Y_d, (hipDeviceptr_t)B_d, Nbytes), memcpy_err);
+    END_CAPTURE_SYNC(memcpy_err);
+
+    if (memcpy_err == hipSuccess) {
+      hipLaunchKernelGGL(HipTest::vectorADD, dim3(1), dim3(1), 0, 0,
+                         static_cast<const TestType*>(X_d), static_cast<const TestType*>(Y_d), Z_d,
+                         NUM_ELM);
+      HIP_CHECK(hipGetLastError());
+      HIP_CHECK(hipMemcpyDtoH(C_h, (hipDeviceptr_t)Z_d, Nbytes));
+      HIP_CHECK(hipDeviceSynchronize());
+      HipTest::checkVectorADD<TestType>(A_h, B_h, C_h, NUM_ELM);
+    }
 
     HipTest::freeArrays<TestType>(A_d, B_d, C_d, A_h, B_h, C_h, false);
     HIP_CHECK(hipFree(X_d));

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

@@ -127,12 +127,17 @@ TEST_CASE("Unit_hipMemcpyToFromSymbol_SyncAndAsync") {
   }
   INFO("Stream :: " << streamType);
 
+  hipError_t memcpy_err = hipSuccess;
+
   SECTION("Singular Value") {
     int set{42};
     int result{0};
     if (streamType == StreamTestType::NoStream) {
-      HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), &set, sizeof(int)));
-      HIP_CHECK(hipMemcpyFromSymbol(&result, HIP_SYMBOL(devSymbol), sizeof(int)));
+      memcpy_err = hipSuccess;
+      BEGIN_CAPTURE_SYNC(memcpy_err, false);
+      HIP_CHECK_ERROR(hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), &set, sizeof(int)), memcpy_err);
+      HIP_CHECK_ERROR(hipMemcpyFromSymbol(&result, HIP_SYMBOL(devSymbol), sizeof(int)), memcpy_err);
+      END_CAPTURE_SYNC(memcpy_err);
     } else {
       HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(devSymbol), &set, sizeof(int), 0,
                                        hipMemcpyHostToDevice, stream));
@@ -141,7 +146,9 @@ TEST_CASE("Unit_hipMemcpyToFromSymbol_SyncAndAsync") {
                                          hipMemcpyDeviceToHost, stream));
       HIP_CHECK(hipStreamSynchronize(stream));
     }
-    REQUIRE(result == set);
+    if (memcpy_err == hipSuccess) {
+      REQUIRE(result == set);
+    }
   }
 
   SECTION("Array Values") {
@@ -149,8 +156,13 @@ TEST_CASE("Unit_hipMemcpyToFromSymbol_SyncAndAsync") {
     int set[size] = {4, 2, 4, 2, 4, 2, 4, 2, 4, 2};
     int result[size] = {0};
     if (streamType == StreamTestType::NoStream) {
-      HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), set, sizeof(int) * size));
-      HIP_CHECK(hipMemcpyFromSymbol(&result, HIP_SYMBOL(devSymbol), sizeof(int) * size));
+      memcpy_err = hipSuccess;
+      BEGIN_CAPTURE_SYNC(memcpy_err, false);
+      HIP_CHECK_ERROR(hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), set, sizeof(int) * size),
+                      memcpy_err);
+      HIP_CHECK_ERROR(hipMemcpyFromSymbol(&result, HIP_SYMBOL(devSymbol), sizeof(int) * size),
+                      memcpy_err);
+      END_CAPTURE_SYNC(memcpy_err);
     } else {
       HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(devSymbol), set, sizeof(int) * size, 0,
                                        hipMemcpyHostToDevice, stream));
@@ -159,8 +171,10 @@ TEST_CASE("Unit_hipMemcpyToFromSymbol_SyncAndAsync") {
                                          hipMemcpyDeviceToHost, stream));
       HIP_CHECK(hipStreamSynchronize(stream));
     }
-    for (size_t i = 0; i < size; i++) {
-      REQUIRE(result[i] == set[i]);
+    if (memcpy_err == hipSuccess) {
+      for (size_t i = 0; i < size; i++) {
+        REQUIRE(result[i] == set[i]);
+      }
     }
   }
 
@@ -170,9 +184,14 @@ TEST_CASE("Unit_hipMemcpyToFromSymbol_SyncAndAsync") {
     int set[size] = {9, 9, 9, 9, 9, 2, 4, 2, 4, 2};
     int result[size] = {0};
     if (streamType == StreamTestType::NoStream) {
-      HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), set, offset));
-      HIP_CHECK(hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), set + 5, offset, offset));
-      HIP_CHECK(hipMemcpyFromSymbol(result, HIP_SYMBOL(devSymbol), sizeof(int) * size));
+      memcpy_err = hipSuccess;
+      BEGIN_CAPTURE_SYNC(memcpy_err, false);
+      HIP_CHECK_ERROR(hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), set, offset), memcpy_err);
+      HIP_CHECK_ERROR(hipMemcpyToSymbol(HIP_SYMBOL(devSymbol), set + 5, offset, offset),
+                      memcpy_err);
+      HIP_CHECK_ERROR(hipMemcpyFromSymbol(result, HIP_SYMBOL(devSymbol), sizeof(int) * size),
+                      memcpy_err);
+      END_CAPTURE_SYNC(memcpy_err);
     } else {
       HIP_CHECK(hipMemcpyToSymbolAsync(HIP_SYMBOL(devSymbol), set, offset, 0, hipMemcpyHostToDevice,
                                        stream));
@@ -184,8 +203,10 @@ TEST_CASE("Unit_hipMemcpyToFromSymbol_SyncAndAsync") {
                                          hipMemcpyDeviceToHost, stream));
       HIP_CHECK(hipStreamSynchronize(stream));
     }
-    for (size_t i = 0; i < size; i++) {
-      REQUIRE(result[i] == set[i]);
+    if (memcpy_err == hipSuccess) {
+      for (size_t i = 0; i < size; i++) {
+        REQUIRE(result[i] == set[i]);
+      }
     }
   }
 }

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

@@ -81,7 +81,14 @@ TEST_CASE("Unit_hipMemcpyHtoA_Positive_ZeroCount") {
                                sizeof(int) * width, 1, hipMemcpyHostToDevice));
   fill_value = 41;
   std::fill_n(host_alloc.host_ptr(), width, fill_value);
-  HIP_CHECK(hipMemcpyHtoA(array_alloc.ptr(), 0, host_alloc.ptr(), 0));
+
+  hipError_t memcpy_err = hipSuccess;
+  BEGIN_CAPTURE_SYNC(memcpy_err, false);
+  HIP_CHECK_ERROR(hipMemcpyHtoA(array_alloc.ptr(), 0, host_alloc.ptr(), 0), memcpy_err);
+  END_CAPTURE_SYNC(memcpy_err);
+  if (memcpy_err = hipErrorStreamCaptureImplicit) {
+    return;
+  }
 
   HIP_CHECK(hipMemcpy2DFromArray(host_alloc.host_ptr(), sizeof(int) * width, array_alloc.ptr(), 0,
                                  0, sizeof(int) * width, 1, hipMemcpyDeviceToHost));
@@ -133,3 +140,18 @@ TEST_CASE("Unit_hipMemcpyHtoA_Negative_Parameters") {
                     hipErrorInvalidValue);
   }
 }
+
+TEST_CASE("Unit_hipMemcpyHtoA_Capture") {
+  CHECK_IMAGE_SUPPORT
+
+  const auto width = 512;
+  const auto size = width * sizeof(int);
+
+  ArrayAllocGuard<int> array_alloc(make_hipExtent(width, 0, 0), hipArrayDefault);
+  LinearAllocGuard<int> host_alloc(LinearAllocs::hipHostMalloc, size);
+
+  hipError_t memcpy_err = hipSuccess;
+  BEGIN_CAPTURE_SYNC(memcpy_err, false);
+  HIP_CHECK_ERROR(hipMemcpyHtoA(array_alloc.ptr(), 0, host_alloc.ptr(), size), memcpy_err);
+  END_CAPTURE_SYNC(memcpy_err);
+}

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

@@ -190,3 +190,25 @@ TEST_CASE("Unit_hipMemcpyParam2D_Negative_Parameters") {
                   dst_alloc.width(), dst_alloc.height(), hipMemcpyDeviceToDevice);
   }
 }
+
+TEST_CASE("Unit_hipMemcpyParam2D_Capture") { 
+  CHECK_IMAGE_SUPPORT
+
+  constexpr size_t cols = 128;
+  constexpr size_t rows = 128;
+
+  LinearAllocGuard2D<int> device_alloc(cols, rows);
+  LinearAllocGuard<int> host_alloc(LinearAllocs::hipHostMalloc, device_alloc.pitch() * rows);
+
+  hip_Memcpy2D params = {};
+  memset(&params, 0x0, sizeof(hip_Memcpy2D));
+
+  InitializeMemcpy2DParams(&params, device_alloc.ptr(), device_alloc.pitch(), host_alloc.ptr(),
+                           device_alloc.pitch(), device_alloc.width(), device_alloc.height(),
+                           hipMemcpyHostToDevice);
+
+  hipError_t memcpy_err = hipSuccess;
+  BEGIN_CAPTURE_SYNC(memcpy_err, false);
+  HIP_CHECK_ERROR(hipMemcpyParam2D(&params), memcpy_err);
+  END_CAPTURE_SYNC(memcpy_err);
+}

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

@@ -77,15 +77,21 @@ TEST_CASE("Unit_hipMemcpyPeer_Positive_Default") {
     VectorSet<<<block_count, thread_count, 0>>>(src_alloc.ptr(), expected_value, element_count);
     HIP_CHECK(hipGetLastError());
 
-    HIP_CHECK(
-        hipMemcpyPeer(dst_alloc.ptr(), dst_device, src_alloc.ptr(), src_device, allocation_size));
+    hipError_t memcpy_err = hipSuccess;
+    BEGIN_CAPTURE_SYNC(memcpy_err, false);
+    HIP_CHECK_ERROR(
+        hipMemcpyPeer(dst_alloc.ptr(), dst_device, src_alloc.ptr(), src_device, allocation_size),
+        memcpy_err);
+    END_CAPTURE_SYNC(memcpy_err);
 
-    HIP_CHECK(
-        hipMemcpy(result.host_ptr(), dst_alloc.ptr(), allocation_size, hipMemcpyDeviceToHost));
+    if (memcpy_err == hipSuccess) {
+      HIP_CHECK(
+          hipMemcpy(result.host_ptr(), dst_alloc.ptr(), allocation_size, hipMemcpyDeviceToHost));
 
-    HIP_CHECK(hipDeviceDisablePeerAccess(dst_device));
+      HIP_CHECK(hipDeviceDisablePeerAccess(dst_device));
 
-    ArrayFindIfNot(result.host_ptr(), expected_value, element_count);
+      ArrayFindIfNot(result.host_ptr(), expected_value, element_count);
+    }
   } else {
     INFO("Peer access cannot be enabled between devices " << src_device << " " << dst_device);
   }
@@ -192,14 +198,21 @@ TEST_CASE("Unit_hipMemcpyPeer_Positive_ZeroSize") {
 
     constexpr int set_value_h = 21;
     std::fill_n(result.host_ptr(), element_count, set_value_h);
-    HIP_CHECK(hipMemcpyPeer(dst_alloc.ptr(), dst_device, src_alloc.ptr(), src_device, 0));
 
-    HIP_CHECK(
-        hipMemcpy(result.host_ptr(), dst_alloc.ptr(), allocation_size, hipMemcpyDeviceToHost));
+    hipError_t memcpy_err = hipSuccess;
+    BEGIN_CAPTURE_SYNC(memcpy_err, false);
+    HIP_CHECK_ERROR(hipMemcpyPeer(dst_alloc.ptr(), dst_device, src_alloc.ptr(), src_device, 0),
+                    memcpy_err);
+    END_CAPTURE_SYNC(memcpy_err);
 
-    HIP_CHECK(hipDeviceDisablePeerAccess(dst_device));
+    if (memcpy_err == hipSuccess) {
+      HIP_CHECK(
+          hipMemcpy(result.host_ptr(), dst_alloc.ptr(), allocation_size, hipMemcpyDeviceToHost));
 
-    ArrayFindIfNot(result.host_ptr(), expected_value, element_count);
+      HIP_CHECK(hipDeviceDisablePeerAccess(dst_device));
+
+      ArrayFindIfNot(result.host_ptr(), expected_value, element_count);
+    }
   } else {
     INFO("Peer access cannot be enabled between devices " << src_device << " " << dst_device);
   }