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 on this 也给出了一个错误代码，虽然是不同的代码：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;
}

并且 re运行ning 给出了预期的输出，没有错误。

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

Answer 1

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

= B();

does not copy the virtual function pointer table。因此，有必要在您从中访问虚函数的任何对象中正确设置虚函数指针 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;
}

这将起作用。我并不是说这是正确的编码方式。我只是用它来表明这里至少有一个问题是 table 的处理。正如 Jerome Richard 在评论中指出的那样，使用底层 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。

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

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

CUDA - Dynamic Shared Memory with Derived Classes

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