c++ - Cuda虚拟类

Question

我想在 cuda 内核中执行一些虚拟方法，但我不想在同一个内核中创建对象，而是想在主机上创建它并将其复制到 gpu 内存。

我正在内核中成功创建对象并调用虚拟方法。复制对象时出现问题。这是有道理的，因为显然虚函数指针是伪造的。发生的只是“Cuda grid launch failed”，至少 Nsight 是这么说的。但是当查看 SASS 时，它会在取消引用虚函数指针时崩溃，这是有道理的。

我当然在使用 Cuda 4.2 以及在适配卡上使用“compute_30”进行编译。

那么推荐的方法是什么？还是根本不支持此功能？

我有想法首先运行一个不同的内核，它创建虚拟对象并提取虚函数指针以在复制它们之前“修补”我的对象。遗憾的是，这并没有真正起作用(还没有弄清楚)，而且这将是一个丑陋的解决方案。

附言这实际上是 this 的重新运行遗憾的是，这个问题从未得到完全回答。

编辑:

所以我找到了一种方法来做我想做的事。但要明确一点:这根本不是答案或解决方案，答案已经提供，这只是一个 hack，只是为了好玩。

所以首先让我们看看Cuda在调用虚方法时做了什么，下面是debug SASS

//R0 is the address of our object
LD.CG R0, [R0];  
IADD R0, R0, 0x4;  
NOP;  
MOV R0, R0;  
LD.CG R0, [R0];
...
IADD R0, RZ, R9;  
MOV R0, R0;  
LDC R0, c[0x2][R0];
...
BRX R0 - 0x5478

因此假设“c[0x2][INDEX]”对于所有内核都是常量，我们可以通过运行内核并执行此操作来获取类的索引，其中 obj 是类的新创建对象:

unsigned int index = *(unsigned int*)(*(unsigned int*)obj + 4);

然后使用这样的东西:

struct entry
{
    unsigned int vfptr;// := &vfref, thats our value to store in an object
    int dummy;// := 1234, great for debugging
    unsigned int vfref;// := &dummy
    unsigned int index;
    char ClassName[256];//use it as a key for a dict
};

将其存储在主机和设备内存中(内存位置是设备内存)，在主机上您可以使用类名查找要“修补”的对象。

但是再说一次:我不会在任何严肃的事情上使用它，因为从性能角度来看，虚函数根本不是很好。

Answer 1

目前，CUDA 编译器和运行时(自 CUDA 5.0 起)不支持您尝试执行的操作。 CUDA C 编程指南 v5.0 的 D.2.6.3 节内容如下:

D.2.6.3 Virtual Functions

When a function in a derived class overrides a virtual function in a base class, the execution space qualifiers (i.e., __host__, __device__) on the overridden and overriding functions must match.

It is not allowed to pass as an argument to a __global__ function an object of a class with virtual functions.

The virtual function table is placed in global or constant memory by the compiler.

我建议您将类的数据与类的功能分开封装。例如，将数据存储在结构中。如果您计划对这些对象的数组进行操作，请将数据存储在数组结构中(为了性能——超出了这个问题的范围)。使用 cudaMalloc 在主机上分配数据结构，然后将数据作为参数传递给内核，而不是通过虚拟方法传递类。

然后在设备上使用虚拟方法构造您的对象。具有虚方法的类的构造函数会将设备指针内核参数作为参数。然后，虚拟设备方法可以对设备数据进行操作。

同样的方法可以在设备的一个内核中分配数据，并在设备的另一个内核中访问它(同样，具有虚函数的类不能作为内核的参数)。