[CUDA] Mangle __host__ __device__ functions differently than __host__ or __device__ functions.

[jlebar]

Summary

This is important because in a later patch, we will allow host
device functions to be overloaded with host / device
functions with the same signature, and we don't want a naming conflict
in this case.

Based on changes from http://reviews.llvm.org/D12453?vs=on&id=33483.

[rsmith]

This makes the "constexpr implies __host__ __device__" patch look slightly questionable: two translation units defining the same constexpr function will mangle that function differently depending on whether the translation unit is built with CUDA support enabled. That will cause you to get duplicates of static locals and the like (but I suppose you do anyway between the host and the device, so maybe that's not much more broken than it would be regardless).

lib/AST/ItaniumMangle.cpp

488–489

According to http://mentorembedded.github.io/cxx-abi/abi.html#mangling-type, order-insensitive attributes should be sorted into *reverse* alphabetic order (alphabetically-first goes nearest to the base type).

Given that enable_if is order-sensitive but host and device are not, I'm not really sure what the Itanium ABI expects us to do here regarding their relative order. John?

[jlebar]

In D18458#383266, @rsmith wrote:

This makes the "constexpr implies __host__ __device__" patch look slightly questionable: two translation units defining the same constexpr function will mangle that function differently depending on whether the translation unit is built with CUDA support enabled. That will cause you to get duplicates of static locals and the like (but I suppose you do anyway between the host and the device, so maybe that's not much more broken than it would be regardless).

Hm. I agree that would be surprising to users.

I suppose we could say that HD can overload D but not H. Then we could apply mangling to all device functions, solving the ambiguity that way. That feels pretty ham-fisted, though. :-/

[jlebar]

! In D18458#383276, @jlebar wrote:

! In D18458#383266, @rsmith wrote:

This makes the "constexpr implies __host__ __device__" patch look slightly questionable: two translation units defining the same constexpr function will mangle that function differently depending on whether the translation unit is built with CUDA support enabled. That will cause you to get duplicates of static locals and the like (but I suppose you do anyway between the host and the device, so maybe that's not much more broken than it would be regardless).

The breakage seems to be worse than this. :( Eigen seems to do the following:

foo.h:
  #ifdef __CUDACC__  // If compiling CUDA code
  #define HOST_DEVICE __host__ __device__
  #else
  #define HOST_DEVICE
  #endif

  HOST_DEVICE void foo();

foo.cc:  // Compiled as CUDA
  HOST_DEVICE void foo() { ... }

bar.cc:  // *Not* compiled as CUDA
  #include "foo.h"
  void bar() { foo(); }

With this patch, foo() has a different mangled name in foo.o and bar.cc, and
we're hosed.

If we think this use-case is reasonable (I think it is?) I think this means
that we cannot mangle host device functions differently when doing host
compilation. That seems to restrict us to saying that H and HD functions with
the same signatures cannot overload. This leaves us with two options:

1. No overloading between HD and H or D functions with the same signature.
   
   I don't see how to do this while still letting constexpr be HD; the issue is that there are constexpr std math functions which we want to overload for device. We could let constexpr be something other than HD, but if that new thing can overload with D, then I think we still have the same problem.
2. No overloading between HD and H, but OK to overload HD and D.
   
   If we did this, we'd still need to give D functions a different mangled name. But we don't have this problem of referencing symbols defined in a file compiled in CUDA mode from a file compiled without CUDA.
   
   tra pointed out a problem with this, which is that if someone (say, nvidia) gave us a C++ library consisting of precompiled device code plus headers, we wouldn't be able to link with it, because we would use different mangling.
   
   I also don't like this because it's inconsistent to say HD can overload D but not H. But that's a minor point at this point.

Richard, what do you think? Maybe you have an alternative idea?

[rsmith]

It seems like we have the following constraint: on host, no attributes must mangle the same as __host__ __device__ and constexpr (and probably __global__?).

Are there any others? What do we need to do to be ABI-compatible with NVCC? (And is that possible if we allow __host__ to overload __host__ __device__?)

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.

[jlebar]

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

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

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.

[rsmith]

In D18458#383755, @jlebar wrote:

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

[jlebar]

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?

[jlebar]

New plan, R2: Let nvcc win.

After much discussion, we're abandoning this because we want to maintain abi compatibility with nvcc. I'm about to upload a revised approach to D18380 that won't require this.

lib/AST/ItaniumMangle.cpp

 void CXXNameMangler::mangleFunctionEncoding(const FunctionDecl *FD) {
   // <encoding> ::= <function name> <bare-function-type>
   mangleName(FD);
 
   // Don't mangle in the type if this isn't a decl we should typically mangle.
   if (!Context.shouldMangleDeclName(FD))
     return;
 
+  // CUDA __host__ __device__ functions co-exist with both __host__ and
+  // __device__ functions, so they need a different mangled name.  We sort
+  // "device", "host", and "enable_if" attrs alphabetically.
+  bool IsCudaHostDevice =
+      FD->hasAttr<CUDADeviceAttr>() && FD->hasAttr<CUDAHostAttr>();
+  if (IsCudaHostDevice)
+    Out << "Ua6device";
+
   if (FD->hasAttr<EnableIfAttr>()) {
     FunctionTypeDepthState Saved = FunctionTypeDepth.push();
     Out << "Ua9enable_ifI";
     // FIXME: specific_attr_iterator iterates in reverse order. Fix that and use
     // it here.
     for (AttrVec::const_reverse_iterator I = FD->getAttrs().rbegin(),
                                          E = FD->getAttrs().rend();
          I != E; ++I) {
       EnableIfAttr *EIA = dyn_cast<EnableIfAttr>(*I);
       if (!EIA)
         continue;
       Out << 'X';
       mangleExpression(EIA->getCond());
       Out << 'E';
     }
     Out << 'E';
     FunctionTypeDepth.pop(Saved);
   }
 
+  if (IsCudaHostDevice)
+    Out << "Ua4host";
+
   // Whether the mangling of a function type includes the return type depends on
   // the context and the nature of the function. The rules for deciding whether
   // the return type is included are:
   //
   //   1. Template functions (names or types) have return types encoded, with
   //   the exceptions listed below.
   //   2. Function types not appearing as part of a function name mangling,
   //   e.g. parameters, pointer types, etc., have return type encoded, with the

test/CodeGenCUDA/convergent.cu

 
 // DEVICE: Function Attrs:
 // DEVICE-SAME: convergent
 // DEVICE-NEXT: define void @_Z3foov
 __device__ void foo() {}
 
 // HOST: Function Attrs:
 // HOST-NOT: convergent
-// HOST-NEXT: define void @_Z3barv
+// HOST-NEXT: define void @_Z3barUa6deviceUa4hostv
 // DEVICE: Function Attrs:
 // DEVICE-SAME: convergent
-// DEVICE-NEXT: define void @_Z3barv
+// DEVICE-NEXT: define void @_Z3barUa6deviceUa4hostv
 __host__ __device__ void baz();
 __host__ __device__ void bar() {
-  // DEVICE: call void @_Z3bazv() [[CALL_ATTR:#[0-9]+]]
+  // DEVICE: call void @_Z3bazUa6deviceUa4hostv() [[CALL_ATTR:#[0-9]+]]
   baz();
 }
 
-// DEVICE: declare void @_Z3bazv() [[BAZ_ATTR:#[0-9]+]]
+// DEVICE: declare void @_Z3bazUa6deviceUa4hostv() [[BAZ_ATTR:#[0-9]+]]
 // DEVICE: attributes [[BAZ_ATTR]] = {
 // DEVICE-SAME: convergent
 // DEVICE-SAME: }
 // DEVICE: attributes [[CALL_ATTR]] = { convergent }
 
-// HOST: declare void @_Z3bazv() [[BAZ_ATTR:#[0-9]+]]
+// HOST: declare void @_Z3bazUa6deviceUa4hostv() [[BAZ_ATTR:#[0-9]+]]
 // HOST: attributes [[BAZ_ATTR]] = {
 // HOST-NOT: convergent
 // NOST-SAME: }

test/CodeGenCUDA/device-var-init.cu

__device__ void df() {
   static __shared__ UC s_uc;
 }
 
 // CHECK:   call void @_ZN2ECC1Ev(%struct.EC* %ec)
 // CHECK:   call void @_ZN3ETCC1IJEEEDpT_(%struct.ETC* %etc)
 // CHECK:   call void @_ZN2UCC1Ev(%struct.UC* %uc)
 // CHECK:   call void @_ZN3ECIC1Ev(%struct.ECI* %eci)
 // CHECK:   call void @_ZN3NECC1Ev(%struct.NEC* %nec)
-// CHECK:   call void @_ZN3NCVC1Ev(%struct.NCV* %ncv)
-// CHECK:   call void @_ZN3NCFC1Ev(%struct.NCF* %ncf)
-// CHECK:   call void @_ZN4NCFSC1Ev(%struct.NCFS* %ncfs)
+// CHECK:   call void @_ZN3NCVC1EUa6deviceUa4hostv(%struct.NCV* %ncv)
+// CHECK:   call void @_ZN3NCFC1EUa6deviceUa4hostv(%struct.NCF* %ncf)
+// CHECK:   call void @_ZN4NCFSC1EUa6deviceUa4hostv(%struct.NCFS* %ncfs)
 // CHECK:   call void @_ZN3UTCC1IJEEEDpT_(%struct.UTC* %utc)
 // CHECK:   call void @_ZN4NETCC1IJEEEDpT_(%struct.NETC* %netc)
 // CHECK:   call void @_ZN7EC_I_ECC1Ev(%struct.EC_I_EC* %ec_i_ec)
 // CHECK:   call void @_ZN8EC_I_EC1C1Ev(%struct.EC_I_EC1* %ec_i_ec1)
-// CHECK:   call void @_ZN5T_V_TC1Ev(%struct.T_V_T* %t_v_t)
+// CHECK:   call void @_ZN5T_V_TC1EUa6deviceUa4hostv(%struct.T_V_T* %t_v_t)
 // CHECK:   call void @_ZN7T_B_NECC1Ev(%struct.T_B_NEC* %t_b_nec)
 // CHECK:   call void @_ZN7T_F_NECC1Ev(%struct.T_F_NEC* %t_f_nec)
 // CHECK:   call void @_ZN8T_FA_NECC1Ev(%struct.T_FA_NEC* %t_fa_nec)
 // CHECK:   call void @_ZN2UCC1Ev(%struct.UC* addrspacecast (%struct.UC addrspace(3)* @_ZZ2dfvE4s_uc to %struct.UC*))
 // CHECK: ret void
 
 // We should not emit global init function.
 // CHECK-NOT: @_GLOBAL__sub_I

test/CodeGenCUDA/function-overload.cu

 __device__
 #else
 __host__
 #endif
 void wrapper() {
   s_cd_dh scddh;
   // CHECK-BOTH: call void @_ZN7s_cd_dhC1Ev(
   s_cd_hd scdhd;
-  // CHECK-BOTH: call void @_ZN7s_cd_hdC1Ev
+  // CHECK-BOTH: call void @_ZN7s_cd_hdC1EUa6deviceUa4hostv
 
-  // CHECK-BOTH: call void @_ZN7s_cd_hdD1Ev(
+  // CHECK-BOTH: call void @_ZN7s_cd_hdD1EUa6deviceUa4hostv(
   // CHECK-BOTH: call void @_ZN7s_cd_dhD1Ev(
 }
 // CHECK-BOTH: ret void
 
 // Now it's time to check what's been generated for the methods we used.
 
 // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_dhC2Ev(
 // CHECK-HOST:   store i32 11,
 // CHECK-DEVICE: store i32 12,
 // CHECK-BOTH: ret void
 
-// CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_hdC2Ev(
+// CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_hdC2EUa6deviceUa4hostv(
 // CHECK-BOTH:   store i32 31,
 // CHECK-BOTH: ret void
 
-// CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_hdD2Ev(
+// CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_hdD2EUa6deviceUa4hostv(
 // CHECK-BOTH: store i32 32,
 // CHECK-BOTH: ret void
 
 // CHECK-BOTH: define linkonce_odr void @_ZN7s_cd_dhD2Ev(
 // CHECK-HOST:   store i32 21,
 // CHECK-DEVICE: store i32 22,
 // CHECK-BOTH: ret void

test/CodeGenCUDA/mangling.cu

+// REQUIRES: x86-registered-target
+// REQUIRES: nvptx-registered-target
+
+// RUN: %clang_cc1 -triple nvptx-nvidia-cuda -fcuda-is-device -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -o - %s | FileCheck %s
+
+#include "Inputs/cuda.h"
+
+// Check that __host__ __device__ function mangled names explicitly contain
+// "host" and "device" attributes.  This is important because HD overloads may
+// coexist with H and D overloads.
+
+// CHECK: define i32 @_Z11host_deviceUa6deviceUa4hostv()
+__host__ __device__ int host_device() { return 0; }
+
+// The enable_if attribute should appear in-between the device and host attrs
+// in the mangled name.
+// CHECK: define i32 @_Z8enableifUa6deviceUa9enable_if{{.*}}Ua4hostv
+__attribute__((enable_if(1, "")))
+__host__ __device__ int enableif() { return 0; }