r190684 - Certain multi-platform languages, such as OpenCL, have the concept of

David Tweed david.tweed at arm.com
Fri Sep 13 05:04:22 PDT 2013


Author: davidtweed
Date: Fri Sep 13 07:04:22 2013
New Revision: 190684

URL: http://llvm.org/viewvc/llvm-project?rev=190684&view=rev
Log:
Certain multi-platform languages, such as OpenCL, have the concept of
address spaces which is both (1) a "semantic" concept and
(2) possibly a hardware level restriction. It is desirable to
be able to discard/merge the LLVM-level address spaces on arguments for which
there is no difference to the current backend while keeping
track of the semantic address spaces in a funciton prototype. To do this
enable addition of the address space into the name-mangling process. Add
some tests to document this behaviour against inadvertent changes.

Patch by Michele Scandale!

Added:
    cfe/trunk/test/CodeGenOpenCL/address-spaces-mangling.cl
Modified:
    cfe/trunk/include/clang/AST/ASTContext.h
    cfe/trunk/include/clang/Basic/LangOptions.def
    cfe/trunk/include/clang/Basic/LangOptions.h
    cfe/trunk/include/clang/Basic/TargetInfo.h
    cfe/trunk/include/clang/Driver/CC1Options.td
    cfe/trunk/lib/AST/ASTContext.cpp
    cfe/trunk/lib/AST/ItaniumMangle.cpp
    cfe/trunk/lib/Basic/TargetInfo.cpp
    cfe/trunk/lib/Basic/Targets.cpp
    cfe/trunk/lib/Frontend/CompilerInvocation.cpp
    cfe/trunk/test/CodeGenOpenCL/local.cl

Modified: cfe/trunk/include/clang/AST/ASTContext.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/AST/ASTContext.h?rev=190684&r1=190683&r2=190684&view=diff
==============================================================================
--- cfe/trunk/include/clang/AST/ASTContext.h (original)
+++ cfe/trunk/include/clang/AST/ASTContext.h Fri Sep 13 07:04:22 2013
@@ -393,6 +393,10 @@ private:
   /// \brief The logical -> physical address space map.
   const LangAS::Map *AddrSpaceMap;
 
+  /// \brief Address space map mangling must be used with language specific 
+  /// address spaces (e.g. OpenCL/CUDA)
+  bool AddrSpaceMapMangling;
+
   friend class ASTDeclReader;
   friend class ASTReader;
   friend class ASTWriter;
@@ -1920,6 +1924,12 @@ public:
       return (*AddrSpaceMap)[AS - LangAS::Offset];
   }
 
+  bool addressSpaceMapManglingFor(unsigned AS) const {
+    return AddrSpaceMapMangling || 
+           AS < LangAS::Offset || 
+           AS >= LangAS::Offset + LangAS::Count;
+  }
+
 private:
   // Helper for integer ordering
   unsigned getIntegerRank(const Type *T) const;

Modified: cfe/trunk/include/clang/Basic/LangOptions.def
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/LangOptions.def?rev=190684&r1=190683&r2=190684&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/LangOptions.def (original)
+++ cfe/trunk/include/clang/Basic/LangOptions.def Fri Sep 13 07:04:22 2013
@@ -142,6 +142,7 @@ LANGOPT(HexagonQdsp6Compat , 1, 0, "hexa
 LANGOPT(ObjCAutoRefCount , 1, 0, "Objective-C automated reference counting")
 LANGOPT(ObjCARCWeak         , 1, 0, "__weak support in the ARC runtime")
 LANGOPT(FakeAddressSpaceMap , 1, 0, "OpenCL fake address space map")
+ENUM_LANGOPT(AddressSpaceMapMangling , AddrSpaceMapMangling, 2, ASMM_Target, "OpenCL address space map mangling mode")
 
 LANGOPT(MRTD , 1, 0, "-mrtd calling convention")
 BENIGN_LANGOPT(DelayedTemplateParsing , 1, 0, "delayed template parsing")

Modified: cfe/trunk/include/clang/Basic/LangOptions.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/LangOptions.h?rev=190684&r1=190683&r2=190684&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/LangOptions.h (original)
+++ cfe/trunk/include/clang/Basic/LangOptions.h Fri Sep 13 07:04:22 2013
@@ -66,6 +66,8 @@ public:
     SOB_Trapping    // -ftrapv
   };
 
+  enum AddrSpaceMapMangling { ASMM_Target, ASMM_On, ASMM_Off };
+
 public:
   clang::ObjCRuntime ObjCRuntime;
 

