Cuda 虚拟类

Question

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

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

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

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

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

附：这实际上是this 问题的重播，遗憾的是从未完全回答。

编辑：

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

那么首先让我们看看调用虚方法时Cuda在做什么，下面是调试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
};

将其存储在主机和设备内存中（内存位置是设备内存），在主机上，您可以使用 ClassName 作为“修补”对象的查找。

但再说一遍：我不会在任何严肃的事情上使用它，因为在性能方面，虚函数根本不是很好。

Answer 1

当前，CUDA 编译器和运行时（从 CUDA 5.0 开始）不支持您尝试执行的操作。 CUDA C Programming Guide v5.0 的 D.2.6.3 部分内容如下：

D.2.6.3 虚拟功能

当派生类中的函数覆盖基类中的虚函数时，被覆盖函数和覆盖函数上的执行空间限定符（即__host__、__device__）必须匹配。

不允许将类的对象作为参数传递给__global__ 函数 带有虚函数。

编译器将虚函数表放置在全局或常量内存中。

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

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

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