         for (int j=0; j<numGPUs; j++)
         {
             int access;
-            if(p2p) {
-                cudaDeviceCanAccessPeer(&access,g[i],g[j]);
-                if (access)
+            if (i != j) {
+                if(p2p) {
+                    cudaDeviceCanAccessPeer(&access,g[i],g[j]);
+                    if (access)
+                    {
+                        cudaSetDevice(g[i]);
+                        cudaDeviceEnablePeerAccess(g[j],0);
+                        cudaCheckError();
+                        cudaSetDevice(g[j]);
+                        cudaDeviceEnablePeerAccess(g[i],0);
+                        cudaCheckError();
+                    }
+                }
+
+                cudaSetDevice(g[i]);
+                cudaDeviceSynchronize();
+                cudaCheckError();
+                delay<<<1,1>>>(NULL);
+                cudaEventRecord(start[i]);
+
+                for (int r=0; r<repeat; r++)
+                {
+                    cudaMemcpyPeerAsync(buffers[i],g[i],buffers[j],g[j],
+                                        sizeof(int)*numElems,stream0[i]);
+                    cudaMemcpyPeerAsync(buffers[j],g[j],buffers[i],g[i],
+                                        sizeof(int)*numElems,stream1[i]);
+                }
+
+                cudaEventRecord(stop[i]);
+                cudaDeviceSynchronize();
+                cudaCheckError();
+
+                float time_ms;
+                cudaEventElapsedTime(&time_ms,start[i],stop[i]);
+                double time_s=time_ms/1e3;
+
+                double gb=2.0*numElems*sizeof(int)*repeat/(double)1e9;
+                bandwidthMatrix[i*numGPUs+j]=gb/time_s;
+                if(p2p && access)
                 {
                     cudaSetDevice(g[i]);
-                    cudaDeviceEnablePeerAccess(g[j],0);
-                    cudaCheckError();
+                    cudaDeviceDisablePeerAccess(g[j]);
                     cudaSetDevice(g[j]);
-                    cudaDeviceEnablePeerAccess(g[i],0);
-                    cudaCheckError();
-                }
-            }
-
-            cudaSetDevice(g[i]);
-            cudaDeviceSynchronize();
-            cudaCheckError();
-            delay<<<1,1>>>(NULL);
-            cudaEventRecord(start[i]);
-
-            for (int r=0; r<repeat; r++)
-            {
-                cudaMemcpyPeerAsync(buffers[i],g[i],buffers[j],g[j],
-                                    sizeof(int)*numElems,stream0[i]);
-                cudaMemcpyPeerAsync(buffers[j],g[j],buffers[i],g[i],
-                                    sizeof(int)*numElems,stream1[i]);
-            }
-
-            cudaEventRecord(stop[i]);
-            cudaDeviceSynchronize();
-            cudaCheckError();
-
-            float time_ms;
-            cudaEventElapsedTime(&time_ms,start[i],stop[i]);
-            double time_s=time_ms/1e3;
-
-            double gb=2.0*numElems*sizeof(int)*repeat/(double)1e9;
-            if(i==j) gb*=2;  //must count both the read and the write here
-            bandwidthMatrix[i*numGPUs+j]=gb/time_s;
-            if(p2p && access)
-            {
-                cudaSetDevice(g[i]);
-                cudaDeviceDisablePeerAccess(g[j]);
-                cudaSetDevice(g[j]);
-                cudaDeviceDisablePeerAccess(g[i]);
+                    cudaDeviceDisablePeerAccess(g[i]);
+                }
             }
         }
     }

        printf("%6d ",g[i]);

        for (int j=0; j<numGPUs; j++)
        {
            printf("%6.02f ", latencyMatrix[i*numGPUs+j]);
        }

        printf("\n");
    }

        printf("%6d ",g[i]);

        for (int j=0; j<numGPUs; j++)
            if (i == j)
                printf("%6c", 'X');
            else
                printf("%6.02f ", bandwidthMatrix[i*numGPUs+j]);

        printf("\n");
    }
