AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-479616 - hipMemcpyParam2D and hipMemcpy2D does not support inter gpu copy without p2p access.

Browse Source 
Skip Unit_hipGetProcAddress_MemoryApisPeerToPeer if p2p access is not supported.

Change-Id: I9f8598b9d1edf1d189583c5600b2e4d937564ea4


[ROCm/hip-tests commit: 067e3836ee]
This commit is contained in:
Jaydeep Patel
2024-09-11 07:01:08 +00:00
committed by Rakesh Roy
parent 149d95ec0c
commit c442ee3990
3 changed files with 62 additions and 55 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip-tests/catch/unit/memory/hipGetProcAddress_Memory_APIs.cc
+62 -51
View File
@@ -3781,36 +3781,39 @@ TEST_CASE("Unit_hipGetProcAddress_MemoryApisMemcpy2DRelated") {
if (deviceCount > 1) { if (deviceCount > 1) {
HIP_CHECK(hipSetDevice(0)); HIP_CHECK(hipSetDevice(0));
int can_access_peer = 0;
char *sDevMem = nullptr; HIP_CHECK(hipDeviceCanAccessPeer(&can_access_peer, 0, 1));
size_t sPitch; if (can_access_peer) {
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&sDevMem), char *sDevMem = nullptr;
size_t sPitch;
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&sDevMem),
&sPitch, width, height)); &sPitch, width, height));
REQUIRE(sDevMem != nullptr); REQUIRE(sDevMem != nullptr);
HIP_CHECK(hipMemset2D(sDevMem, sPitch, value, width, height)); HIP_CHECK(hipMemset2D(sDevMem, sPitch, value, width, height));
HIP_CHECK(hipSetDevice(1)); HIP_CHECK(hipSetDevice(1));
char *dDevMem = nullptr; char *dDevMem = nullptr;
size_t dPitch; size_t dPitch;
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&dDevMem), HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&dDevMem),
&dPitch, width, height)); &dPitch, width, height));
REQUIRE(dDevMem != nullptr); REQUIRE(dDevMem != nullptr);
hip_Memcpy2D desc = {}; hip_Memcpy2D desc = {};
desc.srcMemoryType = hipMemoryTypeDevice; desc.srcMemoryType = hipMemoryTypeDevice;
desc.srcDevice = hipDeviceptr_t(sDevMem); desc.srcDevice = hipDeviceptr_t(sDevMem);
desc.srcPitch = sPitch; desc.srcPitch = sPitch;
desc.dstMemoryType = hipMemoryTypeDevice; desc.dstMemoryType = hipMemoryTypeDevice;
desc.dstDevice = hipDeviceptr_t(dDevMem); desc.dstDevice = hipDeviceptr_t(dDevMem);
desc.dstPitch = dPitch; desc.dstPitch = dPitch;
desc.WidthInBytes = width*sizeof(char); desc.WidthInBytes = width*sizeof(char);
desc.Height = height; desc.Height = height;
HIP_CHECK(dyn_hipMemcpyParam2D_ptr(&desc)); 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(sDevMem));
HIP_CHECK(hipFree(dDevMem)); HIP_CHECK(hipFree(dDevMem));
}
} }
} }
} }
@@ -3974,39 +3977,42 @@ TEST_CASE("Unit_hipGetProcAddress_MemoryApisMemcpy2DRelated") {
if (deviceCount > 1) { if (deviceCount > 1) {
HIP_CHECK(hipSetDevice(0)); HIP_CHECK(hipSetDevice(0));
int can_access_peer = 0;
char *sDevMem = nullptr; HIP_CHECK(hipDeviceCanAccessPeer(&can_access_peer, 0, 1));
size_t sPitch; if (can_access_peer) {
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&sDevMem), char *sDevMem = nullptr;
size_t sPitch;
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&sDevMem),
&sPitch, width, height)); &sPitch, width, height));
REQUIRE(sDevMem != nullptr); REQUIRE(sDevMem != nullptr);
HIP_CHECK(hipMemset2D(sDevMem, sPitch, value, width, height)); HIP_CHECK(hipMemset2D(sDevMem, sPitch, value, width, height));
HIP_CHECK(hipSetDevice(1)); HIP_CHECK(hipSetDevice(1));
char *dDevMem = nullptr; char *dDevMem = nullptr;
size_t dPitch; size_t dPitch;
HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&dDevMem), HIP_CHECK(hipMallocPitch(reinterpret_cast<void**>(&dDevMem),
&dPitch, width, height)); &dPitch, width, height));
REQUIRE(dDevMem != nullptr); REQUIRE(dDevMem != nullptr);
hip_Memcpy2D desc = {}; hip_Memcpy2D desc = {};
desc.srcMemoryType = hipMemoryTypeDevice; desc.srcMemoryType = hipMemoryTypeDevice;
desc.srcDevice = hipDeviceptr_t(sDevMem); desc.srcDevice = hipDeviceptr_t(sDevMem);
desc.srcPitch = sPitch; desc.srcPitch = sPitch;
desc.dstMemoryType = hipMemoryTypeDevice; desc.dstMemoryType = hipMemoryTypeDevice;
desc.dstDevice = hipDeviceptr_t(dDevMem); desc.dstDevice = hipDeviceptr_t(dDevMem);
desc.dstPitch = dPitch; desc.dstPitch = dPitch;
desc.WidthInBytes = width*sizeof(char); desc.WidthInBytes = width*sizeof(char);
desc.Height = height; desc.Height = height;
HIP_CHECK(dyn_hipMemcpyParam2DAsync_ptr(&desc, stream)); 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(hipStreamDestroy(stream));
HIP_CHECK(hipFree(sDevMem)); HIP_CHECK(hipFree(sDevMem));
HIP_CHECK(hipFree(dDevMem)); HIP_CHECK(hipFree(dDevMem));
}
} }
} }
} }
@@ -6729,7 +6735,12 @@ TEST_CASE("Unit_hipGetProcAddress_MemoryApisPeerToPeer") {
int canAccessPeer = 0; int canAccessPeer = 0;
HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, deviceId, peerDeviceId)); 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 N = 16;
const int Nbytes = N * sizeof(int); const int Nbytes = N * sizeof(int);
projects/hip-tests/catch/unit/memory/hipMemcpy2D.cc
-1
View File
@@ -34,7 +34,6 @@ TEST_CASE("Unit_hipMemcpy2D_Positive_Basic") {
SECTION("Device to Host") { Memcpy2DDeviceToHostShell<async>(hipMemcpy2D); } SECTION("Device to Host") { Memcpy2DDeviceToHostShell<async>(hipMemcpy2D); }
SECTION("Device to Device") { SECTION("Device to Device") {
SECTION("Peer access disabled") { Memcpy2DDeviceToDeviceShell<async, false>(hipMemcpy2D); }
SECTION("Peer access enabled") { Memcpy2DDeviceToDeviceShell<async, true>(hipMemcpy2D); } SECTION("Peer access enabled") { Memcpy2DDeviceToDeviceShell<async, true>(hipMemcpy2D); }
} }
projects/hip-tests/catch/unit/memory/hipMemcpyParam2D.cc
-3
View File
@@ -36,9 +36,6 @@ TEST_CASE("Unit_hipMemcpyParam2D_Positive_Basic") {
#endif #endif
SECTION("Device to Device") { SECTION("Device to Device") {
SECTION("Peer access disabled") {
Memcpy2DDeviceToDeviceShell<async, false>(MemcpyParam2DAdapter<async>());
}
SECTION("Peer access enabled") { SECTION("Peer access enabled") {
Memcpy2DDeviceToDeviceShell<async, true>(MemcpyParam2DAdapter<async>()); Memcpy2DDeviceToDeviceShell<async, true>(MemcpyParam2DAdapter<async>());
} }