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

Posted 2023-03-23

【中文标题】如何将 4D 数组的一部分从主机内存复制到设备内存？

【英文标题】：How can I copy a part of 4D array from Host memory to Device memory?

【发布时间】：2021-07-22 01:08:18

【问题描述】：

我在 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;

【问题讨论】：

你所展示的不可能是正确的；传递给cudaMemcpy 的指针没有指针运算。您总是从每个指针复制第一组字节。但抛开这一点，将所有内容复制到主机上的连续缓冲区。大致使用您所概述的内容，除了 memcpy 而不是 cudaMemcpy （使用适当的指针算术/指针偏移量）。然后在单个 cudaMemcpy 调用中将该连续缓冲区复制到设备。

@RobertCrovella 抱歉，我复制了错误的代码版本。您的意思是复制每个可以通过 for 循环创建的起始指针，并使用单个 cudaMemcpy 复制它吗？ CUDA Samples 中有示例代码吗？

【参考方案1】：

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

这些方法存在两个问题：

它们不一定扩展到 4D 外壳。即你可能仍然需要一个循环 这些操作的性能（主机设备传输速度）通常是 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 内核来为您分散数据。

【讨论】：

感谢您提供详细信息。我尝试了您的建议，使连续缓冲区需要 0.0021 毫秒，memcpy 需要 310 毫秒。这是令人印象深刻的！实际上原始的 4d 数组非常大，而我想在显卡中计算的 4d 部分非常小。所以在d_temp很小的情况下，整个过程只需要0.0035毫秒（平均）。当我重新对齐矩阵的一部分并将其传递给设备时，我了解到它非常有效。太有趣了！再次感谢您。