From 7bd95c3da611b139566fd487ff97fad3134a34df Mon Sep 17 00:00:00 2001 From: Sebastian Luzynski Date: Tue, 16 Apr 2024 15:42:56 +0000 Subject: [PATCH] SWDEV-456701 - Fix memcpy D2D test failure This change improves the device to device memcpy test to not rely on delay kernel. It uses blocking context instead. Signed-off-by: Sebastian Luzynski Change-Id: I18090cb7febe806104a343acfd0bb535a96ae04d [ROCm/hip-tests commit: ea96f486eff9813d13c98b202b3075e594db228e] --- .../catch/unit/memory/hipMemcpy3D.cc | 28 +++++++++++++------ 1 file changed, 20 insertions(+), 8 deletions(-) diff --git a/projects/hip-tests/catch/unit/memory/hipMemcpy3D.cc b/projects/hip-tests/catch/unit/memory/hipMemcpy3D.cc index b62f23c8a6..d3211a9113 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemcpy3D.cc +++ b/projects/hip-tests/catch/unit/memory/hipMemcpy3D.cc @@ -60,19 +60,31 @@ TEST_CASE("Unit_hipMemcpy3D_Positive_Synchronization_Behavior") { SECTION("Device to Pinned Host") { Memcpy3DDtoHPinnedSyncBehavior(Memcpy3DWrapper<>, true); } - SECTION("Device to Device") { -#if HT_NVIDIA - Memcpy3DDtoDSyncBehavior(Memcpy3DWrapper<>, false); -#else - Memcpy3DDtoDSyncBehavior(Memcpy3DWrapper<>, true); -#endif - } - #if HT_NVIDIA // Disabled on AMD due to defect - EXSWHTEC-232 SECTION("Host to Host") { Memcpy3DHtoHSyncBehavior(Memcpy3DWrapper<>, true); } #endif } +TEST_CASE("Unit_hipMemcpy3D_Positive_DeviceToDevice_Synchronization_Behavior") { + CHECK_IMAGE_SUPPORT + + LinearAllocGuard3D src_alloc(make_hipExtent(32 * sizeof(int), 32, 8)); + LinearAllocGuard3D dst_alloc(make_hipExtent(32 * sizeof(int), 32, 8)); + HipTest::BlockingContext b_context{nullptr}; + hipStream_t kernel_stream{nullptr}; + + auto parms = GetMemcpy3DParms(dst_alloc.pitched_ptr(), make_hipPos(0, 0, 0), src_alloc.pitched_ptr(), + make_hipPos(0, 0, 0), dst_alloc.extent(), hipMemcpyDeviceToDevice); + + b_context.block_stream(); + REQUIRE(b_context.is_blocked()); + HIP_CHECK(hipMemcpy3D(&parms)); + HIP_CHECK_ERROR(hipStreamQuery(kernel_stream), hipErrorNotReady); + b_context.unblock_stream(); + HIP_CHECK(hipDeviceSynchronize()); + REQUIRE(hipStreamQuery(kernel_stream) == hipSuccess); +} + TEST_CASE("Unit_hipMemcpy3D_Positive_Parameters") { CHECK_IMAGE_SUPPORT