r258734 - [CUDA] Don't generate aliases for static extern "C" functions.
Justin Lebar via cfe-commits
cfe-commits at lists.llvm.org
Mon Jan 25 14:36:38 PST 2016
Author: jlebar
Date: Mon Jan 25 16:36:37 2016
New Revision: 258734
URL: http://llvm.org/viewvc/llvm-project?rev=258734&view=rev
Log:
[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
Added:
cfe/trunk/test/CodeGenCUDA/alias.cu
Modified:
cfe/trunk/lib/CodeGen/CodeGenModule.cpp
Modified: cfe/trunk/lib/CodeGen/CodeGenModule.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CodeGenModule.cpp?rev=258734&r1=258733&r2=258734&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CodeGenModule.cpp (original)
+++ cfe/trunk/lib/CodeGen/CodeGenModule.cpp Mon Jan 25 16:36:37 2016
@@ -3836,6 +3836,10 @@ static void EmitGlobalDeclMetadata(CodeG
/// 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;
Added: cfe/trunk/test/CodeGenCUDA/alias.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/alias.cu?rev=258734&view=auto
==============================================================================
--- cfe/trunk/test/CodeGenCUDA/alias.cu (added)
+++ cfe/trunk/test/CodeGenCUDA/alias.cu Mon Jan 25 16:36:37 2016
@@ -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; }
+}
+}
More information about the cfe-commits
mailing list