如何将 4D 数组的一部分从主机内存复制到设备内存？

Question

我在 Host 数组中展平了 4-D 数组。
我想复制 4-D 数组的一部分（红色区域），如下图所示。

我不知道如何复制未序列化的数组。
我复制一部分数组的原因是因为原始数组大小超过 10GB，我只需要它的 10%。
所以一开始，我用for循环尝试了它。但这花费了太多时间。
有没有更好的主意..？

int main(){
    int nx = 100; ny = 200; nz = 300; nch = 400;
    int idx_x_beg = 50;   int_x_end = 100;
    int idx_y_beg = 100;  int_y_end = 200;
    int idx_z_beg = 150;  int_z_end = 300;
    int idx_ch_beg = 200; int_ch_end = 400;

    double *h_4dArray = (double *)malloc(sizeof(double)*nx*ny*nz*ch);
    double *d_4dArray;
    cudaMalloc((void**)&d_4dArray, (sizeof(cuDoubleReal)*nx*ny*nz*ch));

    for (int temp_ch = 0; temp_ch < (idx_ch_end - idx_ch_beg + 1); temp_ch++) {
        for (int temp_z = 0; temp_z < (idx_z_end - idx_z_beg + 1); temp_z++) {
            for (int temp_y = 0; temp_y < (idx_y_end - idx_y_beg + 1); temp_y++) {
                cudaMemcpy(d_4dArray + temp_ch*idx_z_size*idx_y_size*idx_x_size + temp_z*idx_y_size*idx_x_size + temp_y*idx_x_size
                         , h_4dArray + temp_ch*nz*ny*nx + temp_z*ny*nx + temp_y * nx + idx_x_beg
                         , sizeof(double)*(int_x_end - int_x_beg), cudaMemcpyHostToDevice)
            }
        }
    }

    return 0;
}

Answer 1

对于复制数组的子集，cuda 提供了cudaMemcpy2D（可以复制多维数组的单个 2D 部分）和cudaMemcpy3D（可以复制多维数组的单个 3D 部分）。您可以在 cuda 标签上找到很多问题，以了解如何使用这些问题。

这些方法存在两个问题：

1. 它们不一定扩展到 4D 外壳。即你可能仍然需要一个循环
2. 这些操作的性能（主机设备传输速度）通常是 significantly lower 而不是复制相同字节数的 cudaMemcpy 操作

所以这里没有免费的午餐。我相信最好的建议是在主机上创建一个额外的“连续”缓冲区，将所有切片复制到该缓冲区，然后在单个 cudaMemcpy 调用中将该缓冲区从主机复制到设备。之后，如果您仍然需要设备上的 4D 表示，那么您将需要编写一个设备内核来为您“分散”数据。从概念上讲，与您显示的代码相反。

抱歉，我不会为你编写所有代码。但是，我将使用您显示的代码粗略地完成它的第一部分（将所有内容复制到设备上的单个连续缓冲区）：

int main(){
    int nx = 100; ny = 200; nz = 300; nch = 400;
    int idx_x_beg = 50;   int_x_end = 100;
    int idx_y_beg = 100;  int_y_end = 200;
    int idx_z_beg = 150;  int_z_end = 300;
    int idx_ch_beg = 200; int_ch_end = 400;

    double *h_4dArray = (double *)malloc(sizeof(double)*nx*ny*nz*ch);
    double *d_4dArray, *h_temp, *d_temp;
    size_t temp_sz = (int_x_end - int_x_begin)*(idx_ch_end - idx_ch_beg + 1)*(idx_z_end - idx_z_beg + 1)*(idx_y_end - idx_y_beg + 1);
    h_temp = (double *)malloc(temp_sz*sizeof(double));
    cudaMalloc(&d_temp, temp_sz*sizeof(double));
    cudaMalloc((void**)&d_4dArray, (sizeof(cuDoubleReal)*nx*ny*nz*ch));
    size_t size_tr = 0;
    for (int temp_ch = 0; temp_ch < (idx_ch_end - idx_ch_beg + 1); temp_ch++) {
        for (int temp_z = 0; temp_z < (idx_z_end - idx_z_beg + 1); temp_z++) {
            for (int temp_y = 0; temp_y < (idx_y_end - idx_y_beg + 1); temp_y++) {
                memcpy(h_temp+size_tr
                         , h_4dArray + temp_ch*nz*ny*nx + temp_z*ny*nx + temp_y * nx + idx_x_beg
                         , sizeof(double)*(int_x_end - int_x_beg));
                size_tr += (int_x_end - int_x_beg);
            }
        }
    }
    cudaMemcpy(d_temp, h_temp, temp_sz*sizeof(double), cudaMemcpyHostToDevice);
    // if necessary, put cuda kernel here to scatter data from d_temp to d_4dArray
    return 0;
}

之后，如前所述，如果您需要设备上的 4D 表示，您将需要一个 CUDA 内核来为您分散数据。