CUDA核函数中重写虚类方法时出现地址越界问题

问题背景 #

我在开发项目中定义了带虚方法的基类及多个继承它的派生类，基类的若干虚方法被派生类重写，需要将派生类实例传入CUDA核函数并在其中调用其方法（调用不修改对象）。

正常运行的C++最小示例 #

#include<iostream>

class basefoo{
    public:
    virtual int getiden(int i){return 0;}
};

class derivefoo : public basefoo{
    public:
    int getiden(int i){return i;}
};

int getsq(basefoo *foo, int q){
    int result = (foo->getiden(q))*(foo->getiden(q));
    return result;
}

int main(){
    derivefoo foo;
    basefoo foo2;
    int q=3;
    std::cout<<getsq(&foo, q)<<std::endl;
    std::cout<<getsq(&foo2, q)<<std::endl;
    return 0;
}

预期输出 #

9
0

CUDA最小示例 #

#include <iostream>

class basefoo{
    public:
    basefoo(){}
    __host__ __device__ virtual int getiden(int i){return 0;}
};

class derivefoo : public basefoo{
    public:
    derivefoo(){}
    __host__ __device__ int getiden(int i){return i;}
};

__global__ void get_squares(int* squares, basefoo *foo, int sroot, int nsq){
    const int tid = threadIdx.x + blockDim.x*blockIdx.x;
    if(tid < nsq){
        squares[tid] = 0;
        int result = (foo->getiden(sroot))*(foo->getiden(sroot));
        squares[tid] += result;
    }
}

int main(){
    int sroot = 4;
    int nsquares = 2;
    int *hsquares, *dsquares;
    hsquares = (int *)malloc(nsquares*sizeof(int));
    cudaMalloc(&dsquares, nsquares*sizeof(int));

    basefoo foo, *dfoo;
    derivefoo foo2, *dfoo2;

    cudaMalloc(&dfoo, sizeof(basefoo));
    cudaMemcpy(dfoo, &foo, sizeof(basefoo), cudaMemcpyHostToDevice);
    cudaMalloc(&dfoo2, sizeof(derivefoo));
    cudaMemcpy(dfoo2, &foo2, sizeof(derivefoo), cudaMemcpyHostToDevice);

    get_squares<<<1, 2>>>(dsquares, dfoo, sroot, nsquares);
    cudaMemcpy(hsquares, dsquares, nsquares*sizeof(int), cudaMemcpyDeviceToHost);

    printf("%d %d \n", hsquares[0], hsquares[1]);

    get_squares<<<1, 2>>>(dsquares, dfoo2, sroot, nsquares);
    cudaMemcpy(hsquares, dsquares, nsquares*sizeof(int), cudaMemcpyDeviceToHost);

    printf("%d %d \n", hsquares[0], hsquares[1]);


    return 0;

}

错误信息 #

使用compute-sanitizer --tool memcheck执行时出现多个地址越界错误，第一个错误如下：

========= Invalid __global__ read of size 4 bytes
=========     at 0xc8 in get_squares(int *, basefoo *, int, int)
=========     by thread (0,0,0) in block (0,0,0)
=========     Address 0x555f56399068 is out of bounds
=========     and is 45,900,972,322,712 bytes before the nearest allocation at 0x7f1e7d600000 of size 8 bytes
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame: [0x32e130]
=========                in /lib/x86_64-linux-gnu/libcuda.so.1
=========     Host Frame:libcudart_static_4d8b33a106dceb3c07a56e26de61f2d53bb62a68 [0x1091e]
=========                in /home/arin/Desktop/cuda codes/ASMD/./inhtest
=========     Host Frame:cudaLaunchKernel [0x70b2e]
=========                in /home/arin/Desktop/cuda codes/ASMD/./inhtest
=========     Host Frame:cudaError cudaLaunchKernel<char>(char const*, dim3, dim3, void**, unsigned long, CUstream_st*) [0xb1ba]
=========                in /home/arin/Desktop/cuda codes/ASMD/./inhtest
=========     Host Frame:__device_stub__Z11get_squaresPiP7basefooii(int*, basefoo*, int, int) [0xafec]
=========                in /home/arin/Desktop/cuda codes/ASMD/./inhtest
=========     Host Frame:get_squares(int*, basefoo*, int, int) [0xb036]
=========                in /home/arin/Desktop/cuda codes/ASMD/./inhtest
=========     Host Frame:main [0xad1d]
=========                in /home/arin/Desktop/cuda codes/ASMD/./inhtest
=========     Host Frame:../sysdeps/nptl/libc_start_call_main.h:58:__libc_start_call_main [0x29d90]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========     Host Frame:../csu/libc-start.c:379:__libc_start_main [0x29e40]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========     Host Frame:_start [0xaae5]
=========                in /home/arin/Desktop/cuda codes/ASMD/./inhtest

关键现象：若删除basefoo中getiden方法的virtual关键字，memcheck检测无任何错误，但无法实现方法重写，而这是项目必需的。环境为CUDA 12.2、Ubuntu 22.04.3。

问题原因 #

核心原因是CUDA中主机与设备的虚函数表（vtable）不兼容：

• 带虚函数的类实例在主机端会包含一个指向主机虚函数表的指针（vptr）。当通过cudaMemcpy将主机端的类实例拷贝到设备端时，这个vptr并不会被更新，依然指向主机内存中的vtable。而设备端线程无法访问主机内存空间，因此尝试通过这个vptr调用虚方法时，就会触发地址越界的非法内存访问错误。
• 当去掉virtual关键字后，函数调用变为静态绑定，不需要访问虚函数表，也就不会出现跨内存空间的非法访问，因此memcheck检测不到错误，但同时也失去了多态特性。

此外，即使能让设备访问主机的vtable，主机和设备的vtable布局也可能不同（CUDA编译器对设备代码的编译优化规则与主机不同），依然会导致方法调用错误。

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