Post by David Tweedwhile I think most people agree with the direction things are going there
look to still
be some fiddly details. As one instance, when I run this on a standard OSS
LLVM build I get
a new test failure in test/CodeGenOpenCL/local.cl. Since behaviour is being
made more sophisticated, it seems
it would be good to have add some tests that verify the new behaviour so we
can detect any
modifications that change it. But the patch looks to be progressing.
Hi David,
I've fixed the test and added another test specific for mangling checking. To
simplify testing I've added a command line option (similar to
-ffake-address-space-map).
In attachment the new version of the patch.
Thanks in advance.
Best Regards,
Michele
-------------- next part --------------
diff --git a/include/clang/AST/ASTContext.h b/include/clang/AST/ASTContext.h
index 377282f..fd351ab 100644
--- a/include/clang/AST/ASTContext.h
+++ b/include/clang/AST/ASTContext.h
@@ -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;
diff --git a/include/clang/Basic/LangOptions.def b/include/clang/Basic/LangOptions.def
index a2e94ff..55db34c 100644
--- a/include/clang/Basic/LangOptions.def
+++ b/include/clang/Basic/LangOptions.def
@@ -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")
diff --git a/include/clang/Basic/LangOptions.h b/include/clang/Basic/LangOptions.h
index 21ca7eb..4532054 100644
--- a/include/clang/Basic/LangOptions.h
+++ b/include/clang/Basic/LangOptions.h
@@ -66,6 +66,8 @@ public:
SOB_Trapping // -ftrapv
};
+ enum AddrSpaceMapMangling { ASMM_Target, ASMM_On, ASMM_Off };
+
public:
clang::ObjCRuntime ObjCRuntime;
diff --git a/include/clang/Basic/TargetInfo.h b/include/clang/Basic/TargetInfo.h
index ee3a28d..bda6af3 100644
--- a/include/clang/Basic/TargetInfo.h
+++ b/include/clang/Basic/TargetInfo.h
@@ -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
diff --git a/include/clang/Driver/CC1Options.td b/include/clang/Driver/CC1Options.td
index b74f445..d91001f 100644
--- a/include/clang/Driver/CC1Options.td
+++ b/include/clang/Driver/CC1Options.td
@@ -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">,
diff --git a/lib/AST/ASTContext.cpp b/lib/AST/ASTContext.cpp
index 85ac734..6e77f4e 100644
--- a/lib/AST/ASTContext.cpp
+++ b/lib/AST/ASTContext.cpp
@@ -694,6 +694,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,
@@ -893,6 +906,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);
diff --git a/lib/AST/ItaniumMangle.cpp b/lib/AST/ItaniumMangle.cpp
index c3121c0..e135227 100644
--- a/lib/AST/ItaniumMangle.cpp
+++ b/lib/AST/ItaniumMangle.cpp
@@ -1753,15 +1753,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;
}
diff --git a/lib/Basic/TargetInfo.cpp b/lib/Basic/TargetInfo.cpp
index 3feaf9e..e993055 100644
--- a/lib/Basic/TargetInfo.cpp
+++ b/lib/Basic/TargetInfo.cpp
@@ -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";
diff --git a/lib/Basic/Targets.cpp b/lib/Basic/Targets.cpp
index aa0993d..a6a613d 100644
--- a/lib/Basic/Targets.cpp
+++ b/lib/Basic/Targets.cpp
@@ -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 {
@@ -4575,6 +4577,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,
@@ -5137,6 +5140,7 @@ namespace {
TLSSupported = false;
LongWidth = LongAlign = 64;
AddrSpaceMap = &SPIRAddrSpaceMap;
+ UseAddrSpaceMapMangling = true;
// Define available target features
// These must be defined in sorted order!
NoAsmVariants = true;
diff --git a/lib/Frontend/CompilerInvocation.cpp b/lib/Frontend/CompilerInvocation.cpp
index ce4fdc3..2947041 100644
--- a/lib/Frontend/CompilerInvocation.cpp
+++ b/lib/Frontend/CompilerInvocation.cpp
@@ -1326,6 +1326,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);
diff --git a/test/CodeGenOpenCL/address-spaces-mangling.cl b/test/CodeGenOpenCL/address-spaces-mangling.cl
index e69de29..7ce74d3 100644
--- a/test/CodeGenOpenCL/address-spaces-mangling.cl
+++ b/test/CodeGenOpenCL/address-spaces-mangling.cl
@@ -0,0 +1,22 @@
+// 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
+
+__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
diff --git a/test/CodeGenOpenCL/local.cl b/test/CodeGenOpenCL/local.cl
index 852fa43..b5c67d9 100644
--- a/test/CodeGenOpenCL/local.cl
+++ b/test/CodeGenOpenCL/local.cl
@@ -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;
}