CUDA结构体内存分配报错：数组字段bits分配时出现段错误

z_cu是通过cudaMalloc分配的设备内存指针，主机代码无法直接访问设备内存中的结构体成员。在z_cu_from_z函数里，to指向设备内存，直接执行&to->bits相当于尝试读取设备内存中结构体的bits字段地址，主机端没有权限直接操作设备地址空间，因此触发段错误。

有两种常见的修复思路，任选其一即可：

方案1：主机端构建完整结构体后拷贝到设备 #

先在主机内存中创建z_cu_t实例，完成bits的设备内存分配，再将整个结构体拷贝到设备：

修改main.cu：

#include "z_cu_t.cuh"

int main()
{
    z_t z = z_init(2);
    z_cu_t z_cu_host; // 主机端结构体实例
    z_cu_t* z_cu_dev;

    z_cu_from_z(&z_cu_host, &z); // 先在主机端构建结构体
    cudaMalloc(&z_cu_dev, sizeof(z_cu_t));
    cudaMemcpy(z_cu_dev, &z_cu_host, sizeof(z_cu_t), cudaMemcpyHostToDevice);

    z_free(&z);
    z_cu_free<<<1, 1>>>(z_cu_dev);
    cudaFree(z_cu_dev);
    cudaFree(z_cu_host.bits); // 释放主机结构体持有的设备指针
    
    return 0;
}

修改z_cu_t.cu中的z_cu_from_z函数：

void z_cu_from_z(z_cu_t* to, z_t const* from)
{
    to->size = from->size;
    cudaMalloc(&to->bits, sizeof(int) * from->size);
    cudaMemcpy(to->bits, from->bits, sizeof(int) * from->size, cudaMemcpyHostToDevice);
}

方案2：通过cudaMemcpy修改设备结构体成员 #

保持z_cu为设备指针，先分配bits的设备内存，再将指针值拷贝到设备结构体的bits字段：

修改z_cu_t.cu中的z_cu_from_z函数：

void z_cu_from_z(z_cu_t* to, z_t const* from)
{
    // 拷贝size字段到设备
    cudaMemcpy(&to->size, &from->size, sizeof(size_t), cudaMemcpyHostToDevice);
    
    // 分配bits的设备内存，得到设备指针
    int* d_bits;
    cudaMalloc(&d_bits, sizeof(int) * from->size);
    // 将d_bits指针拷贝到设备结构体的bits字段
    cudaMemcpy(&to->bits, &d_bits, sizeof(int*), cudaMemcpyHostToDevice);
    
    cudaMemcpy(d_bits, from->bits, sizeof(int) * from->size, cudaMemcpyHostToDevice);
}

• 主机指针和设备指针属于不同的地址空间，主机不能直接解引用设备指针访问其成员。
• 释放内存时要确保所有分配的设备内存都被释放，避免内存泄漏。

内容的提问来源于stack exchange，提问作者P'tit Ju