6

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

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

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

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

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

PS这实际上是这个问题的重播，遗憾的是从未完全回答过。

编辑 ：

所以我找到了一种方法来做我想做的事。但要明确一点：这根本不是答案或解决方案，答案已经提供，这只是一个 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 作为“修补”对象的查找。

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

4

1 回答 1

6

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

D.2.6.3 虚拟功能

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

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

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

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

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

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

于 2012-10-03T07:46:34.663 回答