[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


git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@315668 91177308-0d34-0410-b5e6-96231b3b80d8

include/clang/Basic/AddressSpaces.h

@@ -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.

lib/AST/ASTContext.cpp

@@ -707,6 +707,7 @@ static const LangAS::Map *getAddressSpaceMap(const TargetInfo &T,
       1, // opencl_global
       3, // opencl_local
       2, // opencl_constant
+      0, // opencl_private
       4, // opencl_generic
       5, // cuda_device
       6, // cuda_constant

lib/AST/Expr.cpp

@@ -3293,20 +3293,20 @@ Expr::isNullPointerConstant(ASTContext &Ctx,
       // 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;
-        }
-
-        if (IsASValid && !Q.hasQualifiers() &&
-            Pointee->isVoidType() &&                      // to void*
-            CE->getSubExpr()->getType()->isIntegerType()) // from int.
+        // 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 (PointeeHasDefaultAS && Pointee->isVoidType() && // to void*
+            CE->getSubExpr()->getType()->isIntegerType())   // from int.
           return CE->getSubExpr()->isNullPointerConstant(Ctx, NPC);
       }
     }

lib/AST/ItaniumMangle.cpp

@@ -2227,15 +2227,17 @@ void CXXNameMangler::mangleQualifiers(Qualifiers Quals, const DependentAddressSp
     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(Qualifiers Quals, const DependentAddressSp
       case LangAS::cuda_shared:     ASString = "CUshared";   break;
       }
     }
-    mangleVendorQualifier(ASString);
+    if (!ASString.empty())
+      mangleVendorQualifier(ASString);
   }
 
   // The ARC ownership qualifiers start with underscores.

lib/AST/TypePrinter.cpp

@@ -1677,16 +1677,19 @@ void Qualifiers::print(raw_ostream &OS, const PrintingPolicy& Policy,
     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, const PrintingPolicy& Policy,
         OS << "__attribute__((address_space(";
         OS << addrspace - LangAS::FirstTargetAddressSpace;
         OS << ")))";
+      }
     }
   }
   if (Qualifiers::GC gc = getObjCGCAttr()) {

lib/Basic/Targets/AMDGPU.cpp

@@ -47,6 +47,7 @@ static const LangAS::Map AMDGPUPrivIsZeroDefIsGenMap = {
     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 AMDGPUGenIsZeroDefIsGenMap = {
     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 AMDGPUPrivIsZeroDefIsPrivMap = {
     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 AMDGPUGenIsZeroDefIsPrivMap = {
     1, // opencl_global
     3, // opencl_local
     2, // opencl_constant
+    5, // opencl_private
     0, // opencl_generic
     1, // cuda_device
     2, // cuda_constant

lib/Basic/Targets/NVPTX.h

@@ -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

lib/Basic/Targets/SPIR.h

@@ -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

lib/Basic/Targets/TCE.h

@@ -35,6 +35,7 @@ static const unsigned TCEOpenCLAddrSpaceMap[] = {
     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

lib/CodeGen/CGDecl.cpp

@@ -956,7 +956,9 @@ void CodeGenFunction::EmitLifetimeEnd(llvm::Value *Size, llvm::Value *Addr) {
 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);
 

lib/Sema/SemaChecking.cpp

@@ -340,7 +340,7 @@ static bool SemaOpenCLBuiltinNDRangeAndBlock(Sema &S, CallExpr *TheCall) {
 
   // 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 &S, unsigned BuiltinID,
   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)));

lib/Sema/SemaDecl.cpp

@@ -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(VarDecl *NewVD) {
             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 getOpenCLKernelParameterType(Sema &S, QualType PT) {
     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, Declarator &D, DeclContext *DC,
     // 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(DeclContext *DC, SourceLocation StartLoc,
   // 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;

lib/Sema/SemaType.cpp

@@ -4938,7 +4938,6 @@ TypeSourceInfo *Sema::GetTypeForDeclarator(Declarator &D, Scope *S) {
 
   TypeSourceInfo *ReturnTypeInfo = nullptr;
   QualType T = GetDeclSpecTypeForDeclarator(state, ReturnTypeInfo);
-
   if (D.isPrototypeContext() && getLangOpts().ObjCAutoRefCount)
     inferARCWriteback(state, T);
 
@@ -5752,9 +5751,10 @@ static void HandleAddressSpaceTypeAttribute(QualType &Type,
       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(QualType &CurType, const AttributeList &Attr,
   }
 }
 
+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(TypeProcessingState &state, QualType &type,
     }
   }
 
-  // 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) {

test/CodeGenOpenCL/address-spaces-mangling.cl

@@ -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

test/CodeGenOpenCL/address-spaces.cl

@@ -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

test/SemaOpenCL/address-spaces.cl

@@ -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, global int* g2, local int* l, local int* l
   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}}
+}

test/SemaOpenCL/cl20-device-side-enqueue.cl

@@ -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}}

test/SemaOpenCL/extern.cl

@@ -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;
-}

test/SemaOpenCL/storageclass-cl20.cl

@@ -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}}
 }

test/SemaOpenCL/storageclass.cl

@@ -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}}
   }
 
+  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 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}}
+  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}}
 }

test/SemaTemplate/address_space-dependent.cpp

@@ -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;