CUDA - 具有派生类的动态共享内存

Question

我一直在尝试运行其他一些代码，但在动态共享内存方面遇到了一些问题。根据文档 (https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#shared)，我应该分配一个内存数组，然后将指针类型转换到该数组中的特定位置，如下所示：

extern __shared__ float array[];

short* array0 = (short*)array; 
float* array1 = (float*)&array0[128];
int*   array2 =   (int*)&array1[64];

但是，在我自己的代码中，这并不一定总是有效，我也不太明白为什么。

我的基本结构有 2 个类 A 和 B 以及一个错误检查宏

#define cudaCheckError() { \
  cudaError_t err = cudaGetLastError(); \
  if(err != cudaSuccess) { \
    printf("Cuda error: %s:%d: Error code %d, %s\n", __FILE__, __LINE__, err,cudaGetErrorString(err)); \
    exit(1); \
  } \
}

class A {
    public:
    
    __device__ virtual int foo() const = 0;
};

class B : public A {
    public:
    
    __device__ B() {}
    
    __device__ virtual int foo() const override {
        return 1;
    }
};

和我的内核

__global__
void kernel() {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    
    extern __shared__ int shared[];
    B* b_array = (B *) &shared[0];
    
    if (idx == 0) {
        b_array[0] = B();
        
        printf("%i", b_array[0].foo());
    }
    
    __syncthreads();
    
    return;
}

调用具有足够共享内存并指定kernel<<<1, 1, 1000>>> 的内核并检查错误代码会产生错误Error code 700, an illegal memory access was encountered。对此运行 cuda-memcheck 也会给出一个错误代码，虽然是一个不同的代码：Error code 719, unspecified launch failure

将内核更改为：

__global__
void kernel() {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    
    extern __shared__ B shared[];
    
    if (idx == 0) {
        shared[0] = B();
        
        printf("%i", shared[0].foo());
    }
    
    __syncthreads();
    
    return;
}

并且重新运行会给出预期的输出而不会出现错误。

这是 CUDA 中的派生类和类型转换的某种问题吗？我不是在主机和设备之间复制对象，所以这应该不是问题。是不是不可能像我想要的那样强制转换为对象数组？

Answer 1

根据我的经验，一个对象副本：

= B();

does not copy the virtual function pointer table。因此，无论您从哪个对象访问虚函数，都必须正确设置虚函数指针表。

这允许：

extern __shared__ B shared[];

这不是：

extern __shared__ int shared[];

AFAIK 的这方面是特定于实现的； C++ 标准不要求。

作为一个证明点，我们可以在你失败的内核中做这样的事情：

__global__
void kernel() {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    
    extern __shared__ int shared[];
    B* b_array = (B *) &shared[0];
    
    if (idx == 0) {
        B temp = B();
        memcpy(b_array, &temp, sizeof(B));
        
        printf("%i", b_array[0].foo());
    }
    
    __syncthreads();
    
    return;
}

这将起作用。我并不是说这是编码它的正确方法。我只是用它来暗示这里至少有一个问题是桌子的处理。正如 Jerome Richard 在 cmets 中指出的那样，使用底层 int 数组对其他内容进行类型双关可能是非法的，但正如您所指出的，cuda docs 似乎暗示了这一点。

我们还可以按照您的失败示例构建主机代码测试用例：

$ cat t131.cpp
#include <cstdio>

class A {
    public:
     virtual int foo() const = 0;
};

class B : public A {
    public:
     B() {}
     virtual int foo() const override {
        return 3;
    }
};

void k1() {

    int sh1[100];
    B* b_array = (B *) &sh1[0];
        b_array[0] = B();

        printf("k1 %i\n", b_array[0].foo());


    return;
}

int main(){
  k1();
}

$ g++ t131.cpp -o t131
$ ./t131
Segmentation fault (core dumped)
$

这也失败了。

如果您发现我的描述有问题或只是希望此案得到处理，欢迎您file a bug。

这里的确切代码很重要，因此对上面的测试用例稍作更改可能会导致代码正常工作或失败。