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

Richard Smith via cfe-commits cfe-commits at lists.llvm.org
Fri Mar 25 15:58:24 PDT 2016


rsmith added a comment.

In http://reviews.llvm.org/D18458#383755, @jlebar wrote:

> 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.


OK, that makes things pretty easy (though we don't get the answer we might want): unattributed must be mangled the same as H and HD, so we cannot support overloading H and HD.

> > 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.

> 

> That suggests, to your second question, that it's not possible to maintain ABI compatibility if we allow D or H to overload HD.


OK, so the question for you is, how much ABI compatibility with NVCC are you prepared to give up in order to allow HD / D overloading and HD / H overloading?

> > 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.


Yes, that seems like a good argument for mangling H the same as unattributed.


http://reviews.llvm.org/D18458





More information about the cfe-commits mailing list