[Clang] Recognize target address space in superset calculation

[jchlanda]

Summary

Use target's address space map to handle cases when both language and
target address spaces are provided. In such case, attempt language to
target translation and only then perform the calculation.
The main motivation is to be able to use language address spaces as
inputs for builtins, which are defined in terms of target address space
(as discussed here: https://reviews.llvm.org/D112718) and hence the
definition of builtins with generic address space pointers that would
allow any other address space pointers inputs (bar constant).

This patch attempts to find a happy medium between not recognising target
address spaces at all (current state) and allowing all uses of it, based on
the assumption that users must know better. What it does not to is to
provide a bidirectional translation mechanism, which I'm not sure could ever
be done, with the current address space implementation (use of 0, the value
of default, etc).

Based on OpenCL rules, this patch follows the conversion guidelines for
generic and constant address space pointers as described here:
https://www.khronos.org/registry/OpenCL/specs/2.2/html/OpenCL_API.html#_memory_model
similarly it honours SYCL global, global device` and global host semantics.

[jchlanda]

@tra @Naghasan @t4c1 you might find it interesting, a follow up from the discussion here: https://reviews.llvm.org/D112718

[tra]

clang/include/clang/AST/Type.h

486

If A and B are both target AS, we fall through to the code which is dealing with language AS, which would not do us any good. If that's not expected to happen, we should have an assert to ensure it.

Next, I'm not particularly fond of IsSYCLOrOpenCL. Do we need it at all. If we do know that AS maps to OpenCL Constant or Generic, I would assume that those AS would follow the same semantics. Besides, will we ever see OpenCL language AS in non-OpenCL code?

Next, the function *is* OpenCL specific and would not work for CUDA or HIP. I think it needs to be generalized to provide language-specific AS mapping rules.

489

Is the check intended to tell if A is a target AS? If so, we do have isTargetAddressSpace() for that (and it uses '>= LangAS::FirstTargetAddressSpace', which suggests that > may be incorrect, too).

498–499

getLangASFromTargetAS((*ASMap)[static_cast<unsigned>(LangAS::opencl_constant)])

500

if (!IsBTargetAS) would be more directly related to what we're doing here.

508

Is the code above intended to ensure that both A and B are target AS at this point?
If so, then it could be simplified to something like this:

if (ASMap) {
  if (!isTargetAddressSpace(A))
    A = getLangASFromTargetAS((*ASMap)[static_cast<unsigned>(A)]);
  if (!isTargetAddressSpace(B))
    B = getLangASFromTargetAS((*ASMap)[static_cast<unsigned>(B)]);

  Generic = getLangASFromTargetAS((*ASMap)[static_cast<unsigned>(LangAS::opencl_generic)])
  Constant = getLangASFromTargetAS((*ASMap)[static_cast<unsigned>(LangAS::opencl_constant)]);

  // ... proceed inferring whether A is superset of B in target AS.
  return;
}
assert (isTargetAddressSpace(A) && isTargetAddressSpace(B));

clang/lib/Sema/SemaExpr.cpp

9204

Should you pass IsSYCLOrOpenCL to it too? The way isAddressSpaceSupersetOf is implemented now it may give you a different result without it.

Also, it may make sense to plumb ASMap into lhq.isAddressSpaceSupersetOf, too, and just use the old code + couple of new arguments.

9219

Do we need to pass ASMap and IsSYCLOrOpenCL here, too?

[Anastasia]

You should be able to provide an address space of the pointer using the number, see details in:
https://github.com/llvm/llvm-project/blob/main/clang/include/clang/Basic/Builtins.def#L65

However if language address spaces are needed I wonder if the best approach is just to extend
the Builtin definitions with language address spaces similar to other qualifiers?

This patch attempts to find a happy medium between not recognising target
address spaces at all (current state) and allowing all uses of it, based on
the assumption that users must know better. What it does not to is to
provide a bidirectional translation mechanism, which I'm not sure could ever
be done, with the current address space implementation (use of 0, the value
of default, etc).

Can you provide an example of where it could be useful? Note that I feel that
such functionality could be implemented on top of full implementation of
target specific address space proposed in https://reviews.llvm.org/D62574.

[jchlanda]

In D124382#3472888, @Anastasia wrote:

Can you provide an example of where it could be useful? Note that I feel that
such functionality could be implemented on top of full implementation of
target specific address space proposed in https://reviews.llvm.org/D62574.

The use case we had was when calling target builtin (that specifies address
space) from within OpenCL C. Currently this errors out, similarly, the explicit
type cast to address space yields an error
((__attribute__((address_space(3))) int *)&l_woof in the example below). This
is important for libclc, which is implemented in OpenCL C and deals directly
with target builtins.

__kernel void woof() {
  __local int l_woof;
  __nvvm_cp_async_mbarrier_arrive_shared(&l_woof);
}

I wasn't aware of that patch, sorry, I've not had a close look yet, but it
seems worryingly dated.

[Anastasia]

In D124382#3479773, @jchlanda wrote:

In D124382#3472888, @Anastasia wrote:

Can you provide an example of where it could be useful? Note that I feel that
such functionality could be implemented on top of full implementation of
target specific address space proposed in https://reviews.llvm.org/D62574.

The use case we had was when calling target builtin (that specifies address
space) from within OpenCL C. Currently this errors out, similarly, the explicit
type cast to address space yields an error
((__attribute__((address_space(3))) int *)&l_woof in the example below). This
is important for libclc, which is implemented in OpenCL C and deals directly
with target builtins.

__kernel void woof() {
  __local int l_woof;
  __nvvm_cp_async_mbarrier_arrive_shared(&l_woof);
}

I wasn't aware of that patch, sorry, I've not had a close look yet, but it
seems worryingly dated.

Ok, I think the current behavior of builtins is to work with any address space. The way it had worked so far is since the builtins are only intended to be used in toolchains (instead of arbitrary code), the toolchain developers were responsible for making sure that the address spaces are used adequately in those builtins. However the question of extending the clang builtin functions with the notion of language address spaces has popped up before. And I think we could add this feature in a very light way for example by reserving the numbers from the clang LangAS enum to be used with the language address spaces in the prototypes of builtins. Although we could think of more elegant alternatives too. My understanding is there was never a strong enough case to add this functionality.

Just to understand further - do you need the builtins to have a specific address space prototype? And do you need this to improve error handling in libclc code base?

I imagine you could also create some sort of the wrapper functions around the building with the right address spaces, i.e. something like

void __libclc_builtin1(local int* p){
   __builtin1(p);
}

So the prototype in Clang for __builtin1 would still be permissive wrt address space of the pointer but as you only use __libclc_builtin1 in the codebase, you can ensure the correct uses. While this is how this problem has been worked around I think extending Clang builtins definitions might be inevitable to avoid forcing toolchains to create wrapper functions. However if we are aiming for this goal, I think more targeted solutions would make more sense instead of solving this problem indirectly by allowing conversions between target and language address spaces as this for example won’t work for builtins that are shared between targets.

[jchlanda]

@Anastasia @tra apologies for a late reply, I'm catching up with the thread after holidays.

[jchlanda]

In D124382#3480600, @Anastasia wrote:

And I think we could add this feature in a very light way for example by reserving the numbers from the clang LangAS enum to be used with the language address spaces in the prototypes of builtins.

I'm not sure I understand how that would look, could you please elaborate?

Just to understand further - do you need the builtins to have a specific address space prototype? And do you need this to improve error handling in libclc code base?

With the wrappers suggested we could declare all the pointers to be in generic AS and that would get us around the target vs language AS problem. I don't think that would improve the situation, as from llvm perspective/use case all those builtins would be incorrect and there would be no way for users to tell that there is a specific AS requirement on them, nor would the compiler be able to warn/error. Then the only thing making it work would be those wrappers, embedded deeply in the source of libclc, which at the moment is not even shipped with upstream llvm.
The builtins in question are not exclusively OpenCL/SYCL related, say TARGET_BUILTIN(__nvvm_cp_async_ca_shared_global_16, "vv*3vC*1", "", AND(SM_80,PTX70)) would need to take both pointer in a generic address space here. It feels like explicitly providing AS in the prototype is needed.

this for example won’t work for builtins that are shared between targets.

While I agree in principle, I'm not sure if there are any target agnostic and AS specific builtins, sounds like a dangerous thing to introduce. And in any case, it's the target that provides the AS map.

[jchlanda]

clang/include/clang/AST/Type.h

486

I would only like to handle the mixed AS case, it feels like trying to walk back from both HW AS and potentially do the logic of global and constant would be against the intention of users. Asserting on only one HW AS could backfire, as I think it should be allowed to assign between different HW AS.

The reason I added IsSYCLOrOpenCL is because this code is also exercised by checkPointerTypesForAssignment which is not OpenCL specific, so I had to have a way of conditionally enabling the conversion to generic AS.

I agree, it is too restrictive now, especially that the AS map provides values for SYCL, OpenCL and CUDA, so perhaps I should extend IsSYCLOrOpenCL to be an enum specifying which language the function deals with and act accordingly?

489

Yeap, will update to isTargetAddressSpace.

498–499

OK.

500

You are right, will update.

508

Yes at the end of AS map accesses all address spaces have to be expressed in therms of HW values, but I want it to happen only in the case of mixed AS (Language and HW). I will add assert and use helpers, like you suggested in the snippet, but would like to keep the ^ condition.

clang/lib/Sema/SemaExpr.cpp

9204

Yes, will add.

I thought it would introduce much bigger diff, and there was already a handy static version of it, don't mind modifying the member if you'd prefer.

9219

I feel like it would be covered by the first if in this block.

[tra]

clang/include/clang/AST/Type.h

508

would like to keep the ^ condition.

OK. Adding a comment explaining what's going on would be helpful here.

[jchlanda]

Use helper functions when handling address space values.

[jchlanda]

clang/include/clang/AST/Type.h

486

Next, I'm not particularly fond of IsSYCLOrOpenCL. Do we need it at all. If we do know that AS maps to OpenCL Constant or Generic, I would assume that those AS would follow the same semantics. Besides, will we ever see OpenCL language AS in non-OpenCL code?

Next, the function *is* OpenCL specific and would not work for CUDA or HIP. I think it needs to be generalized to provide language-specific AS mapping rules.

I've changed that bool flag to be an enum specifying OpenCL/SYCL/None. The rational here is that the handling of AS values differs slightly (SYCL introduces globa device and global host). It would appear that CUDA follows completely different code-path and by the time isAddressSpaceSuperset called all of the language AS values are stripped and set to 0, which is why (along the fact that we don't have a valid use case for CUDA) I left it out, and only return true for an exact match.

[tra]

My concerns have been addressed. I'll defer the final LGTM to @Anastasia.

[Anastasia]

In D124382#3496417, @jchlanda wrote:

In D124382#3480600, @Anastasia wrote:

And I think we could add this feature in a very light way for example by reserving the numbers from the clang LangAS enum to be used with the language address spaces in the prototypes of builtins.

I'm not sure I understand how that would look, could you please elaborate?

We could reserve the IDs from LangAS enums to be used in https://github.com/llvm/llvm-project/blob/main/clang/include/clang/Basic/Builtins.def#L65, but we could also extend the syntax more naturally i.e. numbers could be used for target address spaces and we could add some letter based syntax for language address spaces.

Just to understand further - do you need the builtins to have a specific address space prototype? And do you need this to improve error handling in libclc code base?

With the wrappers suggested we could declare all the pointers to be in generic AS and that would get us around the target vs language AS problem. I don't think that would improve the situation, as from llvm perspective/use case all those builtins would be incorrect and there would be no way for users to tell that there is a specific AS requirement on them, nor would the compiler be able to warn/error. Then the only thing making it work would be those wrappers, embedded deeply in the source of libclc, which at the moment is not even shipped with upstream llvm.
The builtins in question are not exclusively OpenCL/SYCL related, say TARGET_BUILTIN(__nvvm_cp_async_ca_shared_global_16, "vv*3vC*1", "", AND(SM_80,PTX70)) would need to take both pointer in a generic address space here. It feels like explicitly providing AS in the prototype is needed.

My understanding was that Clang low level builtins are not desirable for uses in the application code directly. They are more targeted at tooling developers and low level libraries uses. Which is why this problem has been worked around by declaring the wrapper overload with correct address spaces in the tooling projects.

My understanding was that you need those builtins in the libclc code?

this for example won’t work for builtins that are shared between targets.

While I agree in principle, I'm not sure if there are any target agnostic and AS specific builtins, sounds like a dangerous thing to introduce. And in any case, it's the target that provides the AS map.

In OpenCL we actually have quite some target agnostic builtins.

[Anastasia]

I feel that to progress further on this change, it would be good to get details about the use cases and the limitations first.

However, if there is sufficient evidence of the need to extend clang builtins with language address spaces I am not convinced the approach here is suitable as:

1. It doesn't allow uses of language builtins with language address spaces generically as mapping to the target address spaces is not portable. In general there can also be generic buitins that are normally mapped to native LLVM intrinsics (not target intrinsics!) shared among multiple targets that might also benefit from having this implemented in a target agnostic way. Such intrinsics could be shares for example between PTX and SPIR-V targets as there is quite some overlap in the functionality.
2. It has much wider impact on the language semantics then just allowing language address spaces being used in builtins i.e. it results in implicit conversions more broadly. This might not be desirable evolution and we might need to reach some consensus with more languages or targets using the address spaces in order to proceed with such change. In fact current title and description doesn't adequately reflect the impact of the change.

Has extending the builtin definition syntax been considered for this problem? That seems like a more natural and fairly localized change. For example the syntax in https://github.com/llvm/llvm-project/blob/main/clang/include/clang/Basic/Builtins.def#L65 could be changed to:

// * -> pointer (optionally followed by an address space number for target address space
//               or by 00 and a number for language address space as it is set in LangAS, if no
//               address space is specified than any address space will be accepted)

Note that we will likely need to set LangAS entry values explicitly which would make maintaing the enum slightly more painful but it doesn't seem like a concern.

If the only use case we have right now is ability to specify generic address space in kernel-like langauges we could also just reserve a special 00 number in address space field of the buitin prototype description for generic address space and leave full support as a future work.

If we understand the use case better we might be able to look at other alternatives too...

clang/include/clang/AST/Type.h

public:
 
   /// Returns true if address space A is equal to or a superset of B.
   /// OpenCL v2.0 defines conversion rules (OpenCLC v2.0 s6.5.5) and notion of
   /// overlapping address spaces.
   /// CL1.1 or CL1.2:
   ///   every address space is a superset of itself.
   /// CL2.0 adds:
   ///   __generic is a superset of any address space except for __constant.
-  static bool isAddressSpaceSupersetOf(LangAS A, LangAS B) {
+  /// If ASMap is provided and address spaces are given in both language and
+  /// target form the function will attempt to convert language to target
+  /// address space.
+  static bool isAddressSpaceSupersetOf(LangAS A, LangAS B,
+                                       const LangASMap *ASMap = nullptr,
+                                       bool IsSYCLOrOpenCL = false) {
+    if (ASMap) {
+      bool IsATargetAS = false;
+      bool IsBTargetAS = false;
+      if (A > LangAS::FirstTargetAddressSpace)
+        IsATargetAS = true;
+      if (B > LangAS::FirstTargetAddressSpace)
+        IsBTargetAS = true;
+      if (IsATargetAS ^ IsBTargetAS) {
+        LangAS Generic = static_cast<LangAS>(
+            (*ASMap)[static_cast<unsigned>(LangAS::opencl_generic)] +
+            static_cast<unsigned>(LangAS::FirstTargetAddressSpace));
+        LangAS Constant = static_cast<LangAS>(
+            (*ASMap)[static_cast<unsigned>(LangAS::opencl_constant)] +
+            static_cast<unsigned>(LangAS::FirstTargetAddressSpace));
+        if (IsATargetAS)
+          B = static_cast<LangAS>(
+              (*ASMap)[static_cast<unsigned>(B)] +
+              static_cast<unsigned>(LangAS::FirstTargetAddressSpace));
+        else
+          A = static_cast<LangAS>(
+              (*ASMap)[static_cast<unsigned>(A)] +
+              static_cast<unsigned>(LangAS::FirstTargetAddressSpace));
+        // When dealing with target AS return true if:
+        // * A is equal to B, or
+        // * in OpenCL or SYCL and A is generic and B is not constant (making
+        // sure that constant and generic are in target address spaces).
+        if (IsSYCLOrOpenCL)
+          return A == B ||
+                 (A == Generic && B != Constant && Generic != Constant);
+        return A == B;
+      }
+    }
+
     // Address spaces must match exactly.
     return A == B ||
            // Otherwise in OpenCLC v2.0 s6.5.5: every address space except
            // for __constant can be used as __generic.
            (A == LangAS::opencl_generic && B != LangAS::opencl_constant) ||
            // We also define global_device and global_host address spaces,
            // to distinguish global pointers allocated on host from pointers
            // allocated on device, which are a subset of __global.
public:
   /// a superset of the address space in the argument qualifiers.
   bool isAddressSpaceSupersetOf(Qualifiers other) const {
     return isAddressSpaceSupersetOf(getAddressSpace(), other.getAddressSpace());
   }
 
   /// Determines if these qualifiers compatibly include another set.
   /// Generally this answers the question of whether an object with the other
   /// qualifiers can be safely used as an object with these qualifiers.
-  bool compatiblyIncludes(Qualifiers other) const {
-    return isAddressSpaceSupersetOf(other) &&
+  bool compatiblyIncludes(Qualifiers other, const LangASMap *ASMap = nullptr,
+                          bool IsSYCLOrOpenCL = false) {
+    return isAddressSpaceSupersetOf(this->getAddressSpace(),
+                                    other.getAddressSpace(), ASMap,
+                                    IsSYCLOrOpenCL) &&
            // ObjC GC qualifiers can match, be added, or be removed, but can't
            // be changed.
            (getObjCGCAttr() == other.getObjCGCAttr() || !hasObjCGCAttr() ||
             !other.hasObjCGCAttr()) &&
            // ObjC lifetime qualifiers must match exactly.
            getObjCLifetime() == other.getObjCLifetime() &&
            // CVR qualifiers may subset.
            (((Mask & CVRMask) | (other.Mask & CVRMask)) == (Mask & CVRMask)) &&

clang/lib/Sema/SemaCast.cpp

void CastOperation::checkAddressSpaceCast(QualType SrcType, QualType DestType) {
   //   local int ** p;
   //   return (generic int **) p;
   // warn even though local -> generic is permitted.
   if (Self.getLangOpts().OpenCL) {
     const Type *DestPtr, *SrcPtr;
     bool Nested = false;
     unsigned DiagID = diag::err_typecheck_incompatible_address_space;
     DestPtr = Self.getASTContext().getCanonicalType(DestType.getTypePtr()),
-    SrcPtr  = Self.getASTContext().getCanonicalType(SrcType.getTypePtr());
+    SrcPtr = Self.getASTContext().getCanonicalType(SrcType.getTypePtr());
+    const LangASMap &ASMap =
+        Self.getASTContext().getTargetInfo().getAddressSpaceMap();
 
     while (isa<PointerType>(DestPtr) && isa<PointerType>(SrcPtr)) {
       const PointerType *DestPPtr = cast<PointerType>(DestPtr);
       const PointerType *SrcPPtr = cast<PointerType>(SrcPtr);
       QualType DestPPointee = DestPPtr->getPointeeType();
       QualType SrcPPointee = SrcPPtr->getPointeeType();
-      if (Nested
-              ? DestPPointee.getAddressSpace() != SrcPPointee.getAddressSpace()
-              : !DestPPointee.isAddressSpaceOverlapping(SrcPPointee)) {
+      LangAS DestAS = DestPPointee.getAddressSpace();
+      LangAS SrcAS = SrcPPointee.getAddressSpace();
+      const bool OverlappingAS =
+          Qualifiers::isAddressSpaceSupersetOf(DestAS, SrcAS, &ASMap, true) ||
+          Qualifiers::isAddressSpaceSupersetOf(SrcAS, DestAS, &ASMap, true);
+      if (Nested ? DestAS != SrcAS : !OverlappingAS) {
         Self.Diag(OpRange.getBegin(), DiagID)
             << SrcType << DestType << Sema::AA_Casting
             << SrcExpr.get()->getSourceRange();
         if (!Nested)
           SrcExpr = ExprError();
         return;
       }
 

clang/lib/Sema/SemaExpr.cpp

checkPointerTypesForAssignment(Sema &S, QualType LHSType, QualType RHSType) {
   // As a special case, 'non-__weak A *' -> 'non-__weak const *' is okay.
   if (lhq.getObjCLifetime() != rhq.getObjCLifetime() &&
       lhq.compatiblyIncludesObjCLifetime(rhq)) {
     // Ignore lifetime for further calculation.
     lhq.removeObjCLifetime();
     rhq.removeObjCLifetime();
   }
 
-  if (!lhq.compatiblyIncludes(rhq)) {
+  const bool IsSYCLOrOpenCL =
+      S.getLangOpts().OpenCL || S.getLangOpts().SYCLIsDevice;
+  const LangASMap &ASMap = S.Context.getTargetInfo().getAddressSpaceMap();
+  if (!lhq.compatiblyIncludes(rhq, &ASMap)) {
+    const bool AddressSpaceSuperset = Qualifiers::isAddressSpaceSupersetOf(
+        lhq.getAddressSpace(), rhq.getAddressSpace(), &ASMap, IsSYCLOrOpenCL);
+
     // Treat address-space mismatches as fatal.
-    if (!lhq.isAddressSpaceSupersetOf(rhq))
+    if (!AddressSpaceSuperset)
       return Sema::IncompatiblePointerDiscardsQualifiers;
 
+    // In OpenCL/SYCL don't issue discard qualifier warning if address spaces
+    // overlap.
+    else if (AddressSpaceSuperset && IsSYCLOrOpenCL)
+      ; // keep Compatible
+
     // It's okay to add or remove GC or lifetime qualifiers when converting to
     // and from void*.
-    else if (lhq.withoutObjCGCAttr().withoutObjCLifetime()
-                        .compatiblyIncludes(
-                                rhq.withoutObjCGCAttr().withoutObjCLifetime())
-             && (lhptee->isVoidType() || rhptee->isVoidType()))
+    else if (lhq.withoutObjCGCAttr().withoutObjCLifetime().compatiblyIncludes(
+                 rhq.withoutObjCGCAttr().withoutObjCLifetime()) &&
+             (lhptee->isVoidType() || rhptee->isVoidType()))
       ; // keep old
 
     // Treat lifetime mismatches as fatal.
     else if (lhq.getObjCLifetime() != rhq.getObjCLifetime())
       ConvTy = Sema::IncompatiblePointerDiscardsQualifiers;
 
     // For GCC/MS compatibility, other qualifier mismatches are treated
     // as still compatible in C.

clang/test/Sema/address_space_type_casts_amdgpu.cl

+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -verify -pedantic -fsyntax-only %s
+
+void __builtins_AS_3(__attribute__((address_space(3))) int *);
+
+// Check calling a function using address space 3 (local for AMD) pointer works
+// with __local.
+__kernel void ker(__local int *IL) {
+  __builtins_AS_3(IL);
+}
+
+// Check casting __local to address space 3 (local for AMD) pointer works.
+__kernel void ker_2(__global int *Array, int N) {
+  __local int IL;
+  __attribute__((address_space(3))) int *I3;
+  I3 = (__attribute__((address_space(3))) int *)&IL;
+  Array[N] = *I3;
+}
+
+// Check casting __local to address space 5 (private for AMD) pointer errors.
+__kernel void ker_3(__global int *Array, int N) {
+  __local int IP;
+  __attribute__((address_space(5))) int *I5;
+  I5 = (__attribute__((address_space(5))) int *)&IP; // expected-error {{casting '__local int *' to type '__attribute__((address_space(5))) int *' changes address space of pointer}}
+  Array[N] = *I5;
+}
+
+// Check casting of address_space(3) to __generic pointer works.
+__kernel void ker_4(__global int *Array, int N, __attribute__((address_space(3))) int *AS3_ptr) {
+  __generic int *IG;
+  IG = AS3_ptr;
+}
+
+// Check casting of address_space(4) (__constant) to __generic pointer fails.
+__kernel void ker_5(__global int *Array, int N, __attribute__((address_space(4))) int *AS4_ptr) {
+  __generic int *IG;
+  IG = AS4_ptr; // expected-error {{assigning '__attribute__((address_space(4))) int *__private' to '__generic int *__private' changes address space of pointer}}
+}

clang/test/Sema/address_space_type_casts_default.cl

+// REQUIRES: x86-registered-target
+// RUN: %clang_cc1 -cl-std=CL2.0 -verify -pedantic -fsyntax-only %s
+
+// The same as address_space_type_cast_amdgpu.cl, but as x86 does not provide
+// ASMap all cases should error out.
+
+void __builtins_AS_3(__attribute__((address_space(3))) int *); // expected-note {{passing argument to parameter here}}
+
+// No relatioship between address_space(3) and __local on x86.
+__kernel void ker(__local int *IL) {
+  __builtins_AS_3(IL); // expected-error {{passing '__local int *__private' to parameter of type '__attribute__((address_space(3))) int *' changes address space of pointer}}
+}
+
+// No relatioship between address_space(3) and __local on x86.
+__kernel void ker_2(__global int *Array, int N) {
+  __local int IL;
+  __attribute__((address_space(3))) int *I3;
+  I3 = (__attribute__((address_space(3))) int *)&IL; // expected-error {{casting '__local int *' to type '__attribute__((address_space(3))) int *' changes address space of pointer}}
+  Array[N] = *I3;
+}
+
+// No relatioship between address_space(5) and __private on x86.
+__kernel void ker_3(__global int *Array, int N) {
+  __private int IP;
+  __attribute__((address_space(5))) int *I5;
+  I5 = (__attribute__((address_space(5))) int *)&IP; // expected-error {{casting '__private int *' to type '__attribute__((address_space(5))) int *' changes address space of pointer}}
+  Array[N] = *I5;
+}
+
+// Without ASMap compiler can't tell if address_space(3) is not equal to __constant, fail.
+__kernel void ker_4(__global int *Array, int N, __attribute__((address_space(3))) int *AS3_ptr) {
+  __generic int *IG;
+  IG = AS3_ptr; // expected-error {{assigning '__attribute__((address_space(3))) int *__private' to '__generic int *__private' changes address space of pointer}}
+}

clang/test/SemaOpenCL/atomic-ops.cl

void f(atomic_int *i, const atomic_int *ci,
   __opencl_atomic_fetch_min(i, 1, memory_order_seq_cst, memory_scope_work_group);
   __opencl_atomic_fetch_max(i, 1, memory_order_seq_cst, memory_scope_work_group);
   __opencl_atomic_fetch_min(f, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to atomic integer ('__generic atomic_float *' (aka '__generic _Atomic(float) *') invalid)}}
   __opencl_atomic_fetch_max(f, 1, memory_order_seq_cst, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to atomic integer ('__generic atomic_float *' (aka '__generic _Atomic(float) *') invalid)}}
 
   bool cmpexch_1 = __opencl_atomic_compare_exchange_strong(i, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
   bool cmpexch_2 = __opencl_atomic_compare_exchange_strong(p, P, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
   bool cmpexch_3 = __opencl_atomic_compare_exchange_strong(f, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{incompatible pointer types passing '__generic int *__private' to parameter of type '__generic float *'}}
-  (void)__opencl_atomic_compare_exchange_strong(i, CI, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{passing 'const __generic int *__private' to parameter of type '__generic int *' discards qualifiers}}
+  (void)__opencl_atomic_compare_exchange_strong(i, CI, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
 
   bool cmpexchw_1 = __opencl_atomic_compare_exchange_weak(i, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
   bool cmpexchw_2 = __opencl_atomic_compare_exchange_weak(p, P, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
   bool cmpexchw_3 = __opencl_atomic_compare_exchange_weak(f, I, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{incompatible pointer types passing '__generic int *__private' to parameter of type '__generic float *'}}
-  (void)__opencl_atomic_compare_exchange_weak(i, CI, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group); // expected-warning {{passing 'const __generic int *__private' to parameter of type '__generic int *' discards qualifiers}}
+  (void)__opencl_atomic_compare_exchange_weak(i, CI, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
 
   // Pointers to different address spaces are allowed.
   bool cmpexch_10 = __opencl_atomic_compare_exchange_strong((global atomic_int *)0x308, (constant int *)0x309, 1, memory_order_seq_cst, memory_order_seq_cst, memory_scope_work_group);
 
   __opencl_atomic_init(ci, 0); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const __generic atomic_int *' (aka 'const __generic _Atomic(int) *') invalid)}}
   __opencl_atomic_store(ci, 0, memory_order_release, memory_scope_work_group); // expected-error {{address argument to atomic operation must be a pointer to non-const _Atomic type ('const __generic atomic_int *' (aka 'const __generic _Atomic(int) *') invalid)}}
   __opencl_atomic_load(ci, memory_order_acquire, memory_scope_work_group);
 

clang/test/SemaOpenCL/numbered-address-space.cl

 // REQUIRES: amdgpu-registered-target
 // RUN: %clang_cc1 -cl-std=CL2.0 -triple amdgcn-amd-amdhsa -verify -pedantic -fsyntax-only %s
 
 void test_numeric_as_to_generic_implicit_cast(__attribute__((address_space(3))) int *as3_ptr, float src) {
-  generic int* generic_ptr = as3_ptr; // FIXME: This should error
+  generic int *generic_ptr = as3_ptr;
+}
+
+// AS 4 is constant on AMDGPU, casting it to generic is illegal.
+void test_numeric_as_const_to_generic_implicit_cast(__attribute__((address_space(4))) int *as4_ptr, float src) {
+  generic int *generic_ptr = as4_ptr; // expected-error{{initializing '__generic int *__private' with an expression of type '__attribute__((address_space(4))) int *__private' changes address space of pointer}}
 }
 
 void test_numeric_as_to_generic_explicit_cast(__attribute__((address_space(3))) int *as3_ptr, float src) {
-  generic int* generic_ptr = (generic int*) as3_ptr; // Should maybe be valid?
+  generic int *generic_ptr = (generic int *)as3_ptr;
 }
 
 void test_generic_to_numeric_as_implicit_cast(void) {
   generic int* generic_ptr = 0;
   __attribute__((address_space(3))) int *as3_ptr = generic_ptr; // expected-error{{initializing '__attribute__((address_space(3))) int *__private' with an expression of type '__generic int *__private' changes address space of pointer}}
 }
 
 void test_generic_to_numeric_as_explicit_cast(void) {
   generic int* generic_ptr = 0;
   __attribute__((address_space(3))) int *as3_ptr = (__attribute__((address_space(3))) int *)generic_ptr;
 }
 
 void test_generic_as_to_builtin_parameter_explicit_cast_numeric(__attribute__((address_space(3))) int *as3_ptr, float src) {
-  generic int* generic_ptr = as3_ptr; // FIXME: This should error
-  volatile float result = __builtin_amdgcn_ds_fmaxf((__attribute__((address_space(3))) float*) generic_ptr, src, 0, 0, false); // expected-error {{passing '__attribute__((address_space(3))) float *' to parameter of type '__local float *' changes address space of pointer}}
+  generic int *generic_ptr = as3_ptr;
+  // This is legal, as address_space(3) corresponds to local on amdgpu.
+  volatile float result = __builtin_amdgcn_ds_fmaxf((__attribute__((address_space(3))) float *)generic_ptr, src, 0, 0, false);
 }
 
 void test_generic_as_to_builtin_parameterimplicit_cast_numeric(__attribute__((address_space(3))) int *as3_ptr, float src) {
-  generic int* generic_ptr = as3_ptr;
+  generic int *generic_ptr = as3_ptr;
   volatile float result = __builtin_amdgcn_ds_fmaxf(generic_ptr, src, 0, 0, false); // expected-error {{passing '__generic int *__private' to parameter of type '__local float *' changes address space of pointer}}
 }
-

clang/test/SemaOpenCL/predefined-expr.cl

 // RUN: %clang_cc1 %s -verify
 // RUN: %clang_cc1 %s -verify -cl-std=CL2.0
 
 void f() {
-  char *f1 = __func__;          //expected-error-re{{initializing '{{__generic|__private}} char *__private' with an expression of type 'const __constant char *' changes address space of pointer}}
-  constant char *f2 = __func__; //expected-warning{{initializing '__constant char *__private' with an expression of type 'const __constant char[2]' discards qualifiers}}
+  char *f1 = __func__; // expected-error-re{{initializing '{{__generic|__private}} char *__private' with an expression of type 'const __constant char *' changes address space of pointer}}
+  constant char *f2 = __func__;
   constant const char *f3 = __func__;
 }

clang/test/SemaOpenCL/vector-conv.cl

void vector_conv_invalid(const global int4 *const_global_ptr) {
   int4 e = (int4)u; // expected-error{{invalid conversion between ext-vector type 'int4' (vector of 4 'int' values) and 'uint4' (vector of 4 'unsigned int' values)}}
 
   uint3 u4 = (uint3)u; // expected-error{{invalid conversion between ext-vector type 'uint3' (vector of 3 'unsigned int' values) and 'uint4' (vector of 4 'unsigned int' values)}}
 
   e = (const int4)i;
   e = (constant int4)i;
   e = (private int4)i;
 
-  private int4 *private_ptr = (const private int4 *)const_global_ptr; // expected-error{{casting 'const __global int4 *' to type 'const __private int4 *' changes address space of pointer}}
-  global int4 *global_ptr = const_global_ptr;                 // expected-warning {{initializing '__global int4 *__private' with an expression of type 'const __global int4 *__private' discards qualifiers}}
+private
+  int4 *private_ptr = (const private int4 *)const_global_ptr; // expected-error{{casting 'const __global int4 *' to type 'const __private int4 *' changes address space of pointer}}
+  global int4 *global_ptr = const_global_ptr;
   global_ptr = (global int4 *)const_global_ptr;
 }