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

        for (int j=0; j<numGPUs; j++)
        {
            int access;
            if(p2p) {
                cudaDeviceCanAccessPeer(&access,g[i],g[j]);
                if (access)
                {
                    cudaDeviceEnablePeerAccess(g[j],0 );
                    cudaCheckError();
                }
            }

            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);
            }

            cudaEventRecord(stop[i]);
            cudaDeviceSynchronize();
            cudaCheckError();

            float time_ms;
            cudaEventElapsedTime(&time_ms,start[i],stop[i]);
            double time_s=time_ms/1e3;

            double gb=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)
            {
                cudaDeviceDisablePeerAccess(g[j]);
                cudaCheckError();
            }
        }
    }

    printf("   D\\D");

    for (int j=0; j<numGPUs; j++)
    {
        printf("%6d ", g[j]);
    }

    printf("\n");


        for (int j=0; j<numGPUs; j++)
        {
            int 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]);
                    cudaDeviceDisablePeerAccess(g[j]);
                    cudaSetDevice(g[j]);
                    cudaDeviceDisablePeerAccess(g[i]);
                }
            }
        }
    }
