Ah, I saw you already sent a fix in r315678. Thank you for the fix! No revert! :)
On Fri, Oct 13, 2017 at 4:10 PM Eric Liu <ioe...@google.com> wrote: > 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
