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())
+  if (!getTargetCodeGenInfo().shouldEmitStaticExternCAliases())
     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));
   }
 }
 
 bool CodeGenModule::lookupRepresentativeDecl(StringRef MangledName,
                                              GlobalDecl &Result) const {
   auto Res = Manglings.find(MangledName);

lib/CodeGen/TargetInfo.h

public:
   /// a wrapper for the block invoke function with target-specific calling
   /// convention and ABI as an OpenCL kernel. The wrapper function accepts
   /// block context and block arguments in target-specific way and calls
   /// the original block invoke function.
   virtual llvm::Function *
   createEnqueuedBlockKernel(CodeGenFunction &CGF,
                             llvm::Function *BlockInvokeFunc,
                             llvm::Value *BlockLiteral) const;
+
+  /// \return true if the target supports alias from the unmangled name to the
+  /// mangled name of functions declared within an extern "C" region and marked
+  /// as 'used', and having internal linkage.
+  virtual bool shouldEmitStaticExternCAliases() const { return true; }
 };
 
 } // namespace CodeGen
 } // namespace clang
 
 #endif // LLVM_CLANG_LIB_CODEGEN_TARGETINFO_H

lib/CodeGen/TargetInfo.cpp

 
 class NVPTXTargetCodeGenInfo : public TargetCodeGenInfo {
 public:
   NVPTXTargetCodeGenInfo(CodeGenTypes &CGT)
     : TargetCodeGenInfo(new NVPTXABIInfo(CGT)) {}
 
   void setTargetAttributes(const Decl *D, llvm::GlobalValue *GV,
                            CodeGen::CodeGenModule &M) const override;
+  bool shouldEmitStaticExternCAliases() const override;
 
 private:
   // Adds a NamedMDNode with F, Name, and Operand as operands, and adds the
   // resulting MDNode to the nvvm.annotations MDNode.
   static void addNVVMMetadata(llvm::Function *F, StringRef Name, int Operand);
 };
 
 ABIArgInfo NVPTXABIInfo::classifyReturnType(QualType RetTy) const {

 
   llvm::Metadata *MDVals[] = {
       llvm::ConstantAsMetadata::get(F), llvm::MDString::get(Ctx, Name),
       llvm::ConstantAsMetadata::get(
           llvm::ConstantInt::get(llvm::Type::getInt32Ty(Ctx), Operand))};
   // Append metadata to nvvm.annotations
   MD->addOperand(llvm::MDNode::get(Ctx, MDVals));
 }
+
+bool NVPTXTargetCodeGenInfo::shouldEmitStaticExternCAliases() const {
+  return false;
+}
 }
 
 //===----------------------------------------------------------------------===//
 // SystemZ ABI Implementation
 //===----------------------------------------------------------------------===//
 
 namespace {
 

   LangAS getGlobalVarAddressSpace(CodeGenModule &CGM,
                                   const VarDecl *D) const override;
   llvm::SyncScope::ID getLLVMSyncScopeID(SyncScope S,
                                          llvm::LLVMContext &C) const override;
   llvm::Function *
   createEnqueuedBlockKernel(CodeGenFunction &CGF,
                             llvm::Function *BlockInvokeFunc,
                             llvm::Value *BlockLiteral) const override;
+  bool shouldEmitStaticExternCAliases() const override;
 };
 }
 
 void AMDGPUTargetCodeGenInfo::setTargetAttributes(
     const Decl *D, llvm::GlobalValue *GV, CodeGen::CodeGenModule &M) const {
   if (GV->isDeclaration())
     return;
   const FunctionDecl *FD = dyn_cast_or_null<FunctionDecl>(D);

     Name = "";
     break;
   case SyncScope::OpenCLSubGroup:
     Name = "subgroup";
   }
   return C.getOrInsertSyncScopeID(Name);
 }
 
+bool AMDGPUTargetCodeGenInfo::shouldEmitStaticExternCAliases() const {
+  return false;
+}
+
 //===----------------------------------------------------------------------===//
 // SPARC v8 ABI Implementation.
 // Based on the SPARC Compliance Definition version 2.4.1.
 //
 // Ensures that complex values are passed in registers.
 //
 namespace {
 class SparcV8ABIInfo : public DefaultABIInfo {

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; }
 }
 }