diff options
| -rw-r--r-- | clang/lib/CodeGen/CodeGenModule.cpp | 4 | ||||
| -rw-r--r-- | clang/test/CodeGenCUDA/alias.cu | 17 |
2 files changed, 21 insertions, 0 deletions
diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 5a85de7e234..219f95e635e 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/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/clang/test/CodeGenCUDA/alias.cu b/clang/test/CodeGenCUDA/alias.cu new file mode 100644 index 00000000000..6efff6b92aa --- /dev/null +++ b/clang/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; } +} +} |
