如何在cuda中使用虚拟课堂
How to use virtual class in cuda?
我写了一个简单的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, ¶ms, 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
)也是如此。
相关文章:
- 虚拟决赛作为安全
- PowerPC ppc64le上的Gcc Woverloaded虚拟错误
- 如何在C++中获得"静态纯虚拟"功能?
- C++无法定义虚拟函数 OUTER 类和头文件
- 用常见虚拟函数实现的任意组合来实现派生类的正确方法是什么
- 在模板基类中为继承类中的可选重写生成虚拟方法
- 尝试将unique_ptrs推送到向量时使用纯虚拟函数错误
- 有没有比在库中添加一个并非由所有派生类实现的新虚拟函数更好的设计实践
- 大小虚拟继承中的派生类
- 链接器找不到在虚拟类 c++ 中访问的静态字段的符号
- 使用 C++ 和 i2c 工具从虚拟 i2c 写入和读取
- 重载 -> shared_ptr 个实例中的箭头运算符<interface>,接口中没有纯虚拟析构函数
- 如果整个应用程序是虚拟映射的,为什么 new 会进行系统调用?
- 跨 DLL 边界访问虚拟方法是否安全/可能?
- std::is_trivially_copyable_v 关于虚拟功能
- 删除C++继承中虚拟类成员的代码重复
- 尝试模拟纯虚拟课堂
- 使用虚拟成员从课堂制作 POD
- Cuda虚拟课堂
- 如何在cuda中使用虚拟课堂