如何解决有什么方法可以将vtable从主机复制到设备CUDA和C ++
由于某些与“虚拟表”或“虚拟指针”相关的原因,Cuda似乎不允许我“将从虚拟基类派生的类的对象传递给__global__
函数”。
我想知道是否可以通过某种方式手动设置“虚拟指针”,以便可以使用多态性?
解决方法
有什么方法可以将vtable从主机复制到设备
您不想将vtable从主机复制到设备。主机上的vtable(即在主机上创建的对象中)在vtable中具有一组主机函数指针。将此类对象复制到设备时,vtable不会被更改或“固定”,因此最终在设备上出现一个对象,该对象的vtable充满了主机指针。
如果您随后尝试调用这些虚拟功能之一(使用设备代码中的设备上的对象),则会发生不良情况。 vtable中列出的数字函数入口点是在设备代码中没有任何意义的地址。
以便我可以使用多态性
我建议在设备代码中使用多态的方法是在设备上创建对象。这使用一组设备函数指针(而不是主机函数指针)设置了vtable,并且诸如this之类的问题证明了它可以工作。一阶近似,如果您有一种方法可以在宿主代码中创建一组多态对象,那么我不知道您为什么不能在设备代码中使用类似方法。问题确实与互操作性有关-在主机和设备之间移动此类对象-the stated limitations in the programming guide所指的是这种情况。
我想知道是否可以通过某种方式手动设置“虚拟指针”
可能有。为了共享知识,我将概述一种方法。但是,我对C ++的了解还不够多,很难说这是否可以接受/合法。我唯一可以说的是在非常有限的测试中,它似乎可以正常工作。 但是我认为这是不合法的,因此我不建议您将这种方法用于实验以外的其他用途。即使我们不确定该方法是否合法,也已经有明确规定CUDA限制(如上所述),您不应尝试在主机和设备之间传递具有虚拟功能的对象。因此,我仅提供它作为观察,这可能对实验或研究很有趣。我不建议将其用于生产代码。
this thread中概述了基本思想。它基于这样的想法,即普通的对象复制似乎并不复制虚拟函数指针表,这对我来说很有意义,但是整个对象确实包含该表。因此,如果我们使用这样的方法:
template<typename T>
__device__ void fixVirtualPointers(T *other) {
T temp = T(*other); // object-copy moves the "guts" of the object w/o changing vtable
memcpy(other,&temp,sizeof(T)); // pointer copy seems to move vtable
}
似乎可以获取给定的对象,创建该类型的新“虚拟”对象,然后通过对对象进行基于指针的复制来“修复” vtable(考虑整个对象的大小)而不是“典型的”对象副本。使用此方法后果自负。 This blog也许也很有趣,尽管我不能保证那里任何陈述的正确性。
除此之外,cuda
标签上还有许多其他建议,您不妨查看them。
我想提供一种不同的方法来修复vtable,它不依赖于在对象之间复制vtable。这个想法是在设备上使用new放置,以使编译器生成适当的vtable。但是,这种方法也违反了编程指南中规定的限制。
#include <cstdio>
struct A{
__host__ __device__
virtual void foo(){
printf("A\n");
}
};
struct B : public A{
B(int i = 13) : data(i){}
__host__ __device__
virtual void foo() override{
printf("B %d\n",data);
}
int data;
};
template<class T>
__global__
void fixKernel(T* ptr){
T tmp(*ptr);
new (ptr) T(tmp);
}
__global__
void useKernel(A* ptr){
ptr->foo();
}
int main(){
A a;
a.foo();
B b(7);
b.foo();
A* ab = new B();
ab->foo();
A* d_a;
cudaMalloc(&d_a,sizeof(A));
cudaMemcpy(d_a,&a,sizeof(A),cudaMemcpyHostToDevice);
B* d_b;
cudaMalloc(&d_b,sizeof(B));
cudaMemcpy(d_b,&b,sizeof(B),cudaMemcpyHostToDevice);
fixKernel<<<1,1>>>(d_a);
useKernel<<<1,1>>>(d_a);
fixKernel<<<1,1>>>(d_b);
useKernel<<<1,1>>>(d_b);
cudaDeviceSynchronize();
cudaFree(d_b);
cudaFree(d_a);
delete ab;
}
版权声明:本文内容由互联网用户自发贡献,该文观点与技术仅代表作者本人。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如发现本站有涉嫌侵权/违法违规的内容, 请发送邮件至 dio@foxmail.com 举报,一经查实,本站将立刻删除。