CUDA全局函数无法调用对象虚函数报错700，如何修复？

问题原因 #

当把Foo::print()改为虚函数后触发CUDA Error 700（非法内存访问），本质原因是虚函数的多态依赖虚表（vtable）实现，而Host和Device端的虚表完全独立：

• 在Host侧创建的Bar对象，其虚表指针指向Host进程空间内的虚表地址
• 直接将Host对象拷贝到Device后，Device端对象仍持有Host的虚表指针，而Device无法访问Host内存空间，调用虚函数时会触发非法内存访问。

同时原代码存在隐性问题：cudaMalloc分配了sizeof(Foo)大小的内存，但Bar作为含虚表指针的子类，内存大小可能大于Foo，拷贝时会截断对象，进一步加剧错误。

修复方案 #

核心思路是让Device端对象持有Device侧的虚表指针，最可靠的方式是在Device端构造对象。

方案1：Device端构造对象（推荐） #

不在Host侧创建子类对象，而是在Device端分配内存后，通过内核函数构造Bar对象，这样对象的虚表指针会指向Device侧的虚表。

修改后的完整代码示例：

#include <iostream>

#define checkCudaErrors(val) check_cuda((val), #val, __FILE__, __LINE__)

void check_cuda(cudaError_t result, char const *const func, const char *const file, int const line) {
    if (result) {
        std::cerr << "CUDA error = " << static_cast<unsigned int>(result) << " at " << file << ":" << line << " '"
             << func << "' \n";
        cudaDeviceReset();
        exit(-1);
    }
}

struct Foo {
    __host__ __device__ virtual void print() const {
        printf("Foo: print()\n");
    }
    __host__ __device__ virtual ~Foo() = default; // 必须定义虚析构函数，确保Device端正确销毁对象
};

struct Bar: Foo {
    __host__ __device__ void print() const override {
        printf("Bar: print()\n"); // 可选：重写print方法验证多态效果
    }
};

__global__
void test(const Foo* foo) {
    foo->print();
}

// 用于在Device端构造Bar对象的内核
__global__ void createBar(Foo** ptr) {
    *ptr = new Bar();
}

// 用于在Device端销毁对象的内核
__global__ void destroyFoo(Foo* ptr) {
    delete ptr;
}

int main()
{
    Foo * gpu_bar;
    // 分配Bar大小的内存，确保空间足够
    checkCudaErrors(cudaMalloc(&gpu_bar, sizeof(Bar)));

    // 在Device端构造Bar对象
    createBar<<<1,1>>>(&gpu_bar);
    checkCudaErrors(cudaGetLastError());

    // 调用测试内核
    test<<<2,3>>>(gpu_bar);
    checkCudaErrors(cudaGetLastError());
    checkCudaErrors(cudaDeviceSynchronize());

    // 销毁Device端对象并释放内存
    destroyFoo<<<1,1>>>(gpu_bar);
    checkCudaErrors(cudaGetLastError());
    checkCudaErrors(cudaFree(gpu_bar));

    return 0;
}

编译注意事项 #

使用nvcc编译时需要启用C11及以上标准，因为Device端的new/delete需要C11支持：

nvcc -std=c++11 your_code.cu -o your_program

额外说明 #

• 必须给基类Foo定义虚析构函数，否则Device端销毁对象时无法正确调用子类的析构函数，会导致内存泄漏。
• 不要尝试通过直接拷贝Host对象到Device来解决问题，即使拷贝完整的Bar对象，其虚表指针仍然指向Host侧的虚表，Device无法访问。

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