From: Justin Lebar <jlebar@google.com>
Date: Mon, 25 Jan 2016 22:36:37 +0000 (+0000)
Subject: [CUDA] Don't generate aliases for static extern "C" functions.
X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=bf6ef01db4fd57637e17c1523740552605c4844a;p=clang

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

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

Reviewers: tra

Subscribers: cfe-commits, jhen, echristo

Differential Revision: http://reviews.llvm.org/D16501

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@258734 91177308-0d34-0410-b5e6-96231b3b80d8
---

diff --git a/lib/CodeGen/CodeGenModule.cpp b/lib/CodeGen/CodeGenModule.cpp
index 5a85de7e23..219f95e635 100644
--- a/lib/CodeGen/CodeGenModule.cpp
+++ b/lib/CodeGen/CodeGenModule.cpp
@@ -3836,6 +3836,10 @@ static void EmitGlobalDeclMetadata(CodeGenModule &CGM,
 /// 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;
diff --git a/test/CodeGenCUDA/alias.cu b/test/CodeGenCUDA/alias.cu
new file mode 100644
index 0000000000..6efff6b92a
--- /dev/null
+++ b/test/CodeGenCUDA/alias.cu
@@ -0,0 +1,17 @@
+// 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; }
+}
+}