CUDA中结构体数组的数组构建问题：非法内存访问排查

你踩的这个坑其实是CUDA新手很容易犯的——直接把主机内存里的结构体数组传给了内核。内核运行在GPU设备上，根本读不到主机内存里的h_a结构体，哪怕结构体里的arr是设备指针也没用，因为内核连结构体本身都访问不了，自然会触发非法内存访问错误。

你的理解缺口在哪？ #

• 你搞混了「主机结构体」和「设备结构体」的生存环境：主机上的h_a存放在主机内存，GPU内核无法直接读取这块内存。结构体里的d_arr虽然是设备内存地址，但内核要拿到这个地址，得先能读到结构体本身才行。
• 直接传递主机指针给内核是无效操作：CUDA内核的参数如果是主机指针，设备端会把它当成设备内存地址来解析，这显然是错误的，必然导致内存访问异常。

怎么修正实现你的需求？ #

核心思路是：把结构体数组也搬到设备内存里，让内核能直接访问设备上的结构体，进而通过结构体里的指针访问设备数组。具体步骤如下：

1. 在设备上分配结构体数组的内存
2. 将主机端的结构体内容（已经包含设备数组指针）拷贝到设备端结构体
3. 内核接收设备端的结构体数组指针，而非主机端的
4. 执行内核后同步设备，确保操作完成再做后续拷贝

下面是修正后的完整代码：

#include <stdio.h>
#include <stdlib.h>
#define N 10

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

typedef struct StructA {
    int* arr;
} StructA;

__global__ void kernel2(StructA *in) {
    int idx = threadIdx.x;
    if (idx < N) { // 加个边界检查更安全
        in[0].arr[idx] = 0;
        in[1].arr[idx] = 1;
        printf("d_arr[%d] = %d , d_arr2[%d] = %d \n", idx, in[0].arr[idx], idx, in[1].arr[idx]);
    }
}

int main(){
    int* h_arr = (int*)malloc(N*sizeof(int));
    int* h_arr2 = (int*)malloc(N*sizeof(int));
    StructA *h_a = (StructA*)malloc(sizeof(StructA) * 2);
    
    // 初始化主机数组
    for(int i=0; i<N; i++){
        h_arr[i] = i+1;
        h_arr2[i] = i+1;
    }

    // 1. 分配设备数组内存并拷贝数据
    int *d_arr, *d_arr2;
    gpuErrchk(cudaMalloc((void**) &d_arr, sizeof(int)*N));
    gpuErrchk(cudaMalloc((void**) &d_arr2, sizeof(int)*N));
    gpuErrchk(cudaMemcpy(d_arr, h_arr, sizeof(int)*N, cudaMemcpyHostToDevice));
    gpuErrchk(cudaMemcpy(d_arr2, h_arr2, sizeof(int)*N, cudaMemcpyHostToDevice));

    // 2. 主机结构体指向设备数组
    h_a[0].arr = d_arr;
    h_a[1].arr = d_arr2;

    // 3. 分配设备结构体数组内存，并将主机结构体拷贝到设备
    StructA *d_a;
    gpuErrchk(cudaMalloc((void**)&d_a, sizeof(StructA)*2));
    gpuErrchk(cudaMemcpy(d_a, h_a, sizeof(StructA)*2, cudaMemcpyHostToDevice));

    // 4. 调用内核，传递设备端的结构体数组指针
    kernel2<<<1,N>>>(d_a);
    gpuErrchk(cudaPeekAtLastError());
    gpuErrchk(cudaDeviceSynchronize()); // 必须同步，不然printf和后续拷贝可能拿不到最新结果

    // 5. 将设备数组数据拷贝回主机
    gpuErrchk(cudaMemcpy(h_arr, d_arr, sizeof(int)*N, cudaMemcpyDeviceToHost));
    gpuErrchk(cudaMemcpy(h_arr2, d_arr2, sizeof(int)*N, cudaMemcpyDeviceToHost));

    // 打印结果
    printf("\nh_arr after kernel: ");
    for(int i=0; i<N; i++) printf("%d ", h_arr[i]);
    printf("\nh_arr2 after kernel: ");
    for(int i=0; i<N; i++) printf("%d ", h_arr2[i]);
    printf("\n");

    // 清理资源
    free(h_arr);
    free(h_arr2);
    free(h_a);
    cudaFree(d_arr);
    cudaFree(d_arr2);
    cudaFree(d_a);

    return 0;
}

关键要点总结 #

• 对于包含指针的结构体，要遵循「设备数组分配→主机结构体绑定设备指针→设备结构体分配→主机结构体拷贝到设备→内核访问设备结构体」的流程
• 内核只能直接访问设备内存中的数据，所有要被内核访问的结构体、数组都必须先拷贝到设备内存
• 记得添加cudaDeviceSynchronize()，不然内核还没执行完，主机就开始读取数据，会得到错误结果

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