OpenCL address space and mangling
Michele Scandale
michele.scandale at gmail.com
Thu Sep 12 04:35:06 PDT 2013
On 09/11/2013 11:55 AM, David Tweed wrote:
> while 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;
}
More information about the cfe-commits
mailing list