CUDA设备内存副本: cudaMemcpyDeviceToDevice与复制内核
我正在编写一个cuda内核来将一个数组复制到另一个。它们都在GPU内存中。我不想使用cudamemcpyDeviceToDevice，因为它的性能很差。
天真的内核：
__global__ void GpuCopy( float* des , float* __restrict__ sour ,const int M , const int N )
{
int tx=blockIdx.x*blockDim.x+threadIdx.x;
if(tx<N*M)
des[tx]=sour[tx];
} 我认为天真的内核不会获得高性能，所以我尝试使用__shared__内存，但它看起来不太好：
__shared__ float TILE[tile];
int tid=threadIdx.x;
for(int i=0; i<M*N/tile;i++)
{
TILE[tid]=sour[i*tile+tid]
des[i*tile+tid]=TILE[tid]
}前一个代码片段将全局内存复制到des[]，而后者将全局内存复制到__shared__，然后将__shared__复制到des[]。我认为后者比前者慢。
那么，如何编写__shared__代码来复制内存呢？另一个问题是，如果我想使用__const__内存，并且数组(已经在GPU中)比常量内存大，如何使用__const__将其复制到另一个GPU内存中。
Stack Overflow用户
发布于 2014-03-12 14:26:12
对于普通的线性到线性内存复制，共享内存不会给你带来任何好处。你天真的内核应该没问题。可能会有一些小的优化，可以用较少的线程块来运行，但是调优将在某种程度上取决于您的特定GPU。
共享内存可以在执行某种修改复制(例如转置操作)的内核中使用。在这些情况下，通过共享内存进行旅行的成本被改进的合并性能所抵消。但是，使用幼稚的内核，读和写都应该结合在一起。
对于单个大型副本操作，cudaMemcpyDeviceToDevice应该提供非常好的性能，因为单个调用的开销是在整个数据移动中摊销的。也许您应该对这两种方法进行计时--使用nvprof很容易。注释中引用的讨论提到了一个特定的用例，其中矩阵象限正在被交换。在这种情况下，NxN矩阵需要1.5N的cudaMemcpy操作，但与单个内核调用相比较。在这种情况下，API调用设置的开销将开始成为一个重要因素。但是，当将单个cudaMemcpy操作与单个等效的内核调用进行比较时，cudaMemcpy操作应该是快速的。
__constant__内存不能被设备代码修改，因此您必须使用基于cudaMemcpyFromSymbol和cudaMemcpyToSymbol的主机代码。
https://stackoverflow.com/questions/22345391
复制相似问题
腾讯云开发者