Modified: cfe/trunk/include/clang/Basic/TargetInfo.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/TargetInfo.h?rev=190684&r1=190683&r2=190684&view=diff
==============================================================================
--- cfe/trunk/include/clang/Basic/TargetInfo.h (original)
+++ cfe/trunk/include/clang/Basic/TargetInfo.h Fri Sep 13 07:04:22 2013
@@ -202,6 +202,10 @@ protected:
   /// zero length bitfield, regardless of the zero length bitfield type.
   unsigned ZeroLengthBitfieldBoundary;
 
+  /// \brief Specify if mangling based on address space map should be used or
+  /// not for language specific address spaces
+  bool UseAddrSpaceMapMangling;
+
 public:
   IntType getSizeType() const { return SizeType; }
   IntType getIntMaxType() const { return IntMaxType; }
@@ -431,6 +435,12 @@ public:
     return ComplexLongDoubleUsesFP2Ret;
   }
 
+  /// \brief Specify if mangling based on address space map should be used or
+  /// not for language specific address spaces
+  bool useAddressSpaceMapMangling() const {
+    return UseAddrSpaceMapMangling;
+  }
+
   ///===---- Other target property query methods --------------------------===//
 
   /// \brief Appends the target-specific \#define values for this

Modified: cfe/trunk/include/clang/Driver/CC1Options.td
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Driver/CC1Options.td?rev=190684&r1=190683&r2=190684&view=diff
==============================================================================
--- cfe/trunk/include/clang/Driver/CC1Options.td (original)
+++ cfe/trunk/include/clang/Driver/CC1Options.td Fri Sep 13 07:04:22 2013
@@ -460,6 +460,8 @@ def fno_bitfield_type_align : Flag<["-"]
   HelpText<"Ignore bit-field types when aligning structures">;
 def ffake_address_space_map : Flag<["-"], "ffake-address-space-map">,
   HelpText<"Use a fake address space map; OpenCL testing purposes only">;
+def faddress_space_map_mangling_EQ : Joined<["-"], "faddress-space-map-mangling=">, MetaVarName<"<yes|no|target>">,
+  HelpText<"Set the mode for address space map based mangling; OpenCL testing purposes only">;
 def funknown_anytype : Flag<["-"], "funknown-anytype">,
   HelpText<"Enable parser support for the __unknown_anytype type; for testing purposes only">;
 def fdebugger_support : Flag<["-"], "fdebugger-support">,

Modified: cfe/trunk/lib/AST/ASTContext.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/ASTContext.cpp?rev=190684&r1=190683&r2=190684&view=diff
==============================================================================
--- cfe/trunk/lib/AST/ASTContext.cpp (original)
+++ cfe/trunk/lib/AST/ASTContext.cpp Fri Sep 13 07:04:22 2013
@@ -695,6 +695,19 @@ static const LangAS::Map *getAddressSpac
   }
 }
 
