[clang] a29aa47 - [OpenCL] Move addr space deduction to Sema.
Anastasia Stulova via cfe-commits
cfe-commits at lists.llvm.org
Wed Nov 27 04:45:07 PST 2019
Author: Anastasia Stulova
Date: 2019-11-27T12:44:42Z
New Revision: a29aa47106205ec95c12e0ebac4260c5de878a6a
URL: https://github.com/llvm/llvm-project/commit/a29aa47106205ec95c12e0ebac4260c5de878a6a
DIFF: https://github.com/llvm/llvm-project/commit/a29aa47106205ec95c12e0ebac4260c5de878a6a.diff
LOG: [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
Added:
clang/test/SemaOpenCLCXX/addrspace-auto.cl
Modified:
clang/include/clang/AST/Type.h
clang/include/clang/Sema/Sema.h
clang/lib/AST/Expr.cpp
clang/lib/Sema/SemaDecl.cpp
clang/lib/Sema/SemaTemplateInstantiate.cpp
clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
clang/lib/Sema/SemaType.cpp
clang/lib/Sema/TreeTransform.h
clang/test/SemaOpenCL/event_t.cl
clang/test/SemaOpenCL/invalid-block.cl
clang/test/SemaOpenCL/invalid-pipes-cl2.0.cl
clang/test/SemaOpenCL/sampler_t.cl
clang/test/SemaOpenCLCXX/address-space-deduction.cl
clang/test/SemaOpenCLCXX/restricted.cl
Removed:
################################################################################
diff --git a/clang/include/clang/AST/Type.h b/clang/include/clang/AST/Type.h
index ecbbd73e19fb..05e78aa78236 100644
--- a/clang/include/clang/AST/Type.h
+++ b/clang/include/clang/AST/Type.h
@@ -2069,6 +2069,8 @@ class alignas(8) Type : public ExtQualsTypeCommonBase {
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();
diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h
index 0a6f58a484ae..ac5a4953e00d 100644
--- a/clang/include/clang/Sema/Sema.h
+++ b/clang/include/clang/Sema/Sema.h
@@ -8760,6 +8760,8 @@ class Sema final {
bool CheckARCMethodDecl(ObjCMethodDecl *method);
bool inferObjCARCLifetime(ValueDecl *decl);
+ void deduceOpenCLAddressSpace(ValueDecl *decl);
+
ExprResult
HandleExprPropertyRefExpr(const ObjCObjectPointerType *OPT,
Expr *BaseExpr,
diff --git a/clang/lib/AST/Expr.cpp b/clang/lib/AST/Expr.cpp
index 3f722f8fd541..322b3a7fa740 100644
--- a/clang/lib/AST/Expr.cpp
+++ b/clang/lib/AST/Expr.cpp
@@ -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();
}
diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index 6ea4923dc2ba..dffb460cedc9 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -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;
}
diff --git a/clang/lib/Sema/SemaTemplateInstantiate.cpp b/clang/lib/Sema/SemaTemplateInstantiate.cpp
index 0daa33cfbef5..d75be4be988c 100644
--- a/clang/lib/Sema/SemaTemplateInstantiate.cpp
+++ b/clang/lib/Sema/SemaTemplateInstantiate.cpp
@@ -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
diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
index a2fd8a92dd61..9a6c7b5277b5 100644
--- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
+++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
@@ -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;
diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp
index b87978035ad8..2f5fdfb1f912 100644
--- a/clang/lib/Sema/SemaType.cpp
+++ b/clang/lib/Sema/SemaType.cpp
@@ -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) {
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index 28c5738eb266..47bd98a85030 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -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()) {
diff --git a/clang/test/SemaOpenCL/event_t.cl b/clang/test/SemaOpenCL/event_t.cl
index e7daf88576cc..ab7f09170e9c 100644
--- a/clang/test/SemaOpenCL/event_t.cl
+++ b/clang/test/SemaOpenCL/event_t.cl
@@ -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'}}
diff --git a/clang/test/SemaOpenCL/invalid-block.cl b/clang/test/SemaOpenCL/invalid-block.cl
index 5d6dc380a37a..7cbcea96d0ac 100644
--- a/clang/test/SemaOpenCL/invalid-block.cl
+++ b/clang/test/SemaOpenCL/invalid-block.cl
@@ -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}}
}
diff --git a/clang/test/SemaOpenCL/invalid-pipes-cl2.0.cl b/clang/test/SemaOpenCL/invalid-pipes-cl2.0.cl
index 69fa2b6da823..de1b4f8858fa 100644
--- a/clang/test/SemaOpenCL/invalid-pipes-cl2.0.cl
+++ b/clang/test/SemaOpenCL/invalid-pipes-cl2.0.cl
@@ -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}}
}
diff --git a/clang/test/SemaOpenCL/sampler_t.cl b/clang/test/SemaOpenCL/sampler_t.cl
index fe9d997c8960..888e973cc31d 100644
--- a/clang/test/SemaOpenCL/sampler_t.cl
+++ b/clang/test/SemaOpenCL/sampler_t.cl
@@ -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 at -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
diff --git a/clang/test/SemaOpenCLCXX/address-space-deduction.cl b/clang/test/SemaOpenCLCXX/address-space-deduction.cl
index ac6b2cabbd0c..9bffeafb1c2d 100644
--- a/clang/test/SemaOpenCLCXX/address-space-deduction.cl
+++ b/clang/test/SemaOpenCLCXX/address-space-deduction.cl
@@ -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 = ⅈ
+ 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() {
diff --git a/clang/test/SemaOpenCLCXX/addrspace-auto.cl b/clang/test/SemaOpenCLCXX/addrspace-auto.cl
new file mode 100644
index 000000000000..56fd9eb58ddc
--- /dev/null
+++ b/clang/test/SemaOpenCLCXX/addrspace-auto.cl
@@ -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 *'}}
+}
diff --git a/clang/test/SemaOpenCLCXX/restricted.cl b/clang/test/SemaOpenCLCXX/restricted.cl
index fc4938df5bf1..c00c634073fe 100644
--- a/clang/test/SemaOpenCLCXX/restricted.cl
+++ b/clang/test/SemaOpenCLCXX/restricted.cl
@@ -32,12 +32,14 @@ B *test_dynamic_cast(B *p) {
__constant _Thread_local int a = 1;
// expected-error at -1 {{C++ for OpenCL version 1.0 does not support the '_Thread_local' storage class specifier}}
// expected-warning at -2 {{'_Thread_local' is a C11 extension}}
-
+// expected-error at -3 {{thread-local storage is not supported for the current target}}
__constant __thread int b = 2;
// expected-error at -1 {{C++ for OpenCL version 1.0 does not support the '__thread' storage class specifier}}
+// expected-error at -2 {{thread-local storage is not supported for the current target}}
kernel void test_storage_classes() {
register int x;
// expected-error at -1 {{C++ for OpenCL version 1.0 does not support the 'register' storage class specifier}}
thread_local int y;
// expected-error at -1 {{C++ for OpenCL version 1.0 does not support the 'thread_local' storage class specifier}}
+ // expected-error at -2 {{thread-local storage is not supported for the current target}}
}
More information about the cfe-commits
mailing list