[PR42707][OpenCL] Fix addr space deduction for auto

[Anastasia]

Summary

A regression was introduced in D64400 because auto is using the same logic as templates. However, deduction of addr spaces wasn't working correctly even before.

Here are the rules that are implemented in this patch:

• For non-reference and non-pointer types deduction can be done early because addr space is not going to be taken from init expr
• For ref or ptr auto types we should prevent deducing addr space before the deduction of the whole type (incl its pointee addr spaces) from init expr

[rjmccall]

Isn't the general rule for template argument deduction (which this devolves to) just to ignore top-level qualifiers? And then you can substitute in the substituted type and end up with a properly qualified type for the parameter / variable, and you can add extra qualifiers as necessary. Why are special rules for pointers and references required?

[Anastasia]

In D65744#1626563, @rjmccall wrote:

Isn't the general rule for template argument deduction (which this devolves to) just to ignore top-level qualifiers? And then you can substitute in the substituted type and end up with a properly qualified type for the parameter / variable, and you can add extra qualifiers as necessary. Why are special rules for pointers and references required?

The issue I am trying to solve is that the addr space qualifier for the pointee type in pointers and references is supposed to default to generic if none were provided either in the parameter or argument of template/auto type. But for all regular types the deduction happens early during parsing. So I need to prevent the early deduction and make sure the deduction happens once the type is being known (i.e. deduced).

[rjmccall]

I see. Is the deduction rule recursive or something, where a pointer to pointer is inferred to point to the same address space as the pointee? I still don't understand why pointers and references are handled differently here, instead of having the rule be "don't infer if the type is opaquely dependent or an auto placeholder".

[Anastasia]

In D65744#1627589, @rjmccall wrote:

I see. Is the deduction rule recursive or something, where a pointer to pointer is inferred to point to the same address space as the pointee?

It is recursive indeed and we currently deduce all pointees to generic AS.

I still don't understand why pointers and references are handled differently here, instead of having the rule be "don't infer if the type is opaquely dependent or an auto placeholder".

Because we currently only infer pointers and references in TreeTransforms. Everything else is inferred earlier. That is mainly because we need some context info to infer in most of other cases and therefore it's not easy to do in TreeTransforms. There is some more information on this review: https://reviews.llvm.org/D64400 where the inference for pointers and references was added originally to TreeTransforms.

[mantognini]

I think this looks good. Maybe the tests should be extended to test auto as function return type, and if there's some special handling around decltype(auto), then it should be tested too, but I'm not sure it's actually needed here. What do you think?

lib/Sema/SemaType.cpp

7441 ↗

