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