#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"
llvm::OwningPtr<CXXABI> ABI;
CXXABI *createCXXABI(const TargetInfo &T);
+ /// \brief The logical -> physical address space map.
+ const LangAS::Map &AddrSpaceMap;
+
friend class ASTDeclReader;
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;
--- /dev/null
+//===--- 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
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
--- /dev/null
+//===--- 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
#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>
const llvm::fltSemantics *FloatFormat, *DoubleFormat, *LongDoubleFormat;
unsigned char RegParmMax, SSERegParmMax;
TargetCXXABI CXXABI;
+ const LangAS::Map *AddrSpaceMap;
unsigned HasAlignMac68kSupport : 1;
unsigned RealTypeUsesObjCFPRet : 3;
virtual const char *getStaticInitSectionSpecifier() const {
return 0;
}
+
+ const LangAS::Map &getAddressSpaceMap() const {
+ return *AddrSpaceMap;
+ }
+
protected:
virtual uint64_t getPointerWidthV(unsigned AddrSpace) const {
return PointerWidth;
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)
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)
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
void ParseMicrosoftTypeAttributes(ParsedAttributes &attrs);
void ParseBorlandTypeAttributes(ParsedAttributes &attrs);
void ParseOpenCLAttributes(ParsedAttributes &attrs);
+ void ParseOpenCLQualifiers(DeclSpec &DS);
void ParseTypeofSpecifier(DeclSpec &DS);
void ParseDecltypeSpecifier(DeclSpec &DS);
#include "llvm/Support/Allocator.h"
#include "clang/Sema/Ownership.h"
#include "clang/Basic/SourceLocation.h"
+#include "clang/AST/Expr.h"
#include <cassert>
namespace clang {
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.
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.
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,
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),
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;
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;
// 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);
//
//===----------------------------------------------------------------------===//
+#include "clang/Basic/AddressSpaces.h"
#include "clang/Basic/TargetInfo.h"
#include "clang/Basic/LangOptions.h"
#include "llvm/ADT/APFloat.h"
#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
// 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.
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;
}
// 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
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());
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
// 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.
// 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);
// 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.
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: {
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);
// 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());
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: {
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: {
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,
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);
#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"
}
}
+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;
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,
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)
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;
}
}
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;
}
}
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))
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;
}
}
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:
.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)
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:
#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"
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
"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.
--- /dev/null
+// 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
+
--- /dev/null
+// 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) { }
projects / clang / commitdiff