[PATCH] D18458: [CUDA] Mangle __host__ __device__ functions differently than __host__ or __device__ functions.
Justin Lebar via cfe-commits
cfe-commits at lists.llvm.org
Fri Mar 25 14:38:45 PDT 2016
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.
More information about the cfe-commits