Anastasia updated this revision to Diff 219903.
Anastasia added a comment.
- Move addr space deduction to a later phase.
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D65744/new/
https://reviews.llvm.org/D65744
Files:
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/addrspace-auto.cl
clang/test/SemaOpenCLCXX/restricted.cl
Index: clang/test/SemaOpenCLCXX/restricted.cl
===================================================================
--- clang/test/SemaOpenCLCXX/restricted.cl
+++ clang/test/SemaOpenCLCXX/restricted.cl
@@ -32,12 +32,14 @@
__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}}
+// 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@-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}}
}
Index: clang/test/SemaOpenCLCXX/addrspace-auto.cl
===================================================================
--- /dev/null
+++ 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 *'}}
+}
Index: clang/test/SemaOpenCLCXX/address-space-deduction.cl
===================================================================
--- clang/test/SemaOpenCLCXX/address-space-deduction.cl
+++ clang/test/SemaOpenCLCXX/address-space-deduction.cl
@@ -65,18 +65,30 @@
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
Index: clang/test/SemaOpenCL/sampler_t.cl
===================================================================
--- clang/test/SemaOpenCL/sampler_t.cl
+++ clang/test/SemaOpenCL/sampler_t.cl
@@ -48,6 +48,9 @@
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 @@
}
#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
Index: clang/test/SemaOpenCL/invalid-pipes-cl2.0.cl
===================================================================
--- clang/test/SemaOpenCL/invalid-pipes-cl2.0.cl
+++ 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 'write_only pipe int ({{(void)?}})' can only be used as a function parameter in OpenCL}}
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}}
}
Index: clang/test/SemaOpenCL/invalid-block.cl
===================================================================
--- clang/test/SemaOpenCL/invalid-block.cl
+++ clang/test/SemaOpenCL/invalid-block.cl
@@ -58,11 +58,11 @@
: 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}}
}
Index: clang/test/SemaOpenCL/event_t.cl
===================================================================
--- clang/test/SemaOpenCL/event_t.cl
+++ clang/test/SemaOpenCL/event_t.cl
@@ -10,7 +10,7 @@
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'}}
Index: clang/lib/Sema/TreeTransform.h
===================================================================
--- clang/lib/Sema/TreeTransform.h
+++ clang/lib/Sema/TreeTransform.h
@@ -4535,14 +4535,6 @@
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.
-static 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) {
@@ -4551,9 +4543,6 @@
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
@@ -4592,9 +4581,6 @@
if (PointeeType.isNull())
return QualType();
- if (SemaRef.getLangOpts().OpenCL)
- deduceOpenCLPointeeAddrSpace(SemaRef, PointeeType);
-
QualType Result = TL.getType();
if (getDerived().AlwaysRebuild() ||
PointeeType != TL.getPointeeLoc().getType()) {
@@ -4624,9 +4610,6 @@
if (PointeeType.isNull())
return QualType();
- if (SemaRef.getLangOpts().OpenCL)
- deduceOpenCLPointeeAddrSpace(SemaRef, PointeeType);
-
QualType Result = TL.getType();
if (getDerived().AlwaysRebuild() ||
PointeeType != T->getPointeeTypeAsWritten()) {
Index: clang/lib/Sema/SemaType.cpp
===================================================================
--- clang/lib/Sema/SemaType.cpp
+++ clang/lib/Sema/SemaType.cpp
@@ -1962,6 +1962,19 @@
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() &&
+ !isa<ParenType>(PointeeType) && !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.
@@ -1998,6 +2011,9 @@
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);
}
@@ -2058,6 +2074,9 @@
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);
@@ -2626,6 +2645,9 @@
if (checkQualifiedFunction(*this, T, Loc, QFK_BlockPointer))
return QualType();
+ if (getLangOpts().OpenCL)
+ T = deduceOpenCLPointeeAddrSpace(*this, T);
+
return Context.getBlockPointerType(T);
}
@@ -4350,6 +4372,18 @@
bool ExpectNoDerefChunk =
state.getCurrentAttributes().hasAttribute(ParsedAttr::AT_NoDeref);
+ auto DeduceOpenCLPointeeAddrSpaceForParenType = [&] {
+ auto AdjustedCI = state.getCurrentChunkIndex() - 1;
+ // 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 ||
+ D.getTypeObject(AdjustedCI).Kind == DeclaratorChunk::BlockPointer)
+ T = deduceOpenCLPointeeAddrSpace(S, T);
+ };
+
// Walk the DeclTypeInfo, building the recursive type as we go.
// DeclTypeInfos are ordered from the identifier out, which is
// opposite of what we want :).
@@ -4362,6 +4396,12 @@
case DeclaratorChunk::Paren:
if (i == 0)
warnAboutRedundantParens(S, D, T);
+
+ if (LangOpts.OpenCL && state.getCurrentChunkIndex() > 0) {
+ // Look the tokens ahead to see whether we are parsing a pointee
+ // in which case we might need to deduce addr space.
+ DeduceOpenCLPointeeAddrSpaceForParenType();
+ }
T = S.BuildParenType(T);
break;
case DeclaratorChunk::BlockPointer:
@@ -4508,8 +4548,15 @@
checkNullabilityConsistency(S, SimplePointerKind::Array, DeclType.Loc);
}
+ if (LangOpts.OpenCL && state.getCurrentChunkIndex() > 1) {
+ // Look the tokens ahead and see if we are parsing a pointer or a
+ // reference to array.
+ DeduceOpenCLPointeeAddrSpaceForParenType();
+ }
+
T = S.BuildArrayType(T, ASM, ArraySize, ATI.TypeQuals,
SourceRange(DeclType.Loc, DeclType.EndLoc), Name);
+
break;
}
case DeclaratorChunk::Function: {
@@ -7358,137 +7405,6 @@
}
}
-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) {
@@ -7718,8 +7634,6 @@
if (!state.getSema().getLangOpts().OpenCL ||
type.getAddressSpace() != LangAS::Default)
return;
-
- deduceOpenCLImplicitAddrSpace(state, type, TAL);
}
void Sema::completeExprArrayBound(Expr *E) {
Index: clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
===================================================================
--- clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
+++ clang/lib/Sema/SemaTemplateInstantiateDecl.cpp
@@ -837,6 +837,9 @@
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;
Index: clang/lib/Sema/SemaTemplateInstantiate.cpp
===================================================================
--- clang/lib/Sema/SemaTemplateInstantiate.cpp
+++ clang/lib/Sema/SemaTemplateInstantiate.cpp
@@ -1461,8 +1461,12 @@
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
Index: clang/lib/Sema/SemaDecl.cpp
===================================================================
--- clang/lib/Sema/SemaDecl.cpp
+++ clang/lib/Sema/SemaDecl.cpp
@@ -5990,6 +5990,21 @@
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 (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.
@@ -6371,95 +6386,6 @@
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());
@@ -6792,6 +6718,100 @@
}
}
+ if (getLangOpts().OpenCL) {
+
+ deduceOpenCLAddressSpace(NewVD);
+
+ // 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 (NewVD->getType().getAddressSpace() == LangAS::opencl_local ||
+ NewVD->getType().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() &&
+ !(NewVD->getType().getAddressSpace() == LangAS::opencl_constant ||
+ NewVD->getType().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 (NewVD->getType().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;
+ }
+ }
+
// Handle attributes prior to checking for duplicates in MergeVarDecl
ProcessDeclAttributes(S, NewVD, D);
@@ -11078,6 +11098,9 @@
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()) {
@@ -12611,6 +12634,10 @@
if (New->hasAttr<BlocksAttr>()) {
Diag(New->getLocation(), diag::err_block_on_nonlocal);
}
+
+ if (getLangOpts().OpenCL)
+ deduceOpenCLAddressSpace(New);
+
return New;
}
Index: clang/lib/AST/Expr.cpp
===================================================================
--- clang/lib/AST/Expr.cpp
+++ clang/lib/AST/Expr.cpp
@@ -1805,7 +1805,7 @@
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();
}
Index: clang/include/clang/Sema/Sema.h
===================================================================
--- clang/include/clang/Sema/Sema.h
+++ clang/include/clang/Sema/Sema.h
@@ -8505,6 +8505,8 @@
bool CheckARCMethodDecl(ObjCMethodDecl *method);
bool inferObjCARCLifetime(ValueDecl *decl);
+ void deduceOpenCLAddressSpace(ValueDecl *decl);
+
ExprResult
HandleExprPropertyRefExpr(const ObjCObjectPointerType *OPT,
Expr *BaseExpr,
Index: clang/include/clang/AST/Type.h
===================================================================
--- clang/include/clang/AST/Type.h
+++ clang/include/clang/AST/Type.h
@@ -2044,6 +2044,8 @@
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;
@@ -6475,6 +6477,10 @@
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();
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
