From 08e2b6a15df3a4e96d42a08aba558135836bdda3 Mon Sep 17 00:00:00 2001 From: Vladana Stojiljkovic Date: Fri, 27 Sep 2024 12:53:31 +0200 Subject: [PATCH] SWDEV-486969 - Reuse existing tests to verify behavior of capturing sync Memcpy APis Change-Id: I7e94743d6957ffaadae9cff297e3f0d93f9ff806 [ROCm/hip-tests commit: 6c5468f802976bb592f4c2e810e04dc0ea00d346] --- .../unit/memory/array_memcpy_tests_common.hh | 1 - .../catch/unit/memory/hipDrvMemcpy3D.cc | 26 +++++++++-- .../catch/unit/memory/hipMemcpy2D.cc | 18 ++++++++ .../catch/unit/memory/hipMemcpy2DFromArray.cc | 28 ++++++++++-- .../catch/unit/memory/hipMemcpy2DToArray.cc | 30 ++++++++++--- .../catch/unit/memory/hipMemcpy3D.cc | 36 ++++++++++++--- .../catch/unit/memory/hipMemcpyAtoA.cc | 18 +++++--- .../catch/unit/memory/hipMemcpyAtoD.cc | 20 ++++++--- .../catch/unit/memory/hipMemcpyAtoH.cc | 28 +++++++++++- .../catch/unit/memory/hipMemcpyDtoD.cc | 27 ++++++----- .../catch/unit/memory/hipMemcpyFromSymbol.cc | 45 ++++++++++++++----- .../catch/unit/memory/hipMemcpyHtoA.cc | 24 +++++++++- .../catch/unit/memory/hipMemcpyParam2D.cc | 22 +++++++++ .../catch/unit/memory/hipMemcpyPeer.cc | 35 ++++++++++----- 14 files changed, 291 insertions(+), 67 deletions(-) diff --git a/projects/hip-tests/catch/unit/memory/array_memcpy_tests_common.hh b/projects/hip-tests/catch/unit/memory/array_memcpy_tests_common.hh index 71a02a8aca..0e33413ea3 100644 --- a/projects/hip-tests/catch/unit/memory/array_memcpy_tests_common.hh +++ b/projects/hip-tests/catch/unit/memory/array_memcpy_tests_common.hh @@ -172,7 +172,6 @@ template 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); diff --git a/projects/hip-tests/catch/unit/memory/hipDrvMemcpy3D.cc b/projects/hip-tests/catch/unit/memory/hipDrvMemcpy3D.cc index 6e33de8c0a..6de7dae744 100644 --- a/projects/hip-tests/catch/unit/memory/hipDrvMemcpy3D.cc +++ b/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(DrvMemcpy3DWrapper); + Memcpy3DZeroWidthHeightDepth(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(DrvMemcpy3DWrapper); } - SECTION("Array from/to Device") { DrvMemcpy3DArrayDeviceShell(DrvMemcpy3DWrapper); } + SECTION("Array from/to Host") { DrvMemcpy3DArrayHostShell(DrvMemcpy3DWrapper<>); } + SECTION("Array from/to Device") { DrvMemcpy3DArrayDeviceShell(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 device_alloc(extent); + LinearAllocGuard 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(¶ms), memcpy_err); + END_CAPTURE_SYNC(memcpy_err); } \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/memory/hipMemcpy2D.cc b/projects/hip-tests/catch/unit/memory/hipMemcpy2D.cc index 014605d267..7a1e995d81 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemcpy2D.cc +++ b/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 device_alloc(width, height); + LinearAllocGuard 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); +} diff --git a/projects/hip-tests/catch/unit/memory/hipMemcpy2DFromArray.cc b/projects/hip-tests/catch/unit/memory/hipMemcpy2DFromArray.cc index 013aa39222..983e53471f 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemcpy2DFromArray.cc +++ b/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( - std::bind(hipMemcpy2DFromArray, _1, _2, _3, 0, 0, 0, height, hipMemcpyDeviceToHost), - width, height); + Memcpy2DFromArrayZeroWidthHeight(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 A_d(make_hipExtent(width, height, 0), hipArrayDefault); + LinearAllocGuard A_h(LinearAllocs::hipHostMalloc, size); + LinearAllocGuard 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); +} diff --git a/projects/hip-tests/catch/unit/memory/hipMemcpy2DToArray.cc b/projects/hip-tests/catch/unit/memory/hipMemcpy2DToArray.cc index 544578557f..5333b06b9f 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemcpy2DToArray.cc +++ b/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( - std::bind(hipMemcpy2DToArray, _1, 0, 0, _2, _3, 0, height, hipMemcpyHostToDevice), width, - height); + Memcpy2DToArrayZeroWidthHeight(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( - std::bind(hipMemcpy2DToArray, _1, 0, 0, _2, _3, 0, height, hipMemcpyDeviceToDevice), - width, height); + Memcpy2DToArrayZeroWidthHeight(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 array_alloc(make_hipExtent(width, height, 0), hipArrayDefault); + LinearAllocGuard 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); +} diff --git a/projects/hip-tests/catch/unit/memory/hipMemcpy3D.cc b/projects/hip-tests/catch/unit/memory/hipMemcpy3D.cc index d3211a9113..7d75bbcfbd 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemcpy3D.cc +++ b/projects/hip-tests/catch/unit/memory/hipMemcpy3D.cc @@ -41,7 +41,9 @@ TEST_CASE("Unit_hipMemcpy3D_Positive_Basic") { SECTION("Peer access disabled") { Memcpy3DDeviceToDeviceShell(Memcpy3DWrapper<>); } - SECTION("Peer access enabled") { Memcpy3DDeviceToDeviceShell(Memcpy3DWrapper<>); } + SECTION("Peer access enabled") { + Memcpy3DDeviceToDeviceShell(Memcpy3DWrapper<>); + } } SECTION("Host to Device") { Memcpy3DHostToDeviceShell(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(Memcpy3DWrapper); + Memcpy3DZeroWidthHeightDepth(Memcpy3DWrapper<>); } TEST_CASE("Unit_hipMemcpy3D_Positive_Array") { CHECK_IMAGE_SUPPORT constexpr bool async = false; - SECTION("Array from/to Host") { Memcpy3DArrayHostShell(Memcpy3DWrapper); } + SECTION("Array from/to Host") { Memcpy3DArrayHostShell(Memcpy3DWrapper<>); } #if HT_NVIDIA // Disabled on AMD due to defect - EXSWHTEC-238 - SECTION("Array from/to Device") { Memcpy3DArrayDeviceShell(Memcpy3DWrapper); } + SECTION("Array from/to Device") { Memcpy3DArrayDeviceShell(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 dev_alloc(extent); + LinearAllocGuard 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(¶ms), memcpy_err); + END_CAPTURE_SYNC(memcpy_err); +} diff --git a/projects/hip-tests/catch/unit/memory/hipMemcpyAtoA.cc b/projects/hip-tests/catch/unit/memory/hipMemcpyAtoA.cc index 5bb28de94e..5e23c33fff 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemcpyAtoA.cc +++ b/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)); diff --git a/projects/hip-tests/catch/unit/memory/hipMemcpyAtoD.cc b/projects/hip-tests/catch/unit/memory/hipMemcpyAtoD.cc index 1f41d35942..ea6e1dff2a 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemcpyAtoD.cc +++ b/projects/hip-tests/catch/unit/memory/hipMemcpyAtoD.cc @@ -62,12 +62,20 @@ TEST_CASE("Unit_hipMemcpyAtoD_Basic") { hipChannelFormatDesc desc = hipCreateChannelDesc(); 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)); diff --git a/projects/hip-tests/catch/unit/memory/hipMemcpyAtoH.cc b/projects/hip-tests/catch/unit/memory/hipMemcpyAtoH.cc index 4933454297..0df76407cf 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemcpyAtoH.cc +++ b/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(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 array_alloc(make_hipExtent(width, height, 0), flag); + LinearAllocGuard 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); +} diff --git a/projects/hip-tests/catch/unit/memory/hipMemcpyDtoD.cc b/projects/hip-tests/catch/unit/memory/hipMemcpyDtoD.cc index 6eb6c5342f..4c01b80715 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemcpyDtoD.cc +++ b/projects/hip-tests/catch/unit/memory/hipMemcpyDtoD.cc @@ -77,19 +77,22 @@ TEMPLATE_TEST_CASE("Unit_hipMemcpyDtoD_Basic", "", HipTest::checkVectorADD(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(X_d), - static_cast(Y_d), Z_d, NUM_ELM); - HIP_CHECK(hipGetLastError()); - HIP_CHECK(hipMemcpyDtoH(C_h, (hipDeviceptr_t)Z_d, Nbytes)); - HIP_CHECK(hipDeviceSynchronize()); - HipTest::checkVectorADD(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(X_d), static_cast(Y_d), Z_d, + NUM_ELM); + HIP_CHECK(hipGetLastError()); + HIP_CHECK(hipMemcpyDtoH(C_h, (hipDeviceptr_t)Z_d, Nbytes)); + HIP_CHECK(hipDeviceSynchronize()); + HipTest::checkVectorADD(A_h, B_h, C_h, NUM_ELM); + } HipTest::freeArrays(A_d, B_d, C_d, A_h, B_h, C_h, false); HIP_CHECK(hipFree(X_d)); diff --git a/projects/hip-tests/catch/unit/memory/hipMemcpyFromSymbol.cc b/projects/hip-tests/catch/unit/memory/hipMemcpyFromSymbol.cc index e124ab19b3..0cf024684a 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemcpyFromSymbol.cc +++ b/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]); + } } } } diff --git a/projects/hip-tests/catch/unit/memory/hipMemcpyHtoA.cc b/projects/hip-tests/catch/unit/memory/hipMemcpyHtoA.cc index 6243b452f1..c01769f953 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemcpyHtoA.cc +++ b/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 array_alloc(make_hipExtent(width, 0, 0), hipArrayDefault); + LinearAllocGuard 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); +} diff --git a/projects/hip-tests/catch/unit/memory/hipMemcpyParam2D.cc b/projects/hip-tests/catch/unit/memory/hipMemcpyParam2D.cc index 72cfdebf84..4630317f94 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemcpyParam2D.cc +++ b/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 device_alloc(cols, rows); + LinearAllocGuard host_alloc(LinearAllocs::hipHostMalloc, device_alloc.pitch() * rows); + + hip_Memcpy2D params = {}; + memset(¶ms, 0x0, sizeof(hip_Memcpy2D)); + + InitializeMemcpy2DParams(¶ms, 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(¶ms), memcpy_err); + END_CAPTURE_SYNC(memcpy_err); +} diff --git a/projects/hip-tests/catch/unit/memory/hipMemcpyPeer.cc b/projects/hip-tests/catch/unit/memory/hipMemcpyPeer.cc index 6daf341e6b..26f1b70761 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemcpyPeer.cc +++ b/projects/hip-tests/catch/unit/memory/hipMemcpyPeer.cc @@ -77,15 +77,21 @@ TEST_CASE("Unit_hipMemcpyPeer_Positive_Default") { VectorSet<<>>(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); }