diff --git a/tests/src/hipPeerToPeer_simple.cpp b/tests/src/hipPeerToPeer_simple.cpp index aa305f92fa..4aaa6a452b 100644 --- a/tests/src/hipPeerToPeer_simple.cpp +++ b/tests/src/hipPeerToPeer_simple.cpp @@ -26,7 +26,9 @@ THE SOFTWARE. #include "hip_runtime.h" #include "test_common.h" -bool p_memcpyWithPeer = false; +bool p_memcpyWithPeer = false; // use the peer device for the P2P copy +bool p_mirrorPeers = false; // in addition to mapping current to peer space, map peer to current space. +int p_peerDevice = -1; // explicly specify which peer to use, else use p_gpuDevice + 1. void parseMyArguments(int argc, char *argv[]) { @@ -37,12 +39,19 @@ void parseMyArguments(int argc, char *argv[]) if (!strcmp(arg, "--memcpyWithPeer")) { p_memcpyWithPeer = true; + } else if (!strcmp(arg, "--mirrorPeers")) { + p_mirrorPeers = true; + } else if (!strcmp(arg, "--peerDevice")) { + if (++i >= argc || !HipTest::parseInt(argv[i], &p_peerDevice)) { + failed("Bad peerDevice argument"); + } } else { failed("Bad argument '%s'", arg); } }; }; + int main(int argc, char *argv[]) { parseMyArguments(argc, argv); @@ -50,23 +59,33 @@ int main(int argc, char *argv[]) int deviceCnt; HIPCHECK(hipGetDeviceCount(&deviceCnt)); - HIPCHECK(hipSetDevice(p_gpuDevice)); - int peerDevice = ((p_gpuDevice + 1) % deviceCnt); + int currentDevice = p_gpuDevice; + int peerDevice = (p_peerDevice == -1) ? ((currentDevice + 1) % deviceCnt) : p_peerDevice; - printf ("N=%zu device=%d peerDevice=%d (%d devices total)\n", N, p_gpuDevice, peerDevice, deviceCnt); + printf ("N=%zu device=%d peerDevice=%d (%d devices total)\n", N, currentDevice, peerDevice, deviceCnt); // Must be on a multi-gpu system: - assert (p_gpuDevice != peerDevice); + assert (currentDevice != peerDevice); int canAccessPeer; - HIPCHECK(hipDeviceCanAccessPeer(&canAccessPeer, p_gpuDevice, peerDevice)); - printf ("dev#%d canAccessPeer:#%d=%d\n", p_gpuDevice, peerDevice, canAccessPeer); + HIPCHECK(hipDeviceCanAccessPeer(&canAccessPeer, currentDevice, peerDevice)); + printf ("dev#%d canAccessPeer:#%d=%d\n", currentDevice, peerDevice, canAccessPeer); assert(canAccessPeer); + HIPCHECK(hipSetDevice(currentDevice)); HIPCHECK(hipDeviceEnablePeerAccess(peerDevice, 0)); + if (p_mirrorPeers) { + int canAccessPeer; + HIPCHECK(hipDeviceCanAccessPeer(&canAccessPeer, peerDevice, currentDevice)); + assert(canAccessPeer); + + HIPCHECK(hipSetDevice(peerDevice)); + HIPCHECK(hipDeviceEnablePeerAccess(currentDevice, 0)); + } + size_t Nbytes = N*sizeof(char); char *A_d0, *A_d1; @@ -75,7 +94,7 @@ int main(int argc, char *argv[]) A_h = (char*)malloc(Nbytes); // allocate and initialize memory on device0 - HIPCHECK (hipSetDevice(p_gpuDevice)); + HIPCHECK (hipSetDevice(currentDevice)); HIPCHECK (hipMalloc(&A_d0, Nbytes) ); HIPCHECK ( hipMemset(A_d0, memsetval, Nbytes) ); @@ -87,7 +106,7 @@ int main(int argc, char *argv[]) // Device0 push to device1, using P2P: - HIPCHECK (hipSetDevice(p_memcpyWithPeer ? peerDevice : p_gpuDevice)); + HIPCHECK (hipSetDevice(p_memcpyWithPeer ? peerDevice : currentDevice)); HIPCHECK (hipMemcpy(A_d1, A_d0, Nbytes, hipMemcpyDefault)); // Copy data back to host: diff --git a/tests/src/test_common.cpp b/tests/src/test_common.cpp index 332c2856d3..35e3d6d3f2 100644 --- a/tests/src/test_common.cpp +++ b/tests/src/test_common.cpp @@ -111,7 +111,7 @@ int parseStandardArguments(int argc, char *argv[], bool failOnUndefinedArg) failed("Bad iterations argument"); } - } else if (!strcmp(arg, "--gpu") || (!strcmp(arg, "-g"))) { + } else if (!strcmp(arg, "--gpu") || (!strcmp(arg, "-gpuDevice")) || (!strcmp(arg, "-g"))) { if (++i >= argc || !HipTest::parseInt(argv[i], &p_gpuDevice)) { failed("Bad gpuDevice argument"); }