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:12:06 PDT 2017
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 <ioeric at 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 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/abc71be9/attachment-0001.html>
More information about the cfe-commits
mailing list