如何在 cuda 中使用虚拟类？

Question

我编写了一个简单的 cuda 代码来测试是否可以将指针向量复制到 GPU 的类。

这就是我所拥有的：

test.hpp：

class Class {
public:
    Class() {};
    virtual ~Class() {};

    __host__ __device__ int print() { return 42; };
    __host__ __device__ virtual int getClass() const = 0;
};                                                                                                                                                                                       

class AlphaClass : public Class
{
public:
    AlphaClass() {
      className = "Alpha";
      classAvgGrade = 90;
      classSize = 100;
      for(int i = 0; i < classSize; i++){
          classScores.push_back((90+i)%100);
      } 
    };
    ~AlphaClass() { };

    std::string className;
    int classAvgGrade;
    int classSize; 
    std::vector<int> classScores;

    __host__ __device__ void incClassSize(){
        classSize++;
    };

    __host__ __device__ virtual int getClass() const{
        return 0;
    }; 
};  


class BetaClass : public Class
{                                                                                                                                                                                        
public:
    BetaClass() {
      className = "Beta";
      classAvgGrade = 80;
      classSize = 120;
      for(int i = 0; i < classSize; i++){
          classScores.push_back((80+i)%100);
      } 
    }
    ~BetaClass() { };
    std::string className;
    int classAvgGrade;
    int classSize;
    std::vector<int> classScores;
    
    __host__ __device__ void incClassSize(){
        classSize++;
    }   

    __host__ __device__ virtual int getClass() const{
        return 1;
    };

};
    
    
class CudaClass : public Class
{     
public:
    CudaClass() {
      className = "Cuda";
      classAvgGrade = 70;
      classSize = 110;
      for(int i = 0; i < classSize; i++){
          classScores.push_back(70);
      }   
    };
    ~CudaClass() {
        //do nothing
    };
    
    std::string className;
    int classAvgGrade;
    int classSize;
    std::vector<int> classScores;
    
    __host__ __device__ void incClassSize(){
        classSize++;
    };
 };

test.cpp：

struct GlobalConstants {
    Class** classList;
};  

__constant__ GlobalConstants cuConstRaytracerParams;


   __global__ void useClass()
  {

    Class** deviceClassList = cuConstRaytracerParams.classList;
    AlphaClass* alpha = (AlphaClass*) deviceClassList[0];
    BetaClass* beta = (BetaClass*) deviceClassList[1];
    CudaClass* cuda = (CudaClass*) deviceClassList[2];

    printf("%s\n", alpha->className);
    printf("%s\n", beta->className); 
    printf("%s\n", cuda->ClassName);

    printf("alpha avg = %d\n", alpha->classAvgGrade);
    printf("beta avg = %d\n", beta->classAvgGrade);
    printf("cuda avg = %d\n", cuda->classAvgGrade);

  };  


...

  

    AlphaClass *alpha;
    alpha = new AlphaClass();
    BetaClass *beta;
    beta = new BetaClass();
    CudaClass *cuda;
    cuda = new CudaClass();
    std::vector<Class*> classes;
    classes.push_back(alpha);
    classes.push_back(beta);
    classes.push_back(cuda);
    
    AlphaClass* alpha_ptr;
    BetaClass* beta_ptr;
    CudaClass* cuda_ptr;
        
    // copy classes to cuda
    thrust::device_vector<Class*> deviceClassList;
    for(int i = 0; i < classes.size(); i++){
        if(classes[i]->getClass() == 0){
            cudaMalloc(&alpha_ptr, sizeof(AlphaClass));
            cudaMemcpy(alpha_ptr, &classes[i],sizeof(AlphaClass), cudaMemcpyHostToDevice);
            deviceClassList.push_back(alpha_ptr);

        }else if(classes[i]->getClass() == 1){
            cudaMalloc(&beta_ptr, sizeof(BetaClass));
            cudaMemcpy(beta_ptr, &classes[i],sizeof(BetaClass), cudaMemcpyHostToDevice);
            deviceClassList.push_back(beta_ptr);

            
        }else if(classes[i]->getClass() == 2){
            cudaMalloc(&cuda_ptr, sizeof(CudaClass));
            cudaMemcpy(cuda_ptr, &classes[i],sizeof(CudaClass), cudaMemcpyHostToDevice);
            deviceClassList.push_back(cuda_ptr);

        }else{
            //do nothing
        }
    }
    Class** class_ptr = thrust::raw_pointer_cast(&deviceClassList[0]);
          
    //ADD IT TO CUDA PARAM STRUCT
    GlobalConstants params;
    params.classList = class_ptr;
    cudaMemcpyToSymbol(cuConstRaytracerParams, &params, sizeof(GlobalConstants));

    useClass<<<1,1>>>();
    cudaDeviceSynchronize();

    ...cleanup code

