"Selective usage" of __device__ in template class

Dear All,

I ran into an issue while experimenting with some non-trivial code designs. Maybe you’ll have more of an idea about this than I do right now…

I am trying to create a templated class that would provide the same user interface in host and device code “for what it does”. (That’s not super relevant to this issue.) The setup is basically that I would use the class in host code to set up some non-trivial data in memory, that the top level object is operating on. I then copy this non-trivial data to the GPU’s global memory, and in my kernel I would construct a thin UI on top of that data that provides smart access to that data.

To do this, this smart class receives a “vector type” as a template parameter. Which in the host code would be std::vector, but in the device code it would be a simple type that provides an std::vector like interface on top of the array allocated in the global memory. I do this with code like:

template< template< typename > class vector_type >
class SmartClass {
   __host__ __device__
   __host__ __device___
   SmartClass( const SmartClassData& );

   SmartClassData data() const;

Of course the whole setup is a bit more complicated, but this is the general way in which I pass a “vector type” to the class. Then, in a piece of test code I do something like:

__global__ someKernel( SmartClassData data ) {
   SmartClass< device_vector > deviceObject( data );
int main() {
   SmartClass< std::vector > hostObject;
   auto deviceData = hostObject.data();
   someKernel<<<...>>>( deviceData );

Now… This actually works. I can generate binary code with this that behaves as I would expect it to. But I’m getting an annoying, but understandable warning from the compilation:

/data/software/cuda/11.2.0/x86_64-ubuntu1804/bin/../targets/x86_64-linux/include/vector_types.h(421): warning: calling a __host__ function("std::vector< ::detray::dummy_surface, ::std::allocator< ::detray::dummy_surface> > ::vector") from a __host__ __device__ function("std::vector< ::detray::dummy_surface, ::std::allocator< ::detray::dummy_surface> > ::vector [subobject]") is not allowed

/data/software/cuda/11.2.0/x86_64-ubuntu1804/bin/../targets/x86_64-linux/include/vector_types.h(421): warning: calling a __host__ function("std::vector< ::detray::dummy_volume<    ::device_volume_vector> , ::std::allocator< ::detray::dummy_volume<    ::device_volume_vector> > > ::vector") from a __host__ __device__ function("std::vector< ::detray::dummy_volume<    ::device_volume_vector> , ::std::allocator< ::detray::dummy_volume<    ::device_volume_vector> > > ::vector [subobject]") is not allowed

Since nvcc sees that I instantiate the type with std::vector in the host code, it complains that in that instantiation __device__ is not valid for the default constructor.

Now… I could just do the simple thing, and never allow using the default constructor of this type in device code. For this particular class that could actually be made to work. But I was wondering if there may be a better solution to this issue. Where I could control which function is marked with __device__ and __host__ using some “template magic”. ;-)