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(CanonicalType); } +inline bool Type::isUndeducedAutoType() const { + return isa(CanonicalType); +} + inline bool Type::isObjCQualifiedIdType() const { if (const auto *OPT = getAs()) return OPT->isObjCQualifiedIdType(); 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/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/lib/Sema/SemaDecl.cpp =================================================================== --- clang/lib/Sema/SemaDecl.cpp +++ clang/lib/Sema/SemaDecl.cpp @@ -5990,6 +5990,22 @@ return false; } +void Sema::deduceOpenCLAddressSpace(ValueDecl *Decl) { + if (Decl->getType().getQualifiers().hasAddressSpace()) + return; + if (VarDecl *Var = dyn_cast(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. @@ -6348,6 +6364,105 @@ 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, @@ -6371,95 +6486,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 +6818,13 @@ } } + 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); @@ -11078,6 +11111,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 +12647,10 @@ if (New->hasAttr()) { Diag(New->getLocation(), diag::err_block_on_nonlocal); } + + if (getLangOpts().OpenCL) + deduceOpenCLAddressSpace(New); + return New; } Index: clang/lib/Sema/SemaTemplateInstantiate.cpp =================================================================== --- clang/lib/Sema/SemaTemplateInstantiate.cpp +++ clang/lib/Sema/SemaTemplateInstantiate.cpp @@ -1461,8 +1461,12 @@ int indexAdjustment, Optional 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/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/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() && + !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); } @@ -7358,137 +7380,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 = - // 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 +7609,6 @@ if (!state.getSema().getLangOpts().OpenCL || type.getAddressSpace() != LangAS::Default) return; - - deduceOpenCLImplicitAddrSpace(state, type, TAL); } void Sema::completeExprArrayBound(Expr *E) { 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 QualType TreeTransform::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()) { // 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/test/SemaOpenCL/event_t.cl =================================================================== --- clang/test/SemaOpenCL/event_t.cl +++ 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 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/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/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 '__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}} } 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/SemaOpenCLCXX/address-space-deduction.cl =================================================================== --- clang/test/SemaOpenCLCXX/address-space-deduction.cl +++ clang/test/SemaOpenCLCXX/address-space-deduction.cl @@ -65,30 +65,42 @@ x3::x3(const x3 &t) {} template -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() { 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/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}} 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}} }