The CUDA programming guide explicitly prohibits passing an object of any class with virtual functions to a __global__ function [1]. To my knowledge, Nvidia hasn't supplied an official explanation for why this isn't possible, but I surmise that it's because the vtable lookup results in an indirect jump, which isn't possible in device code [2].
In my project, I have a set of classes with virtual functions that rely on dynamic polymorphism in host code. I'd like to pass instances of these classes to the device but supply type information to the compiler via a template instead of using dynamic polymorphism. Even though I don't require dynamic polymorphism for my device code, I would like to be able to re-use my original classes with virtual functions instead of re-implementing them to remove the virtual functions.
Is it possible to write my code such that a vtable lookup is avoided in this case? If so, is it safe to rely on this behavior or should I be wary of using this in my project?
In the following minimal example, I'm passing a derived class object with virtual functions to a kernel function, but there is no reference to the base class or need for dynamic polymorphism. The compiler is aware of the type of the argument at compile-time and need not generate code that performs a vtable lookup.
// example.cu
#include <cstdio>
class Base {
public:
__host__ __device__ virtual void foo() const = 0;
virtual ~Base() = default;
};
class Derived : public Base {
public:
__host__ __device__ void foo() const override
{
std::printf("Called Derived::foo()\n");
}
};
template<class T>
__global__ void bar(const T derived)
{
derived.foo();
}
int main()
{
Derived d;
bar<Derived><<<1, 1>>>(d);
if (cudaPeekAtLastError() != cudaSuccess) return 1;
if (cudaDeviceSynchronize() != cudaSuccess) return 2;
return 0;
}
$ nvcc --run example.cu && echo $?
Called Derived::foo()
0
This example successfully compiles and runs for me with nvcc v9.0 (with gcc v6.1 as the host compiler). Is it possible to safely use this idiom, or am I relying on undefined behavior?
