Hi Yaxun, this is causing failures in ppc build bots ( http://lab.llvm.org:8011/builders/clang-ppc64be-linux-multistage/builds/5486). I'll revert the patch for now. Please take a look. Thanks!
On Fri, Oct 13, 2017 at 5:37 AM Yaxun Liu via cfe-commits < cfe-commits@lists.llvm.org> wrote: > Author: yaxunl > Date: Thu Oct 12 20:37:48 2017 > New Revision: 315668 > > URL: http://llvm.org/viewvc/llvm-project?rev=315668&view=rev > Log: > [OpenCL] Add LangAS::opencl_private to represent private address space in > AST > > Currently Clang uses default address space (0) to represent private > address space for OpenCL > in AST. There are two issues with this: > > Multiple address spaces including private address space cannot be > diagnosed. > There is no mangling for default address space. For example, if private > int* is emitted as > i32 addrspace(5)* in IR. It is supposed to be mangled as PUAS5i but it is > mangled as > Pi instead. > > This patch attempts to represent OpenCL private address space explicitly > in AST. It adds > a new enum LangAS::opencl_private and adds it to the variable types which > are implicitly > private: > > automatic variables without address space qualifier > > function parameter > > pointee type without address space qualifier (OpenCL 1.2 and below) > > Differential Revision: https://reviews.llvm.org/D35082 > > Removed: > cfe/trunk/test/SemaOpenCL/extern.cl > Modified: > cfe/trunk/include/clang/Basic/AddressSpaces.h > cfe/trunk/lib/AST/ASTContext.cpp > cfe/trunk/lib/AST/Expr.cpp > cfe/trunk/lib/AST/ItaniumMangle.cpp > cfe/trunk/lib/AST/TypePrinter.cpp > cfe/trunk/lib/Basic/Targets/AMDGPU.cpp > cfe/trunk/lib/Basic/Targets/NVPTX.h > cfe/trunk/lib/Basic/Targets/SPIR.h > cfe/trunk/lib/Basic/Targets/TCE.h > cfe/trunk/lib/CodeGen/CGDecl.cpp > cfe/trunk/lib/Sema/SemaChecking.cpp > cfe/trunk/lib/Sema/SemaDecl.cpp > cfe/trunk/lib/Sema/SemaType.cpp > cfe/trunk/test/CodeGenOpenCL/address-spaces-mangling.cl > cfe/trunk/test/CodeGenOpenCL/address-spaces.cl > cfe/trunk/test/SemaOpenCL/address-spaces.cl > cfe/trunk/test/SemaOpenCL/cl20-device-side-enqueue.cl > cfe/trunk/test/SemaOpenCL/storageclass-cl20.cl > cfe/trunk/test/SemaOpenCL/storageclass.cl > cfe/trunk/test/SemaTemplate/address_space-dependent.cpp > > Modified: cfe/trunk/include/clang/Basic/AddressSpaces.h > URL: > http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/AddressSpaces.h?rev=315668&r1=315667&r2=315668&view=diff > > ============================================================================== > --- cfe/trunk/include/clang/Basic/AddressSpaces.h (original) > +++ cfe/trunk/include/clang/Basic/AddressSpaces.h Thu Oct 12 20:37:48 2017 > @@ -25,16 +25,17 @@ namespace LangAS { > /// > enum ID { > // The default value 0 is the value used in QualType for the the > situation > - // where there is no address space qualifier. For most languages, this > also > - // corresponds to the situation where there is no address space > qualifier in > - // the source code, except for OpenCL, where the address space value 0 > in > - // QualType represents private address space in OpenCL source code. > + // where there is no address space qualifier. > Default = 0, > > // OpenCL specific address spaces. > + // In OpenCL each l-value must have certain non-default address space, > each > + // r-value must have no address space (i.e. the default address space). > The > + // pointee of a pointer must have non-default address space. > opencl_global, > opencl_local, > opencl_constant, > + opencl_private, > opencl_generic, > > // CUDA specific address spaces. > > Modified: cfe/trunk/lib/AST/ASTContext.cpp > URL: > http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/ASTContext.cpp?rev=315668&r1=315667&r2=315668&view=diff > > ============================================================================== > --- cfe/trunk/lib/AST/ASTContext.cpp (original) > +++ cfe/trunk/lib/AST/ASTContext.cpp Thu Oct 12 20:37:48 2017 > @@ -707,6 +707,7 @@ static const LangAS::Map *getAddressSpac > 1, // opencl_global > 3, // opencl_local > 2, // opencl_constant > + 0, // opencl_private > 4, // opencl_generic > 5, // cuda_device > 6, // cuda_constant > > Modified: cfe/trunk/lib/AST/Expr.cpp > URL: > http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/Expr.cpp?rev=315668&r1=315667&r2=315668&view=diff > > ============================================================================== > --- cfe/trunk/lib/AST/Expr.cpp (original) > +++ cfe/trunk/lib/AST/Expr.cpp Thu Oct 12 20:37:48 2017 > @@ -3293,20 +3293,20 @@ Expr::isNullPointerConstant(ASTContext & > // Check that it is a cast to void*. > if (const PointerType *PT = CE->getType()->getAs<PointerType>()) { > QualType Pointee = PT->getPointeeType(); > - Qualifiers Q = Pointee.getQualifiers(); > - // In OpenCL v2.0 generic address space acts as a placeholder > - // and should be ignored. > - bool IsASValid = true; > - if (Ctx.getLangOpts().OpenCLVersion >= 200) { > - if (Pointee.getAddressSpace() == LangAS::opencl_generic) > - Q.removeAddressSpace(); > - else > - IsASValid = false; > - } > + // Only (void*)0 or equivalent are treated as nullptr. If pointee > type > + // has non-default address space it is not treated as nullptr. > + // (__generic void*)0 in OpenCL 2.0 should not be treated as > nullptr > + // since it cannot be assigned to a pointer to constant address > space. > + bool PointeeHasDefaultAS = > + Pointee.getAddressSpace() == LangAS::Default || > + (Ctx.getLangOpts().OpenCLVersion >= 200 && > + Pointee.getAddressSpace() == LangAS::opencl_generic) || > + (Ctx.getLangOpts().OpenCL && > + Ctx.getLangOpts().OpenCLVersion < 200 && > + Pointee.getAddressSpace() == LangAS::opencl_private); > > - if (IsASValid && !Q.hasQualifiers() && > - Pointee->isVoidType() && // to void* > - CE->getSubExpr()->getType()->isIntegerType()) // from int. > + if (PointeeHasDefaultAS && Pointee->isVoidType() && // to void* > + CE->getSubExpr()->getType()->isIntegerType()) // from int. > return CE->getSubExpr()->isNullPointerConstant(Ctx, NPC); > } > } > > Modified: cfe/trunk/lib/AST/ItaniumMangle.cpp > URL: > http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/ItaniumMangle.cpp?rev=315668&r1=315667&r2=315668&view=diff > > ============================================================================== > --- cfe/trunk/lib/AST/ItaniumMangle.cpp (original) > +++ cfe/trunk/lib/AST/ItaniumMangle.cpp Thu Oct 12 20:37:48 2017 > @@ -2227,15 +2227,17 @@ void CXXNameMangler::mangleQualifiers(Qu > if (Context.getASTContext().addressSpaceMapManglingFor(AS)) { > // <target-addrspace> ::= "AS" <address-space-number> > unsigned TargetAS = > Context.getASTContext().getTargetAddressSpace(AS); > - ASString = "AS" + llvm::utostr(TargetAS); > + if (TargetAS != 0) > + ASString = "AS" + llvm::utostr(TargetAS); > } else { > switch (AS) { > default: llvm_unreachable("Not a language specific address space"); > - // <OpenCL-addrspace> ::= "CL" [ "global" | "local" | "constant | > - // "generic" ] > + // <OpenCL-addrspace> ::= "CL" [ "global" | "local" | "constant" | > + // "private"| "generic" ] > case LangAS::opencl_global: ASString = "CLglobal"; break; > case LangAS::opencl_local: ASString = "CLlocal"; break; > case LangAS::opencl_constant: ASString = "CLconstant"; break; > + case LangAS::opencl_private: ASString = "CLprivate"; break; > case LangAS::opencl_generic: ASString = "CLgeneric"; break; > // <CUDA-addrspace> ::= "CU" [ "device" | "constant" | "shared" ] > case LangAS::cuda_device: ASString = "CUdevice"; break; > @@ -2243,7 +2245,8 @@ void CXXNameMangler::mangleQualifiers(Qu > case LangAS::cuda_shared: ASString = "CUshared"; break; > } > } > - mangleVendorQualifier(ASString); > + if (!ASString.empty()) > + mangleVendorQualifier(ASString); > } > > // The ARC ownership qualifiers start with underscores. > > Modified: cfe/trunk/lib/AST/TypePrinter.cpp > URL: > http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/TypePrinter.cpp?rev=315668&r1=315667&r2=315668&view=diff > > ============================================================================== > --- cfe/trunk/lib/AST/TypePrinter.cpp (original) > +++ cfe/trunk/lib/AST/TypePrinter.cpp Thu Oct 12 20:37:48 2017 > @@ -1677,16 +1677,19 @@ void Qualifiers::print(raw_ostream &OS, > addSpace = true; > } > if (unsigned addrspace = getAddressSpace()) { > - if (addSpace) > - OS << ' '; > - addSpace = true; > - switch (addrspace) { > + if (addrspace != LangAS::opencl_private) { > + if (addSpace) > + OS << ' '; > + addSpace = true; > + switch (addrspace) { > case LangAS::opencl_global: > OS << "__global"; > break; > case LangAS::opencl_local: > OS << "__local"; > break; > + case LangAS::opencl_private: > + break; > case LangAS::opencl_constant: > case LangAS::cuda_constant: > OS << "__constant"; > @@ -1705,6 +1708,7 @@ void Qualifiers::print(raw_ostream &OS, > OS << "__attribute__((address_space("; > OS << addrspace - LangAS::FirstTargetAddressSpace; > OS << ")))"; > + } > } > } > if (Qualifiers::GC gc = getObjCGCAttr()) { > > Modified: cfe/trunk/lib/Basic/Targets/AMDGPU.cpp > URL: > http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/Targets/AMDGPU.cpp?rev=315668&r1=315667&r2=315668&view=diff > > ============================================================================== > --- cfe/trunk/lib/Basic/Targets/AMDGPU.cpp (original) > +++ cfe/trunk/lib/Basic/Targets/AMDGPU.cpp Thu Oct 12 20:37:48 2017 > @@ -47,6 +47,7 @@ static const LangAS::Map AMDGPUPrivIsZer > 1, // opencl_global > 3, // opencl_local > 2, // opencl_constant > + 0, // opencl_private > 4, // opencl_generic > 1, // cuda_device > 2, // cuda_constant > @@ -58,6 +59,7 @@ static const LangAS::Map AMDGPUGenIsZero > 1, // opencl_global > 3, // opencl_local > 2, // opencl_constant > + 5, // opencl_private > 0, // opencl_generic > 1, // cuda_device > 2, // cuda_constant > @@ -69,6 +71,7 @@ static const LangAS::Map AMDGPUPrivIsZer > 1, // opencl_global > 3, // opencl_local > 2, // opencl_constant > + 0, // opencl_private > 4, // opencl_generic > 1, // cuda_device > 2, // cuda_constant > @@ -80,6 +83,7 @@ static const LangAS::Map AMDGPUGenIsZero > 1, // opencl_global > 3, // opencl_local > 2, // opencl_constant > + 5, // opencl_private > 0, // opencl_generic > 1, // cuda_device > 2, // cuda_constant > > Modified: cfe/trunk/lib/Basic/Targets/NVPTX.h > URL: > http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/Targets/NVPTX.h?rev=315668&r1=315667&r2=315668&view=diff > > ============================================================================== > --- cfe/trunk/lib/Basic/Targets/NVPTX.h (original) > +++ cfe/trunk/lib/Basic/Targets/NVPTX.h Thu Oct 12 20:37:48 2017 > @@ -28,6 +28,7 @@ static const unsigned NVPTXAddrSpaceMap[ > 1, // opencl_global > 3, // opencl_local > 4, // opencl_constant > + 0, // opencl_private > // FIXME: generic has to be added to the target > 0, // opencl_generic > 1, // cuda_device > > Modified: cfe/trunk/lib/Basic/Targets/SPIR.h > URL: > http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/Targets/SPIR.h?rev=315668&r1=315667&r2=315668&view=diff > > ============================================================================== > --- cfe/trunk/lib/Basic/Targets/SPIR.h (original) > +++ cfe/trunk/lib/Basic/Targets/SPIR.h Thu Oct 12 20:37:48 2017 > @@ -27,6 +27,7 @@ static const unsigned SPIRAddrSpaceMap[] > 1, // opencl_global > 3, // opencl_local > 2, // opencl_constant > + 0, // opencl_private > 4, // opencl_generic > 0, // cuda_device > 0, // cuda_constant > > Modified: cfe/trunk/lib/Basic/Targets/TCE.h > URL: > http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/Targets/TCE.h?rev=315668&r1=315667&r2=315668&view=diff > > ============================================================================== > --- cfe/trunk/lib/Basic/Targets/TCE.h (original) > +++ cfe/trunk/lib/Basic/Targets/TCE.h Thu Oct 12 20:37:48 2017 > @@ -35,6 +35,7 @@ static const unsigned TCEOpenCLAddrSpace > 3, // opencl_global > 4, // opencl_local > 5, // opencl_constant > + 0, // opencl_private > // FIXME: generic has to be added to the target > 0, // opencl_generic > 0, // cuda_device > > Modified: cfe/trunk/lib/CodeGen/CGDecl.cpp > URL: > http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGDecl.cpp?rev=315668&r1=315667&r2=315668&view=diff > > ============================================================================== > --- cfe/trunk/lib/CodeGen/CGDecl.cpp (original) > +++ cfe/trunk/lib/CodeGen/CGDecl.cpp Thu Oct 12 20:37:48 2017 > @@ -956,7 +956,9 @@ void CodeGenFunction::EmitLifetimeEnd(ll > CodeGenFunction::AutoVarEmission > CodeGenFunction::EmitAutoVarAlloca(const VarDecl &D) { > QualType Ty = D.getType(); > - assert(Ty.getAddressSpace() == LangAS::Default); > + assert( > + Ty.getAddressSpace() == LangAS::Default || > + (Ty.getAddressSpace() == LangAS::opencl_private && > getLangOpts().OpenCL)); > > AutoVarEmission emission(D); > > > Modified: cfe/trunk/lib/Sema/SemaChecking.cpp > URL: > http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaChecking.cpp?rev=315668&r1=315667&r2=315668&view=diff > > ============================================================================== > --- cfe/trunk/lib/Sema/SemaChecking.cpp (original) > +++ cfe/trunk/lib/Sema/SemaChecking.cpp Thu Oct 12 20:37:48 2017 > @@ -340,7 +340,7 @@ static bool SemaOpenCLBuiltinNDRangeAndB > > // First argument is an ndrange_t type. > Expr *NDRangeArg = TheCall->getArg(0); > - if (NDRangeArg->getType().getAsString() != "ndrange_t") { > + if (NDRangeArg->getType().getUnqualifiedType().getAsString() != > "ndrange_t") { > S.Diag(NDRangeArg->getLocStart(), > diag::err_opencl_builtin_expected_type) > << TheCall->getDirectCallee() << "'ndrange_t'"; > @@ -784,8 +784,11 @@ static bool SemaOpenCLBuiltinToAddr(Sema > case Builtin::BIto_local: > Qual.setAddressSpace(LangAS::opencl_local); > break; > + case Builtin::BIto_private: > + Qual.setAddressSpace(LangAS::opencl_private); > + break; > default: > - Qual.removeAddressSpace(); > + llvm_unreachable("Invalid builtin function"); > } > Call->setType(S.Context.getPointerType(S.Context.getQualifiedType( > RT.getUnqualifiedType(), Qual))); > > Modified: cfe/trunk/lib/Sema/SemaDecl.cpp > URL: > http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDecl.cpp?rev=315668&r1=315667&r2=315668&view=diff > > ============================================================================== > --- cfe/trunk/lib/Sema/SemaDecl.cpp (original) > +++ cfe/trunk/lib/Sema/SemaDecl.cpp Thu Oct 12 20:37:48 2017 > @@ -6324,7 +6324,7 @@ NamedDecl *Sema::ActOnVariableDeclarator > // The event type cannot be used with the __local, __constant and > __global > // address space qualifiers. > if (R->isEventT()) { > - if (R.getAddressSpace()) { > + if (R.getAddressSpace() != LangAS::opencl_private) { > Diag(D.getLocStart(), diag::err_event_t_addr_space_qual); > D.setInvalidType(); > } > @@ -7427,7 +7427,7 @@ void Sema::CheckVariableDeclarationType( > return; > } > } > - } else if (T.getAddressSpace() != LangAS::Default) { > + } else if (T.getAddressSpace() != LangAS::opencl_private) { > // Do not allow other address spaces on automatic variable. > Diag(NewVD->getLocation(), diag::err_as_qualified_auto_decl) << 1; > NewVD->setInvalidDecl(); > @@ -8062,7 +8062,8 @@ static OpenCLParamType getOpenCLKernelPa > if (PointeeType->isPointerType()) > return PtrPtrKernelParam; > if (PointeeType.getAddressSpace() == LangAS::opencl_generic || > - PointeeType.getAddressSpace() == 0) > + PointeeType.getAddressSpace() == LangAS::opencl_private || > + PointeeType.getAddressSpace() == LangAS::Default) > return InvalidAddrSpacePtrKernelParam; > return PtrKernelParam; > } > @@ -8832,9 +8833,7 @@ Sema::ActOnFunctionDeclarator(Scope *S, > // OpenCL v1.1 s6.5: Using an address space qualifier in a function > return > // type declaration will generate a compilation error. > unsigned AddressSpace = NewFD->getReturnType().getAddressSpace(); > - if (AddressSpace == LangAS::opencl_local || > - AddressSpace == LangAS::opencl_global || > - AddressSpace == LangAS::opencl_constant) { > + if (AddressSpace != LangAS::Default) { > Diag(NewFD->getLocation(), > diag::err_opencl_return_value_with_address_space); > NewFD->setInvalidDecl(); > @@ -11939,13 +11938,13 @@ ParmVarDecl *Sema::CheckParameter(DeclCo > // duration shall not be qualified by an address-space qualifier." > // Since all parameters have automatic store duration, they can not have > // an address space. > - if (T.getAddressSpace() != 0) { > - // OpenCL allows function arguments declared to be an array of a type > - // to be qualified with an address space. > - if (!(getLangOpts().OpenCL && T->isArrayType())) { > - Diag(NameLoc, diag::err_arg_with_address_space); > - New->setInvalidDecl(); > - } > + if (T.getAddressSpace() != LangAS::Default && > + // OpenCL allows function arguments declared to be an array of a > type > + // to be qualified with an address space. > + !(getLangOpts().OpenCL && > + (T->isArrayType() || T.getAddressSpace() == > LangAS::opencl_private))) { > + Diag(NameLoc, diag::err_arg_with_address_space); > + New->setInvalidDecl(); > } > > return New; > > Modified: cfe/trunk/lib/Sema/SemaType.cpp > URL: > http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaType.cpp?rev=315668&r1=315667&r2=315668&view=diff > > ============================================================================== > --- cfe/trunk/lib/Sema/SemaType.cpp (original) > +++ cfe/trunk/lib/Sema/SemaType.cpp Thu Oct 12 20:37:48 2017 > @@ -4938,7 +4938,6 @@ TypeSourceInfo *Sema::GetTypeForDeclarat > > TypeSourceInfo *ReturnTypeInfo = nullptr; > QualType T = GetDeclSpecTypeForDeclarator(state, ReturnTypeInfo); > - > if (D.isPrototypeContext() && getLangOpts().ObjCAutoRefCount) > inferARCWriteback(state, T); > > @@ -5752,9 +5751,10 @@ static void HandleAddressSpaceTypeAttrib > ASIdx = LangAS::opencl_constant; break; > case AttributeList::AT_OpenCLGenericAddressSpace: > ASIdx = LangAS::opencl_generic; break; > + case AttributeList::AT_OpenCLPrivateAddressSpace: > + ASIdx = LangAS::opencl_private; break; > default: > - assert(Attr.getKind() == > AttributeList::AT_OpenCLPrivateAddressSpace); > - ASIdx = 0; break; > + llvm_unreachable("Invalid address space"); > } > > Type = S.Context.getAddrSpaceQualType(Type, ASIdx); > @@ -6986,6 +6986,92 @@ static void HandleOpenCLAccessAttr(QualT > } > } > > +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 alwasy 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 > + // definining 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::BlockPointer); > + 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. > + (D.getContext() == Declarator::MemberContext && !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)) > + return; > + > + unsigned ImpAddr; > + // 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) { > + 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() == Declarator::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 processTypeAttrs(TypeProcessingState &state, QualType &type, > TypeAttrLocation TAL, AttributeList *attrs) { > // Scan through and apply attributes to this type where it makes > sense. Some > @@ -7157,39 +7243,11 @@ static void processTypeAttrs(TypeProcess > } > } > > - // 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 (state.getSema().getLangOpts().OpenCLVersion >= 200 && > - !hasOpenCLAddressSpace && type.getAddressSpace() == 0 && > - (TAL == TAL_DeclSpec || TAL == TAL_DeclChunk)) { > - Declarator &D = state.getDeclarator(); > - if (state.getCurrentChunkIndex() > 0 && > - (D.getTypeObject(state.getCurrentChunkIndex() - 1).Kind == > - DeclaratorChunk::Pointer || > - D.getTypeObject(state.getCurrentChunkIndex() - 1).Kind == > - DeclaratorChunk::BlockPointer)) { > - type = state.getSema().Context.getAddrSpaceQualType( > - type, LangAS::opencl_generic); > - } else if (state.getCurrentChunkIndex() == 0 && > - D.getContext() == Declarator::FileContext && > - !D.isFunctionDeclarator() && !D.isFunctionDefinition() && > - D.getDeclSpec().getStorageClassSpec() != > DeclSpec::SCS_typedef && > - !type->isSamplerT()) > - type = state.getSema().Context.getAddrSpaceQualType( > - type, LangAS::opencl_global); > - else if (state.getCurrentChunkIndex() == 0 && > - D.getContext() == Declarator::BlockContext && > - D.getDeclSpec().getStorageClassSpec() == > DeclSpec::SCS_static) > - type = state.getSema().Context.getAddrSpaceQualType( > - type, LangAS::opencl_global); > - } > + if (!state.getSema().getLangOpts().OpenCL || > + type.getAddressSpace() != LangAS::Default) > + return; > + > + deduceOpenCLImplicitAddrSpace(state, type, TAL); > } > > void Sema::completeExprArrayBound(Expr *E) { > > Modified: cfe/trunk/test/CodeGenOpenCL/address-spaces-mangling.cl > URL: > http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/address-spaces-mangling.cl?rev=315668&r1=315667&r2=315668&view=diff > > ============================================================================== > --- cfe/trunk/test/CodeGenOpenCL/address-spaces-mangling.cl (original) > +++ cfe/trunk/test/CodeGenOpenCL/address-spaces-mangling.cl Thu Oct 12 > 20:37:48 2017 > @@ -1,5 +1,7 @@ > -// RUN: %clang_cc1 %s -ffake-address-space-map > -faddress-space-map-mangling=yes -triple %itanium_abi_triple -emit-llvm -o > - | FileCheck -check-prefix=ASMANG %s > -// RUN: %clang_cc1 %s -ffake-address-space-map > -faddress-space-map-mangling=no -triple %itanium_abi_triple -emit-llvm -o - > | FileCheck -check-prefix=NOASMANG %s > +// RUN: %clang_cc1 %s -ffake-address-space-map > -faddress-space-map-mangling=yes -triple %itanium_abi_triple -emit-llvm -o > - | FileCheck -check-prefixes=ASMANG,ASMAN10 %s > +// RUN: %clang_cc1 %s -cl-std=CL2.0 -ffake-address-space-map > -faddress-space-map-mangling=yes -triple %itanium_abi_triple -emit-llvm -o > - | FileCheck -check-prefixes=ASMANG,ASMAN20 %s > +// RUN: %clang_cc1 %s -ffake-address-space-map > -faddress-space-map-mangling=no -triple %itanium_abi_triple -emit-llvm -o - > | FileCheck -check-prefixes=NOASMANG,NOASMAN10 %s > +// RUN: %clang_cc1 %s -cl-std=CL2.0 -ffake-address-space-map > -faddress-space-map-mangling=no -triple %itanium_abi_triple -emit-llvm -o - > | FileCheck -check-prefixes=NOASMANG,NOASMAN20 %s > > // We check that the address spaces are mangled the same in both version > of OpenCL > // RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=CL2.0 > -emit-llvm -o - | FileCheck -check-prefix=OCL-20 %s > @@ -10,15 +12,17 @@ > // warnings, but we do want it for comparison purposes. > __attribute__((overloadable)) > void ff(int *arg) { } > -// ASMANG: @_Z2ffPi > -// NOASMANG: @_Z2ffPi > +// ASMANG10: @_Z2ffPi > +// ASMANG20: @_Z2ffPU3AS4i > +// NOASMANG10: @_Z2ffPi > +// NOASMANG20: @_Z2ffPU9CLgenerici > // OCL-20-DAG: @_Z2ffPU3AS4i > // OCL-12-DAG: @_Z2ffPi > > __attribute__((overloadable)) > void f(private int *arg) { } > // ASMANG: @_Z1fPi > -// NOASMANG: @_Z1fPi > +// NOASMANG: @_Z1fPU9CLprivatei > // OCL-20-DAG: @_Z1fPi > // OCL-12-DAG: @_Z1fPi > > @@ -42,3 +46,11 @@ void f(constant int *arg) { } > // NOASMANG: @_Z1fPU10CLconstanti > // OCL-20-DAG: @_Z1fPU3AS2i > // OCL-12-DAG: @_Z1fPU3AS2i > + > +#if __OPENCL_C_VERSION__ >= 200 > +__attribute__((overloadable)) > +void f(generic int *arg) { } > +// ASMANG20: @_Z1fPU3AS4i > +// NOASMANG20: @_Z1fPU9CLgenerici > +// OCL-20-DAG: @_Z1fPU3AS4i > +#endif > > Modified: cfe/trunk/test/CodeGenOpenCL/address-spaces.cl > URL: > http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenOpenCL/address-spaces.cl?rev=315668&r1=315667&r2=315668&view=diff > > ============================================================================== > --- cfe/trunk/test/CodeGenOpenCL/address-spaces.cl (original) > +++ cfe/trunk/test/CodeGenOpenCL/address-spaces.cl Thu Oct 12 20:37:48 > 2017 > @@ -7,6 +7,24 @@ > // RUN: %clang_cc1 %s -O0 -triple amdgcn-mesa-mesa3d -emit-llvm -o - | > FileCheck --check-prefixes=CHECK,SPIR %s > // RUN: %clang_cc1 %s -O0 -triple r600-- -emit-llvm -o - | FileCheck > --check-prefixes=CHECK,SPIR %s > > +// SPIR: %struct.S = type { i32, i32, i32* } > +// CL20SPIR: %struct.S = type { i32, i32, i32 addrspace(4)* } > +struct S { > + int x; > + int y; > + int *z; > +}; > + > +// CL20-DAG: @g_extern_var = external addrspace(1) global float > +// CL20-DAG: @l_extern_var = external addrspace(1) global float > +// CL20-DAG: @test_static.l_static_var = internal addrspace(1) global > float 0.000000e+00 > +// CL20-DAG: @g_static_var = internal addrspace(1) global float > 0.000000e+00 > + > +#ifdef CL20 > +// CL20-DAG: @g_s = common addrspace(1) global %struct.S zeroinitializer > +struct S g_s; > +#endif > + > // SPIR: i32* %arg > // GIZ: i32 addrspace(5)* %arg > void f__p(__private int *arg) {} > @@ -58,3 +76,52 @@ void f(int *arg) { > // CL20-DAG: @f.ii = internal addrspace(1) global i32 0 > #endif > } > + > +typedef int int_td; > +typedef int *intp_td; > +// SPIR: define void @test_typedef(i32 addrspace(1)* %x, i32 > addrspace(2)* %y, i32* %z) > +void test_typedef(global int_td *x, constant int_td *y, intp_td z) { > + *x = *y; > + *z = 0; > +} > + > +// SPIR: define void @test_struct() > +void test_struct() { > + // SPIR: %ps = alloca %struct.S* > + // CL20SPIR: %ps = alloca %struct.S addrspace(4)* > + struct S *ps; > + // SPIR: store i32 0, i32* %x > + // CL20SPIR: store i32 0, i32 addrspace(4)* %x > + ps->x = 0; > +#ifdef CL20 > + // CL20SPIR: store i32 0, i32 addrspace(1)* getelementptr inbounds > (%struct.S, %struct.S addrspace(1)* @g_s, i32 0, i32 0) > + g_s.x = 0; > +#endif > +} > + > +// SPIR-LABEL: define void @test_void_par() > +void test_void_par(void) {} > + > +// SPIR-LABEL: define i32 @test_func_return_type() > +int test_func_return_type(void) { > + return 0; > +} > + > +#ifdef CL20 > +extern float g_extern_var; > + > +// CL20-LABEL: define {{.*}}void @test_extern( > +kernel void test_extern(global float *buf) { > + extern float l_extern_var; > + buf[0] += g_extern_var + l_extern_var; > +} > + > +static float g_static_var; > + > +// CL20-LABEL: define {{.*}}void @test_static( > +kernel void test_static(global float *buf) { > + static float l_static_var; > + buf[0] += g_static_var + l_static_var; > +} > + > +#endif > > Modified: cfe/trunk/test/SemaOpenCL/address-spaces.cl > URL: > http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCL/address-spaces.cl?rev=315668&r1=315667&r2=315668&view=diff > > ============================================================================== > --- cfe/trunk/test/SemaOpenCL/address-spaces.cl (original) > +++ cfe/trunk/test/SemaOpenCL/address-spaces.cl Thu Oct 12 20:37:48 2017 > @@ -1,4 +1,5 @@ > // RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only > +// RUN: %clang_cc1 %s -cl-std=CL2.0 -verify -pedantic -fsyntax-only > > __constant int ci = 1; > > @@ -7,9 +8,15 @@ __kernel void foo(__global int *gip) { > __local int lj = 2; // expected-error {{'__local' variable cannot have > an initializer}} > > int *ip; > +#if __OPENCL_C_VERSION__ < 200 > ip = gip; // expected-error {{assigning '__global int *' to 'int *' > changes address space of pointer}} > ip = &li; // expected-error {{assigning '__local int *' to 'int *' > changes address space of pointer}} > ip = &ci; // expected-error {{assigning '__constant int *' to 'int *' > changes address space of pointer}} > +#else > + ip = gip; > + ip = &li; > + ip = &ci; // expected-error {{assigning '__constant int *' to > '__generic int *' changes address space of pointer}} > +#endif > } > > void explicit_cast(global int* g, local int* l, constant int* c, private > int* p, const constant int *cc) > @@ -40,3 +47,19 @@ void ok_explicit_casts(global int *g, gl > l = (local int*) l2; > p = (private int*) p2; > } > + > +__private int func_return_priv(void); //expected-error {{return > value cannot be qualified with address space}} > +__global int func_return_global(void); //expected-error {{return > value cannot be qualified with address space}} > +__local int func_return_local(void); //expected-error {{return > value cannot be qualified with address space}} > +__constant int func_return_constant(void); //expected-error {{return > value cannot be qualified with address space}} > +#if __OPENCL_C_VERSION__ >= 200 > +__generic int func_return_generic(void); //expected-error {{return > value cannot be qualified with address space}} > +#endif > + > +void func_multiple_addr(void) { > + typedef __private int private_int_t; > + __local __private int var1; // expected-error {{multiple address > spaces specified for type}} > + __local __private int *var2; // expected-error {{multiple address > spaces specified for type}} > + __local private_int_t var3; // expected-error {{multiple address > spaces specified for type}} > + __local private_int_t *var4; // expected-error {{multiple address > spaces specified for type}} > +} > > Modified: cfe/trunk/test/SemaOpenCL/cl20-device-side-enqueue.cl > URL: > http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCL/cl20-device-side-enqueue.cl?rev=315668&r1=315667&r2=315668&view=diff > > ============================================================================== > --- cfe/trunk/test/SemaOpenCL/cl20-device-side-enqueue.cl (original) > +++ cfe/trunk/test/SemaOpenCL/cl20-device-side-enqueue.cl Thu Oct 12 > 20:37:48 2017 > @@ -222,7 +222,7 @@ kernel void foo(global int *buf) > > kernel void bar(global int *buf) > { > - ndrange_t n; > + __private ndrange_t n; > buf[0] = get_kernel_sub_group_count_for_ndrange(n, ^(){}); > buf[0] = get_kernel_sub_group_count_for_ndrange(0, ^(){}); // > expected-error{{illegal call to 'get_kernel_sub_group_count_for_ndrange', > expected 'ndrange_t' argument type}} > buf[0] = get_kernel_sub_group_count_for_ndrange(n, 1); // > expected-error{{illegal call to 'get_kernel_sub_group_count_for_ndrange', > expected block argument type}} > > Removed: cfe/trunk/test/SemaOpenCL/extern.cl > URL: > http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCL/extern.cl?rev=315667&view=auto > > ============================================================================== > --- cfe/trunk/test/SemaOpenCL/extern.cl (original) > +++ cfe/trunk/test/SemaOpenCL/extern.cl (removed) > @@ -1,9 +0,0 @@ > -// RUN: %clang_cc1 -x cl -cl-opt-disable -cl-std=CL1.2 -emit-llvm > -ffake-address-space-map %s -o - -verify | FileCheck %s > -// expected-no-diagnostics > - > -// CHECK: @foo = external addrspace(2) constant float > -extern constant float foo; > - > -kernel void test(global float* buf) { > - buf[0] += foo; > -} > > Modified: cfe/trunk/test/SemaOpenCL/storageclass-cl20.cl > URL: > http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCL/storageclass-cl20.cl?rev=315668&r1=315667&r2=315668&view=diff > > ============================================================================== > --- cfe/trunk/test/SemaOpenCL/storageclass-cl20.cl (original) > +++ cfe/trunk/test/SemaOpenCL/storageclass-cl20.cl Thu Oct 12 20:37:48 > 2017 > @@ -1,21 +1,41 @@ > // RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only -cl-std=CL2.0 > > -static constant int G1 = 0; > int G2 = 0; > global int G3 = 0; > local int G4 = 0; // expected-error{{program scope variable > must reside in global or constant address space}} > > -void kernel foo() { > - static int S1 = 5; > - static global int S2 = 5; > - static private int S3 = 5; // expected-error{{static local variable > must reside in global or constant address space}} > +static float g_implicit_static_var = 0; > +static constant float g_constant_static_var = 0; > +static global float g_global_static_var = 0; > +static local float g_local_static_var = 0; // expected-error > {{program scope variable must reside in global or constant address space}} > +static private float g_private_static_var = 0; // expected-error > {{program scope variable must reside in global or constant address space}} > +static generic float g_generic_static_var = 0; // expected-error > {{program scope variable must reside in global or constant address space}} > + > +extern float g_implicit_extern_var; > +extern constant float g_constant_extern_var; > +extern global float g_global_extern_var; > +extern local float g_local_extern_var; // expected-error {{extern > variable must reside in global or constant address space}} > +extern private float g_private_extern_var; // expected-error {{extern > variable must reside in global or constant address space}} > +extern generic float g_generic_extern_var; // expected-error {{extern > variable must reside in global or constant address space}} > > +void kernel foo() { > constant int L1 = 0; > local int L2; > global int L3; // expected-error{{function > scope variable cannot be declared in global address space}} > generic int L4; // > expected-error{{automatic variable qualified with an invalid address space}} > __attribute__((address_space(100))) int L5; // > expected-error{{automatic variable qualified with an invalid address space}} > > - extern global int G5; > - extern int G6; // expected-error{{extern variable must reside in global > or constant address space}} > + static float l_implicit_static_var = 0; > + static constant float l_constant_static_var = 0; > + static global float l_global_static_var = 0; > + static local float l_local_static_var = 0; // expected-error > {{static local variable must reside in global or constant address space}} > + static private float l_private_static_var = 0; // expected-error > {{static local variable must reside in global or constant address space}} > + static generic float l_generic_static_var = 0; // expected-error > {{static local variable must reside in global or constant address space}} > + > + extern float l_implicit_extern_var; > + extern constant float l_constant_extern_var; > + extern global float l_global_extern_var; > + extern local float l_local_extern_var; // expected-error {{extern > variable must reside in global or constant address space}} > + extern private float l_private_extern_var; // expected-error {{extern > variable must reside in global or constant address space}} > + extern generic float l_generic_extern_var; // expected-error {{extern > variable must reside in global or constant address space}} > } > > Modified: cfe/trunk/test/SemaOpenCL/storageclass.cl > URL: > http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCL/storageclass.cl?rev=315668&r1=315667&r2=315668&view=diff > > ============================================================================== > --- cfe/trunk/test/SemaOpenCL/storageclass.cl (original) > +++ cfe/trunk/test/SemaOpenCL/storageclass.cl Thu Oct 12 20:37:48 2017 > @@ -5,6 +5,20 @@ constant int G2 = 0; > int G3 = 0; // expected-error{{program scope variable must reside > in constant address space}} > global int G4 = 0; // expected-error{{program scope variable must reside > in constant address space}} > > +static float g_implicit_static_var = 0; // expected-error {{program scope > variable must reside in constant address space}} > +static constant float g_constant_static_var = 0; > +static global float g_global_static_var = 0; // expected-error > {{program scope variable must reside in constant address space}} > +static local float g_local_static_var = 0; // expected-error > {{program scope variable must reside in constant address space}} > +static private float g_private_static_var = 0; // expected-error > {{program scope variable must reside in constant address space}} > +static generic float g_generic_static_var = 0; // expected-error{{OpenCL > version 1.2 does not support the 'generic' type qualifier}} // > expected-error {{program scope variable must reside in constant address > space}} > + > +extern float g_implicit_extern_var; // expected-error {{extern variable > must reside in constant address space}} > +extern constant float g_constant_extern_var; > +extern global float g_global_extern_var; // expected-error {{extern > variable must reside in constant address space}} > +extern local float g_local_extern_var; // expected-error {{extern > variable must reside in constant address space}} > +extern private float g_private_extern_var; // expected-error {{extern > variable must reside in constant address space}} > +extern generic float g_generic_extern_var; // expected-error{{OpenCL > version 1.2 does not support the 'generic' type qualifier}} // > expected-error {{extern variable must reside in constant address space}} > + > void kernel foo(int x) { > // static is not allowed at local scope before CL2.0 > static int S1 = 5; // expected-error{{variables in function > scope cannot be declared static}} > @@ -45,10 +59,17 @@ void f() { > __attribute__((address_space(100))) int L4; // > expected-error{{automatic variable qualified with an invalid address space}} > } > > - > - extern constant float L5; > - extern local float L6; // expected-error{{extern variable must reside > in constant address space}} > - > - static int L7 = 0; // expected-error{{variables in function scope > cannot be declared static}} > - static int L8; // expected-error{{variables in function scope > cannot be declared static}} > + static float l_implicit_static_var = 0; // expected-error > {{variables in function scope cannot be declared static}} > + static constant float l_constant_static_var = 0; // expected-error > {{variables in function scope cannot be declared static}} > + static global float l_global_static_var = 0; // expected-error > {{variables in function scope cannot be declared static}} > + static local float l_local_static_var = 0; // expected-error > {{variables in function scope cannot be declared static}} > + static private float l_private_static_var = 0; // expected-error > {{variables in function scope cannot be declared static}} > + static generic float l_generic_static_var = 0; // > expected-error{{OpenCL version 1.2 does not support the 'generic' type > qualifier}} // expected-error {{variables in function scope cannot be > declared static}} > + > + extern float l_implicit_extern_var; // expected-error {{extern variable > must reside in constant address space}} > + extern constant float l_constant_extern_var; > + extern global float l_global_extern_var; // expected-error {{extern > variable must reside in constant address space}} > + extern local float l_local_extern_var; // expected-error {{extern > variable must reside in constant address space}} > + extern private float l_private_extern_var; // expected-error {{extern > variable must reside in constant address space}} > + extern generic float l_generic_extern_var; // expected-error{{OpenCL > version 1.2 does not support the 'generic' type qualifier}} // > expected-error {{extern variable must reside in constant address space}} > } > > Modified: cfe/trunk/test/SemaTemplate/address_space-dependent.cpp > URL: > http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaTemplate/address_space-dependent.cpp?rev=315668&r1=315667&r2=315668&view=diff > > ============================================================================== > --- cfe/trunk/test/SemaTemplate/address_space-dependent.cpp (original) > +++ cfe/trunk/test/SemaTemplate/address_space-dependent.cpp Thu Oct 12 > 20:37:48 2017 > @@ -43,7 +43,7 @@ void neg() { > > template <long int I> > void tooBig() { > - __attribute__((address_space(I))) int *bounds; // expected-error > {{address space is larger than the maximum supported (8388599)}} > + __attribute__((address_space(I))) int *bounds; // expected-error > {{address space is larger than the maximum supported (8388598)}} > } > > template <long int I> > @@ -101,7 +101,7 @@ int main() { > car<1, 2, 3>(); // expected-note {{in instantiation of function > template specialization 'car<1, 2, 3>' requested here}} > HasASTemplateFields<1> HASTF; > neg<-1>(); // expected-note {{in instantiation of function template > specialization 'neg<-1>' requested here}} > - correct<0x7FFFF7>(); > + correct<0x7FFFF6>(); > tooBig<8388650>(); // expected-note {{in instantiation of function > template specialization 'tooBig<8388650>' requested here}} > > __attribute__((address_space(1))) char *x; > > > _______________________________________________ > cfe-commits mailing list > cfe-commits@lists.llvm.org > http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits >
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
