Possible to automatically create __device__ (rvalue/copy reference/...) constructors?

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi,

I’m working on an application where I need to run an rvalue constructor on the device for many classes. The default rvalue constructor, by default, compiles for __host__, as we all know. I need to use the default rvalue constructor, but on the device. From what I can tell, this entails writing down the whole default rvalue constructor (std::move applied to each of the constituent pieces of data in the class), but prefixing it with __device__ __host__. Clearly, this is not a very sustainable software development technique since if someone adds a piece of data to any class (there are around 100 classes in this piece of software), they have to also modify the custom constructor.

Is it possible to generate the default constructors, but enabling them for __device__? I tried this:

struct DeviceConstructible
{
  __device__ __host__ DeviceConstructible(DeviceConstructible&& other) = default;
};

But sadly doesn’t give the effect I’m looking for. Looks like i may have to cook up some nasty macro or clang script to generate these automatically, huh? Or am I missing some CUDA feature?

Thanks,

Gavin

2

AFAIK for explicit default constructors/destructors:

… = default;

CUDA should automatically decorate these for you. What happens if you just drop the __device__ __host__ decoration?

Perhaps I am not understanding the intricacies here.

3

Oh, wow, this is fantastic to hear this works. I’ve tested your suggestion and it works perfectly… but… but… the manual says (B.1.3):

It is equivalent to declare a function with only the host execution space specifier or todeclare it without any of the host, device, or global execution space specifier;in either case the function is compiled for the host only.

I never thought to try that because of this. Thank you for pointing out this exception to the manual’s stated rule.

4

Doc support:

G.4.9.2. Implicitly-declared and explicitly-defaulted functions

Let F denote a function that is either implicitly-declared or is explicitly-defaulted on its first declaration The execution space specifiers (host, device) for F are the union of the execution space specifiers of all the functions that invoke it (note that a __global__ caller will be treated as a __device__ caller for this analysis).

(emphasis added)

1 Like