+static bool isAddrSpaceMapManglingEnabled(const TargetInfo &TI,
+                                          const LangOptions &LangOpts) {
+  switch (LangOpts.getAddressSpaceMapMangling()) {
+  default: return false;
+  case LangOptions::ASMM_Target:
+    return TI.useAddressSpaceMapMangling();
+  case LangOptions::ASMM_On:
+    return true;
+  case LangOptions::ASMM_Off:
+    return false;
+  }
+}
+
 ASTContext::ASTContext(LangOptions& LOpts, SourceManager &SM,
                        const TargetInfo *t,
                        IdentifierTable &idents, SelectorTable &sels,
@@ -900,6 +913,7 @@ void ASTContext::InitBuiltinTypes(const
   
   ABI.reset(createCXXABI(Target));
   AddrSpaceMap = getAddressSpaceMap(Target, LangOpts);
+  AddrSpaceMapMangling = isAddrSpaceMapManglingEnabled(Target, LangOpts);
   
   // C99 6.2.5p19.
   InitBuiltinType(VoidTy,              BuiltinType::Void);

Modified: cfe/trunk/lib/AST/ItaniumMangle.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/ItaniumMangle.cpp?rev=190684&r1=190683&r2=190684&view=diff
==============================================================================
--- cfe/trunk/lib/AST/ItaniumMangle.cpp (original)
+++ cfe/trunk/lib/AST/ItaniumMangle.cpp Fri Sep 13 07:04:22 2013
@@ -1755,15 +1755,33 @@ void CXXNameMangler::mangleQualifiers(Qu
     Out << 'K';
 
   if (Quals.hasAddressSpace()) {
-    // Extension:
+    // Address space extension:
     //
-    //   <type> ::= U <address-space-number>
-    // 
-    // where <address-space-number> is a source name consisting of 'AS' 
-    // followed by the address space <number>.
+    //   <type> ::= U <target-addrspace>
+    //   <type> ::= U <OpenCL-addrspace>
+    //   <type> ::= U <CUDA-addrspace>
+
     SmallString<64> ASString;
-    ASString = "AS" + llvm::utostr_32(
-        Context.getASTContext().getTargetAddressSpace(Quals.getAddressSpace()));
+    unsigned AS = Quals.getAddressSpace();
+    bool IsLangAS = (LangAS::Offset <= AS) && (AS < LangAS::Last);
+
+    if (Context.getASTContext().addressSpaceMapManglingFor(AS)) {
+      //  <target-addrspace> ::= "AS" <address-space-number>
+      unsigned TargetAS = Context.getASTContext().getTargetAddressSpace(AS);
+      ASString = "AS" + llvm::utostr_32(TargetAS);
+    } else {
+      switch (AS) {
+      default: llvm_unreachable("Not a language specific address space");
+      //  <OpenCL-addrspace> ::= "CL" [ "global" | "local" | "constant" ]
+      case LangAS::opencl_global:   ASString = "CLglobal";   break;
+      case LangAS::opencl_local:    ASString = "CLlocal";    break;
+      case LangAS::opencl_constant: ASString = "CLconstant"; break;
+      //  <CUDA-addrspace> ::= "CU" [ "device" | "constant" | "shared" ]
+      case LangAS::cuda_device:     ASString = "CUdevice";   break;
+      case LangAS::cuda_constant:   ASString = "CUconstant"; break;
+      case LangAS::cuda_shared:     ASString = "CUshared";   break;
+      }
+    }
     Out << 'U' << ASString.size() << ASString;
   }
   

Modified: cfe/trunk/lib/Basic/TargetInfo.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/TargetInfo.cpp?rev=190684&r1=190683&r2=190684&view=diff
==============================================================================
--- cfe/trunk/lib/Basic/TargetInfo.cpp (original)
+++ cfe/trunk/lib/Basic/TargetInfo.cpp Fri Sep 13 07:04:22 2013
@@ -88,6 +88,7 @@ TargetInfo::TargetInfo(const llvm::Tripl
 
   // Default to an empty address space map.
   AddrSpaceMap = &DefaultAddrSpaceMap;
+  UseAddrSpaceMapMangling = false;
 
   // Default to an unknown platform name.
   PlatformName = "unknown";

Modified: cfe/trunk/lib/Basic/Targets.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/Targets.cpp?rev=190684&r1=190683&r2=190684&view=diff
==============================================================================
--- cfe/trunk/lib/Basic/Targets.cpp (original)
+++ cfe/trunk/lib/Basic/Targets.cpp Fri Sep 13 07:04:22 2013
@@ -1266,6 +1266,7 @@ namespace {
       TLSSupported = false;
       LongWidth = LongAlign = 64;
       AddrSpaceMap = &NVPTXAddrSpaceMap;
+      UseAddrSpaceMapMangling = true;
       // Define available target features
       // These must be defined in sorted order!
       NoAsmVariants = true;
@@ -1424,6 +1425,7 @@ public:
       : TargetInfo(Triple), GPU(GK_R600) {
     DescriptionString = DescriptionStringR600;
     AddrSpaceMap = &R600AddrSpaceMap;
+    UseAddrSpaceMapMangling = true;
   }
 
   virtual const char * getClobbers() const {
@@ -4577,6 +4579,7 @@ namespace {
                           "f32:32:32-f64:32:32-v64:32:32-"
                           "v128:32:32-a0:0:32-n32";
       AddrSpaceMap = &TCEOpenCLAddrSpaceMap;
+      UseAddrSpaceMapMangling = true;
     }
 
     virtual void getTargetDefines(const LangOptions &Opts,
@@ -5139,6 +5142,7 @@ namespace {
       TLSSupported = false;
       LongWidth = LongAlign = 64;
       AddrSpaceMap = &SPIRAddrSpaceMap;
+      UseAddrSpaceMapMangling = true;
       // Define available target features
       // These must be defined in sorted order!
       NoAsmVariants = true;

Modified: cfe/trunk/lib/Frontend/CompilerInvocation.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Frontend/CompilerInvocation.cpp?rev=190684&r1=190683&r2=190684&view=diff
==============================================================================
--- cfe/trunk/lib/Frontend/CompilerInvocation.cpp (original)
+++ cfe/trunk/lib/Frontend/CompilerInvocation.cpp Fri Sep 13 07:04:22 2013
@@ -1329,6 +1329,28 @@ static void ParseLangArgs(LangOptions &O
   Opts.ApplePragmaPack = Args.hasArg(OPT_fapple_pragma_pack);
   Opts.CurrentModule = Args.getLastArgValue(OPT_fmodule_name);
 
+  if (Arg *A = Args.getLastArg(OPT_faddress_space_map_mangling_EQ)) {
+    switch (llvm::StringSwitch<unsigned>(A->getValue())
+      .Case("target", LangOptions::ASMM_Target)
+      .Case("no", LangOptions::ASMM_Off)
+      .Case("yes", LangOptions::ASMM_On)
+      .Default(255)) {
+    default:
+      Diags.Report(diag::err_drv_invalid_value) 
+        << "-faddress-space-map-mangling=" << A->getValue();
+      break;
+    case LangOptions::ASMM_Target:
+      Opts.setAddressSpaceMapMangling(LangOptions::ASMM_Target);
+      break;
+    case LangOptions::ASMM_On:
+      Opts.setAddressSpaceMapMangling(LangOptions::ASMM_On);
+      break;
+    case LangOptions::ASMM_Off:
+      Opts.setAddressSpaceMapMangling(LangOptions::ASMM_Off);
+      break;
+    }
+  }
+
   // Check if -fopenmp is specified.
   Opts.OpenMP = Args.hasArg(OPT_fopenmp);
 

Added: cfe/trunk/test/CodeGenOpenCL/address-spaces-mangling.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/address-spaces-mangling.cl?rev=190684&view=auto
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/address-spaces-mangling.cl (added)
+++ cfe/trunk/test/CodeGenOpenCL/address-spaces-mangling.cl Fri Sep 13 07:04:22 2013
@@ -0,0 +1,30 @@
+// RUN: %clang_cc1 %s -ffake-address-space-map -faddress-space-map-mangling=yes -emit-llvm -o - | FileCheck -check-prefix=ASMANG %s
+// RUN: %clang_cc1 %s -ffake-address-space-map -faddress-space-map-mangling=no -emit-llvm -o - | FileCheck -check-prefix=NOASMANG %s
+
+// We can't name this f as private is equivalent to default
+// no specifier given address space so we get multiple definition
+// warnings, but we do want it for comparison purposes.
+__attribute__((overloadable))
+void ff(int *arg) { }
+// ASMANG: @_Z2ffPi
+// NOASMANG: @_Z2ffPi
+
+__attribute__((overloadable))
+void f(private int *arg) { }
+// ASMANG: @_Z1fPi
+// NOASMANG: @_Z1fPi
+
+__attribute__((overloadable))
+void f(global int *arg) { }
+// ASMANG: @_Z1fPU3AS1i
+// NOASMANG: @_Z1fPU8CLglobali
+
+__attribute__((overloadable))
+void f(local int *arg) { }
+// ASMANG: @_Z1fPU3AS2i
+// NOASMANG: @_Z1fPU7CLlocali
+
+__attribute__((overloadable))
+void f(constant int *arg) { }
+// ASMANG: @_Z1fPU3AS3i
+// NOASMANG: @_Z1fPU10CLconstanti

Modified: cfe/trunk/test/CodeGenOpenCL/local.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/local.cl?rev=190684&r1=190683&r2=190684&view=diff
==============================================================================
--- cfe/trunk/test/CodeGenOpenCL/local.cl (original)
+++ cfe/trunk/test/CodeGenOpenCL/local.cl Fri Sep 13 07:04:22 2013
@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 %s -ffake-address-space-map -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 %s -ffake-address-space-map -faddress-space-map-mangling=no -emit-llvm -o - | FileCheck %s
 
 __kernel void foo(void) {
   // CHECK: @foo.i = internal addrspace(2)
@@ -6,7 +6,7 @@ __kernel void foo(void) {
   ++i;
 }
 
-// CHECK-LABEL: define void @_Z3barPU3AS2i
+// CHECK-LABEL: define void @_Z3barPU7CLlocali
 __kernel void __attribute__((__overloadable__)) bar(local int *x) {
   *x = 5;
 }





More information about the cfe-commits mailing list