[PATCH] D18380: [CUDA] Make unattributed constexpr functions (usually) implicitly host+device.

[Justin Lebar via cfe-commits]

jlebar added a comment.

In http://reviews.llvm.org/D18380#385240, @tra wrote:

> What if instead of permanently sticking HD attributes on the constexpr function, we instead postpone decision to the point of overload resolution and figure out effective attributes or call preference based on contents of the whole overload set regardless of the order the decls were added to the set.


The problem we were trying to prevent by requiring that the __device__ overload come first is:

  constexpr int foo();
  __device__ void bar() { foo(); }
  __device__ int foo();
  __device__ void baz() { foo(); }

In this example, we're forced to instantiate both versions of foo() on the device.  Being lazy about making the first foo HD doesn't help, because at the time we see bar, it's the only option available.

(Instantiating both foos is a problem if they have the same mangling.  And we want them to have the same mangling so we maintain ABI compatibility with nvcc.)


http://reviews.llvm.org/D18380