[CUDA] Don't generate aliases for static extern "C" functions.

[jlebar]

Summary

These aliases are done to support inline asm, but there's nothing we can
do: NVPTX doesn't support aliases.

[tra]

Failing silently is not a good idea. At the very least there should produce an error.
The right thing to do here, IMO, would be to generate a stub with alias name that just jumps to or calls aliasee.

[jlebar]

In D16501#335225, @tra wrote:

Failing silently is not a good idea. At the very least there should produce an error.
The right thing to do here, IMO, would be to generate a stub with alias name that just jumps to or calls aliasee.

Discussed IRL. We're not breaking anyone who's not already broken -- if you currently have "device attribute((used)) static int foo()", that's just not going to compile for you without this patch. With this patch, at least this will work for some people; we can always add this stub function later if we want (I have concerns about potentially surprising performance implications, but we can postpone worrying about that).

[tra]

OK. If someone attempts to rely on this feature in CUDA it will be obviously broken due to the missing C-style mangled name. We can deal with it then.

lib/CodeGen/CodeGenModule.cpp

 }
 
 /// For each function which is declared within an extern "C" region and marked
 /// as 'used', but has internal linkage, create an alias from the unmangled
 /// name to the mangled name if possible. People expect to be able to refer
 /// to such functions with an unmangled name from inline assembly within the
 /// same translation unit.
 void CodeGenModule::EmitStaticExternCAliases() {
+  // Don't do anything if we're generating CUDA device code -- the NVPTX
+  // assembly target doesn't support aliases.
+  if (Context.getTargetInfo().getTriple().isNVPTX())
+    return;
   for (auto &I : StaticExternCValues) {
     IdentifierInfo *Name = I.first;
     llvm::GlobalValue *Val = I.second;
     if (Val && !getModule().getNamedValue(Name->getName()))
       addUsedGlobal(llvm::GlobalAlias::create(Name->getName(), Val));
   }
 }
 

test/CodeGenCUDA/alias.cu

+// REQUIRES: x86-registered-target
+// REQUIRES: nvptx-registered-target
+
+// RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -emit-llvm \
+// RUN:   -o - %s | FileCheck %s
+
+#include "Inputs/cuda.h"
+
+// Check that we don't generate an alias from "foo" to the mangled name for
+// ns::foo() -- nvptx doesn't support aliases.
+
+namespace ns {
+extern "C" {
+// CHECK-NOT: @foo = internal alias
+__device__ __attribute__((used)) static int foo() { return 0; }
+}
+}