在 CUDA 中提取和设置数组的部分答案

【问题标题】：Extract and Set portions of an array in CUDA在 CUDA 中提取和设置数组的部分
【发布时间】：2012-04-02 12:19:06
【问题描述】：

我必须提取数组的部分并将块设置为另一个数组。

例如，我有一个二维数组（一维格式），如 A[32 X 32];还有另一个数组 B[64 X 64]，我想复制一个 8X8 的 B 块，从 B 的 (0,8) 开始，并将其放在 A 的 (8,8) 中。

目前，我可能会使用类似于下面的内核的东西，以便在传递偏移量时获取部分数据。类似的也可以用于将块设置为更大的数组。

__global__ void get_chunk (double *data, double *sub, int xstart, int ystart, int rows, int cols, int subset)
 {
    int i,j;
    i = blockIdx.x * blockDim.x + threadIdx.x;

    for (j = 0; j < subset; j++)
            sub[i*subset+j] = data[i*cols + (xstart*cols + ystart)+j];

 }

我认为使用 cudamemCpy* 的变体（也许是 cudamemCpyArray(...)）也可以做到这一点，但我不知道该怎么做。我需要一些代码示例，或者一些关于如何完成的指导。

PS 我在 nvidia 论坛上遇到了完全相同的问题，没有得到回复，所以在这里尝试。 http://forums.nvidia.com/index.php?showtopic=223386

谢谢。

【问题讨论】：

标签： cuda memcpy

【解决方案1】：

如果您只想将数据从设备上的一个数组复制到另一个数组，则不需要内核。

如果您的源数据中有设备指针，并且主机代码中有分配的目标指针：

伪代码：

//source and target device pointers
double * source_d, target_d;

//get offseted source pointer
double * offTarget_d + offset * sizeof(double);

//copy n elements from offseted source data to target device pointer
cudaMemcpy(offTarget_d, source_d, n * sizeof(double), cudaMemcpyDeviceToDevice);

不清楚是只想复制一维数组的范围，还是要将二维数组中每一行的范围复制到另一个二维数组的目标行中

【讨论】：

@djmj：for 循环实际上是个好主意，但不是一个单位步长。每个线程的多个副本将比一个副本更有效 - 它有助于隐藏线程调度开销和设置成本。
是的，但是线程调度的性能成本是多少？，大多数消息来源指出，gpu 应该保持忙碌，使用尽可能多的线程，同时保持 100% 的占用率。（我假设一个简单的一维数组复制内核）。内存延迟在这里是一个更大的问题，可以使用 100% 的占用率来隐藏它。查看性能图表：cs.berkeley.edu/~volkov/volkov10-GTC.pdf#page=30
@djmj：我不确定我是否跟随。您链接到的性能图表与您在评论中写的完全矛盾。他们表明，最高的 memcpy 带宽是通过使用更少的线程和每个线程更多的并行工作来实现的。达到的最高性能是 4% 占用率，而不是 100%......
是的，你是对的，我 100% 的陈述是不正确的，在给定的示例中，即使收益似乎也很小（并且希望看到展开的版本！）。我只是想指出，内存延迟在这里会更重要，我假设动态大循环的执行速度要慢于展开的小循环。在我上次编辑时，我以某种方式从我之前的评论中删除了该展开声明。
最初的问题是关于复制二维数组的二维子集。这将需要一堆不相交的cudaMemcpy 调用，这将是低效的。如果你有很多小的 2D 块要复制，我建议你编写一个内核并让一组 warp 大小的线程独立地执行每个复制。