通过 CUDA 线程复制全局内存答案

【问题标题】：copy global memory by CUDA threads通过 CUDA 线程复制全局内存
【发布时间】：2011-07-28 09:20:48
【问题描述】：

我需要通过 CUDA 线程（不是来自主机）将全局内存中的一个数组复制到全局内存中的另一个数组。

我的代码如下：

__global__ void copy_kernel(int *g_data1, int *g_data2, int n)
{
  int idx = blockIdx.x * blockDim.x + threadIdx.x;
  int start, end;
  start = some_func(idx);
  end = another_func(idx);
  unsigned int i;
  for (i = start; i < end; i++) {
      g_data2[i] = g_data1[idx];
  }
}

这是非常低效的，因为对于某些 idx，[start, end] 区域非常大，这使得该线程发出过多的复制命令。有什么方法可以有效地实现吗？

谢谢，

郑

【问题讨论】：

标签： cuda gpgpu

【解决方案1】：

你写它的方式，我猜每个线程都试图将整个“开始”写入“结束”块。这真的真的效率低下。

你需要做这样的事情。

___shared___ unsigned sm_start[BLOCK_SIZE];
___shared___ unsigned sm_end[BLOCK_SIZE];
sm_start[threadIdx.x] = start;
sm_end[threadIdx.y] = end;
__syncthreads();
for (int n = 0; n < blockdDim.x; n++) {
g_data2 += sm_start[n];
unsigned lim = sm_end[n] - sm_start[n];
  for (int i = threadIdx.x; i < lim; i += blockDim.x) {
      g_data2[i] = g_data1[idx];
  }
}

【讨论】：

【解决方案2】：

试试这个：

CUresult cuMemcpyDtoD(
    CUdeviceptr dst,
    CUdeviceptr src,
    unsigned int bytes   
)

更新：

你是对的：http://forums.nvidia.com/index.php?showtopic=88745

没有有效的方法来正确地执行此操作，因为 CUDA 的设计希望您只使用内核中的少量数据。

【讨论】：

和qba的回答一样的问题：这个函数只能从宿主函数调用，不能从__global__函数调用。
因此，基于 nvidia 论坛主题，看起来最好的方法是让您的主题使用您的设备可以原生操作的最大数据类型（这样每次迭代只会产生一个副本指令），最佳线程数必须凭经验确定，因为它取决于设备的内存架构。