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

[Richard Smith via cfe-commits]

rsmith added a comment.

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


================
Comment at: lib/AST/ItaniumMangle.cpp:488-489
@@ +487,4 @@
+  // 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 =
----------------
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?


http://reviews.llvm.org/D18458