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

[Justin Lebar via cfe-commits]

jlebar added a comment.

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


At the moment, getting this feature to work seems more important than maintaining ABI compatibility with NVCC.  But I cannot confidently assign a probability to how likely it will be at some point in the future that we'll want this ABI compatibility.  I really don't know.

So, that's one option.  Here's another:

The motivation behind this one is, we have this pie-in-the-sky notion that, morally, device code should be able to call anything it wants.  Only if we cannot codegen for device a function transitively invoked by a device function will we error out.  constexpr-is-implicitly-HD is a step towards this more ambitious goal.

Setting aside the constexpr bit, it seems to me that when we codegen an unattributed function for device, we should mark the function as having internal linkage (or whatever the thing is called such that it's not visible from other TUs).  The reason is, other TUs cannot rely on this function being present in the first object file, because the function is only generated on-demand.  If you want to call an HD function defined in another .cu file, then the header in both files needs to explicitly define it as HD.

If that is true -- that unattributed functions which we codegen for device can/should be made internal -- then the mangling of those names has no bearing on ABI compatibility.  So we could say, no explicit-HD / D or explicit-HD / H overloading, but *implicit*-HD / D overloading is OK, and we will mangle implicit-HD functions differently to allow this.

Does that sound like it might work?


http://reviews.llvm.org/D18458