From: Peter Collingbourne Date: Fri, 18 Mar 2011 22:38:29 +0000 (+0000) Subject: Add support for language-specific address spaces. On top of that, X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=207f4d8543529221932af82836016a2ef066c917;p=clang 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 ad69f0f35a..68a12bba24 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 @@ class ASTContext { llvm::OwningPtr 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 @@ public: 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 0000000000..d44a9c3b03 --- /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 ebdcddeb35..433f42ecf3 100644 --- a/include/clang/Basic/LangOptions.h +++ b/include/clang/Basic/LangOptions.h @@ -124,6 +124,8 @@ public: 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 0000000000..6f9785f256 --- /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 b9087f2c47..fd9e373b73 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 #include #include @@ -78,6 +79,7 @@ protected: 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 @@ public: 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 c7d9524b41..6fa4bf9304 100644 --- a/include/clang/Basic/TokenKinds.def +++ b/include/clang/Basic/TokenKinds.def @@ -252,7 +252,7 @@ KEYWORD(mutable , KEYCXX) 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(__forceinline , KEYALL) 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 bb8759e270..418823482b 100644 --- a/include/clang/Driver/CC1Options.td +++ b/include/clang/Driver/CC1Options.td @@ -513,6 +513,8 @@ def fno_bitfield_type_align : Flag<"-fno-bitfield-type-align">, 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 bd39cdb452..b9af05e58f 100644 --- a/include/clang/Parse/Parser.h +++ b/include/clang/Parse/Parser.h @@ -1546,6 +1546,7 @@ private: 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 10b859e9f9..382e50494e 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 namespace clang { @@ -76,6 +77,13 @@ public: 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 @@ public: 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 d20c25851e..fd5fd0e54b 100644 --- a/lib/AST/ASTContext.cpp +++ b/lib/AST/ASTContext.cpp @@ -190,6 +190,22 @@ CXXABI *ASTContext::createCXXABI(const TargetInfo &T) { 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 @@ ASTContext::ASTContext(const LangOptions& LOpts, SourceManager &SM, 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 @@ ASTContext::getTypeInfo(const Type *T) const { Align = Target.getPointerAlign(0); break; case Type::BlockPointer: { - unsigned AS = cast(T)->getPointeeType().getAddressSpace(); + unsigned AS = getTargetAddressSpace( + cast(T)->getPointeeType()); Width = Target.getPointerWidth(AS); Align = Target.getPointerAlign(AS); break; @@ -815,13 +833,14 @@ ASTContext::getTypeInfo(const Type *T) const { case Type::RValueReference: { // alignof and sizeof should never enter this code path here, so we go // the pointer route. - unsigned AS = cast(T)->getPointeeType().getAddressSpace(); + unsigned AS = getTargetAddressSpace( + cast(T)->getPointeeType()); Width = Target.getPointerWidth(AS); Align = Target.getPointerAlign(AS); break; } case Type::Pointer: { - unsigned AS = cast(T)->getPointeeType().getAddressSpace(); + unsigned AS = getTargetAddressSpace(cast(T)->getPointeeType()); Width = Target.getPointerWidth(AS); Align = Target.getPointerAlign(AS); break; @@ -1468,7 +1487,7 @@ QualType ASTContext::getConstantArrayType(QualType EltTy, // 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 a9eeb8b4cc..1696cae840 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 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 @@ TargetInfo::TargetInfo(const std::string &T) : Triple(T) { // 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 15ab6b080b..7a1600f166 100644 --- a/lib/CodeGen/CGCall.cpp +++ b/lib/CodeGen/CGCall.cpp @@ -627,7 +627,8 @@ CodeGenTypes::GetFunctionType(const CGFunctionInfo &FI, bool IsVariadic, 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 b23b392222..bf5e7a5d1a 100644 --- a/lib/CodeGen/CGDebugInfo.cpp +++ b/lib/CodeGen/CGDebugInfo.cpp @@ -462,8 +462,8 @@ llvm::DIType CGDebugInfo::CreatePointerLikeType(unsigned Tag, // 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 2cd4c87d3f..792865668b 100644 --- a/lib/CodeGen/CGDecl.cpp +++ b/lib/CodeGen/CGDecl.cpp @@ -179,7 +179,8 @@ CodeGenFunction::CreateStaticVarDecl(const VarDecl &D, 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 @@ CodeGenFunction::AddInitializerToStaticVarDecl(const VarDecl &D, 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 @@ void CodeGenFunction::EmitStaticVarDecl(const VarDecl &D, // 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 @@ CodeGenFunction::EmitAutoVarAlloca(const VarDecl &D) { // 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 99ea773fa0..34a231cea1 100644 --- a/lib/CodeGen/CGExpr.cpp +++ b/lib/CodeGen/CGExpr.cpp @@ -726,7 +726,7 @@ RValue CodeGenFunction::EmitLoadOfBitfieldLValue(LValue LV, // 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 a735c3f409..11b08c0b04 100644 --- a/lib/CodeGen/CGExprConstant.cpp +++ b/lib/CodeGen/CGExprConstant.cpp @@ -800,7 +800,7 @@ public: 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 11e85870d1..ea74abff54 100644 --- a/lib/CodeGen/CodeGenModule.cpp +++ b/lib/CodeGen/CodeGenModule.cpp @@ -1060,7 +1060,7 @@ llvm::Constant *CodeGenModule::GetAddrOfGlobalVar(const VarDecl *D, 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 @@ void CodeGenModule::EmitGlobalVarDefinition(const VarDecl *D) { // 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 5254922f13..b9acbc4e9b 100644 --- a/lib/CodeGen/CodeGenTypes.cpp +++ b/lib/CodeGen/CodeGenTypes.cpp @@ -270,14 +270,16 @@ const llvm::Type *CodeGenTypes::ConvertNewType(QualType T) { 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(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 llvm::Type *CodeGenTypes::ConvertNewType(QualType T) { const QualType FTy = cast(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 9d9c99611c..2b5c289cbf 100644 --- a/lib/Frontend/CompilerInvocation.cpp +++ b/lib/Frontend/CompilerInvocation.cpp @@ -672,6 +672,8 @@ static void LangOptsToArgs(const LangOptions &Opts, 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 @@ static void ParseLangArgs(LangOptions &Opts, ArgList &Args, InputKind IK, 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 42083803b6..f27ff79f19 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::ParseOpenCLAttributes(ParsedAttributes &attrs) { } } +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 @@ void Parser::ParseDeclarationSpecifiers(DeclSpec &DS, 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 "" as a synonym for // "id". This is hopelessly old fashioned and dangerous, @@ -1697,6 +1762,20 @@ bool Parser::ParseOptionalTypeSpecifier(DeclSpec &DS, bool& isInvalid, 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 @@ void Parser::ParseEnumBody(SourceLocation StartLoc, Decl *EnumDecl) { 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 @@ bool Parser::isTypeSpecifierQualifier() { 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 @@ bool Parser::isDeclarationSpecifier(bool DisambiguatingWithExpression) { 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 @@ bool Parser::isDeclarationSpecifier(bool DisambiguatingWithExpression) { 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 @@ void Parser::ParseTypeQualifierListOpt(DeclSpec &DS, 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 1663fd58a9..a715989c68 100644 --- a/lib/Sema/AttributeList.cpp +++ b/lib/Sema/AttributeList.cpp @@ -94,6 +94,7 @@ AttributeList::Kind AttributeList::getKind(const IdentifierInfo *Name) { .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 5f6de1409e..3c1127784d 100644 --- a/lib/Sema/SemaDeclAttr.cpp +++ b/lib/Sema/SemaDeclAttr.cpp @@ -2724,6 +2724,7 @@ static void ProcessInheritableDeclAttr(Scope *scope, Decl *D, 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 b78c21bbb6..dc2d3ac518 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 @@ static bool handleFunctionTypeAttr(TypeProcessingState &state, 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(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(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 @@ static void processTypeAttrs(TypeProcessingState &state, QualType &type, "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 0000000000..e030c77d3e --- /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 0000000000..313587c1d2 --- /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) { }