r315668 - [OpenCL] Add LangAS::opencl_private to represent private address space in AST
Yaxun Liu via cfe-commits
cfe-commits at lists.llvm.org
Thu Oct 12 20:37:48 PDT 2017
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;
More information about the cfe-commits
mailing list