(On Diff #213345)

Shouldn't the parentheses around IsAutoPointee be removed for style consistency?

7441 ↗

(On Diff #213345)

With the if statement introduced above, IsAutoPointee can be true only in C++ mode. Could it be an issue to not guard (T->isAutoType() && IsPointee) for non-C++ mode? (I guess not, but I couldn't convince myself.)

lib/Sema/TreeTransform.h

4550 ↗

(On Diff #213345)

Nitpicking: there are two spaces between } and while.

[rjmccall]

In D65744#1629055, @Anastasia wrote:

In D65744#1627589, @rjmccall wrote:

I see. Is the deduction rule recursive or something, where a pointer to pointer is inferred to point to the same address space as the pointee?

It is recursive indeed and we currently deduce all pointees to generic AS.

Is that likely to change? If all pointees are inferred to be generic (in the absence of an explicit qualifier) when building a pointer type, this seems very over-thought vs. just adding the generic AS whenever you build a pointer to an unqualified but non-dependent type.

[Anastasia]

missing test case

1 auto* accumulator()
2 {
3   __global int * glptr;
4   return glptr;
5 }
6 int i;
7 auto* ai = &i;
8 int* ii = &i;

In D65744#1646348, @rjmccall wrote:

In D65744#1629055, @Anastasia wrote:

In D65744#1627589, @rjmccall wrote:

I see. Is the deduction rule recursive or something, where a pointer to pointer is inferred to point to the same address space as the pointee?

It is recursive indeed and we currently deduce all pointees to generic AS.

Is that likely to change? If all pointees are inferred to be generic (in the absence of an explicit qualifier) when building a pointer type, this seems very over-thought vs. just adding the generic AS whenever you build a pointer to an unqualified but non-dependent type.

I don't think this is likely to change. Are you suggesting to move the deduction logic for pointee of pointers, references and block pointers into ASTContext helper that creates a pointer/reference/block pointer type? I guess it should work if the only way to construct the pointer is by using this helper. I can certainly prototype this change and see.

lib/Sema/SemaType.cpp

7441 ↗

(On Diff #213345)

I think TreeTransforms will only be used in C++ mode. But also isAutoType should only be true in C++ mode. So I think we should be fine.

[rjmccall]

In D65744#1652355, @Anastasia wrote:

I don't think this is likely to change. Are you suggesting to move the deduction logic for pointee of pointers, references and block pointers into ASTContext helper that creates a pointer/reference/block pointer type?

No. I'm suggesting that the deduction logic should be much more straightforward, just some sort of "is the type non-dependent and lacking a qualifier", and it should be applied in the two basic places we build these types in Sema, i.e. in the type-from-declarator logic and in the Build{Pointer,Reference}Type logic. Instead we have something very elaborate that apparently recursively looks through pointer types and is contingent on the exact spelling, e.g. trying to find auto types, which seems both brittle and unnecessary.

include/clang/AST/Type.h

6509 ↗

(On Diff #213345)

Hmm. So this method, confusingly, will not return true for a deduced auto, unless the deduced type is itself an undeduced auto (which I'm not sure can happen). I think it at least needs a different name; isUndeducedAutoType() would be okay if the latter case is not possible. But it might be better if we can just define away the need for the method entirely.

lib/Sema/SemaType.cpp

7441 ↗

(On Diff #213345)

I don't think TreeTransform is expected to be C++-only, but I agree that isAutoType is good enough.

[Anastasia]

Moved addr space of pointee inference into Build* helpers of Sema.

[Anastasia]

In D65744#1653081, @rjmccall wrote:

In D65744#1652355, @Anastasia wrote:

I don't think this is likely to change. Are you suggesting to move the deduction logic for pointee of pointers, references and block pointers into ASTContext helper that creates a pointer/reference/block pointer type?

No. I'm suggesting that the deduction logic should be much more straightforward, just some sort of "is the type non-dependent and lacking a qualifier", and it should be applied in the two basic places we build these types in Sema, i.e. in the type-from-declarator logic and in the Build{Pointer,Reference}Type logic.

Ok thanks, I made this change now.

Instead we have something very elaborate that apparently recursively looks through pointer types and is contingent on the exact spelling, e.g. trying to find auto types, which seems both brittle and unnecessary.

I realized that I didn't have to do this actually. It was done by mistake earlier.

[Anastasia]

include/clang/AST/Type.h

6509 ↗

(On Diff #213345)

I feel I still need this helper in order to detect the auto type. Alternatively I can create a static function for this instead of publicly exposing this functionality. Or maybe you have other ideas in mind...

[Anastasia]

In D65744#1639175, @mantognini wrote:

I think this looks good. Maybe the tests should be extended to test auto as function return type, and if there's some special handling around decltype(auto), then it should be tested too, but I'm not sure it's actually needed here. What do you think?

Do you mean anything specific to test?

[Anastasia]

• Added forgotten test

[rjmccall]

clang/lib/Sema/SemaType.cpp

7460

Okay, I understand why you're doing this now, and it makes sense. I would like to propose changing the entire way deduceOpenCLImplicitAddrSpace works. Why don't we do it more like how ObjC ARC infers its implicit ownership qualifiers:

• In SemaType, we add the default address space to non-qualified, non-dependent, non-undeduced-auto pointees when parsing a pointer or reference type.
• In SemaType, we add the default address space to non-qualified pointees when building a pointer or reference type.
• We add the default address space at the top level when when building a variable.

Then all of this context-specific logic where we're looking at different declarator chunks and trying to infer the relationship of the current chunk to the overall type being parsed can just go away or get pushed into a more appropriate position.

[Anastasia]

• Move addr space deduction to a later phase.

[Anastasia]

clang/lib/Sema/SemaDecl.cpp

5998

I don't seem to need a check for dependent or auto types because the substitution happens using type info rather than getting type from the declaration. Not sure if I should explain it here or add the checks just in case?

clang/lib/Sema/SemaType.cpp

7460

Ok, it mainly works, however I still need a bit of parsing horribleness when deducing addr space of declarations with parenthesis in GetFullTypeForDeclarator. This is the case for block pointers or pointers/references to arrays. It is incorrect to add address spaces on ParenType while building a pointer or references so I have to detect this as special case.

[rjmccall]

clang/lib/Sema/SemaDecl.cpp

5998

If you have adequate test-case coverage (both inside and out of templates) then I don't think you need further explanation.

6721

Since you're moving this code anyway, can this be split into its own function? I'm not sure if it's actually important that some of these failures return immediately and some of them fall through to later checks.

clang/lib/Sema/SemaType.cpp

7460

You can't add an address space outside a ParenType? That seems odd; what problems are you seeing exactly?

If it's really just specific to ParenType, you could simply drill through them and then rebuild the ParenTypes.

[Anastasia]

clang/lib/Sema/SemaDecl.cpp

6721

I'm not sure if it's actually important that some of these failures return immediately and some of them fall through to later checks.

Yes, it looks a bit random. Do we have any guideline whether we should return or keep diagnosing as long as possible?

clang/lib/Sema/SemaType.cpp

7460

You can't add an address space outside a ParenType? That seems odd; what problems are you seeing exactly?

When we add addr space for a pointee and it's a ParenType the addr space should actually be added to a first non-ParenType but not ParenType itself. For example is we declare a reference to an array we want the array type to have an address space not ParenType.

If it's really just specific to ParenType, you could simply drill through them and then rebuild the ParenTypes.

Ok, in the case I explained above we would have to add address space to the first non-ParenTypes and then rebuild all ParenTypes. I will try that.

[rjmccall]

clang/lib/Sema/SemaDecl.cpp

6721

If the user is likely to have made multiple independent errors, it's good to diagnose as many of them as possible. But if it's just as likely that the user messed up in a single way that should've meant we didn't take this code path, then it's better to bail out early.

In this case, most of these diagnostics are looking for different special OpenCL types, and those are probably all independent of each other.

...it does occur to me to wonder if more of these checks should be looking through array types the way that the check for half does. Presumably you shouldn't be able to declare global arrays of images or pipes if you can't declare non-array variables of them.

clang/lib/Sema/SemaType.cpp

7460

ParenType is just a sugar node, and qualifiers on it should be treated identically to qualifiers on the inner type, just like a qualifier on a typedef name (i.e. const int32_t) would. So unless you're seeing a real problem I wouldn't worry about it.

I could totally believe that it prints poorly, though.

[Anastasia]

• Factored OpenCL diagnostics out in a separate helper function
• Removed special case for ParenType

clang/lib/Sema/SemaDecl.cpp

6721

...it does occur to me to wonder if more of these checks should be looking through array types the way that the check for half does. Presumably you shouldn't be able to declare global arrays of images or pipes if you can't declare non-array variables of them.

We actually provide dedicated diagnostic for all other types in Sema::BuildArrayType. No idea why half is handled here I will try to refactor that in a separate patch.

[rjmccall]

clang/lib/Sema/SemaDecl.cpp

6721

Well, half is a perfectly reasonable type to have an array of. I don't know why that's not equally true of images or pipes; are they assumed to have implicit trailing storage? Anyway, if OpenCL says arrays of them are forbidden, we don't need to look through arrays; that's probably worth a comment, though.

clang/include/clang/AST/Type.h

   bool isAlignValT() const;                     // C++17 std::align_val_t
   bool isStdByteType() const;                   // C++17 std::byte
   bool isAtomicType() const;                    // C11 _Atomic()
+  bool isUndeducedAutoType() const;             // C++11 auto or
+                                                // C++14 decltype(auto)
 
 #define IMAGE_TYPE(ImgType, Id, SingletonId, Access, Suffix) \
   bool is##Id##Type() const;

   return isa<AtomicType>(CanonicalType);
 }
 
+inline bool Type::isUndeducedAutoType() const {
+  return isa<AutoType>(CanonicalType);
+}
+
 inline bool Type::isObjCQualifiedIdType() const {
   if (const auto *OPT = getAs<ObjCObjectPointerType>())
     return OPT->isObjCQualifiedIdType();

clang/include/clang/Sema/Sema.h

   bool CheckARCMethodDecl(ObjCMethodDecl *method);
   bool inferObjCARCLifetime(ValueDecl *decl);
 
+  void deduceOpenCLAddressSpace(ValueDecl *decl);
+
   ExprResult
   HandleExprPropertyRefExpr(const ObjCObjectPointerType *OPT,
                             Expr *BaseExpr,

clang/lib/AST/Expr.cpp

     auto Ty = getType();
     auto SETy = getSubExpr()->getType();
     assert(getValueKindForType(Ty) == Expr::getValueKindForType(SETy));
-    if (/*isRValue()*/ !Ty->getPointeeType().isNull()) {
+    if (isRValue()) {
       Ty = Ty->getPointeeType();
       SETy = SETy->getPointeeType();
     }

clang/lib/Sema/SemaDecl.cpp

   return false;
 }
 
+void Sema::deduceOpenCLAddressSpace(ValueDecl *Decl) {
+  if (Decl->getType().getQualifiers().hasAddressSpace())
+    return;
+  if (VarDecl *Var = dyn_cast<VarDecl>(Decl)) {
+    QualType Type = Var->getType();
+    if (Type->isSamplerT() || Type->isVoidType())
+      return;
+    LangAS ImplAS = LangAS::opencl_private;
+    if ((getLangOpts().OpenCLCPlusPlus || getLangOpts().OpenCLVersion >= 200) &&
+        Var->hasGlobalStorage())
+      ImplAS = LangAS::opencl_global;
+    Type = Context.getAddrSpaceQualType(Type, ImplAS);
+    Decl->setType(Type);
+  }
+}
+
 static void checkAttributesAfterMerging(Sema &S, NamedDecl &ND) {
   // Ensure that an auto decl is deduced otherwise the checks below might cache
   // the wrong linkage.

 
   llvm_unreachable("Unknown type of decl!");
 }
+/// Returns true if there hasn't been any invalid type diagnosed.
+static bool diagnoseOpenCLTypes(Scope *S, Sema &Se, Declarator &D,
+                                DeclContext *DC, QualType R) {
+  // OpenCL v2.0 s6.9.b - Image type can only be used as a function argument.
+  // OpenCL v2.0 s6.13.16.1 - Pipe type can only be used as a function
+  // argument.
+  if (R->isImageType() || R->isPipeType()) {
+    Se.Diag(D.getIdentifierLoc(),
+            diag::err_opencl_type_can_only_be_used_as_function_parameter)
+        << R;
+    D.setInvalidType();
+    return false;
+  }
+
+  // OpenCL v1.2 s6.9.r:
+  // The event type cannot be used to declare a program scope variable.
+  // OpenCL v2.0 s6.9.q:
+  // The clk_event_t and reserve_id_t types cannot be declared in program
+  // scope.
+  if (NULL == S->getParent()) {
+    if (R->isReserveIDT() || R->isClkEventT() || R->isEventT()) {
+      Se.Diag(D.getIdentifierLoc(),
+              diag::err_invalid_type_for_program_scope_var)
+          << R;
+      D.setInvalidType();
+      return false;
+    }
+  }
+
+  // OpenCL v1.0 s6.8.a.3: Pointers to functions are not allowed.
+  QualType NR = R;
+  while (NR->isPointerType()) {
+    if (NR->isFunctionPointerType()) {
+      Se.Diag(D.getIdentifierLoc(), diag::err_opencl_function_pointer);
+      D.setInvalidType();
+      return false;
+    }
+    NR = NR->getPointeeType();
+  }
+
+  if (!Se.getOpenCLOptions().isEnabled("cl_khr_fp16")) {
+    // OpenCL v1.2 s6.1.1.1: reject declaring variables of the half and
+    // half array type (unless the cl_khr_fp16 extension is enabled).
+    if (Se.Context.getBaseElementType(R)->isHalfType()) {
+      Se.Diag(D.getIdentifierLoc(), diag::err_opencl_half_declaration) << R;
+      D.setInvalidType();
+      return false;
+    }
+  }
+
+  // OpenCL v1.2 s6.9.r:
+  // The event type cannot be used with the __local, __constant and __global
+  // address space qualifiers.
+  if (R->isEventT()) {
+    if (R.getAddressSpace() != LangAS::opencl_private) {
+      Se.Diag(D.getBeginLoc(), diag::err_event_t_addr_space_qual);
+      D.setInvalidType();
+      return false;
+    }
+  }
+
+  // C++ for OpenCL does not allow the thread_local storage qualifier.
+  // OpenCL C does not support thread_local either, and
+  // also reject all other thread storage class specifiers.
+  DeclSpec::TSCS TSC = D.getDeclSpec().getThreadStorageClassSpec();
+  if (TSC != TSCS_unspecified) {
+    bool IsCXX = Se.getLangOpts().OpenCLCPlusPlus;
+    Se.Diag(D.getDeclSpec().getThreadStorageClassSpecLoc(),
+            diag::err_opencl_unknown_type_specifier)
+        << IsCXX << Se.getLangOpts().getOpenCLVersionTuple().getAsString()
+        << DeclSpec::getSpecifierName(TSC) << 1;
+    D.setInvalidType();
+    return false;
+  }
+
+  if (R->isSamplerT()) {
+    // OpenCL v1.2 s6.9.b p4:
+    // The sampler type cannot be used with the __local and __global address
+    // space qualifiers.
+    if (R.getAddressSpace() == LangAS::opencl_local ||
+        R.getAddressSpace() == LangAS::opencl_global) {
+      Se.Diag(D.getIdentifierLoc(), diag::err_wrong_sampler_addressspace);
+      D.setInvalidType();
+    }
+
+    // OpenCL v1.2 s6.12.14.1:
+    // A global sampler must be declared with either the constant address
+    // space qualifier or with the const qualifier.
+    if (DC->isTranslationUnit() &&
+        !(R.getAddressSpace() == LangAS::opencl_constant ||
+          R.isConstQualified())) {
+      Se.Diag(D.getIdentifierLoc(), diag::err_opencl_nonconst_global_sampler);
+      D.setInvalidType();
+    }
+    if (D.isInvalidType())
+      return false;
+  }
+  return true;
+}
 
 NamedDecl *Sema::ActOnVariableDeclarator(
     Scope *S, Declarator &D, DeclContext *DC, TypeSourceInfo *TInfo,

     return nullptr;
   }
 
-  if (getLangOpts().OpenCL) {
-    // OpenCL v2.0 s6.9.b - Image type can only be used as a function argument.
-    // OpenCL v2.0 s6.13.16.1 - Pipe type can only be used as a function
-    // argument.
-    if (R->isImageType() || R->isPipeType()) {
-      Diag(D.getIdentifierLoc(),
-           diag::err_opencl_type_can_only_be_used_as_function_parameter)
-          << R;
-      D.setInvalidType();
-      return nullptr;
-    }
-
-    // OpenCL v1.2 s6.9.r:
-    // The event type cannot be used to declare a program scope variable.
-    // OpenCL v2.0 s6.9.q:
-    // The clk_event_t and reserve_id_t types cannot be declared in program scope.
-    if (NULL == S->getParent()) {
-      if (R->isReserveIDT() || R->isClkEventT() || R->isEventT()) {
-        Diag(D.getIdentifierLoc(),
-             diag::err_invalid_type_for_program_scope_var) << R;
-        D.setInvalidType();
-        return nullptr;
-      }
-    }
-
-    // OpenCL v1.0 s6.8.a.3: Pointers to functions are not allowed.
-    QualType NR = R;
-    while (NR->isPointerType()) {
-      if (NR->isFunctionPointerType()) {
-        Diag(D.getIdentifierLoc(), diag::err_opencl_function_pointer);
-        D.setInvalidType();
-        break;
-      }
-      NR = NR->getPointeeType();
-    }
-
-    if (!getOpenCLOptions().isEnabled("cl_khr_fp16")) {
-      // OpenCL v1.2 s6.1.1.1: reject declaring variables of the half and
-      // half array type (unless the cl_khr_fp16 extension is enabled).
-      if (Context.getBaseElementType(R)->isHalfType()) {
-        Diag(D.getIdentifierLoc(), diag::err_opencl_half_declaration) << R;
-        D.setInvalidType();
-      }
-    }
-
-    if (R->isSamplerT()) {
-      // OpenCL v1.2 s6.9.b p4:
-      // The sampler type cannot be used with the __local and __global address
-      // space qualifiers.
-      if (R.getAddressSpace() == LangAS::opencl_local ||
-          R.getAddressSpace() == LangAS::opencl_global) {
-        Diag(D.getIdentifierLoc(), diag::err_wrong_sampler_addressspace);
-      }
-
-      // OpenCL v1.2 s6.12.14.1:
-      // A global sampler must be declared with either the constant address
-      // space qualifier or with the const qualifier.
-      if (DC->isTranslationUnit() &&
-          !(R.getAddressSpace() == LangAS::opencl_constant ||
-          R.isConstQualified())) {
-        Diag(D.getIdentifierLoc(), diag::err_opencl_nonconst_global_sampler);
-        D.setInvalidType();
-      }
-    }
-
-    // OpenCL v1.2 s6.9.r:
-    // The event type cannot be used with the __local, __constant and __global
-    // address space qualifiers.
-    if (R->isEventT()) {
-      if (R.getAddressSpace() != LangAS::opencl_private) {
-        Diag(D.getBeginLoc(), diag::err_event_t_addr_space_qual);
-        D.setInvalidType();
-      }
-    }
-
-    // C++ for OpenCL does not allow the thread_local storage qualifier.
-    // OpenCL C does not support thread_local either, and
-    // also reject all other thread storage class specifiers.
-    DeclSpec::TSCS TSC = D.getDeclSpec().getThreadStorageClassSpec();
-    if (TSC != TSCS_unspecified) {
-      bool IsCXX = getLangOpts().OpenCLCPlusPlus;
-      Diag(D.getDeclSpec().getThreadStorageClassSpecLoc(),
-           diag::err_opencl_unknown_type_specifier)
-          << IsCXX << getLangOpts().getOpenCLVersionTuple().getAsString()
-          << DeclSpec::getSpecifierName(TSC) << 1;
-      D.setInvalidType();
-      return nullptr;
-    }
-  }
 
   DeclSpec::SCS SCSpec = D.getDeclSpec().getStorageClassSpec();
   StorageClass SC = StorageClassSpecToVarDeclStorageClass(D.getDeclSpec());

     }
   }
 
+  if (getLangOpts().OpenCL) {
+
+    deduceOpenCLAddressSpace(NewVD);
+
+    diagnoseOpenCLTypes(S, *this, D, DC, NewVD->getType());
+  }
+
   // Handle attributes prior to checking for duplicates in MergeVarDecl
   ProcessDeclAttributes(S, NewVD, D);
 

   if (getLangOpts().ObjCAutoRefCount && inferObjCARCLifetime(VDecl))
     VDecl->setInvalidDecl();
 
+  if (getLangOpts().OpenCL)
+    deduceOpenCLAddressSpace(VDecl);
+
   // If this is a redeclaration, check that the type we just deduced matches
   // the previously declared type.
   if (VarDecl *Old = VDecl->getPreviousDecl()) {

   if (New->hasAttr<BlocksAttr>()) {
     Diag(New->getLocation(), diag::err_block_on_nonlocal);
   }
+
+  if (getLangOpts().OpenCL)
+    deduceOpenCLAddressSpace(New);
+
   return New;
 }
 

clang/lib/Sema/SemaTemplateInstantiate.cpp

                                                  int indexAdjustment,
                                                Optional<unsigned> NumExpansions,
                                                  bool ExpectParameterPack) {
-  return SemaRef.SubstParmVarDecl(OldParm, TemplateArgs, indexAdjustment,
-                                  NumExpansions, ExpectParameterPack);
+  auto NewParm =
+      SemaRef.SubstParmVarDecl(OldParm, TemplateArgs, indexAdjustment,
+                               NumExpansions, ExpectParameterPack);
+  if (NewParm && SemaRef.getLangOpts().OpenCL)
+    SemaRef.deduceOpenCLAddressSpace(NewParm);
+  return NewParm;
 }
 
 QualType

clang/lib/Sema/SemaTemplateInstantiateDecl.cpp

       SemaRef.inferObjCARCLifetime(Var))
     Var->setInvalidDecl();
 
+  if (SemaRef.getLangOpts().OpenCL)
+    SemaRef.deduceOpenCLAddressSpace(Var);
+
   // Substitute the nested name specifier, if any.
   if (SubstQualifier(D, Var))
     return nullptr;

clang/lib/Sema/SemaType.cpp

   return true;
 }
 
+// Helper to deduce addr space of a pointee type in OpenCL mode.
+static QualType deduceOpenCLPointeeAddrSpace(Sema &S, QualType PointeeType) {
+  if (!PointeeType->isUndeducedAutoType() && !PointeeType->isDependentType() &&
+      !PointeeType->isSamplerT() &&
+      !PointeeType.getQualifiers().hasAddressSpace())
+    PointeeType = S.getASTContext().getAddrSpaceQualType(
+        PointeeType,
+        S.getLangOpts().OpenCLCPlusPlus || S.getLangOpts().OpenCLVersion == 200
+            ? LangAS::opencl_generic
+            : LangAS::opencl_private);
+  return PointeeType;
+}
+
 /// Build a pointer type.
 ///
 /// \param T The type to which we'll be building a pointer.

   if (getLangOpts().ObjCAutoRefCount)
     T = inferARCLifetimeForPointee(*this, T, Loc, /*reference*/ false);
 
+  if (getLangOpts().OpenCL)
+    T = deduceOpenCLPointeeAddrSpace(*this, T);
+
   // Build the pointer type.
   return Context.getPointerType(T);
 }

   if (getLangOpts().ObjCAutoRefCount)
     T = inferARCLifetimeForPointee(*this, T, Loc, /*reference*/ true);
 
+  if (getLangOpts().OpenCL)
+    T = deduceOpenCLPointeeAddrSpace(*this, T);
+
   // Handle restrict on references.
   if (LValueRef)
     return Context.getLValueReferenceType(T, SpelledAsLValue);

   if (checkQualifiedFunction(*this, T, Loc, QFK_BlockPointer))
     return QualType();
 
+  if (getLangOpts().OpenCL)
+    T = deduceOpenCLPointeeAddrSpace(*this, T);
+
   return Context.getBlockPointerType(T);
 }
 

   }
 }
 
