From c442ee3990a2c8563977fdec170bfc41ce35ecee Mon Sep 17 00:00:00 2001 From: Jaydeep Patel Date: Wed, 11 Sep 2024 07:01:08 +0000 Subject: [PATCH] SWDEV-479616 - hipMemcpyParam2D and hipMemcpy2D does not support inter gpu copy without p2p access. Skip Unit_hipGetProcAddress_MemoryApisPeerToPeer if p2p access is not supported. Change-Id: I9f8598b9d1edf1d189583c5600b2e4d937564ea4 [ROCm/hip-tests commit: 067e3836ee0f57b87c736d476d9a06edabc3348a] --- .../memory/hipGetProcAddress_Memory_APIs.cc | 113 ++++++++++-------- .../catch/unit/memory/hipMemcpy2D.cc | 1 - .../catch/unit/memory/hipMemcpyParam2D.cc | 3 - 3 files changed, 62 insertions(+), 55 deletions(-) diff --git a/projects/hip-tests/catch/unit/memory/hipGetProcAddress_Memory_APIs.cc b/projects/hip-tests/catch/unit/memory/hipGetProcAddress_Memory_APIs.cc index 17cc396410..7570a1c523 100644 --- a/projects/hip-tests/catch/unit/memory/hipGetProcAddress_Memory_APIs.cc +++ b/projects/hip-tests/catch/unit/memory/hipGetProcAddress_Memory_APIs.cc @@ -3781,36 +3781,39 @@ TEST_CASE("Unit_hipGetProcAddress_MemoryApisMemcpy2DRelated") { if (deviceCount > 1) { HIP_CHECK(hipSetDevice(0)); - - char *sDevMem = nullptr; - size_t sPitch; - HIP_CHECK(hipMallocPitch(reinterpret_cast(&sDevMem), + int can_access_peer = 0; + HIP_CHECK(hipDeviceCanAccessPeer(&can_access_peer, 0, 1)); + if (can_access_peer) { + char *sDevMem = nullptr; + size_t sPitch; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&sDevMem), &sPitch, width, height)); - REQUIRE(sDevMem != nullptr); - HIP_CHECK(hipMemset2D(sDevMem, sPitch, value, width, height)); + REQUIRE(sDevMem != nullptr); + HIP_CHECK(hipMemset2D(sDevMem, sPitch, value, width, height)); - HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipSetDevice(1)); - char *dDevMem = nullptr; - size_t dPitch; - HIP_CHECK(hipMallocPitch(reinterpret_cast(&dDevMem), + char *dDevMem = nullptr; + size_t dPitch; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&dDevMem), &dPitch, width, height)); - REQUIRE(dDevMem != nullptr); - hip_Memcpy2D desc = {}; - desc.srcMemoryType = hipMemoryTypeDevice; - desc.srcDevice = hipDeviceptr_t(sDevMem); - desc.srcPitch = sPitch; - desc.dstMemoryType = hipMemoryTypeDevice; - desc.dstDevice = hipDeviceptr_t(dDevMem); - desc.dstPitch = dPitch; - desc.WidthInBytes = width*sizeof(char); - desc.Height = height; - HIP_CHECK(dyn_hipMemcpyParam2D_ptr(&desc)); + REQUIRE(dDevMem != nullptr); + hip_Memcpy2D desc = {}; + desc.srcMemoryType = hipMemoryTypeDevice; + desc.srcDevice = hipDeviceptr_t(sDevMem); + desc.srcPitch = sPitch; + desc.dstMemoryType = hipMemoryTypeDevice; + desc.dstDevice = hipDeviceptr_t(dDevMem); + desc.dstPitch = dPitch; + desc.WidthInBytes = width*sizeof(char); + desc.Height = height; + HIP_CHECK(dyn_hipMemcpyParam2D_ptr(&desc)); - REQUIRE(validateCharDeviceArray(dDevMem, N, value) == true); + REQUIRE(validateCharDeviceArray(dDevMem, N, value) == true); - HIP_CHECK(hipFree(sDevMem)); - HIP_CHECK(hipFree(dDevMem)); + HIP_CHECK(hipFree(sDevMem)); + HIP_CHECK(hipFree(dDevMem)); + } } } } @@ -3974,39 +3977,42 @@ TEST_CASE("Unit_hipGetProcAddress_MemoryApisMemcpy2DRelated") { if (deviceCount > 1) { HIP_CHECK(hipSetDevice(0)); - - char *sDevMem = nullptr; - size_t sPitch; - HIP_CHECK(hipMallocPitch(reinterpret_cast(&sDevMem), + int can_access_peer = 0; + HIP_CHECK(hipDeviceCanAccessPeer(&can_access_peer, 0, 1)); + if (can_access_peer) { + char *sDevMem = nullptr; + size_t sPitch; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&sDevMem), &sPitch, width, height)); - REQUIRE(sDevMem != nullptr); - HIP_CHECK(hipMemset2D(sDevMem, sPitch, value, width, height)); + REQUIRE(sDevMem != nullptr); + HIP_CHECK(hipMemset2D(sDevMem, sPitch, value, width, height)); - HIP_CHECK(hipSetDevice(1)); + HIP_CHECK(hipSetDevice(1)); - char *dDevMem = nullptr; - size_t dPitch; - HIP_CHECK(hipMallocPitch(reinterpret_cast(&dDevMem), + char *dDevMem = nullptr; + size_t dPitch; + HIP_CHECK(hipMallocPitch(reinterpret_cast(&dDevMem), &dPitch, width, height)); - REQUIRE(dDevMem != nullptr); - hip_Memcpy2D desc = {}; - desc.srcMemoryType = hipMemoryTypeDevice; - desc.srcDevice = hipDeviceptr_t(sDevMem); - desc.srcPitch = sPitch; - desc.dstMemoryType = hipMemoryTypeDevice; - desc.dstDevice = hipDeviceptr_t(dDevMem); - desc.dstPitch = dPitch; - desc.WidthInBytes = width*sizeof(char); - desc.Height = height; - HIP_CHECK(dyn_hipMemcpyParam2DAsync_ptr(&desc, stream)); + REQUIRE(dDevMem != nullptr); + hip_Memcpy2D desc = {}; + desc.srcMemoryType = hipMemoryTypeDevice; + desc.srcDevice = hipDeviceptr_t(sDevMem); + desc.srcPitch = sPitch; + desc.dstMemoryType = hipMemoryTypeDevice; + desc.dstDevice = hipDeviceptr_t(dDevMem); + desc.dstPitch = dPitch; + desc.WidthInBytes = width*sizeof(char); + desc.Height = height; + HIP_CHECK(dyn_hipMemcpyParam2DAsync_ptr(&desc, stream)); - HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipStreamSynchronize(stream)); - REQUIRE(validateCharDeviceArray(dDevMem, N, value) == true); + REQUIRE(validateCharDeviceArray(dDevMem, N, value) == true); - HIP_CHECK(hipStreamDestroy(stream)); - HIP_CHECK(hipFree(sDevMem)); - HIP_CHECK(hipFree(dDevMem)); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipFree(sDevMem)); + HIP_CHECK(hipFree(dDevMem)); + } } } } @@ -6729,7 +6735,12 @@ TEST_CASE("Unit_hipGetProcAddress_MemoryApisPeerToPeer") { int canAccessPeer = 0; HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, deviceId, peerDeviceId)); - REQUIRE(canAccessPeer == 1); + if (!canAccessPeer) { + std::string msg = "Skipped as peer access cannot be enabled between devices " + + std::to_string(deviceId) + " " + std::to_string(peerDeviceId); + HipTest::HIP_SKIP_TEST(msg.c_str()); + return; + } const int N = 16; const int Nbytes = N * sizeof(int); diff --git a/projects/hip-tests/catch/unit/memory/hipMemcpy2D.cc b/projects/hip-tests/catch/unit/memory/hipMemcpy2D.cc index 611cf13363..014605d267 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemcpy2D.cc +++ b/projects/hip-tests/catch/unit/memory/hipMemcpy2D.cc @@ -34,7 +34,6 @@ TEST_CASE("Unit_hipMemcpy2D_Positive_Basic") { SECTION("Device to Host") { Memcpy2DDeviceToHostShell(hipMemcpy2D); } SECTION("Device to Device") { - SECTION("Peer access disabled") { Memcpy2DDeviceToDeviceShell(hipMemcpy2D); } SECTION("Peer access enabled") { Memcpy2DDeviceToDeviceShell(hipMemcpy2D); } } diff --git a/projects/hip-tests/catch/unit/memory/hipMemcpyParam2D.cc b/projects/hip-tests/catch/unit/memory/hipMemcpyParam2D.cc index ec391c5a51..72cfdebf84 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemcpyParam2D.cc +++ b/projects/hip-tests/catch/unit/memory/hipMemcpyParam2D.cc @@ -36,9 +36,6 @@ TEST_CASE("Unit_hipMemcpyParam2D_Positive_Basic") { #endif SECTION("Device to Device") { - SECTION("Peer access disabled") { - Memcpy2DDeviceToDeviceShell(MemcpyParam2DAdapter()); - } SECTION("Peer access enabled") { Memcpy2DDeviceToDeviceShell(MemcpyParam2DAdapter()); }