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

[Justin Lebar via cfe-commits]

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

(At the moment we treat all device functions as internal, but changing
that is definitely on the radar.)

It does not seem that external declarations of functions implicitly
converted to HD is compatible with nvcc's ABI, if we allow the
implicit-HD functions to overload explicit-D functions.

The reason is, suppose I have two TUs:

  a.cu:
    void foo();  // implicitly HD
  b.cu:
    __device__ void foo();

To be ABI compatible with nvcc, both overloads must have the same
mangled name, which means this is an ODR violation, yes?  But saying
that this program is an error means that we don't allow overloading of
implicit-HD with explicit-D.  And anyway that seems pretty wrong to me
-- why does the presence of some unattributed function you don't care
about make your __device__ function decl invalid?  IOW the nvcc
behavior seems broken.

So I guess we're left with breaking ABI compatibility with nvcc.  In
which case the only remaining questions are:

a) How much overloading of HD do we want to allow?  One side of the
spectrum is to allow all HD / D and HD / H overloading.  The other
side is to allow only implicit-HD / D overloading.  And there are
points in between.

b) Do we want to mangle all HD functions differently, or just the HD
functions which can be overloaded?



On Mon, Mar 28, 2016 at 9:23 AM, Richard Smith <richard at metafoo.co.uk> wrote:
> 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