maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

[OpenCL] Move addr space deduction to Sema. 

In order to simplify implementation we are moving add space
deduction into Sema while constructing variable declaration
and on template instantiation. Pointee are deduced to generic
addr space during creation of types.

This commit also
- fixed addr space dedution for auto type;
- factors out in a separate helper function OpenCL specific
  logic from type diagnostics in var decl.

Tags: #clang

Differential Revision: https://reviews.llvm.org/D65744
This commit is contained in:
Anastasia Stulova 2019-11-27 11:03:11 +00:00
parent 3edf2eb897
commit a29aa47106
15 changed files with 236 additions and 257 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
clang/include/clang/AST/Type.h
View File

@ -2069,6 +2069,8 @@ public:
bool isAlignValT() const; // C++17 std::align_val_t
bool isStdByteType() const; // C++17 std::byte
bool isAtomicType() const; // C11 _Atomic()
bool isUndeducedAutoType() const; // C++11 auto or
// C++14 decltype(auto)
#define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \
bool is##Id##Type() const;
@ -6509,6 +6511,10 @@ inline bool Type::isAtomicType() const {
return isa<AtomicType>(CanonicalType);
}
inline bool Type::isUndeducedAutoType() const {
return isa<AutoType>(CanonicalType);
}
inline bool Type::isObjCQualifiedIdType() const {
if (const auto *OPT = getAs<ObjCObjectPointerType>())
return OPT->isObjCQualifiedIdType();

2
clang/include/clang/Sema/Sema.h
View File

@ -8760,6 +8760,8 @@ public:
bool CheckARCMethodDecl(ObjCMethodDecl *method);
bool inferObjCARCLifetime(ValueDecl *decl);
void deduceOpenCLAddressSpace(ValueDecl *decl);
ExprResult
HandleExprPropertyRefExpr(const ObjCObjectPointerType *OPT,
Expr *BaseExpr,

2
clang/lib/AST/Expr.cpp
View File

@ -1814,7 +1814,7 @@ bool CastExpr::CastConsistency() const {
auto Ty = getType();
auto SETy = getSubExpr()->getType();
assert(getValueKindForType(Ty) == Expr::getValueKindForType(SETy));
if (/*isRValue()*/ !Ty->getPointeeType().isNull()) {
if (isRValue()) {
Ty = Ty->getPointeeType();
SETy = SETy->getPointeeType();
}

218
clang/lib/Sema/SemaDecl.cpp
View File

@ -6117,6 +6117,22 @@ bool Sema::inferObjCARCLifetime(ValueDecl *decl) {
return false;
}
void Sema::deduceOpenCLAddressSpace(ValueDecl *Decl) {
if (Decl->getType().getQualifiers().hasAddressSpace())
return;
if (VarDecl *Var = dyn_cast<VarDecl>(Decl)) {
QualType Type = Var->getType();
if (Type->isSamplerT() || Type->isVoidType())
return;
LangAS ImplAS = LangAS::opencl_private;
if ((getLangOpts().OpenCLCPlusPlus || getLangOpts().OpenCLVersion >= 200) &&
Var->hasGlobalStorage())
ImplAS = LangAS::opencl_global;
Type = Context.getAddrSpaceQualType(Type, ImplAS);
Decl->setType(Type);
}
}
static void checkAttributesAfterMerging(Sema &S, NamedDecl &ND) {
// Ensure that an auto decl is deduced otherwise the checks below might cache
// the wrong linkage.
@ -6474,6 +6490,105 @@ static bool isDeclExternC(const Decl *D) {
llvm_unreachable("Unknown type of decl!");
}
/// Returns true if there hasn't been any invalid type diagnosed.
static bool diagnoseOpenCLTypes(Scope *S, Sema &Se, Declarator &D,
DeclContext *DC, QualType R) {
// OpenCL v2.0 s6.9.b - Image type can only be used as a function argument.
// OpenCL v2.0 s6.13.16.1 - Pipe type can only be used as a function
// argument.
if (R->isImageType() || R->isPipeType()) {
Se.Diag(D.getIdentifierLoc(),
diag::err_opencl_type_can_only_be_used_as_function_parameter)
<< R;
D.setInvalidType();
return false;
}
// OpenCL v1.2 s6.9.r:
// The event type cannot be used to declare a program scope variable.
// OpenCL v2.0 s6.9.q:
// The clk_event_t and reserve_id_t types cannot be declared in program
// scope.
if (NULL == S->getParent()) {
if (R->isReserveIDT() || R->isClkEventT() || R->isEventT()) {
Se.Diag(D.getIdentifierLoc(),
diag::err_invalid_type_for_program_scope_var)
<< R;
D.setInvalidType();
return false;
}
}
// OpenCL v1.0 s6.8.a.3: Pointers to functions are not allowed.
QualType NR = R;
while (NR->isPointerType()) {
if (NR->isFunctionPointerType()) {
Se.Diag(D.getIdentifierLoc(), diag::err_opencl_function_pointer);
D.setInvalidType();
return false;
}
NR = NR->getPointeeType();
}
if (!Se.getOpenCLOptions().isEnabled("cl_khr_fp16")) {
// OpenCL v1.2 s6.1.1.1: reject declaring variables of the half and
// half array type (unless the cl_khr_fp16 extension is enabled).
if (Se.Context.getBaseElementType(R)->isHalfType()) {
Se.Diag(D.getIdentifierLoc(), diag::err_opencl_half_declaration) << R;
D.setInvalidType();
return false;
}
}
// OpenCL v1.2 s6.9.r:
// The event type cannot be used with the __local, __constant and __global
// address space qualifiers.
if (R->isEventT()) {
if (R.getAddressSpace() != LangAS::opencl_private) {
Se.Diag(D.getBeginLoc(), diag::err_event_t_addr_space_qual);
D.setInvalidType();
return false;
}
}
// C++ for OpenCL does not allow the thread_local storage qualifier.
// OpenCL C does not support thread_local either, and
// also reject all other thread storage class specifiers.
DeclSpec::TSCS TSC = D.getDeclSpec().getThreadStorageClassSpec();
if (TSC != TSCS_unspecified) {
bool IsCXX = Se.getLangOpts().OpenCLCPlusPlus;
Se.Diag(D.getDeclSpec().getThreadStorageClassSpecLoc(),
diag::err_opencl_unknown_type_specifier)
<< IsCXX << Se.getLangOpts().getOpenCLVersionTuple().getAsString()
<< DeclSpec::getSpecifierName(TSC) << 1;
D.setInvalidType();
return false;
}
if (R->isSamplerT()) {
// OpenCL v1.2 s6.9.b p4:
// The sampler type cannot be used with the __local and __global address
// space qualifiers.
if (R.getAddressSpace() == LangAS::opencl_local ||
R.getAddressSpace() == LangAS::opencl_global) {
Se.Diag(D.getIdentifierLoc(), diag::err_wrong_sampler_addressspace);
D.setInvalidType();
}
// OpenCL v1.2 s6.12.14.1:
// A global sampler must be declared with either the constant address
// space qualifier or with the const qualifier.
if (DC->isTranslationUnit() &&
!(R.getAddressSpace() == LangAS::opencl_constant ||
R.isConstQualified())) {
Se.Diag(D.getIdentifierLoc(), diag::err_opencl_nonconst_global_sampler);
D.setInvalidType();
}
if (D.isInvalidType())
return false;
}
return true;
}
NamedDecl *Sema::ActOnVariableDeclarator(
Scope *S, Declarator &D, DeclContext *DC, TypeSourceInfo *TInfo,
@ -6497,95 +6612,6 @@ NamedDecl *Sema::ActOnVariableDeclarator(
return nullptr;
}
if (getLangOpts().OpenCL) {
// OpenCL v2.0 s6.9.b - Image type can only be used as a function argument.
// OpenCL v2.0 s6.13.16.1 - Pipe type can only be used as a function
// argument.
if (R->isImageType() || R->isPipeType()) {
Diag(D.getIdentifierLoc(),
diag::err_opencl_type_can_only_be_used_as_function_parameter)
<< R;
D.setInvalidType();
return nullptr;
}
// OpenCL v1.2 s6.9.r:
// The event type cannot be used to declare a program scope variable.
// OpenCL v2.0 s6.9.q:
// The clk_event_t and reserve_id_t types cannot be declared in program scope.
if (NULL == S->getParent()) {
if (R->isReserveIDT() || R->isClkEventT() || R->isEventT()) {
Diag(D.getIdentifierLoc(),
diag::err_invalid_type_for_program_scope_var) << R;
D.setInvalidType();
return nullptr;
}
}
// OpenCL v1.0 s6.8.a.3: Pointers to functions are not allowed.
QualType NR = R;
while (NR->isPointerType()) {
if (NR->isFunctionPointerType()) {
Diag(D.getIdentifierLoc(), diag::err_opencl_function_pointer);
D.setInvalidType();
break;
}
NR = NR->getPointeeType();
}
if (!getOpenCLOptions().isEnabled("cl_khr_fp16")) {
// OpenCL v1.2 s6.1.1.1: reject declaring variables of the half and
// half array type (unless the cl_khr_fp16 extension is enabled).
if (Context.getBaseElementType(R)->isHalfType()) {
Diag(D.getIdentifierLoc(), diag::err_opencl_half_declaration) << R;
D.setInvalidType();
}
}
if (R->isSamplerT()) {
// OpenCL v1.2 s6.9.b p4:
// The sampler type cannot be used with the __local and __global address
// space qualifiers.
if (R.getAddressSpace() == LangAS::opencl_local ||
R.getAddressSpace() == LangAS::opencl_global) {
Diag(D.getIdentifierLoc(), diag::err_wrong_sampler_addressspace);
}
// OpenCL v1.2 s6.12.14.1:
// A global sampler must be declared with either the constant address
// space qualifier or with the const qualifier.
if (DC->isTranslationUnit() &&
!(R.getAddressSpace() == LangAS::opencl_constant ||
R.isConstQualified())) {
Diag(D.getIdentifierLoc(), diag::err_opencl_nonconst_global_sampler);
D.setInvalidType();
}
}
// OpenCL v1.2 s6.9.r:
// The event type cannot be used with the __local, __constant and __global
// address space qualifiers.
if (R->isEventT()) {
if (R.getAddressSpace() != LangAS::opencl_private) {
Diag(D.getBeginLoc(), diag::err_event_t_addr_space_qual);
D.setInvalidType();
}
}
// C++ for OpenCL does not allow the thread_local storage qualifier.
// OpenCL C does not support thread_local either, and
// also reject all other thread storage class specifiers.
DeclSpec::TSCS TSC = D.getDeclSpec().getThreadStorageClassSpec();
if (TSC != TSCS_unspecified) {
bool IsCXX = getLangOpts().OpenCLCPlusPlus;
Diag(D.getDeclSpec().getThreadStorageClassSpecLoc(),
diag::err_opencl_unknown_type_specifier)
<< IsCXX << getLangOpts().getOpenCLVersionTuple().getAsString()
<< DeclSpec::getSpecifierName(TSC) << 1;
D.setInvalidType();
return nullptr;
}
}
DeclSpec::SCS SCSpec = D.getDeclSpec().getStorageClassSpec();
StorageClass SC = StorageClassSpecToVarDeclStorageClass(D.getDeclSpec());
@ -6942,6 +6968,13 @@ NamedDecl *Sema::ActOnVariableDeclarator(
}
}
if (getLangOpts().OpenCL) {
deduceOpenCLAddressSpace(NewVD);
diagnoseOpenCLTypes(S, *this, D, DC, NewVD->getType());
}
// Handle attributes prior to checking for duplicates in MergeVarDecl
ProcessDeclAttributes(S, NewVD, D);
@ -11285,6 +11318,9 @@ bool Sema::DeduceVariableDeclarationType(VarDecl *VDecl, bool DirectInit,
if (getLangOpts().ObjCAutoRefCount && inferObjCARCLifetime(VDecl))
VDecl->setInvalidDecl();
if (getLangOpts().OpenCL)
deduceOpenCLAddressSpace(VDecl);
// If this is a redeclaration, check that the type we just deduced matches
// the previously declared type.
if (VarDecl *Old = VDecl->getPreviousDecl()) {
@ -13107,6 +13143,10 @@ Decl *Sema::ActOnParamDeclarator(Scope *S, Declarator &D) {
if (New->hasAttr<BlocksAttr>()) {
Diag(New->getLocation(), diag::err_block_on_nonlocal);
}
if (getLangOpts().OpenCL)
deduceOpenCLAddressSpace(New);
return New;
}

8
clang/lib/Sema/SemaTemplateInstantiate.cpp
View File

@ -1514,8 +1514,12 @@ TemplateInstantiator::TransformFunctionTypeParam(ParmVarDecl *OldParm,
int indexAdjustment,
Optional<unsigned> NumExpansions,
bool ExpectParameterPack) {
return SemaRef.SubstParmVarDecl(OldParm, TemplateArgs, indexAdjustment,
NumExpansions, ExpectParameterPack);
auto NewParm =
SemaRef.SubstParmVarDecl(OldParm, TemplateArgs, indexAdjustment,
NumExpansions, ExpectParameterPack);
if (NewParm && SemaRef.getLangOpts().OpenCL)
SemaRef.deduceOpenCLAddressSpace(NewParm);
return NewParm;
}
QualType

3
clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
View File

@ -930,6 +930,9 @@ Decl *TemplateDeclInstantiator::VisitVarDecl(VarDecl *D,
SemaRef.inferObjCARCLifetime(Var))
Var->setInvalidDecl();
if (SemaRef.getLangOpts().OpenCL)
SemaRef.deduceOpenCLAddressSpace(Var);
// Substitute the nested name specifier, if any.
if (SubstQualifier(D, Var))
return nullptr;

155
clang/lib/Sema/SemaType.cpp
View File

@ -1976,6 +1976,19 @@ bool Sema::CheckQualifiedFunctionForTypeId(QualType T, SourceLocation Loc) {
return true;
}
// Helper to deduce addr space of a pointee type in OpenCL mode.
static QualType deduceOpenCLPointeeAddrSpace(Sema &S, QualType PointeeType) {
if (!PointeeType->isUndeducedAutoType() && !PointeeType->isDependentType() &&
!PointeeType->isSamplerT() &&
!PointeeType.getQualifiers().hasAddressSpace())
PointeeType = S.getASTContext().getAddrSpaceQualType(
PointeeType,
S.getLangOpts().OpenCLCPlusPlus || S.getLangOpts().OpenCLVersion == 200
? LangAS::opencl_generic
: LangAS::opencl_private);
return PointeeType;
}
/// Build a pointer type.
///
/// \param T The type to which we'll be building a pointer.
@ -2012,6 +2025,9 @@ QualType Sema::BuildPointerType(QualType T,
if (getLangOpts().ObjCAutoRefCount)
T = inferARCLifetimeForPointee(*this, T, Loc, /*reference*/ false);
if (getLangOpts().OpenCL)
T = deduceOpenCLPointeeAddrSpace(*this, T);
// Build the pointer type.
return Context.getPointerType(T);
}
@ -2072,6 +2088,9 @@ QualType Sema::BuildReferenceType(QualType T, bool SpelledAsLValue,
if (getLangOpts().ObjCAutoRefCount)
T = inferARCLifetimeForPointee(*this, T, Loc, /*reference*/ true);
if (getLangOpts().OpenCL)
T = deduceOpenCLPointeeAddrSpace(*this, T);
// Handle restrict on references.
if (LValueRef)
return Context.getLValueReferenceType(T, SpelledAsLValue);
@ -2655,6 +2674,9 @@ QualType Sema::BuildBlockPointerType(QualType T,
if (checkQualifiedFunction(*this, T, Loc, QFK_BlockPointer))
return QualType();
if (getLangOpts().OpenCL)
T = deduceOpenCLPointeeAddrSpace(*this, T);
return Context.getBlockPointerType(T);
}
@ -7369,137 +7391,6 @@ static void HandleOpenCLAccessAttr(QualType &CurType, const ParsedAttr &Attr,
}
}
static void deduceOpenCLImplicitAddrSpace(TypeProcessingState &State,
QualType &T, TypeAttrLocation TAL) {
Declarator &D = State.getDeclarator();
// Handle the cases where address space should not be deduced.
//
// The pointee type of a pointer type is always deduced since a pointer always
// points to some memory location which should has an address space.
//
// There are situations that at the point of certain declarations, the address
// space may be unknown and better to be left as default. For example, when
// defining a typedef or struct type, they are not associated with any
// specific address space. Later on, they may be used with any address space
// to declare a variable.
//
// The return value of a function is r-value, therefore should not have
// address space.
//
// The void type does not occupy memory, therefore should not have address
// space, except when it is used as a pointee type.
//
// Since LLVM assumes function type is in default address space, it should not
// have address space.
auto ChunkIndex = State.getCurrentChunkIndex();
bool IsPointee =
ChunkIndex > 0 &&
(D.getTypeObject(ChunkIndex - 1).Kind == DeclaratorChunk::Pointer ||
D.getTypeObject(ChunkIndex - 1).Kind == DeclaratorChunk::Reference ||
D.getTypeObject(ChunkIndex - 1).Kind == DeclaratorChunk::BlockPointer);
// For pointers/references to arrays the next chunk is always an array
// followed by any number of parentheses.
if (!IsPointee && ChunkIndex > 1) {
auto AdjustedCI = ChunkIndex - 1;
if (D.getTypeObject(AdjustedCI).Kind == DeclaratorChunk::Array)
AdjustedCI--;
// Skip over all parentheses.
while (AdjustedCI > 0 &&
D.getTypeObject(AdjustedCI).Kind == DeclaratorChunk::Paren)
AdjustedCI--;
if (D.getTypeObject(AdjustedCI).Kind == DeclaratorChunk::Pointer ||
D.getTypeObject(AdjustedCI).Kind == DeclaratorChunk::Reference)
IsPointee = true;
}
bool IsFuncReturnType =
ChunkIndex > 0 &&
D.getTypeObject(ChunkIndex - 1).Kind == DeclaratorChunk::Function;
bool IsFuncType =
ChunkIndex < D.getNumTypeObjects() &&
D.getTypeObject(ChunkIndex).Kind == DeclaratorChunk::Function;
if ( // Do not deduce addr space for function return type and function type,
// otherwise it will fail some sema check.
IsFuncReturnType || IsFuncType ||
// Do not deduce addr space for member types of struct, except the pointee
// type of a pointer member type or static data members.
(D.getContext() == DeclaratorContext::MemberContext &&
(!IsPointee &&
D.getDeclSpec().getStorageClassSpec() != DeclSpec::SCS_static)) ||
// Do not deduce addr space of non-pointee in type alias because it
// doesn't define any object.
(D.getContext() == DeclaratorContext::AliasDeclContext && !IsPointee) ||
// Do not deduce addr space for types used to define a typedef and the
// typedef itself, except the pointee type of a pointer type which is used
// to define the typedef.
(D.getDeclSpec().getStorageClassSpec() == DeclSpec::SCS_typedef &&
!IsPointee) ||
// Do not deduce addr space of the void type, e.g. in f(void), otherwise
// it will fail some sema check.
(T->isVoidType() && !IsPointee) ||
// Do not deduce addr spaces for dependent types because they might end
// up instantiating to a type with an explicit address space qualifier.
// Except for pointer or reference types because the addr space in
// template argument can only belong to a pointee.
(T->isDependentType() && !T->isPointerType() && !T->isReferenceType()) ||
// Do not deduce addr space of decltype because it will be taken from
// its argument.
T->isDecltypeType() ||
// OpenCL spec v2.0 s6.9.b:
// The sampler type cannot be used with the __local and __global address
// space qualifiers.
// OpenCL spec v2.0 s6.13.14:
// Samplers can also be declared as global constants in the program
// source using the following syntax.
// const sampler_t <sampler name> = <value>
// In codegen, file-scope sampler type variable has special handing and
// does not rely on address space qualifier. On the other hand, deducing
// address space of const sampler file-scope variable as global address
// space causes spurious diagnostic about __global address space
// qualifier, therefore do not deduce address space of file-scope sampler
// type variable.
(D.getContext() == DeclaratorContext::FileContext && T->isSamplerT()))
return;
LangAS ImpAddr = LangAS::Default;
// Put OpenCL automatic variable in private address space.
// OpenCL v1.2 s6.5:
// The default address space name for arguments to a function in a
// program, or local variables of a function is __private. All function
// arguments shall be in the __private address space.
if (State.getSema().getLangOpts().OpenCLVersion <= 120 &&
!State.getSema().getLangOpts().OpenCLCPlusPlus) {
ImpAddr = LangAS::opencl_private;
} else {
// If address space is not set, OpenCL 2.0 defines non private default
// address spaces for some cases:
// OpenCL 2.0, section 6.5:
// The address space for a variable at program scope or a static variable
// inside a function can either be __global or __constant, but defaults to
// __global if not specified.
// (...)
// Pointers that are declared without pointing to a named address space
// point to the generic address space.
if (IsPointee) {
ImpAddr = LangAS::opencl_generic;
} else {
if (D.getContext() == DeclaratorContext::TemplateArgContext) {
// Do not deduce address space for non-pointee type in template arg.
} else if (D.getContext() == DeclaratorContext::FileContext) {
ImpAddr = LangAS::opencl_global;
} else {
if (D.getDeclSpec().getStorageClassSpec() == DeclSpec::SCS_static ||
D.getDeclSpec().getStorageClassSpec() == DeclSpec::SCS_extern) {
ImpAddr = LangAS::opencl_global;
} else {
ImpAddr = LangAS::opencl_private;
}
}
}
}
T = State.getSema().Context.getAddrSpaceQualType(T, ImpAddr);
}
static void HandleLifetimeBoundAttr(TypeProcessingState &State,
QualType &CurType,
ParsedAttr &Attr) {
@ -7729,8 +7620,6 @@ static void processTypeAttrs(TypeProcessingState &state, QualType &type,
if (!state.getSema().getLangOpts().OpenCL ||
type.getAddressSpace() != LangAS::Default)
return;
deduceOpenCLImplicitAddrSpace(state, type, TAL);
}
void Sema::completeExprArrayBound(Expr *E) {

17
clang/lib/Sema/TreeTransform.h
View File

@ -4579,14 +4579,6 @@ QualType TreeTransform<Derived>::TransformDecayedType(TypeLocBuilder &TLB,
return Result;
}
/// Helper to deduce addr space of a pointee type in OpenCL mode.
/// If the type is updated it will be overwritten in PointeeType param.
inline void deduceOpenCLPointeeAddrSpace(Sema &SemaRef, QualType &PointeeType) {
if (PointeeType.getAddressSpace() == LangAS::Default)
PointeeType = SemaRef.Context.getAddrSpaceQualType(PointeeType,
LangAS::opencl_generic);
}
template<typename Derived>
QualType TreeTransform<Derived>::TransformPointerType(TypeLocBuilder &TLB,
PointerTypeLoc TL) {
@ -4595,9 +4587,6 @@ QualType TreeTransform<Derived>::TransformPointerType(TypeLocBuilder &TLB,
if (PointeeType.isNull())
return QualType();
if (SemaRef.getLangOpts().OpenCL)
deduceOpenCLPointeeAddrSpace(SemaRef, PointeeType);
QualType Result = TL.getType();
if (PointeeType->getAs<ObjCObjectType>()) {
// A dependent pointer type 'T *' has is being transformed such
@ -4636,9 +4625,6 @@ TreeTransform<Derived>::TransformBlockPointerType(TypeLocBuilder &TLB,
if (PointeeType.isNull())
return QualType();
if (SemaRef.getLangOpts().OpenCL)
deduceOpenCLPointeeAddrSpace(SemaRef, PointeeType);
QualType Result = TL.getType();
if (getDerived().AlwaysRebuild() ||
PointeeType != TL.getPointeeLoc().getType()) {
@ -4668,9 +4654,6 @@ TreeTransform<Derived>::TransformReferenceType(TypeLocBuilder &TLB,
if (PointeeType.isNull())
return QualType();
if (SemaRef.getLangOpts().OpenCL)
deduceOpenCLPointeeAddrSpace(SemaRef, PointeeType);
QualType Result = TL.getType();
if (getDerived().AlwaysRebuild() ||
PointeeType != T->getPointeeTypeAsWritten()) {

4
clang/test/SemaOpenCL/event_t.cl
View File

@ -1,6 +1,6 @@
// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
event_t glb_evt; // expected-error {{the 'event_t' type cannot be used to declare a program scope variable}}
event_t glb_evt; // expected-error {{the 'event_t' type cannot be used to declare a program scope variable}} expected-error{{program scope variable must reside in constant address space}}
constant struct evt_s {
event_t evt; // expected-error {{the 'event_t' type cannot be used to declare a structure or union field}}
@ -10,7 +10,7 @@ void foo(event_t evt); // expected-note {{passing argument to parameter 'evt' he
void kernel ker(event_t argevt) { // expected-error {{'event_t' cannot be used as the type of a kernel parameter}}
event_t e;
constant event_t const_evt; // expected-error {{the event_t type can only be used with __private address space qualifier}}
constant event_t const_evt; // expected-error {{the event_t type can only be used with __private address space qualifier}} expected-error{{variable in constant address space must be initialized}}
foo(e);
foo(0);
foo(5); // expected-error {{passing 'int' to parameter of incompatible type 'event_t'}}

4
clang/test/SemaOpenCL/invalid-block.cl
View File

@ -58,11 +58,11 @@ void f5(int i) {
: bl2(i); // expected-error {{block type cannot be used as expression in ternary expression in OpenCL}}
}
// A block pointer type and all pointer operations are disallowed
void f6(bl2_t *bl_ptr) { // expected-error{{pointer to type '__generic bl2_t' (aka 'int (__generic ^const __generic)(int)') is invalid in OpenCL}}
void f6(bl2_t *bl_ptr) { // expected-error{{pointer to type 'bl2_t' (aka 'int (__generic ^const)(int)') is invalid in OpenCL}}
bl2_t bl = ^(int i) {
return 1;
};
bl2_t *p; // expected-error {{pointer to type '__generic bl2_t' (aka 'int (__generic ^const __generic)(int)') is invalid in OpenCL}}
bl2_t *p; // expected-error {{pointer to type 'bl2_t' (aka 'int (__generic ^const)(int)') is invalid in OpenCL}}
*bl; // expected-error {{invalid argument type 'bl2_t' (aka 'int (__generic ^const)(int)') to unary expression}}
&bl; // expected-error {{invalid argument type 'bl2_t' (aka 'int (__generic ^const)(int)') to unary expression}}
}

2
clang/test/SemaOpenCL/invalid-pipes-cl2.0.cl
View File

@ -4,7 +4,7 @@
global pipe int gp; // expected-error {{type '__global read_only pipe int' can only be used as a function parameter in OpenCL}}
global reserve_id_t rid; // expected-error {{the '__global reserve_id_t' type cannot be used to declare a program scope variable}}
extern pipe write_only int get_pipe(); // expected-error-re{{type '__global write_only pipe int ({{(void)?}})' can only be used as a function parameter in OpenCL}}
extern pipe write_only int get_pipe(); // expected-error-re{{type '__global write_only pipe int ({{(void)?}})' can only be used as a function parameter in OpenCL}} expected-error{{'write_only' attribute only applies to parameters and typedefs}}
kernel void test_invalid_reserved_id(reserve_id_t ID) { // expected-error {{'reserve_id_t' cannot be used as the type of a kernel parameter}}
}

5
clang/test/SemaOpenCL/sampler_t.cl
View File

@ -48,6 +48,9 @@ constant struct sampler_s {
sampler_t bad(void); //expected-error{{declaring function return value of type 'sampler_t' is not allowed}}
sampler_t global_nonconst_smp = 0; // expected-error {{global sampler requires a const or constant address space qualifier}}
#ifdef CHECK_SAMPLER_VALUE
// expected-warning@-2{{sampler initializer has invalid Filter Mode bits}}
#endif
const sampler_t glb_smp10 = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_NORMALIZED_COORDS_TRUE | CLK_FILTER_LINEAR;
const constant sampler_t glb_smp11 = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_NORMALIZED_COORDS_TRUE | CLK_FILTER_LINEAR;
@ -62,7 +65,7 @@ void kernel ker(sampler_t argsmp) {
}
#if __OPENCL_C_VERSION__ == 200
void bad(sampler_t*); // expected-error{{pointer to type '__generic sampler_t' is invalid in OpenCL}}
void bad(sampler_t *); // expected-error{{pointer to type 'sampler_t' is invalid in OpenCL}}
#else
void bad(sampler_t*); // expected-error{{pointer to type 'sampler_t' is invalid in OpenCL}}
#endif

28
clang/test/SemaOpenCLCXX/address-space-deduction.cl
View File

@ -65,30 +65,42 @@ template <typename T>
x3<T>::x3(const x3<T> &t) {}
template <class T>
T xxx(T *in) {
T xxx(T *in1, T in2) {
// This pointer can't be deduced to generic because addr space
// will be taken from the template argument.
//CHECK: `-VarDecl {{.*}} i 'T *' cinit
T *i = in;
T *i = in1;
T ii;
__private T *ptr = &ii;
ptr = &in2;
return *i;
}
__kernel void test() {
int foo[10];
xxx(&foo[0]);
xxx<__private int>(&foo[0], foo[0]);
// FIXME: Template param deduction fails here because
// temporaries are not in the __private address space.
// It is probably reasonable to put them in __private
// considering that stack and function params are
// implicitly in __private.
// However, if temporaries are left in default addr
// space we should at least pretty print the __private
// addr space. Otherwise diagnostic apprears to be
// confusing.
//xxx(&foo[0], foo[0]);
}
// Addr space for pointer/reference to an array
//CHECK: FunctionDecl {{.*}} t1 'void (const __generic float (&)[2])'
//CHECK: FunctionDecl {{.*}} t1 'void (const float (__generic &)[2])'
void t1(const float (&fYZ)[2]);
//CHECK: FunctionDecl {{.*}} t2 'void (const __generic float (*)[2])'
//CHECK: FunctionDecl {{.*}} t2 'void (const float (__generic *)[2])'
void t2(const float (*fYZ)[2]);
//CHECK: FunctionDecl {{.*}} t3 'void (__generic float (((*)))[2])'
//CHECK: FunctionDecl {{.*}} t3 'void (float (((__generic *)))[2])'
void t3(float(((*fYZ)))[2]);
//CHECK: FunctionDecl {{.*}} t4 'void (__generic float (((*__generic *)))[2])'
//CHECK: FunctionDecl {{.*}} t4 'void (float (((__generic *__generic *)))[2])'
void t4(float(((**fYZ)))[2]);
//CHECK: FunctionDecl {{.*}} t5 'void (__generic float (*__generic (*))[2])'
//CHECK: FunctionDecl {{.*}} t5 'void (float (__generic *(__generic *))[2])'
void t5(float (*(*fYZ))[2]);
__kernel void k() {

35
clang/test/SemaOpenCLCXX/addrspace-auto.cl Normal file
View File

@ -0,0 +1,35 @@
//RUN: %clang_cc1 %s -cl-std=clc++ -pedantic -ast-dump -verify | FileCheck %s
__constant int i = 1;
//CHECK: |-VarDecl {{.*}} ai '__global int':'__global int'
auto ai = i;
kernel void test() {
int i;
//CHECK: VarDecl {{.*}} ai 'int':'int'
auto ai = i;
constexpr int c = 1;
//CHECK: VarDecl {{.*}} used cai '__constant int':'__constant int'
__constant auto cai = c;
//CHECK: VarDecl {{.*}} aii 'int':'int'
auto aii = cai;
//CHECK: VarDecl {{.*}} ref 'int &'
auto &ref = i;
//CHECK: VarDecl {{.*}} ptr 'int *'
auto *ptr = &i;
//CHECK: VarDecl {{.*}} ref_c '__constant int &'
auto &ref_c = cai;
//CHECK: VarDecl {{.*}} ptrptr 'int *__generic *'
auto **ptrptr = &ptr;
//CHECK: VarDecl {{.*}} refptr 'int *__generic &'
auto *&refptr = ptr;
//CHECK: VarDecl {{.*}} invalid gref '__global auto &'
__global auto &gref = i; //expected-error{{variable 'gref' with type '__global auto &' has incompatible initializer of type 'int'}}
__local int *ptr_l;
//CHECK: VarDecl {{.*}} invalid gptr '__global auto *'
__global auto *gptr = ptr_l; //expected-error{{variable 'gptr' with type '__global auto *' has incompatible initializer of type '__local int *'}}
}

4
clang/test/SemaOpenCLCXX/restricted.cl
View File

@ -32,12 +32,14 @@ B *test_dynamic_cast(B *p) {
__constant _Thread_local int a = 1;
// expected-error@-1 {{C++ for OpenCL version 1.0 does not support the '_Thread_local' storage class specifier}}
// expected-warning@-2 {{'_Thread_local' is a C11 extension}}
// expected-error@-3 {{thread-local storage is not supported for the current target}}
__constant __thread int b = 2;
// expected-error@-1 {{C++ for OpenCL version 1.0 does not support the '__thread' storage class specifier}}
// expected-error@-2 {{thread-local storage is not supported for the current target}}
kernel void test_storage_classes() {
register int x;
// expected-error@-1 {{C++ for OpenCL version 1.0 does not support the 'register' storage class specifier}}
thread_local int y;
// expected-error@-1 {{C++ for OpenCL version 1.0 does not support the 'thread_local' storage class specifier}}
// expected-error@-2 {{thread-local storage is not supported for the current target}}
}