cudagpus是什么_如何在cuda中复制不同gpus之间的内存

将数据从一个GPU传输到另一个GPU通常需要通过主机内存进行“分段”.例外情况是GPU和系统拓扑支持对等(P2P)访问并且已明确启用P2P.在这种情况下,数据传输可以直接通过PCIE总线从一个GPU流向另一个GPU.

在任何一种情况下(有或没有P2P可用/启用),典型的cuda runtime API call将是cudaMemcpyPeer / cudaMemcpyPeerAsync,如cuda p2pBandwidthLatencyTest sample code中所示.

在Windows上,P2P的一个要求是在TCC模式下驱动程序支持两个设备.大多数情况下,TCC模式不是GeForce GPU的可用选项(最近,使用CUDA 7.5RC工具包中提供的驱动程序和运行时对GeForce Titan系列GPU进行了例外处理.)

因此,在Windows上,这些GPU将无法利用直接P2P传输.然而,几乎相同的序列可用于传输数据. CUDA运行时将检测传输的性质,并在“引擎盖下”执行分配以创建临时缓冲区.然后,传输将分两部分完成：从始发设备到登台缓冲区的传输,以及从登台缓冲区到目标设备的传输.

以下是一个完整的示例,展示了如何将数据从一个GPU传输到另一个GPU,同时利用P2P访问(如果可用)：

$cat t850.cu

#include

#include

#define SRC_DEV 0

#define DST_DEV 1

#define DSIZE (8*1048576)

#define cudaCheckErrors(msg) \

do { \

cudaError_t __err = cudaGetLastError(); \

if (__err != cudaSuccess) { \

fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \

msg, cudaGetErrorString(__err), \

__FILE__, __LINE__); \

fprintf(stderr, "*** FAILED - ABORTING\n"); \

exit(1); \

} \

} while (0)

int main(int argc, char *argv[]){

int disablePeer = 0;

if (argc > 1) disablePeer = 1;

int devcount;

cudaGetDeviceCount(&devcount);

cudaCheckErrors("cuda failure");

int srcdev = SRC_DEV;

int dstdev = DST_DEV;

if (devcount <= max(srcdev,dstdev)) {printf("not enough cuda devices for the requested operation\n"); return 1;}

int *d_s, *d_d, *h;

int dsize = DSIZE*sizeof(int);

h = (int *)malloc(dsize);

if (h == NULL) {printf("malloc fail\n"); return 1;}

for (int i = 0; i < DSIZE; i++) h[i] = i;

int canAccessPeer = 0;

if (!disablePeer) cudaDeviceCanAccessPeer(&canAccessPeer, srcdev, dstdev);

cudaSetDevice(srcdev);

cudaMalloc(&d_s, dsize);

cudaMemcpy(d_s, h, dsize, cudaMemcpyHostToDevice);

if (canAccessPeer) cudaDeviceEnablePeerAccess(dstdev,0);

cudaSetDevice(dstdev);

cudaMalloc(&d_d, dsize);

cudaMemset(d_d, 0, dsize);

if (canAccessPeer) cudaDeviceEnablePeerAccess(srcdev,0);

cudaCheckErrors("cudaMalloc/cudaMemset fail");

if (canAccessPeer) printf("Timing P2P transfer");

else printf("Timing ordinary transfer");

printf(" of %d bytes\n", dsize);

cudaEvent_t start, stop;

cudaEventCreate(&start); cudaEventCreate(&stop);

cudaEventRecord(start);

cudaMemcpyPeer(d_d, dstdev, d_s, srcdev, dsize);

cudaCheckErrors("cudaMemcpyPeer fail");

cudaEventRecord(stop);

cudaEventSynchronize(stop);

float et;

cudaEventElapsedTime(&et, start, stop);

cudaSetDevice(dstdev);

cudaMemcpy(h, d_d, dsize, cudaMemcpyDeviceToHost);

cudaCheckErrors("cudaMemcpy fail");

for (int i = 0; i < DSIZE; i++) if (h[i] != i) {printf("transfer failure\n"); return 1;}

printf("transfer took %fms\n", et);

return 0;

}

$nvcc -arch=sm_20 -o t850 t850.cu

$./t850

Timing P2P transfer of 33554432 bytes

transfer took 5.135680ms

$./t850 disable

Timing ordinary transfer of 33554432 bytes

transfer took 7.274336ms

$

笔记：

>传递任何命令行参数将禁用P2P,即使它可用.>上述结果适用于可以进行P2P访问的系统,两个GPU都通过PCIE Gen2链路连接,单向传输带宽约为6GB / s. P2P传输时间与此一致(32MB / 5ms~ = 6GB / s).非P2P传输时间较长,但不是两倍.这是因为对于到/来自登台缓冲区的传输,在将一些数据传输到登台缓冲区之后,可以开始传出传输.驱动程序/运行时利用此功能来部分重叠数据传输.