如何在cuda中使用虚拟课堂

How to use virtual class in cuda?

本文关键字:虚拟 课堂 cuda      更新时间:2023-10-16

我写了一个简单的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("%sn", alpha->className);
    printf("%sn", beta->className); 
    printf("%sn", cuda->ClassName);
    printf("alpha avg = %dn", alpha->classAvgGrade);
    printf("beta avg = %dn", beta->classAvgGrade);
    printf("cuda avg = %dn", 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

当我运行这个时,我没有得到正确的值,得到以下结果:

alpha avg = 39696816

beta avg = 70

cuda avg = 0

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

OP提出了几个问题。主要问题在标题"如何在cuda中使用虚拟课堂?"一个不相关的问题是如何在cuda代码中使用字符串。我将主要关注题目中的问题。

根据cuda c编程指南,您可以使用虚函数,但有限制。您遇到的限制是

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

在您的示例代码中,您试图通过常量内存传递对象(设备指针数组)来避免这种情况。然而,我认为编程指南在这里并不精确。我认为不可能将从虚拟基类派生的类的对象复制到设备。问题是(据我所知)您将复制主机虚拟函数表到设备。

示例代码太复杂(并且有其他问题),无法演示该行为。下面的简化代码展示了在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(): %dn", superClass->getNumber() );
}
__global__ void validClassKernel()
{
    Class* classVector[2];
    classVector[0] = new ClassA();
    classVector[1] = new ClassB();
    printf( "classVector[0]->getNumber(): %dn", classVector[0]->getNumber() );
    printf( "classVector[1]->getNumber(): %dn", 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: %sn", cudaGetErrorString(error) );
    }
    invalidClassKernel<<<1,1>>>( devClassA );
    cudaDeviceSynchronize();
    error = cudaGetLastError();
    if(error!=cudaSuccess)
    {
        fprintf(stderr,"ERROR: invalidClassKernel: %sn", cudaGetErrorString(error) );
    }
}

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

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


其他问题:

  • 你不能在设备代码中使用std::string,参见这个问题我们能在内核中使用c++中的字符串数据类型吗?作为一种解决方案,您可以使用固定大小的char数组。std::vector (classScores)也是如此。