Disable emitting static extern C aliases for amdgcn target for CUDA

[yaxunl]

Summary

Patch by Greg Rodgers.
Revised and lit test added by Yaxun Liu.

[rjmccall]

lib/CodeGen/CodeGenModule.cpp

4684

Please add a target hook for this instead of building a list of targets that don't support it.

[yaxunl]

Add a target hook as suggested by John.

[rjmccall]

LGTM.

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())
+  // and AMDGCN assembly target doesn't support aliases.
+  if (Context.getTargetInfo().getTriple().isNVPTX() ||
+      (Context.getTargetInfo().getTriple().getArch() == llvm::Triple::amdgcn))
     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
+// REQUIRES: amdgpu-registered-target
 
 // RUN: %clang_cc1 -fcuda-is-device -triple nvptx-nvidia-cuda -emit-llvm \
 // RUN:   -o - %s | FileCheck %s
+// RUN: %clang_cc1 -fcuda-is-device -triple amdgcn -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; }
 }
 }