c++ - 使用 cudaMemcpy3D 进行高效零填充

Question

我想将存储在主机线性内存中的 3d 数组传输到设备上的更大 (3D) 数组中。作为示例(见下文)，我尝试将 (3x3x3) 数组转换为 (5x5x3) 数组。我希望在主机上获得具有以下模式的 2D 切片:

x x x 0 0
x x x 0 0
x x x 0 0
0 0 0 0 0 
0 0 0 0 0

其中 x 是我的数组的值。但是，我得到类似的结果，其中 y 是下一个 2D 切片的值:

x x x 0 0
x x x 0 0
x x x 0 0
y y y 0 0 
y y y 0 0

根据cudaMemcpy3D documentation我本以为 extent 参数会考虑垂直轴上的填充，但显然没有。我对文档的理解有误吗？如果是，还有其他方法可以执行此操作吗？要传输的数组的最终大小将是 60x60x900 到大小为 1100x1500x900 的数组。我使用零填充来准备傅立叶变换。

这是我使用的简化代码:

cudaError_t cuda_status;
cudaPitchedPtr d_ptr;
cudaExtent d_extent = make_cudaExtent(sizeof(int)*5,sizeof(int)*5,sizeof(int)*3);
cudaExtent h_extent = make_cudaExtent(sizeof(int)*3,sizeof(int)*3,sizeof(int)*3);

int* h_array = (int*) malloc(27*sizeof(int));
int* h_result = (int*) malloc(512*sizeof(int)*5*3);
for (int i = 0; i<27; i++)
{
    h_array[i] = i;
}
cuda_status = cudaMalloc3D(&d_ptr, d_extent);
cout << cudaGetErrorString(cuda_status) << endl;

cudaMemcpy3DParms myParms = {0};
myParms.extent = h_extent;
myParms.srcPtr.ptr = h_array;
myParms.srcPtr.pitch = 3*sizeof(int);
myParms.srcPtr.xsize = 3*sizeof(int);
myParms.srcPtr.ysize = 3*sizeof(int);
myParms.dstPtr = d_ptr;
myParms.kind = cudaMemcpyHostToDevice;

cuda_status = cudaMemcpy3D(&myParms);
cout << cudaGetErrorString(cuda_status) << endl;

cout << "Pitch: " << d_ptr.pitch << " / xsize:" << d_ptr.xsize << " / ysize:" << d_ptr.ysize << endl; // returns Pitch: 512 / xsize:20 / ysize:20 which is as expected

// Copy array to host to be able to print the values - may not be necessary
cout << cudaMemcpy(h_result, (int*) d_ptr.ptr, 512*5*3, cudaMemcpyDeviceToHost) << endl;
cout << h_result[128] << " " << h_result[3*128] << " " << h_result[5*128] << " " << endl; // output : 3 9 15 / expected 3 0 9

Answer 1

这里的问题与您的范围和大小有关。

当一个范围与 cudaMemcpy3D 一起用于非 cudaArray 情况时，它旨在提供 the size of the region in bytes .考虑这一点的一种方法是范围的 3 个维度的乘积应该产生以字节为单位的区域大小。

然而，您所做的是按元素大小缩放 3 个维度中的每一个，这是不正确的:

cudaExtent h_extent = make_cudaExtent(sizeof(int)*3,sizeof(int)*3,sizeof(int)*3);
                                      ^^^^^^^^^^^
                                      this is the only element scaling expected

你在这里犯了类似的错误:

myParms.srcPtr.xsize = 3*sizeof(int); // correct
myParms.srcPtr.ysize = 3*sizeof(int); // incorrect

我们仅按元素大小缩放 x(宽度)维度，不缩放 y(高度)或 z(深度)维度。

我还没有完全验证您的代码，但是通过这 2 处更改，您的代码会产生您指定的预期输出:

$ cat t1593.cu
#include <iostream>
using namespace std;
int main(){

cudaError_t cuda_status;
cudaPitchedPtr d_ptr;
cudaExtent d_extent = make_cudaExtent(sizeof(int)*5,5,3);
cudaExtent h_extent = make_cudaExtent(sizeof(int)*3,3,3);

int* h_array = (int*) malloc(27*sizeof(int));
int* h_result = (int*) malloc(512*sizeof(int)*5*3);
for (int i = 0; i<27; i++)
{
    h_array[i] = i;
}
cuda_status = cudaMalloc3D(&d_ptr, d_extent);
cout << cudaGetErrorString(cuda_status) << endl;

cudaMemcpy3DParms myParms = {0};
myParms.extent = h_extent;
myParms.srcPtr.ptr = h_array;
myParms.srcPtr.pitch = 3*sizeof(int);
myParms.srcPtr.xsize = 3*sizeof(int);
myParms.srcPtr.ysize = 3;
myParms.dstPtr = d_ptr;
myParms.kind = cudaMemcpyHostToDevice;

cuda_status = cudaMemcpy3D(&myParms);
cout << cudaGetErrorString(cuda_status) << endl;

cout << "Pitch: " << d_ptr.pitch << " / xsize:" << d_ptr.xsize << " / ysize:" << d_ptr.ysize << endl; // returns Pitch: 512 / xsize:20 / ysize:20 wich is as expected

// Copy array to host to be able to print the values - may not be necessary
cout << cudaMemcpy(h_result, (int*) d_ptr.ptr, d_ptr.pitch*5*3, cudaMemcpyDeviceToHost) << endl;
cout << h_result[128] << " " << h_result[3*128] << " " << h_result[5*128] << " " << endl; // output : 3 9 15 / expected 3 0 9

}
$ nvcc -o t1593 t1593.cu
$ cuda-memcheck ./t1593
========= CUDA-MEMCHECK
no error
no error
Pitch: 512 / xsize:20 / ysize:5
0
3 0 9
========= ERROR SUMMARY: 0 errors
$

我还应该指出，CUDA 中的跨步 memcpy 操作(例如 cudaMemcpy2D、cudaMemcpy3D)不一定是执行此类传输的最快方法。您可以在 SO cuda 标签上的有关 cudaMemcpy2D 的各种问题中找到有关此特性的文章。

总而言之，以无跨度、无填充的线性传输方式将数据传输到设备可能会更快，然后编写一个 CUDA 内核来获取设备上现在的数据，并将其放入数组中感兴趣的，具有适当的步幅/填充。