r315668 - [OpenCL] Add LangAS::opencl_private to represent private address space in AST
Eric Liu via cfe-commits
cfe-commits at lists.llvm.org
Fri Oct 13 07:10:27 PDT 2017
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 at 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 at lists.llvm.org
> http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
>
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20171013/81cb3c04/attachment-0001.html>
More information about the cfe-commits
mailing list