cuda程序无法传数据到显存

1、CPU和GPU之间

1）CPU->GPU

从CPU向GPU传输数据，最为人熟知的就是cudaMemcpy了。

默认情况下，数据是从系统的分页内存先到锁页内存，然后再到GPU显存。因此如果显式指定使用锁页内存，是可以加快数据传输速度的。

（锁页内存，在cuda编程里使用CudaHostMalloc分配。实质上和linux的mlock系统调用一样，就是给内存页打上标记，不让 *** 作系统将其从物理内存交换到硬盘）

至于为什么cuda要这样设计，个人理解是为了实现的方便。因为 *** 作系统已经处理了硬盘和物理内存间的页交换等情况，显卡驱动只需要实现物理内存到GPU显存这一种数据传输即可，不需要把 *** 作系统内存管理的事情再做一遍。

2) GPU->CPU

GPU向CPU拷贝数据时，锁页内存同样比分页内存快

值得一提的是，适当使用pinned memory显然可以加快IO速度。但是并不是越多越好，因为锁页内存是完全独占住了物理内存， *** 作系统无法调度，可能会影响系统整体性能。

3)同一张GPU卡内部

同一张卡内两块显存对拷，实测P40上高达~285GB/s。也比较接近于GPU卡本身的访存速度

4）数据拷贝的overhead

在上面的测试数据中，可以看到传输数据量从1M->32M增长的过程中，测得的传输带宽是有逐渐增加的。

这是因为每次调用cuda api进行数据传输都有overhead，在数据量小的时候这个overhead在数据传输时间中的占比就显得很高。这也提示我们尽量合并小数据的传输

2、同机的GPU之间

一般可以通过cudaMemcpyPeer/cudaMemcpyPeerAsync函数进行显存拷贝

1）cudaMemcpyPeer withoutP2P

/********代码示例*******/

cudaSetDevice(1)

cudaMalloc((int**)&dest, bytes)

cudaSetDevice(2)

cudaMalloc((int**)&dsrc, bytes)

cudaMemcpyPeer(dest, 1, dsrc, 2, bytes)

通过nvprof+nvpp可以看到：禁用GPU P2P时，数据是先从GPU2拷贝到系统内存（DtoH），然后再从系统内存拷贝到GPU1（HtoD）

当然，这里是在一个进程内做GPU之间的数据拷贝。如果是2个进程分别运行在GPU1和GPU2上，那在CPU上这2个进程间可以通过共享内存或者socket通信来完成数据的拷贝。

2）cudaMemcpyPeer withP2P

/********代码示例*******/

cudaSetDevice(1)

cudaMalloc((int**)&dest, bytes)

cudaSetDevice(2)

cudaMalloc((int**)&dsrc, bytes)

cudaDeviceEnablePeerAccess(1,0)

cudaDeviceEnablePeerAccess(2,0)

cudaMemcpyPeer(dest, 1, dsrc, 2, bytes)

启用GPU P2P时，数据直接从GPU2拷贝到了GPU1，不再经过系统内存。

3)通过变量赋值方式传输数据

深度学习中，卡之间传递的数据其实很多都是参数数值，因此也可以直接用一个GPU内的变量给另一个GPU上的变量赋值来进行数据传输

/********代码示例*******/

cudaOccupancyMaxPotentialBlockSize(&numBlocks, &blockSize, copyp2p_float)

copyp2p_float>>(

(float *)dest, (float *)src, num_elems)

__global__ void copyp2p_float(float *__restrict__ dest, float const *__restrict__ src,

size_t num_elems) {undefined

size_t globalId = blockIdx.x * blockDim.x + threadIdx.x

size_t gridSize = blockDim.x * gridDim.x

#pragma unroll(5)

for (size_t i = globalIdi

dest[i] = src[i]

}

}