使用cudaMalloc指针引用时出现非法内存访问的解决方法

问题背景 #

我希望在A类中创建一个P对象，并让B类持有P的引用。相关代码定义如下：

P结构体定义 #

struct P
{
    int* p_;

    P() { CC(cudaMalloc(&p_, sizeof(int))); }

    __device__ void print() { printf("p: %d\n", p_[0]); }
};

A、B类定义 #

struct A
{
    P a_;
};

struct B
{
    B(A& a) : b_(a.a_) {}
    P& b_;
};

设备端核函数实现 #

__global__ void Ka(A a)
{
    a.a_.p_[0] = 399;
    a.a_.print();
}

__global__ void Kb(B b)
{
    b.b_.print();
}

主机端调用方式 #

A a;
Ka <<<1, 1>>> (a);
CC(cudaDeviceSynchronize());
CC(cudaGetLastError());

B b(a);
Kb <<<1, 1>>> (b);
CC(cudaDeviceSynchronize());
CC(cudaGetLastError());

调试显示二者内存中是同一个对象，但第二个Kb核函数调用时出现非法内存访问。需要在无需复制P（实际场景中复制成本极高）的前提下实现类似功能。

问题原因 #

核心问题是主机端的引用无法直接传递到设备端：

• 主机端创建的B对象中，b_引用指向的是主机端A对象里的a_；
• 当把B对象作为参数传递给核函数时，传递的是B的副本，其中的引用地址仍然是主机端内存地址；
• 设备端线程无法直接访问主机端内存，因此这个引用在设备端是无效地址，触发非法内存访问。

解决方案 #

方案1：用指针替代引用，传递设备端对象地址 #

引用本质是语法糖，在CUDA设备端场景下，指针更适合传递设备内存地址。需要先将A对象拷贝到设备端，再让B持有设备端P对象的指针：

修改后代码

struct P
{
    int* p_;

    P() { CC(cudaMalloc(&p_, sizeof(int))); }
    // 添加析构函数避免内存泄漏
    ~P() { if (p_) CC(cudaFree(p_)); }

    __device__ void print() { printf("p: %d\n", p_[0]); }
};

struct A
{
    P a_;
};

struct B
{
    P* b_;
    // 接收设备端A的指针，直接指向其内部的P对象
    B(A* dev_a) : b_(&dev_a->a_) {}
};

__global__ void Ka(A* a)
{
    a->a_.p_[0] = 399;
    a->a_.print();
}

__global__ void Kb(B b)
{
    b.b_->print();
}

主机端调用

A host_a;
A* dev_a;
// 分配设备端内存并拷贝A对象
CC(cudaMalloc(&dev_a, sizeof(A)));
CC(cudaMemcpy(dev_a, &host_a, sizeof(A), cudaMemcpyHostToDevice));

// 调用核函数修改设备端P对象
Ka<<<1, 1>>>(dev_a);
CC(cudaDeviceSynchronize());
CC(cudaGetLastError());

// 创建B对象，持有设备端P的指针
B host_b(dev_a);
Kb<<<1, 1>>>(host_b);
CC(cudaDeviceSynchronize());
CC(cudaGetLastError());

// 清理资源
CC(cudaFree(dev_a));

方案2：使用统一内存（Unified Memory）简化管理 #

如果你的CUDA版本支持统一内存（Compute Capability 3.0及以上），可以利用统一内存让主机和设备共享对象地址，无需手动拷贝：

修改后代码

struct P
{
    int* p_;

    // 使用cudaMallocManaged分配统一内存
    P() { CC(cudaMallocManaged(&p_, sizeof(int))); }
    ~P() { if (p_) CC(cudaFree(p_)); }

    __device__ void print() { printf("p: %d\n", p_[0]); }
};

struct A
{
    P a_;
};

struct B
{
    P& b_;
    B(A& a) : b_(a.a_) {}
};

__global__ void Ka(A* a)
{
    a->a_.p_[0] = 399;
    a->a_.print();
}

__global__ void Kb(B b)
{
    b.b_.print();
}

主机端调用

// 分配统一内存的A对象
A* dev_a;
CC(cudaMallocManaged(&dev_a, sizeof(A)));

Ka<<<1, 1>>>(dev_a);
CC(cudaDeviceSynchronize());
CC(cudaGetLastError());

// B直接引用统一内存中的P对象
B host_b(*dev_a);
Kb<<<1, 1>>>(host_b);
CC(cudaDeviceSynchronize());
CC(cudaGetLastError());

// 清理资源
CC(cudaFree(dev_a));

统一内存会自动处理主机与设备间的内存同步，对象地址在两端一致，因此引用可以正常在设备端使用。

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