add simple negative P2P tests

[ROCm/hip commit: 4a0c0479d4]

projects/hip/tests/src/CMakeLists.txt

@@ -224,7 +224,7 @@ make_named_test (hipMultiThreadDevice "hipMultiThreadDevice-pyramid" --tests 0x4
 make_named_test (hipMultiThreadDevice "hipMultiThreadDevice-nearzero" --tests 0x10) 
 
 if (${HIP_MULTI_GPU})
-    make_test(hipPeerToPeer_simple )                  # use current device for copy, this fails. 
+    make_test(hipPeerToPeer_simple " ")                  # use current device for copy, this fails. 
     make_test(hipPeerToPeer_simple --memcpyWithPeer)  
     make_test(hipPeerToPeer_simple --mirrorPeers)    # mirror mapping: test to ensure mirror doesn't destroy orig mapping.
 

projects/hip/tests/src/hipMultiThreadDevice.cpp

@@ -116,12 +116,12 @@ int main(int argc, char *argv[])
     /*disable, this takess a while and if the next one works then no need to run serial*/
     if (1 && (p_tests & 0x2)) {
         printf ("\ntest 0x2 : serialized multiThread_pyramid(1) \n");
-        multiThread_pyramid(true, 10);
+        multiThread_pyramid(true, 3);
     }
 
     if (p_tests & 0x4) {
         printf ("\ntest 0x4 : parallel multiThread_pyramid(1) \n");
-        multiThread_pyramid(false, 10);
+        multiThread_pyramid(false, 3);
     }
 
     //if (p_tests & 0x8) {

projects/hip/tests/src/hipPeerToPeer_simple.cpp

@@ -29,6 +29,10 @@ 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.
 
+
+int g_currentDevice;
+int g_peerDevice;
+
 void parseMyArguments(int argc, char *argv[])
 {
     int more_argc = HipTest::parseStandardArguments(argc, argv, false);
@@ -51,45 +55,50 @@ void parseMyArguments(int argc, char *argv[])
 };
 
 
+// Sets globals g_currentDevice, g_peerDevice
+void setupPeerTests() 
+{
+    int deviceCnt;
+
+    HIPCHECK(hipGetDeviceCount(&deviceCnt));
+
+    g_currentDevice = p_gpuDevice;
+    g_peerDevice    = (p_peerDevice == -1) ? ((g_currentDevice + 1) % deviceCnt) : p_peerDevice;
+
+    printf ("N=%zu  device=%d peerDevice=%d (%d devices total)\n", N, g_currentDevice, g_peerDevice, deviceCnt);
+
+    // Must be on a multi-gpu system:
+    assert (g_currentDevice != g_peerDevice);
+
+    int canAccessPeer;
+    HIPCHECK(hipDeviceCanAccessPeer(&canAccessPeer, g_currentDevice, g_peerDevice));
+    printf ("dev#%d canAccessPeer:#%d=%d\n", g_currentDevice, g_peerDevice, canAccessPeer);
+
+    assert(canAccessPeer);
+
+    HIPCHECK (hipSetDevice(g_currentDevice));
+    HIPCHECK(hipDeviceReset());
+    HIPCHECK (hipSetDevice(g_peerDevice));
+    HIPCHECK(hipDeviceReset());
+}
+
 //---
 // Test which enables peer2peer first, then allocates the memory.
 void enablePeerFirst()
 {
     printf ("\n==testing: %s\n", __func__);
 
-    int deviceCnt;
 
-    HIPCHECK(hipGetDeviceCount(&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, currentDevice, peerDevice, deviceCnt);
-
-    // Must be on a multi-gpu system:
-    assert (currentDevice != peerDevice);
-
-    int canAccessPeer;
-    HIPCHECK(hipDeviceCanAccessPeer(&canAccessPeer, currentDevice, peerDevice));
-    printf ("dev#%d canAccessPeer:#%d=%d\n", currentDevice, peerDevice, canAccessPeer);
-
-    assert(canAccessPeer);
-
-    HIPCHECK (hipSetDevice(currentDevice));
-    HIPCHECK(hipDeviceReset());
-    HIPCHECK (hipSetDevice(peerDevice));
-    HIPCHECK(hipDeviceReset());
-
-    HIPCHECK(hipSetDevice(currentDevice));
-    HIPCHECK(hipDeviceEnablePeerAccess(peerDevice, 0));
+    HIPCHECK(hipSetDevice(g_currentDevice));
+    HIPCHECK(hipDeviceEnablePeerAccess(g_peerDevice, 0));
 
     if (p_mirrorPeers) {
         int canAccessPeer;
-        HIPCHECK(hipDeviceCanAccessPeer(&canAccessPeer, peerDevice, currentDevice));
+        HIPCHECK(hipDeviceCanAccessPeer(&canAccessPeer, g_peerDevice, g_currentDevice));
         assert(canAccessPeer);
 
-        HIPCHECK(hipSetDevice(peerDevice));
-            HIPCHECK(hipDeviceEnablePeerAccess(currentDevice, 0));
+        HIPCHECK(hipSetDevice(g_peerDevice));
+            HIPCHECK(hipDeviceEnablePeerAccess(g_currentDevice, 0));
         }
 
         size_t Nbytes = N*sizeof(char);
@@ -100,23 +109,23 @@ void enablePeerFirst()
         A_h = (char*)malloc(Nbytes);
 
         // allocate and initialize memory on device0
-        HIPCHECK (hipSetDevice(currentDevice));
+        HIPCHECK (hipSetDevice(g_currentDevice));
         HIPCHECK (hipMalloc(&A_d0, Nbytes) );
         HIPCHECK ( hipMemset(A_d0, memsetval, Nbytes) ); 
 
         // allocate and initialize memory on peer device
-        HIPCHECK (hipSetDevice(peerDevice));
+        HIPCHECK (hipSetDevice(g_peerDevice));
         HIPCHECK (hipMalloc(&A_d1, Nbytes) );
         HIPCHECK ( hipMemset(A_d1, 0x13, Nbytes) ); 
 
 
 
         // Device0 push to device1, using P2P:
-        HIPCHECK (hipSetDevice(p_memcpyWithPeer ? peerDevice : currentDevice));
+        HIPCHECK (hipSetDevice(p_memcpyWithPeer ? g_peerDevice : g_currentDevice));
         HIPCHECK (hipMemcpy(A_d1, A_d0, Nbytes, hipMemcpyDefault));
 
         // Copy data back to host:
-        HIPCHECK (hipSetDevice(peerDevice));
+        HIPCHECK (hipSetDevice(g_peerDevice));
         HIPCHECK (hipMemcpy(A_h, A_d1, Nbytes, hipMemcpyDeviceToHost));
 
         // Check host data:
@@ -128,35 +137,14 @@ void enablePeerFirst()
     }
 
 
-    //---
-    // Test which allocated memory first, then enables peer2peer.
-    // Enabling peer needs to scan all allocated memory and enable peer access.
-    void allocMemoryFirst()
-    {
-        printf ("\n==testing: %s\n", __func__);
-        int deviceCnt;
-
-        HIPCHECK(hipGetDeviceCount(&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, currentDevice, peerDevice, deviceCnt);
-
-        // Must be on a multi-gpu system:
-        assert (currentDevice != peerDevice);
-
-        int canAccessPeer;
-        HIPCHECK(hipDeviceCanAccessPeer(&canAccessPeer, currentDevice, peerDevice));
-        printf ("dev#%d canAccessPeer:#%d=%d\n", currentDevice, peerDevice, canAccessPeer);
-
-        assert(canAccessPeer);
-
-        HIPCHECK (hipSetDevice(currentDevice));
-    HIPCHECK(hipDeviceReset());
-    HIPCHECK (hipSetDevice(peerDevice));
-    HIPCHECK(hipDeviceReset());
+//---
+// Test which allocated memory first, then enables peer2peer.
+// Enabling peer needs to scan all allocated memory and enable peer access.
+void allocMemoryFirst()
+{
+    printf ("\n==testing: %s\n", __func__);
 
+    setupPeerTests();
 
     size_t Nbytes = N*sizeof(char);
 
@@ -167,39 +155,39 @@ void enablePeerFirst()
 
     //---
     // allocate and initialize memory on device0
-    HIPCHECK (hipSetDevice(currentDevice));
+    HIPCHECK (hipSetDevice(g_currentDevice));
     HIPCHECK (hipMalloc(&A_d0, Nbytes) );
     HIPCHECK ( hipMemset(A_d0, memsetval, Nbytes) ); 
 
     // allocate and initialize memory on peer device
-    HIPCHECK (hipSetDevice(peerDevice));
+    HIPCHECK (hipSetDevice(g_peerDevice));
     HIPCHECK (hipMalloc(&A_d1, Nbytes) );
     HIPCHECK ( hipMemset(A_d1, 0x13, Nbytes) ); 
 
 
     //---
     //Enable peer access, for memory already allocated:
-    HIPCHECK(hipSetDevice(currentDevice));
-    HIPCHECK(hipDeviceEnablePeerAccess(peerDevice, 0));
+    HIPCHECK(hipSetDevice(g_currentDevice));
+    HIPCHECK(hipDeviceEnablePeerAccess(g_peerDevice, 0));
 
     if (p_mirrorPeers) {
         int canAccessPeer;
-        HIPCHECK(hipDeviceCanAccessPeer(&canAccessPeer, peerDevice, currentDevice));
+        HIPCHECK(hipDeviceCanAccessPeer(&canAccessPeer, g_peerDevice, g_currentDevice));
         assert(canAccessPeer);
 
-        HIPCHECK(hipSetDevice(peerDevice));
-        HIPCHECK(hipDeviceEnablePeerAccess(currentDevice, 0));
+        HIPCHECK(hipSetDevice(g_peerDevice));
+        HIPCHECK(hipDeviceEnablePeerAccess(g_currentDevice, 0));
     }
 
 
     //---
     // Copies to test functionality:
     // Device0 push to device1, using P2P:
-    HIPCHECK (hipSetDevice(p_memcpyWithPeer ? peerDevice : currentDevice));
+    HIPCHECK (hipSetDevice(p_memcpyWithPeer ? g_peerDevice : g_currentDevice));
     HIPCHECK (hipMemcpy(A_d1, A_d0, Nbytes, hipMemcpyDefault));
 
     // Copy data back to host:
-    HIPCHECK (hipSetDevice(peerDevice));
+    HIPCHECK (hipSetDevice(g_peerDevice));
     HIPCHECK (hipMemcpy(A_h, A_d1, Nbytes, hipMemcpyDeviceToHost));
 
 
@@ -212,6 +200,40 @@ void enablePeerFirst()
     }
 }
 
+void simpleNegative()
+{
+    printf ("\n==testing: %s\n", __func__);
+
+    setupPeerTests();
+
+    int deviceId;
+    HIPCHECK (hipGetDevice(&deviceId));
+
+    //---
+    //-- self is not a peer
+    int canAccessPeer;
+    hipError_t e = hipDeviceCanAccessPeer(&canAccessPeer, deviceId, deviceId);
+    HIPASSERT( e == hipSuccess);  // no error returned, it doesn't hurt to ask.
+    HIPASSERT (canAccessPeer == 0); // but self is not a peer.
+
+    e = hipSuccess; 
+    //---
+    // Enable same device twice in a row:
+    HIPCHECK(hipSetDevice(g_currentDevice));
+    HIPCHECK(hipDeviceEnablePeerAccess(g_peerDevice, 0));
+    e =(hipDeviceEnablePeerAccess(g_peerDevice, 0));
+    HIPASSERT (e == hipErrorPeerAccessAlreadyEnabled);
+
+    //---
+    // try disabling twice in a row
+    HIPCHECK(hipDeviceDisablePeerAccess(g_peerDevice));
+    e =(hipDeviceDisablePeerAccess(g_peerDevice));
+    HIPASSERT (e == hipErrorPeerAccessNotEnabled);
+
+    
+    // More tests here:
+}
+
 
 
 int main(int argc, char *argv[])
@@ -226,5 +248,9 @@ int main(int argc, char *argv[])
         allocMemoryFirst();
     }
 
+    if (p_tests & 0x4) {
+        simpleNegative();
+    }
+
     passed();
 }