当我运行它时，我没有得到正确的值并得到以下结果：

阿尔法平均值 = 39696816

测试平均 = 70

cuda 平均 = 0

我没有得到任何字符串结果。

Answer 1

OP 提出了几个问题。主要问题在标题“如何在 cuda 中使用虚拟类？”。一个不相关的问题是如何在 cuda 代码中使用字符串。我将主要关注标题中的问题。

根据cuda c programming guide，您可以使用虚函数，但有限制。你遇到的限制是

不允许将派生自虚拟基类的类的对象作为参数传递给 __global__ 函数。

在您的示例代码中，您尝试通过常量内存传递对象（设备指针数组）来避免这种情况。但是我认为编程指南在这里并不精确。我认为不可能将派生自虚拟基类的类的对象复制到设备。问题是（据我了解）您会将 host 虚拟功能表复制到设备。

示例代码过于复杂（并且存在其他问题），无法演示该行为。以下更简化的代码显示了您可以在 cuda 中使用虚函数做什么：

#include <stdio.h>

class Class
{
public:
    __host__ __device__ virtual int getNumber() = 0;
    __host__ __device__ virtual ~Class() {};
};

class ClassA: public Class
{
public:
    int aNumber;
    __host__ __device__ ClassA(): aNumber(0){}

    __host__ __device__ int getNumber()
    {
        return aNumber;
    }
};

class ClassB: public Class
{
public:
    int aNumber;
    int anotherNumber;
    __host__ __device__ ClassB(): aNumber(1), anotherNumber(2){}

    __host__ __device__ int getNumber()
    {
        return aNumber;
    }
};

__global__ void invalidClassKernel( Class* superClass )
{
    printf( "superClass->getNumber(): %d\n", superClass->getNumber() );
}

__global__ void validClassKernel()
{
    Class* classVector[2];
    classVector[0] = new ClassA();
    classVector[1] = new ClassB();

    printf( "classVector[0]->getNumber(): %d\n", classVector[0]->getNumber() );
    printf( "classVector[1]->getNumber(): %d\n", classVector[1]->getNumber() );

    delete classVector[0];
    delete classVector[1];
}

int main()
{
    ClassA hostClassA;
    ClassB hostClassB;

    ClassA* devClassA;
    ClassA* devClassB;
    cudaMalloc( &devClassA, sizeof(ClassA) );
    cudaMalloc( &devClassB, sizeof(ClassB) );
    cudaMemcpy( devClassA, &hostClassA, sizeof( ClassA ), cudaMemcpyHostToDevice );
    cudaMemcpy( devClassB, &hostClassB, sizeof( ClassB ), cudaMemcpyHostToDevice );

    validClassKernel<<<1,1>>>();
    cudaDeviceSynchronize();
    cudaError_t error = cudaGetLastError();
    if(error!=cudaSuccess)
    {
        fprintf(stderr,"ERROR: validClassKernel: %s\n", cudaGetErrorString(error) );
    }

    invalidClassKernel<<<1,1>>>( devClassA );
    cudaDeviceSynchronize();
    error = cudaGetLastError();
    if(error!=cudaSuccess)
    {
        fprintf(stderr,"ERROR: invalidClassKernel: %s\n", cudaGetErrorString(error) );
    }
}

validClassKernel() 展示了如何将派生对象的指针存储在基类指针数组中并访问虚函数getNumber()。在此示例中，对象是在设备代码中创建的。

invalidClassKernel() 表明您不能在设备代码中使用从在主机上创建的虚拟基类派生的对象的副本。代码编译但内核失败并显示an illegal memory access was encountered。这很可能是原始示例代码中的主要问题。

---

其他问题：

• 您不能在设备代码中使用 std::string，请参阅此问题 Can we use the string data type in C++ within kernels。作为一种解决方案，您可以使用 char 恒定大小的数组。这同样适用于 std::vector (classScores)。