[PATCH] D18458: [CUDA] Mangle __host__ __device__ functions differently than __host__ or __device__ functions.

[Justin Lebar via cfe-commits]

jlebar added a comment.

In http://reviews.llvm.org/D18458#383719, @rsmith wrote:

> It seems like we have the following constraint: on host, no attributes must mangle the same as `__host__ __device__` and `constexpr` (and probably `__global__`?).


Yes to `__host__ __device__` and `constexpr`.  Unsure about `__global__`, but let's also say yes for now, to be conservative.

> Are there any others?


An existing assumption is that `__host__` is identical to unattributed.  Probably makes sense to keep that one around for now if we can (modulo changes to unattributed constexpr), as it makes things simpler.

> What do we need to do to be ABI-compatible with NVCC? (And is that possible if we allow `__host__` to overload `__host__ __device__`?)


NVCC doesn't apply any special mangling to D or HD functions, so I think maintaining naming compatibility means, basically, not screwing with mangled names based on attributes.

> One possibility given only that constraint would be to use a different mangling for H functions and D functions, but mangle HD and unattributed functions the same.


I guess using a different mangling for both H and D functions, rather than just for D functions, is in some sense more consistent.  But this  would also be very subtle: We'd be saying, non-constexpr H and unattributed are identical, *except* for their mangled names.


http://reviews.llvm.org/D18458