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

[Richard Smith via cfe-commits]

On 27 Mar 2016 9:56 a.m., "Justin Lebar via cfe-commits" <
cfe-commits at lists.llvm.org> wrote:
>
> 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?

Not completely. That would break the semantics of things like static local
variables in inline functions, uniqueness of local types in inline
functions, etc. For an HD function, that's likely to be fine since HD
already breaks those properties by creating two independent versions of the
function (as noted elsewhere, this seems like a surprising thing to
implicitly do to all constexpr functions, but that's a separate change, and
sidesteps some of these issues since you can't have static locals in a
constexpr function -- yet).

> http://reviews.llvm.org/D18458
>
>
>
> _______________________________________________
> cfe-commits mailing list
> cfe-commits at lists.llvm.org
> http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20160328/59c0f92e/attachment.html>