如何正确将CUDA 3D Surface数据复制到主机端数组？

我编写了一段创建int类型3D Surface的代码，填充数据后希望将内容传回主机端做CPU串行处理，但调用cudaMemcpyFromArray()时返回错误码1（cudaErrorInvalidValue），推测该函数用法有误，求正确实现方式。

__global__
void FillVolume(int width, int height, int depth, cudaSurfaceObject_t volSurface)
{
    unsigned int x = threadIdx.x + blockIdx.x * blockDim.x;
    unsigned int y = threadIdx.y + blockIdx.y * blockDim.y;
    unsigned int z = threadIdx.z + blockIdx.z * blockDim.z;
    if (x < width && y < height && z < depth)
    {
        int val = x + y + z;
        surf3Dwrite(val, volSurface, x * sizeof(int), y, z);
    }
}

void test (int width, int height, int depth)
{   
    cudaExtent vol = { width,height,depth };
    cudaChannelFormatDesc volDesc = cudaCreateChannelDesc<int>();
    cudaArray_t volArray;
    gpuErrchk(cudaMalloc3DArray(&volArray, &volDesc, vol, cudaArraySurfaceLoadStore));
    
    cudaSurfaceObject_t volSurface;
    cudaResourceDesc volDescSurf;
    memset(&volDescSurf, 0, sizeof(cudaResourceDesc));
    volDescSurf.resType = cudaResourceTypeArray;
    volDescSurf.res.array.array = volArray;
    gpuErrchk(cudaCreateSurfaceObject(&volSurface, &volDescSurf));
    gpuErrchk(cudaGetLastError());
    
    dim3 block3D(8, 8, 8);
    dim3 grid3D((vol.width + 7) / 8, (vol.height + 7) / 8, (vol.depth + 7) / 8);
    
    FillVolume <<<grid3D, block3D >>> (width,height,depth,volSurface);  
    gpuErrchk(cudaDeviceSynchronize());
    gpuErrchk(cudaGetLastError());
    
    int extentvol = sizeof(int) * vol.depth * vol.width * vol.height;
    
    int* cPUArray; cPUArray = (int*)malloc(extentvol);
    gpuErrchk(cudaHostRegister(cPUArray, extentvol, 0));
    gpuErrchk(cudaMemcpyFromArray(cPUArray, volArray, 0, 0, extentvol, cudaMemcpyDeviceToHost));
    gpuErrchk(cudaDeviceSynchronize());
    gpuErrchk(cudaGetLastError());
    //  Check results here...
    gpuErrchk(cudaHostUnregister(cPUArray));
    gpuErrchk(cudaDestroySurfaceObject(volSurface));
    gpuErrchk(cudaFreeArray(volArray));
    free(cPUArray);

}

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, const char* file, int line)
{
    if (code != cudaSuccess)
    {
        fprintf(stderr, "GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
        exit(code);
    }
}

错误原因 #

cudaMemcpyFromArray()仅适用于2D数组或线性内存的拷贝，无法直接处理3D cudaArray——3D cudaArray的内存布局是硬件优化的非连续结构，直接用该函数会触发cudaErrorInvalidValue错误。

修复方法 #

改用cudaMemcpy3D()来实现3D数组到主机线性内存的拷贝，该函数专门用于处理多维内存的拷贝操作。

修改后的拷贝代码片段 #

替换原代码中cudaMemcpyFromArray相关的部分为以下内容：

// 配置3D拷贝参数
cudaMemcpy3DParms copyParams = {0};
copyParams.srcArray = volArray;
copyParams.srcPos = make_cudaPos(0, 0, 0);
copyParams.dstPtr = make_cudaPitchedPtr(cPUArray, vol.width * sizeof(int), vol.width, vol.height);
copyParams.dstPos = make_cudaPos(0, 0, 0);
copyParams.extent = vol;
copyParams.kind = cudaMemcpyDeviceToHost;

gpuErrchk(cudaMemcpy3D(&copyParams));

完整修改后的test函数 #

void test (int width, int height, int depth)
{   
    cudaExtent vol = { width,height,depth };
    cudaChannelFormatDesc volDesc = cudaCreateChannelDesc<int>();
    cudaArray_t volArray;
    gpuErrchk(cudaMalloc3DArray(&volArray, &volDesc, vol, cudaArraySurfaceLoadStore));
    
    cudaSurfaceObject_t volSurface;
    cudaResourceDesc volDescSurf;
    memset(&volDescSurf, 0, sizeof(cudaResourceDesc));
    volDescSurf.resType = cudaResourceTypeArray;
    volDescSurf.res.array.array = volArray;
    gpuErrchk(cudaCreateSurfaceObject(&volSurface, &volDescSurf));
    gpuErrchk(cudaGetLastError());
    
    dim3 block3D(8, 8, 8);
    dim3 grid3D((vol.width + 7) / 8, (vol.height + 7) / 8, (vol.depth + 7) / 8);
    
    FillVolume <<<grid3D, block3D >>> (width,height,depth,volSurface);  
    gpuErrchk(cudaDeviceSynchronize());
    gpuErrchk(cudaGetLastError());
    
    int extentvol = sizeof(int) * vol.depth * vol.width * vol.height;
    
    int* cPUArray; cPUArray = (int*)malloc(extentvol);
    gpuErrchk(cudaHostRegister(cPUArray, extentvol, 0));

    // 替换原cudaMemcpyFromArray为cudaMemcpy3D
    cudaMemcpy3DParms copyParams = {0};
    copyParams.srcArray = volArray;
    copyParams.srcPos = make_cudaPos(0, 0, 0);
    copyParams.dstPtr = make_cudaPitchedPtr(cPUArray, vol.width * sizeof(int), vol.width, vol.height);
    copyParams.dstPos = make_cudaPos(0, 0, 0);
    copyParams.extent = vol;
    copyParams.kind = cudaMemcpyDeviceToHost;
    gpuErrchk(cudaMemcpy3D(&copyParams));

    gpuErrchk(cudaDeviceSynchronize());
    gpuErrchk(cudaGetLastError());
    //  Check results here...
    gpuErrchk(cudaHostUnregister(cPUArray));
    gpuErrchk(cudaDestroySurfaceObject(volSurface));
    gpuErrchk(cudaFreeArray(volArray));
    free(cPUArray);

}

额外说明 #

• make_cudaPitchedPtr用于描述主机端的线性内存，参数分别为指针、每行字节数、宽度（元素数）、高度（元素数）
• cudaMemcpy3D会自动处理3D数组的内存布局转换，确保数据正确拷贝到主机端的线性内存中

内容的提问来源于stack exchange，提问作者drG_Bristol