如何在cuda中使用虚拟类？

Question

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

这就是我所拥有的：

测试.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++;
    };
 };

测试.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 编程指南，您可以使用虚拟函数，但有限制。您遇到的限制是

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

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

示例代码太复杂（并且存在其他问题），无法演示该行为。以下更简化的代码显示了您可以使用 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++ inside kernels。作为解决方案，您可以使用char恒定大小的数组。这同样适用于 std::vector ( classScores)。