Add support for language-specific address spaces. On top of that,
add support for the OpenCL __private, __local, __constant and
__global address spaces, as well as the __read_only, _read_write and
__write_only image access specifiers. Patch originally by ARM;
language-specific address space support by myself.
git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@127915 91177308-0d34-0410-b5e6-96231b3b80d8
diff --git a/include/clang/AST/ASTContext.h b/include/clang/AST/ASTContext.h
index ad69f0f..68a12bb 100644
--- a/include/clang/AST/ASTContext.h
+++ b/include/clang/AST/ASTContext.h
@@ -14,6 +14,7 @@
#ifndef LLVM_CLANG_AST_ASTCONTEXT_H
#define LLVM_CLANG_AST_ASTCONTEXT_H
+#include "clang/Basic/AddressSpaces.h"
#include "clang/Basic/IdentifierTable.h"
#include "clang/Basic/LangOptions.h"
#include "clang/Basic/OperatorKinds.h"
@@ -311,6 +312,9 @@
llvm::OwningPtr<CXXABI> ABI;
CXXABI *createCXXABI(const TargetInfo &T);
+ /// \brief The logical -> physical address space map.
+ const LangAS::Map &AddrSpaceMap;
+
friend class ASTDeclReader;
public:
@@ -1295,6 +1299,21 @@
QualType getFloatingTypeOfSizeWithinDomain(QualType typeSize,
QualType typeDomain) const;
+ unsigned getTargetAddressSpace(QualType T) const {
+ return getTargetAddressSpace(T.getQualifiers());
+ }
+
+ unsigned getTargetAddressSpace(Qualifiers Q) const {
+ return getTargetAddressSpace(Q.getAddressSpace());
+ }
+
+ unsigned getTargetAddressSpace(unsigned AS) const {
+ if (AS < LangAS::Offset || AS >= LangAS::Offset + LangAS::Count)
+ return AS;
+ else
+ return AddrSpaceMap[AS - LangAS::Offset];
+ }
+
private:
// Helper for integer ordering
unsigned getIntegerRank(const Type *T) const;
diff --git a/include/clang/Basic/AddressSpaces.h b/include/clang/Basic/AddressSpaces.h
new file mode 100644
index 0000000..d44a9c3
--- /dev/null
+++ b/include/clang/Basic/AddressSpaces.h
@@ -0,0 +1,44 @@
+//===--- AddressSpaces.h - Language-specific address spaces -----*- C++ -*-===//
+//
+// The LLVM Compiler Infrastructure
+//
+// This file is distributed under the University of Illinois Open Source
+// License. See LICENSE.TXT for details.
+//
+//===----------------------------------------------------------------------===//
+//
+// This file provides definitions for the various language-specific address
+// spaces.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_CLANG_BASIC_ADDRESSSPACES_H
+#define LLVM_CLANG_BASIC_ADDRESSSPACES_H
+
+namespace clang {
+
+namespace LangAS {
+
+/// This enum defines the set of possible language-specific address spaces.
+/// It uses a high starting offset so as not to conflict with any address
+/// space used by a target.
+enum ID {
+ Offset = 0xFFFF00,
+
+ opencl_global = Offset,
+ opencl_local,
+ opencl_constant,
+
+ Last,
+ Count = Last-Offset
+};
+
+/// The type of a lookup table which maps from language-specific address spaces
+/// to target-specific ones.
+typedef unsigned Map[Count];
+
+}
+
+}
+
+#endif
diff --git a/include/clang/Basic/LangOptions.h b/include/clang/Basic/LangOptions.h
index ebdcdde..433f42e 100644
--- a/include/clang/Basic/LangOptions.h
+++ b/include/clang/Basic/LangOptions.h
@@ -124,6 +124,8 @@
unsigned DefaultFPContract : 1; // Default setting for FP_CONTRACT
// FIXME: This is just a temporary option, for testing purposes.
unsigned NoBitFieldTypeAlign : 1;
+ unsigned FakeAddressSpaceMap : 1; // Use a fake address space map, for
+ // testing languages such as OpenCL.
unsigned MRTD : 1; // -mrtd calling convention
diff --git a/include/clang/Basic/OpenCL.h b/include/clang/Basic/OpenCL.h
new file mode 100644
index 0000000..6f9785f
--- /dev/null
+++ b/include/clang/Basic/OpenCL.h
@@ -0,0 +1,28 @@
+//===--- OpenCL.h - OpenCL enums --------------------------------*- C++ -*-===//
+//
+// The LLVM Compiler Infrastructure
+//
+// This file is distributed under the University of Illinois Open Source
+// License. See LICENSE.TXT for details.
+//
+//===----------------------------------------------------------------------===//
+//
+// This file defines some OpenCL-specific enums.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_CLANG_BASIC_OPENCL_H
+#define LLVM_CLANG_BASIC_OPENCL_H
+
+namespace clang {
+
+/// Names for the OpenCL image access qualifiers (OpenCL 1.1 6.6).
+enum OpenCLImageAccess {
+ CLIA_read_only = 1,
+ CLIA_write_only = 2,
+ CLIA_read_write = 3
+};
+
+}
+
+#endif
diff --git a/include/clang/Basic/TargetInfo.h b/include/clang/Basic/TargetInfo.h
index b9087f2..fd9e373 100644
--- a/include/clang/Basic/TargetInfo.h
+++ b/include/clang/Basic/TargetInfo.h
@@ -19,6 +19,7 @@
#include "llvm/ADT/StringSwitch.h"
#include "llvm/ADT/Triple.h"
#include "llvm/Support/DataTypes.h"
+#include "clang/Basic/AddressSpaces.h"
#include <cassert>
#include <vector>
#include <string>
@@ -78,6 +79,7 @@
const llvm::fltSemantics *FloatFormat, *DoubleFormat, *LongDoubleFormat;
unsigned char RegParmMax, SSERegParmMax;
TargetCXXABI CXXABI;
+ const LangAS::Map *AddrSpaceMap;
unsigned HasAlignMac68kSupport : 1;
unsigned RealTypeUsesObjCFPRet : 3;
@@ -530,6 +532,11 @@
virtual const char *getStaticInitSectionSpecifier() const {
return 0;
}
+
+ const LangAS::Map &getAddressSpaceMap() const {
+ return *AddrSpaceMap;
+ }
+
protected:
virtual uint64_t getPointerWidthV(unsigned AddrSpace) const {
return PointerWidth;
diff --git a/include/clang/Basic/TokenKinds.def b/include/clang/Basic/TokenKinds.def
index c7d9524..6fa4bf9 100644
--- a/include/clang/Basic/TokenKinds.def
+++ b/include/clang/Basic/TokenKinds.def
@@ -252,7 +252,7 @@
KEYWORD(namespace , KEYCXX)
KEYWORD(new , KEYCXX)
KEYWORD(operator , KEYCXX)
-KEYWORD(private , KEYCXX)
+KEYWORD(private , KEYCXX|KEYOPENCL)
KEYWORD(protected , KEYCXX)
KEYWORD(public , KEYCXX)
KEYWORD(reinterpret_cast , KEYCXX)
@@ -350,6 +350,19 @@
KEYWORD(__kernel , KEYOPENCL)
ALIAS("kernel", __kernel , KEYOPENCL)
KEYWORD(vec_step , KEYOPENCL)
+KEYWORD(__private , KEYOPENCL)
+KEYWORD(__global , KEYOPENCL)
+KEYWORD(__local , KEYOPENCL)
+KEYWORD(__constant , KEYOPENCL)
+ALIAS("global", __global , KEYOPENCL)
+ALIAS("local", __local , KEYOPENCL)
+ALIAS("constant", __constant , KEYOPENCL)
+KEYWORD(__read_only , KEYOPENCL)
+KEYWORD(__write_only , KEYOPENCL)
+KEYWORD(__read_write , KEYOPENCL)
+ALIAS("read_only", __read_only , KEYOPENCL)
+ALIAS("write_only", __write_only , KEYOPENCL)
+ALIAS("read_write", __read_write , KEYOPENCL)
// Borland Extensions.
KEYWORD(__pascal , KEYALL)
diff --git a/include/clang/Driver/CC1Options.td b/include/clang/Driver/CC1Options.td
index bb8759e..4188234 100644
--- a/include/clang/Driver/CC1Options.td
+++ b/include/clang/Driver/CC1Options.td
@@ -513,6 +513,8 @@
HelpText<"Ignore bit-field types when aligning structures">;
def traditional_cpp : Flag<"-traditional-cpp">,
HelpText<"Enable some traditional CPP emulation">;
+def ffake_address_space_map : Flag<"-ffake-address-space-map">,
+ HelpText<"Use a fake address space map; OpenCL testing purposes only">;
//===----------------------------------------------------------------------===//
// Header Search Options
diff --git a/include/clang/Parse/Parser.h b/include/clang/Parse/Parser.h
index bd39cdb..b9af05e 100644
--- a/include/clang/Parse/Parser.h
+++ b/include/clang/Parse/Parser.h
@@ -1546,6 +1546,7 @@
void ParseMicrosoftTypeAttributes(ParsedAttributes &attrs);
void ParseBorlandTypeAttributes(ParsedAttributes &attrs);
void ParseOpenCLAttributes(ParsedAttributes &attrs);
+ void ParseOpenCLQualifiers(DeclSpec &DS);
void ParseTypeofSpecifier(DeclSpec &DS);
void ParseDecltypeSpecifier(DeclSpec &DS);
diff --git a/include/clang/Sema/AttributeList.h b/include/clang/Sema/AttributeList.h
index 10b859e..382e504 100644
--- a/include/clang/Sema/AttributeList.h
+++ b/include/clang/Sema/AttributeList.h
@@ -18,6 +18,7 @@
#include "llvm/Support/Allocator.h"
#include "clang/Sema/Ownership.h"
#include "clang/Basic/SourceLocation.h"
+#include "clang/AST/Expr.h"
#include <cassert>
namespace clang {
@@ -76,6 +77,13 @@
declspec, cxx0x);
return Mem;
}
+
+ AttributeList* CreateIntegerAttribute(ASTContext &C, IdentifierInfo *Name,
+ SourceLocation TokLoc, int Arg) {
+ Expr* IArg = IntegerLiteral::Create(C, llvm::APInt(32, (uint64_t)Arg),
+ C.IntTy, TokLoc);
+ return Create( Name, TokLoc, 0, TokLoc, 0, TokLoc, &IArg, 1, 0);
+ }
};
enum Kind { // Please keep this list alphabetized.
@@ -135,6 +143,7 @@
AT_ns_consumed, // Clang-specific.
AT_ns_consumes_self, // Clang-specific.
AT_objc_gc,
+ AT_opencl_image_access, // OpenCL-specific.
AT_opencl_kernel_function, // OpenCL-specific.
AT_overloadable, // Clang-specific.
AT_ownership_holds, // Clang-specific.
diff --git a/lib/AST/ASTContext.cpp b/lib/AST/ASTContext.cpp
index d20c258..fd5fd0e 100644
--- a/lib/AST/ASTContext.cpp
+++ b/lib/AST/ASTContext.cpp
@@ -190,6 +190,22 @@
return 0;
}
+static const LangAS::Map &getAddressSpaceMap(const TargetInfo &T,
+ const LangOptions &LOpts) {
+ if (LOpts.FakeAddressSpaceMap) {
+ // The fake address space map must have a distinct entry for each
+ // language-specific address space.
+ static const unsigned FakeAddrSpaceMap[] = {
+ 1, // opencl_global
+ 2, // opencl_local
+ 3 // opencl_constant
+ };
+ return FakeAddrSpaceMap;
+ } else {
+ return T.getAddressSpaceMap();
+ }
+}
+
ASTContext::ASTContext(const LangOptions& LOpts, SourceManager &SM,
const TargetInfo &t,
IdentifierTable &idents, SelectorTable &sels,
@@ -204,7 +220,8 @@
sigjmp_bufDecl(0), BlockDescriptorType(0), BlockDescriptorExtendedType(0),
cudaConfigureCallDecl(0),
NullTypeSourceInfo(QualType()),
- SourceMgr(SM), LangOpts(LOpts), ABI(createCXXABI(t)), Target(t),
+ SourceMgr(SM), LangOpts(LOpts), ABI(createCXXABI(t)),
+ AddrSpaceMap(getAddressSpaceMap(t, LOpts)), Target(t),
Idents(idents), Selectors(sels),
BuiltinInfo(builtins),
DeclarationNames(*this),
@@ -806,7 +823,8 @@
Align = Target.getPointerAlign(0);
break;
case Type::BlockPointer: {
- unsigned AS = cast<BlockPointerType>(T)->getPointeeType().getAddressSpace();
+ unsigned AS = getTargetAddressSpace(
+ cast<BlockPointerType>(T)->getPointeeType());
Width = Target.getPointerWidth(AS);
Align = Target.getPointerAlign(AS);
break;
@@ -815,13 +833,14 @@
case Type::RValueReference: {
// alignof and sizeof should never enter this code path here, so we go
// the pointer route.
- unsigned AS = cast<ReferenceType>(T)->getPointeeType().getAddressSpace();
+ unsigned AS = getTargetAddressSpace(
+ cast<ReferenceType>(T)->getPointeeType());
Width = Target.getPointerWidth(AS);
Align = Target.getPointerAlign(AS);
break;
}
case Type::Pointer: {
- unsigned AS = cast<PointerType>(T)->getPointeeType().getAddressSpace();
+ unsigned AS = getTargetAddressSpace(cast<PointerType>(T)->getPointeeType());
Width = Target.getPointerWidth(AS);
Align = Target.getPointerAlign(AS);
break;
@@ -1468,7 +1487,7 @@
// the target.
llvm::APInt ArySize(ArySizeIn);
ArySize =
- ArySize.zextOrTrunc(Target.getPointerWidth(EltTy.getAddressSpace()));
+ ArySize.zextOrTrunc(Target.getPointerWidth(getTargetAddressSpace(EltTy)));
llvm::FoldingSetNodeID ID;
ConstantArrayType::Profile(ID, EltTy, ArySize, ASM, IndexTypeQuals);
diff --git a/lib/Basic/TargetInfo.cpp b/lib/Basic/TargetInfo.cpp
index a9eeb8b..1696cae 100644
--- a/lib/Basic/TargetInfo.cpp
+++ b/lib/Basic/TargetInfo.cpp
@@ -11,6 +11,7 @@
//
//===----------------------------------------------------------------------===//
+#include "clang/Basic/AddressSpaces.h"
#include "clang/Basic/TargetInfo.h"
#include "clang/Basic/LangOptions.h"
#include "llvm/ADT/APFloat.h"
@@ -19,6 +20,8 @@
#include <cstdlib>
using namespace clang;
+static const LangAS::Map DefaultAddrSpaceMap = { 0 };
+
// TargetInfo Constructor.
TargetInfo::TargetInfo(const std::string &T) : Triple(T) {
// Set defaults. Defaults are set for a 32-bit RISC platform, like PPC or
@@ -64,6 +67,9 @@
// Default to using the Itanium ABI.
CXXABI = CXXABI_Itanium;
+
+ // Default to an empty address space map.
+ AddrSpaceMap = &DefaultAddrSpaceMap;
}
// Out of line virtual dtor for TargetInfo.
diff --git a/lib/CodeGen/CGCall.cpp b/lib/CodeGen/CGCall.cpp
index 15ab6b0..7a1600f 100644
--- a/lib/CodeGen/CGCall.cpp
+++ b/lib/CodeGen/CGCall.cpp
@@ -627,7 +627,8 @@
assert(!RetAI.getIndirectAlign() && "Align unused on indirect return.");
ResultType = llvm::Type::getVoidTy(getLLVMContext());
const llvm::Type *STy = ConvertType(RetTy, IsRecursive);
- ArgTys.push_back(llvm::PointerType::get(STy, RetTy.getAddressSpace()));
+ unsigned AS = Context.getTargetAddressSpace(RetTy);
+ ArgTys.push_back(llvm::PointerType::get(STy, AS));
break;
}
diff --git a/lib/CodeGen/CGDebugInfo.cpp b/lib/CodeGen/CGDebugInfo.cpp
index b23b392..bf5e7a5 100644
--- a/lib/CodeGen/CGDebugInfo.cpp
+++ b/lib/CodeGen/CGDebugInfo.cpp
@@ -462,8 +462,8 @@
// Bit size, align and offset of the type.
// Size is always the size of a pointer. We can't use getTypeSize here
// because that does not return the correct value for references.
- uint64_t Size =
- CGM.getContext().Target.getPointerWidth(PointeeTy.getAddressSpace());
+ unsigned AS = CGM.getContext().getTargetAddressSpace(PointeeTy);
+ uint64_t Size = CGM.getContext().Target.getPointerWidth(AS);
uint64_t Align = CGM.getContext().getTypeAlign(Ty);
return
diff --git a/lib/CodeGen/CGDecl.cpp b/lib/CodeGen/CGDecl.cpp
index 2cd4c87..7928656 100644
--- a/lib/CodeGen/CGDecl.cpp
+++ b/lib/CodeGen/CGDecl.cpp
@@ -179,7 +179,8 @@
new llvm::GlobalVariable(CGM.getModule(), LTy,
Ty.isConstant(getContext()), Linkage,
CGM.EmitNullConstant(D.getType()), Name, 0,
- D.isThreadSpecified(), Ty.getAddressSpace());
+ D.isThreadSpecified(),
+ CGM.getContext().getTargetAddressSpace(Ty));
GV->setAlignment(getContext().getDeclAlign(&D).getQuantity());
if (Linkage != llvm::GlobalValue::InternalLinkage)
GV->setVisibility(CurFn->getVisibility());
@@ -222,7 +223,7 @@
OldGV->getLinkage(), Init, "",
/*InsertBefore*/ OldGV,
D.isThreadSpecified(),
- D.getType().getAddressSpace());
+ CGM.getContext().getTargetAddressSpace(D.getType()));
GV->setVisibility(OldGV->getVisibility());
// Steal the name of the old global
@@ -289,7 +290,8 @@
// FIXME: It is really dangerous to store this in the map; if anyone
// RAUW's the GV uses of this constant will be invalid.
const llvm::Type *LTy = CGM.getTypes().ConvertTypeForMem(D.getType());
- const llvm::Type *LPtrTy = LTy->getPointerTo(D.getType().getAddressSpace());
+ const llvm::Type *LPtrTy =
+ LTy->getPointerTo(CGM.getContext().getTargetAddressSpace(D.getType()));
DMEntry = llvm::ConstantExpr::getBitCast(GV, LPtrTy);
// Emit global variable debug descriptor for static vars.
@@ -724,7 +726,8 @@
// Get the element type.
const llvm::Type *LElemTy = ConvertTypeForMem(Ty);
- const llvm::Type *LElemPtrTy = LElemTy->getPointerTo(Ty.getAddressSpace());
+ const llvm::Type *LElemPtrTy =
+ LElemTy->getPointerTo(CGM.getContext().getTargetAddressSpace(Ty));
llvm::Value *VLASize = EmitVLASize(Ty);
diff --git a/lib/CodeGen/CGExpr.cpp b/lib/CodeGen/CGExpr.cpp
index 99ea773..34a231c 100644
--- a/lib/CodeGen/CGExpr.cpp
+++ b/lib/CodeGen/CGExpr.cpp
@@ -726,7 +726,7 @@
// Cast to the access type.
const llvm::Type *PTy = llvm::Type::getIntNPtrTy(getLLVMContext(),
AI.AccessWidth,
- ExprType.getAddressSpace());
+ CGM.getContext().getTargetAddressSpace(ExprType));
Ptr = Builder.CreateBitCast(Ptr, PTy);
// Perform the load.
diff --git a/lib/CodeGen/CGExprConstant.cpp b/lib/CodeGen/CGExprConstant.cpp
index a735c3f..11b08c0 100644
--- a/lib/CodeGen/CGExprConstant.cpp
+++ b/lib/CodeGen/CGExprConstant.cpp
@@ -800,7 +800,7 @@
E->getType().isConstant(CGM.getContext()),
llvm::GlobalValue::InternalLinkage,
C, ".compoundliteral", 0, false,
- E->getType().getAddressSpace());
+ CGM.getContext().getTargetAddressSpace(E->getType()));
return C;
}
case Expr::DeclRefExprClass: {
diff --git a/lib/CodeGen/CodeGenModule.cpp b/lib/CodeGen/CodeGenModule.cpp
index 11e8587..ea74abf 100644
--- a/lib/CodeGen/CodeGenModule.cpp
+++ b/lib/CodeGen/CodeGenModule.cpp
@@ -1060,7 +1060,7 @@
Ty = getTypes().ConvertTypeForMem(ASTTy);
const llvm::PointerType *PTy =
- llvm::PointerType::get(Ty, ASTTy.getAddressSpace());
+ llvm::PointerType::get(Ty, getContext().getTargetAddressSpace(ASTTy));
llvm::StringRef MangledName = getMangledName(D);
return GetOrCreateLLVMGlobal(MangledName, PTy, D);
@@ -1239,7 +1239,8 @@
// from the type of the global (this happens with unions).
if (GV == 0 ||
GV->getType()->getElementType() != InitType ||
- GV->getType()->getAddressSpace() != ASTTy.getAddressSpace()) {
+ GV->getType()->getAddressSpace() !=
+ getContext().getTargetAddressSpace(ASTTy)) {
// Move the old entry aside so that we'll create a new one.
Entry->setName(llvm::StringRef());
diff --git a/lib/CodeGen/CodeGenTypes.cpp b/lib/CodeGen/CodeGenTypes.cpp
index 5254922..b9acbc4 100644
--- a/lib/CodeGen/CodeGenTypes.cpp
+++ b/lib/CodeGen/CodeGenTypes.cpp
@@ -270,14 +270,16 @@
QualType ETy = RTy.getPointeeType();
llvm::OpaqueType *PointeeType = llvm::OpaqueType::get(getLLVMContext());
PointersToResolve.push_back(std::make_pair(ETy, PointeeType));
- return llvm::PointerType::get(PointeeType, ETy.getAddressSpace());
+ unsigned AS = Context.getTargetAddressSpace(ETy);
+ return llvm::PointerType::get(PointeeType, AS);
}
case Type::Pointer: {
const PointerType &PTy = cast<PointerType>(Ty);
QualType ETy = PTy.getPointeeType();
llvm::OpaqueType *PointeeType = llvm::OpaqueType::get(getLLVMContext());
PointersToResolve.push_back(std::make_pair(ETy, PointeeType));
- return llvm::PointerType::get(PointeeType, ETy.getAddressSpace());
+ unsigned AS = Context.getTargetAddressSpace(ETy);
+ return llvm::PointerType::get(PointeeType, AS);
}
case Type::VariableArray: {
@@ -402,7 +404,8 @@
const QualType FTy = cast<BlockPointerType>(Ty).getPointeeType();
llvm::OpaqueType *PointeeType = llvm::OpaqueType::get(getLLVMContext());
PointersToResolve.push_back(std::make_pair(FTy, PointeeType));
- return llvm::PointerType::get(PointeeType, FTy.getAddressSpace());
+ unsigned AS = Context.getTargetAddressSpace(FTy);
+ return llvm::PointerType::get(PointeeType, AS);
}
case Type::MemberPointer: {
diff --git a/lib/Frontend/CompilerInvocation.cpp b/lib/Frontend/CompilerInvocation.cpp
index 9d9c996..2b5c289 100644
--- a/lib/Frontend/CompilerInvocation.cpp
+++ b/lib/Frontend/CompilerInvocation.cpp
@@ -672,6 +672,8 @@
Res.push_back("-fconstant-string-class");
Res.push_back(Opts.ObjCConstantStringClass);
}
+ if (Opts.FakeAddressSpaceMap)
+ Res.push_back("-ffake-address-space-map");
}
static void PreprocessorOptsToArgs(const PreprocessorOptions &Opts,
@@ -1483,6 +1485,7 @@
Opts.FastRelaxedMath = Args.hasArg(OPT_cl_fast_relaxed_math);
Opts.OptimizeSize = 0;
Opts.MRTD = Args.hasArg(OPT_mrtd);
+ Opts.FakeAddressSpaceMap = Args.hasArg(OPT_ffake_address_space_map);
// FIXME: Eliminate this dependency.
unsigned Opt = getOptimizationLevel(Args, IK, Diags);
diff --git a/lib/Parse/ParseDecl.cpp b/lib/Parse/ParseDecl.cpp
index 4208380..f27ff79 100644
--- a/lib/Parse/ParseDecl.cpp
+++ b/lib/Parse/ParseDecl.cpp
@@ -13,6 +13,7 @@
#include "clang/Parse/Parser.h"
#include "clang/Parse/ParseDiagnostic.h"
+#include "clang/Basic/OpenCL.h"
#include "clang/Sema/Scope.h"
#include "clang/Sema/ParsedTemplate.h"
#include "clang/Sema/PrettyDeclStackTrace.h"
@@ -311,6 +312,56 @@
}
}
+void Parser::ParseOpenCLQualifiers(DeclSpec &DS) {
+ SourceLocation Loc = Tok.getLocation();
+ switch(Tok.getKind()) {
+ // OpenCL qualifiers:
+ case tok::kw___private:
+ case tok::kw_private:
+ DS.addAttributes(AttrFactory.CreateIntegerAttribute(
+ Actions.getASTContext(),
+ PP.getIdentifierInfo("address_space"), Loc, 0));
+ break;
+
+ case tok::kw___global:
+ DS.addAttributes(AttrFactory.CreateIntegerAttribute(
+ Actions.getASTContext(),
+ PP.getIdentifierInfo("address_space"), Loc, LangAS::opencl_global));
+ break;
+
+ case tok::kw___local:
+ DS.addAttributes(AttrFactory.CreateIntegerAttribute(
+ Actions.getASTContext(),
+ PP.getIdentifierInfo("address_space"), Loc, LangAS::opencl_local));
+ break;
+
+ case tok::kw___constant:
+ DS.addAttributes(AttrFactory.CreateIntegerAttribute(
+ Actions.getASTContext(),
+ PP.getIdentifierInfo("address_space"), Loc, LangAS::opencl_constant));
+ break;
+
+ case tok::kw___read_only:
+ DS.addAttributes(AttrFactory.CreateIntegerAttribute(
+ Actions.getASTContext(),
+ PP.getIdentifierInfo("opencl_image_access"), Loc, CLIA_read_only));
+ break;
+
+ case tok::kw___write_only:
+ DS.addAttributes(AttrFactory.CreateIntegerAttribute(
+ Actions.getASTContext(),
+ PP.getIdentifierInfo("opencl_image_access"), Loc, CLIA_write_only));
+ break;
+
+ case tok::kw___read_write:
+ DS.addAttributes(AttrFactory.CreateIntegerAttribute(
+ Actions.getASTContext(),
+ PP.getIdentifierInfo("opencl_image_access"), Loc, CLIA_read_write));
+ break;
+ default: break;
+ }
+}
+
void Parser::DiagnoseProhibitedAttributes(ParsedAttributesWithRange &attrs) {
Diag(attrs.Range.getBegin(), diag::err_attributes_not_allowed)
<< attrs.Range;
@@ -1446,6 +1497,20 @@
ParseDecltypeSpecifier(DS);
continue;
+ // OpenCL qualifiers:
+ case tok::kw_private:
+ if (!getLang().OpenCL)
+ goto DoneWithDeclSpec;
+ case tok::kw___private:
+ case tok::kw___global:
+ case tok::kw___local:
+ case tok::kw___constant:
+ case tok::kw___read_only:
+ case tok::kw___write_only:
+ case tok::kw___read_write:
+ ParseOpenCLQualifiers(DS);
+ break;
+
case tok::less:
// GCC ObjC supports types like "<SomeProtocol>" as a synonym for
// "id<SomeProtocol>". This is hopelessly old fashioned and dangerous,
@@ -1697,6 +1762,20 @@
ParseDecltypeSpecifier(DS);
return true;
+ // OpenCL qualifiers:
+ case tok::kw_private:
+ if (!getLang().OpenCL)
+ return false;
+ case tok::kw___private:
+ case tok::kw___global:
+ case tok::kw___local:
+ case tok::kw___constant:
+ case tok::kw___read_only:
+ case tok::kw___write_only:
+ case tok::kw___read_write:
+ ParseOpenCLQualifiers(DS);
+ break;
+
// C++0x auto support.
case tok::kw_auto:
if (!getLang().CPlusPlus0x)
@@ -2269,10 +2348,22 @@
bool Parser::isTypeQualifier() const {
switch (Tok.getKind()) {
default: return false;
+
+ // type-qualifier only in OpenCL
+ case tok::kw_private:
+ return getLang().OpenCL;
+
// type-qualifier
case tok::kw_const:
case tok::kw_volatile:
case tok::kw_restrict:
+ case tok::kw___private:
+ case tok::kw___local:
+ case tok::kw___global:
+ case tok::kw___constant:
+ case tok::kw___read_only:
+ case tok::kw___read_write:
+ case tok::kw___write_only:
return true;
}
}
@@ -2400,7 +2491,19 @@
case tok::kw___w64:
case tok::kw___ptr64:
case tok::kw___pascal:
+
+ case tok::kw___private:
+ case tok::kw___local:
+ case tok::kw___global:
+ case tok::kw___constant:
+ case tok::kw___read_only:
+ case tok::kw___read_write:
+ case tok::kw___write_only:
+
return true;
+
+ case tok::kw_private:
+ return getLang().OpenCL;
}
}
@@ -2413,6 +2516,9 @@
switch (Tok.getKind()) {
default: return false;
+ case tok::kw_private:
+ return getLang().OpenCL;
+
case tok::identifier: // foo::bar
// Unfortunate hack to support "Class.factoryMethod" notation.
if (getLang().ObjC1 && NextToken().is(tok::period))
@@ -2522,6 +2628,15 @@
case tok::kw___ptr64:
case tok::kw___forceinline:
case tok::kw___pascal:
+
+ case tok::kw___private:
+ case tok::kw___local:
+ case tok::kw___global:
+ case tok::kw___constant:
+ case tok::kw___read_only:
+ case tok::kw___read_write:
+ case tok::kw___write_only:
+
return true;
}
}
@@ -2627,6 +2742,21 @@
isInvalid = DS.SetTypeQual(DeclSpec::TQ_restrict, Loc, PrevSpec, DiagID,
getLang());
break;
+
+ // OpenCL qualifiers:
+ case tok::kw_private:
+ if (!getLang().OpenCL)
+ goto DoneWithTypeQuals;
+ case tok::kw___private:
+ case tok::kw___global:
+ case tok::kw___local:
+ case tok::kw___constant:
+ case tok::kw___read_only:
+ case tok::kw___write_only:
+ case tok::kw___read_write:
+ ParseOpenCLQualifiers(DS);
+ break;
+
case tok::kw___w64:
case tok::kw___ptr64:
case tok::kw___cdecl:
diff --git a/lib/Sema/AttributeList.cpp b/lib/Sema/AttributeList.cpp
index 1663fd5..a715989 100644
--- a/lib/Sema/AttributeList.cpp
+++ b/lib/Sema/AttributeList.cpp
@@ -94,6 +94,7 @@
.Case("unavailable", AT_unavailable)
.Case("overloadable", AT_overloadable)
.Case("address_space", AT_address_space)
+ .Case("opencl_image_access", AT_opencl_image_access)
.Case("always_inline", AT_always_inline)
.Case("returns_twice", IgnoredAttribute)
.Case("vec_type_hint", IgnoredAttribute)
diff --git a/lib/Sema/SemaDeclAttr.cpp b/lib/Sema/SemaDeclAttr.cpp
index 5f6de14..3c11277 100644
--- a/lib/Sema/SemaDeclAttr.cpp
+++ b/lib/Sema/SemaDeclAttr.cpp
@@ -2724,6 +2724,7 @@
case AttributeList::AT_IBOutletCollection:
HandleIBOutletCollection(D, Attr, S); break;
case AttributeList::AT_address_space:
+ case AttributeList::AT_opencl_image_access:
case AttributeList::AT_objc_gc:
case AttributeList::AT_vector_size:
case AttributeList::AT_neon_vector_type:
diff --git a/lib/Sema/SemaType.cpp b/lib/Sema/SemaType.cpp
index b78c21b..dc2d3ac 100644
--- a/lib/Sema/SemaType.cpp
+++ b/lib/Sema/SemaType.cpp
@@ -13,6 +13,7 @@
#include "clang/Sema/SemaInternal.h"
#include "clang/Sema/Template.h"
+#include "clang/Basic/OpenCL.h"
#include "clang/AST/ASTContext.h"
#include "clang/AST/CXXInheritance.h"
#include "clang/AST/DeclObjC.h"
@@ -2965,6 +2966,41 @@
return true;
}
+/// Handle OpenCL image access qualifiers: read_only, write_only, read_write
+static void HandleOpenCLImageAccessAttribute(QualType& CurType,
+ const AttributeList &Attr,
+ Sema &S) {
+ // Check the attribute arguments.
+ if (Attr.getNumArgs() != 1) {
+ S.Diag(Attr.getLoc(), diag::err_attribute_wrong_number_arguments) << 1;
+ Attr.setInvalid();
+ return;
+ }
+ Expr *sizeExpr = static_cast<Expr *>(Attr.getArg(0));
+ llvm::APSInt arg(32);
+ if (sizeExpr->isTypeDependent() || sizeExpr->isValueDependent() ||
+ !sizeExpr->isIntegerConstantExpr(arg, S.Context)) {
+ S.Diag(Attr.getLoc(), diag::err_attribute_argument_not_int)
+ << "opencl_image_access" << sizeExpr->getSourceRange();
+ Attr.setInvalid();
+ return;
+ }
+ unsigned iarg = static_cast<unsigned>(arg.getZExtValue());
+ switch (iarg) {
+ case CLIA_read_only:
+ case CLIA_write_only:
+ case CLIA_read_write:
+ // Implemented in a separate patch
+ break;
+ default:
+ // Implemented in a separate patch
+ S.Diag(Attr.getLoc(), diag::err_attribute_invalid_size)
+ << sizeExpr->getSourceRange();
+ Attr.setInvalid();
+ break;
+ }
+}
+
/// HandleVectorSizeAttribute - this attribute is only applicable to integral
/// and float scalars, although arrays, pointers, and function return values are
/// allowed in conjunction with this construct. Aggregates with this attribute
@@ -3119,6 +3155,10 @@
"neon_polyvector_type");
break;
+ case AttributeList::AT_opencl_image_access:
+ HandleOpenCLImageAccessAttribute(type, attr, state.getSema());
+ break;
+
FUNCTION_TYPE_ATTRS_CASELIST:
// Never process function type attributes as part of the
// declaration-specifiers.
diff --git a/test/CodeGenOpenCL/address-spaces.cl b/test/CodeGenOpenCL/address-spaces.cl
new file mode 100644
index 0000000..e030c77
--- /dev/null
+++ b/test/CodeGenOpenCL/address-spaces.cl
@@ -0,0 +1,27 @@
+// RUN: %clang_cc1 %s -ffake-address-space-map -emit-llvm -o - | FileCheck %s
+
+void f__p(__private int *arg) { }
+// CHECK: i32* nocapture %arg
+
+void f__g(__global int *arg) { }
+// CHECK: i32 addrspace(1)* nocapture %arg
+
+void f__l(__local int *arg) { }
+// CHECK: i32 addrspace(2)* nocapture %arg
+
+void f__c(__constant int *arg) { }
+// CHECK: i32 addrspace(3)* nocapture %arg
+
+
+void fp(private int *arg) { }
+// CHECK: i32* nocapture %arg
+
+void fg(global int *arg) { }
+// CHECK: i32 addrspace(1)* nocapture %arg
+
+void fl(local int *arg) { }
+// CHECK: i32 addrspace(2)* nocapture %arg
+
+void fc(constant int *arg) { }
+// CHECK: i32 addrspace(3)* nocapture %arg
+
diff --git a/test/Parser/opencl-image-access.cl b/test/Parser/opencl-image-access.cl
new file mode 100644
index 0000000..313587c
--- /dev/null
+++ b/test/Parser/opencl-image-access.cl
@@ -0,0 +1,16 @@
+// RUN: %clang_cc1 %s -fsyntax-only
+
+typedef void* image2d_t;
+
+__kernel void f__ro(__read_only image2d_t a) { }
+
+__kernel void f__wo(__write_only image2d_t a) { }
+
+__kernel void f__rw(__read_write image2d_t a) { }
+
+
+__kernel void fro(read_only image2d_t a) { }
+
+__kernel void fwo(write_only image2d_t a) { }
+
+__kernel void frw(read_write image2d_t a) { }