Permalink
Browse files

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!


git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@190684 91177308-0d34-0410-b5e6-96231b3b80d8
  • Loading branch information...
1 parent 0909859 commit 1eef85246b411b55c493098266746d0d83c241ea David Tweed committed Sep 13, 2013
@@ -393,6 +393,10 @@ class ASTContext : public RefCountedBase<ASTContext> {
/// \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 @@ class ASTContext : public RefCountedBase<ASTContext> {
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;
@@ -142,6 +142,7 @@ LANGOPT(HexagonQdsp6Compat , 1, 0, "hexagon-qdsp6 backward compatibility")
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")
@@ -66,6 +66,8 @@ class LangOptions : public RefCountedBase<LangOptions>, public LangOptionsBase {
SOB_Trapping // -ftrapv
};
+ enum AddrSpaceMapMangling { ASMM_Target, ASMM_On, ASMM_Off };
+
public:
clang::ObjCRuntime ObjCRuntime;
@@ -202,6 +202,10 @@ class TargetInfo : public RefCountedBase<TargetInfo> {
/// 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 @@ class TargetInfo : public RefCountedBase<TargetInfo> {
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
@@ -460,6 +460,8 @@ def fno_bitfield_type_align : Flag<["-"], "fno-bitfield-type-align">,
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">,
View
@@ -695,6 +695,19 @@ static const LangAS::Map *getAddressSpaceMap(const TargetInfo &T,
}
}
+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 TargetInfo &Target) {
ABI.reset(createCXXABI(Target));
AddrSpaceMap = getAddressSpaceMap(Target, LangOpts);
+ AddrSpaceMapMangling = isAddrSpaceMapManglingEnabled(Target, LangOpts);
// C99 6.2.5p19.
InitBuiltinType(VoidTy, BuiltinType::Void);
View
@@ -1755,15 +1755,33 @@ void CXXNameMangler::mangleQualifiers(Qualifiers Quals) {
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;
}
View
@@ -88,6 +88,7 @@ TargetInfo::TargetInfo(const llvm::Triple &T) : TargetOpts(), Triple(T) {
// Default to an empty address space map.
AddrSpaceMap = &DefaultAddrSpaceMap;
+ UseAddrSpaceMapMangling = false;
// Default to an unknown platform name.
PlatformName = "unknown";
View
@@ -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 @@ class R600TargetInfo : public TargetInfo {
: 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;
@@ -1329,6 +1329,28 @@ static void ParseLangArgs(LangOptions &Opts, ArgList &Args, InputKind IK,
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);
@@ -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
@@ -1,12 +1,12 @@
-// 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)
__local int i;
++i;
}
-// CHECK-LABEL: define void @_Z3barPU3AS2i
+// CHECK-LABEL: define void @_Z3barPU7CLlocali
__kernel void __attribute__((__overloadable__)) bar(local int *x) {
*x = 5;
}

0 comments on commit 1eef852

Please sign in to comment.