-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 always 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
-  // defining 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::Reference ||
-       D.getTypeObject(ChunkIndex - 1).Kind == DeclaratorChunk::BlockPointer);
-  // For pointers/references to arrays the next chunk is always an array
-  // followed by any number of parentheses.
-  if (!IsPointee && ChunkIndex > 1) {
-    auto AdjustedCI = ChunkIndex - 1;
-    if (D.getTypeObject(AdjustedCI).Kind == DeclaratorChunk::Array)
-      AdjustedCI--;
-    // Skip over all parentheses.
-    while (AdjustedCI > 0 &&
-           D.getTypeObject(AdjustedCI).Kind == DeclaratorChunk::Paren)
-      AdjustedCI--;
-    if (D.getTypeObject(AdjustedCI).Kind == DeclaratorChunk::Pointer ||
-        D.getTypeObject(AdjustedCI).Kind == DeclaratorChunk::Reference)
-      IsPointee = true;
-  }
-  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 or static data members.
-      (D.getContext() == DeclaratorContext::MemberContext &&
-       (!IsPointee &&
-        D.getDeclSpec().getStorageClassSpec() != DeclSpec::SCS_static)) ||
-      // Do not deduce addr space of non-pointee in type alias because it
-      // doesn't define any object.
-      (D.getContext() == DeclaratorContext::AliasDeclContext && !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) ||
-      // Do not deduce addr spaces for dependent types because they might end
-      // up instantiating to a type with an explicit address space qualifier.
-      // Except for pointer or reference types because the addr space in
-      // template argument can only belong to a pointee.
-      (T->isDependentType() && !T->isPointerType() && !T->isReferenceType()) ||
-      // Do not deduce addr space of decltype because it will be taken from
-      // its argument.
-      T->isDecltypeType() ||
-      // OpenCL spec v2.0 s6.9.b:
-      // The sampler type cannot be used with the __local and __global address
-      // space qualifiers.
-      // OpenCL spec v2.0 s6.13.14:
-      // Samplers can also be declared as global constants in the program
-      // source using the following syntax.
-      //   const sampler_t <sampler name> = <value>
-      // In codegen, file-scope sampler type variable has special handing and
-      // does not rely on address space qualifier. On the other hand, deducing
-      // address space of const sampler file-scope variable as global address
-      // space causes spurious diagnostic about __global address space
-      // qualifier, therefore do not deduce address space of file-scope sampler
-      // type variable.
-      (D.getContext() == DeclaratorContext::FileContext && T->isSamplerT()))
-    return;
-
-  LangAS ImpAddr = LangAS::Default;
-  // 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 &&
-      !State.getSema().getLangOpts().OpenCLCPlusPlus) {
-    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() == DeclaratorContext::TemplateArgContext) {
-        // Do not deduce address space for non-pointee type in template arg.
-      } else if (D.getContext() == DeclaratorContext::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 HandleLifetimeBoundAttr(TypeProcessingState &State,
                                     QualType &CurType,
                                     ParsedAttr &Attr) {

   if (!state.getSema().getLangOpts().OpenCL ||
       type.getAddressSpace() != LangAS::Default)
     return;
-
-  deduceOpenCLImplicitAddrSpace(state, type, TAL);
 }
 
 void Sema::completeExprArrayBound(Expr *E) {

clang/lib/Sema/TreeTransform.h

   return Result;
 }
 
-/// Helper to deduce addr space of a pointee type in OpenCL mode.
-/// If the type is updated it will be overwritten in PointeeType param.
-static void deduceOpenCLPointeeAddrSpace(Sema &SemaRef, QualType &PointeeType) {
-  if (PointeeType.getAddressSpace() == LangAS::Default)
-    PointeeType = SemaRef.Context.getAddrSpaceQualType(PointeeType,
-                                                       LangAS::opencl_generic);
-}
-
 template<typename Derived>
 QualType TreeTransform<Derived>::TransformPointerType(TypeLocBuilder &TLB,
                                                       PointerTypeLoc TL) {

   if (PointeeType.isNull())
     return QualType();
 
-  if (SemaRef.getLangOpts().OpenCL)
-    deduceOpenCLPointeeAddrSpace(SemaRef, PointeeType);
-
   QualType Result = TL.getType();
   if (PointeeType->getAs<ObjCObjectType>()) {
     // A dependent pointer type 'T *' has is being transformed such

   if (PointeeType.isNull())
     return QualType();
 
-  if (SemaRef.getLangOpts().OpenCL)
-    deduceOpenCLPointeeAddrSpace(SemaRef, PointeeType);
-
   QualType Result = TL.getType();
   if (getDerived().AlwaysRebuild() ||
       PointeeType != TL.getPointeeLoc().getType()) {

   if (PointeeType.isNull())
     return QualType();
 
-  if (SemaRef.getLangOpts().OpenCL)
-    deduceOpenCLPointeeAddrSpace(SemaRef, PointeeType);
-
   QualType Result = TL.getType();
   if (getDerived().AlwaysRebuild() ||
       PointeeType != T->getPointeeTypeAsWritten()) {

clang/test/SemaOpenCL/event_t.cl

 // RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
 
-event_t glb_evt; // expected-error {{the 'event_t' type cannot be used to declare a program scope variable}}
+event_t glb_evt; // expected-error {{the 'event_t' type cannot be used to declare a program scope variable}} expected-error{{program scope variable must reside in constant address space}}
 
 constant struct evt_s {
   event_t evt; // expected-error {{the 'event_t' type cannot be used to declare a structure or union field}}

 
 void kernel ker(event_t argevt) { // expected-error {{'event_t' cannot be used as the type of a kernel parameter}}
   event_t e;
-  constant event_t const_evt; // expected-error {{the event_t type can only be used with __private address space qualifier}}
+  constant event_t const_evt; // expected-error {{the event_t type can only be used with __private address space qualifier}} expected-error{{variable in constant address space must be initialized}}
   foo(e);
   foo(0);
   foo(5); // expected-error {{passing 'int' to parameter of incompatible type 'event_t'}}

clang/test/SemaOpenCL/invalid-block.cl

               : bl2(i);     // expected-error {{block type cannot be used as expression in ternary expression in OpenCL}}
 }
 // A block pointer type and all pointer operations are disallowed
-void f6(bl2_t *bl_ptr) { // expected-error{{pointer to type '__generic bl2_t' (aka 'int (__generic ^const __generic)(int)') is invalid in OpenCL}}
+void f6(bl2_t *bl_ptr) { // expected-error{{pointer to type 'bl2_t' (aka 'int (__generic ^const)(int)') is invalid in OpenCL}}
   bl2_t bl = ^(int i) {
     return 1;
   };
-  bl2_t *p; // expected-error {{pointer to type '__generic bl2_t' (aka 'int (__generic ^const __generic)(int)') is invalid in OpenCL}}
+  bl2_t *p; // expected-error {{pointer to type 'bl2_t' (aka 'int (__generic ^const)(int)') is invalid in OpenCL}}
   *bl;      // expected-error {{invalid argument type 'bl2_t' (aka 'int (__generic ^const)(int)') to unary expression}}
   &bl;      // expected-error {{invalid argument type 'bl2_t' (aka 'int (__generic ^const)(int)') to unary expression}}
 }

clang/test/SemaOpenCL/invalid-pipes-cl2.0.cl

 global pipe int gp;            // expected-error {{type '__global read_only pipe int' can only be used as a function parameter in OpenCL}}
 global reserve_id_t rid;          // expected-error {{the '__global reserve_id_t' type cannot be used to declare a program scope variable}}
 
-extern pipe write_only int get_pipe(); // expected-error-re{{type '__global write_only pipe int ({{(void)?}})' can only be used as a function parameter in OpenCL}}
+extern pipe write_only int get_pipe(); // expected-error-re{{type '__global write_only pipe int ({{(void)?}})' can only be used as a function parameter in OpenCL}} expected-error{{'write_only' attribute only applies to parameters and typedefs}}
 
 kernel void test_invalid_reserved_id(reserve_id_t ID) { // expected-error {{'reserve_id_t' cannot be used as the type of a kernel parameter}}
 }

clang/test/SemaOpenCL/sampler_t.cl

 sampler_t bad(void); //expected-error{{declaring function return value of type 'sampler_t' is not allowed}}
 
 sampler_t global_nonconst_smp = 0; // expected-error {{global sampler requires a const or constant address space qualifier}}
+#ifdef CHECK_SAMPLER_VALUE
+// expected-warning@-2{{sampler initializer has invalid Filter Mode bits}}
+#endif
 
 const sampler_t glb_smp10 = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_NORMALIZED_COORDS_TRUE | CLK_FILTER_LINEAR;
 const constant sampler_t glb_smp11 = CLK_ADDRESS_CLAMP_TO_EDGE | CLK_NORMALIZED_COORDS_TRUE | CLK_FILTER_LINEAR;

 }
 
 #if __OPENCL_C_VERSION__ == 200
-void bad(sampler_t*); // expected-error{{pointer to type '__generic sampler_t' is invalid in OpenCL}}
+void bad(sampler_t *); // expected-error{{pointer to type 'sampler_t' is invalid in OpenCL}}
 #else
 void bad(sampler_t*); // expected-error{{pointer to type 'sampler_t' is invalid in OpenCL}}
 #endif

clang/test/SemaOpenCLCXX/address-space-deduction.cl

 x3<T>::x3(const x3<T> &t) {}
 
 template <class T>
-T xxx(T *in) {
+T xxx(T *in1, T in2) {
   // This pointer can't be deduced to generic because addr space
   // will be taken from the template argument.
   //CHECK: `-VarDecl {{.*}} i 'T *' cinit
-  T *i = in;
+  T *i = in1;
   T ii;
+  __private T *ptr = &ii;
+  ptr = &in2;
   return *i;
 }
 
 __kernel void test() {
   int foo[10];
-  xxx(&foo[0]);
+  xxx<__private int>(&foo[0], foo[0]);
+  // FIXME: Template param deduction fails here because
+  // temporaries are not in the __private address space.
+  // It is probably reasonable to put them in __private
+  // considering that stack and function params are
+  // implicitly in __private.
+  // However, if temporaries are left in default addr
+  // space we should at least pretty print the __private
+  // addr space. Otherwise diagnostic apprears to be
+  // confusing.
+  //xxx(&foo[0], foo[0]);
 }
 
 // Addr space for pointer/reference to an array
-//CHECK: FunctionDecl {{.*}} t1 'void (const __generic float (&)[2])'
+//CHECK: FunctionDecl {{.*}} t1 'void (const float (__generic &)[2])'
 void t1(const float (&fYZ)[2]);
-//CHECK: FunctionDecl {{.*}} t2 'void (const __generic float (*)[2])'
+//CHECK: FunctionDecl {{.*}} t2 'void (const float (__generic *)[2])'
 void t2(const float (*fYZ)[2]);
-//CHECK: FunctionDecl {{.*}} t3 'void (__generic float (((*)))[2])'
+//CHECK: FunctionDecl {{.*}} t3 'void (float (((__generic *)))[2])'
 void t3(float(((*fYZ)))[2]);
-//CHECK: FunctionDecl {{.*}} t4 'void (__generic float (((*__generic *)))[2])'
+//CHECK: FunctionDecl {{.*}} t4 'void (float (((__generic *__generic *)))[2])'
 void t4(float(((**fYZ)))[2]);
-//CHECK: FunctionDecl {{.*}} t5 'void (__generic float (*__generic (*))[2])'
+//CHECK: FunctionDecl {{.*}} t5 'void (float (__generic *(__generic *))[2])'
 void t5(float (*(*fYZ))[2]);
 
 __kernel void k() {

clang/test/SemaOpenCLCXX/addrspace-auto.cl

+//RUN: %clang_cc1 %s -cl-std=clc++ -pedantic -ast-dump -verify | FileCheck %s
+
+__constant int i = 1;
+//CHECK: |-VarDecl {{.*}} ai '__global int':'__global int'
+auto ai = i;
+
+kernel void test() {
+  int i;
+  //CHECK: VarDecl {{.*}} ai 'int':'int'
+  auto ai = i;
+
+  constexpr int c = 1;
+  //CHECK: VarDecl {{.*}} used cai '__constant int':'__constant int'
+  __constant auto cai = c;
+  //CHECK: VarDecl {{.*}} aii 'int':'int'
+  auto aii = cai;
+
+  //CHECK: VarDecl {{.*}} ref 'int &'
+  auto &ref = i;
+  //CHECK: VarDecl {{.*}} ptr 'int *'
+  auto *ptr = &i;
+  //CHECK: VarDecl {{.*}} ref_c '__constant int &'
+  auto &ref_c = cai;
+
+  //CHECK: VarDecl {{.*}} ptrptr 'int *__generic *'
+  auto **ptrptr = &ptr;
+  //CHECK: VarDecl {{.*}} refptr 'int *__generic &'
+  auto *&refptr = ptr;
+
+  //CHECK: VarDecl {{.*}} invalid gref '__global auto &'
+  __global auto &gref = i; //expected-error{{variable 'gref' with type '__global auto &' has incompatible initializer of type 'int'}}
+  __local int *ptr_l;
+  //CHECK: VarDecl {{.*}} invalid gptr '__global auto *'
+  __global auto *gptr = ptr_l; //expected-error{{variable 'gptr' with type '__global auto *' has incompatible initializer of type '__local int *'}}
+}

clang/test/SemaOpenCLCXX/restricted.cl

 __constant _Thread_local int a = 1;
 // expected-error@-1 {{C++ for OpenCL version 1.0 does not support the '_Thread_local' storage class specifier}}
 // expected-warning@-2 {{'_Thread_local' is a C11 extension}}
-
+// expected-error@-3 {{thread-local storage is not supported for the current target}}
 __constant __thread int b = 2;
 // expected-error@-1 {{C++ for OpenCL version 1.0 does not support the '__thread' storage class specifier}}
+// expected-error@-2 {{thread-local storage is not supported for the current target}}
 kernel void test_storage_classes() {
   register int x;
   // expected-error@-1 {{C++ for OpenCL version 1.0 does not support the 'register' storage class specifier}}
   thread_local int y;
   // expected-error@-1 {{C++ for OpenCL version 1.0 does not support the 'thread_local' storage class specifier}}
+  // expected-error@-2 {{thread-local storage is not supported for the current target}}
 }