[PATCH] D47154: Try to make builtin address space declarations not useless

[Anastasia Stulova via Phabricator via cfe-commits]

Anastasia added inline comments.


================
Comment at: include/clang/Basic/BuiltinsAMDGPU.def:49
+
+// FIXME: Need to disallow constant address space.
 BUILTIN(__builtin_amdgcn_div_scale, "dddbb*", "n")
----------------
arsenm wrote:
> Anastasia wrote:
> > Do you plan to provide the support for it later? Or if else perhaps we should elaborate more what's to be done.
> I'm not sure. I don't know how to best enforce this
The only way I guess if we list overloads with each address space explicitly (apart from `constant`)? Or may be with the use of `generic` AS, although that will only work for CL2.0.


================
Comment at: lib/AST/ASTContext.cpp:9093
       unsigned AddrSpace = strtoul(Str, &End, 10);
-      if (End != Str && AddrSpace != 0) {
-        Type = Context.getAddrSpaceQualType(Type,
-                                            getLangASFromTargetAS(AddrSpace));
+      if (End != Str) {
+        // Note AddrSpace == 0 is not the same as an unspecified address space.
----------------
arsenm wrote:
> Anastasia wrote:
> > Could we check against LangAS::Default instead of removing this completely.
> I don't think that really make sense, since that would be leaving this the same. I don't really need it for this patch, but I fundamentally think specifying address space 0 is different from an unspecified address space. According to the description for builtins, if no address space is specified than any address space will be accepted. This is different from a builtin requiring address space 0
I thought `Default` AS was meant to be for the case no AS is specified but I guess it doesn't work the same in Builtins specification syntax.


================
Comment at: lib/CodeGen/CGBuiltin.cpp:3500
+        if (auto *PtrTy = dyn_cast<llvm::PointerType>(PTy)) {
+          if (PtrTy->getAddressSpace() !=
+              ArgValue->getType()->getPointerAddressSpace()) {
----------------
arsenm wrote:
> Anastasia wrote:
> > Would this be correct for OpenCL? Should we use  `isAddressSpaceSupersetOf` helper instead? Would it also sort the issue with constant AS (at least for OpenCL)? 
> The issue I mentioned for the other builtin is that it modifies the memory, and doesn't have to do with the casting.
> 
> At this point the AddrSpaceCast has to be emitted. The checking if the cast is legal I guess would be in the SemaExpr part. I know at one point I was trying to use isAddressSpaceSupersetOf in rewriteBuiltinFunctionDecl, but there was some problem with that. I think it didn't make sense with the magic where the builtin without an address space is supposed to accept any address space or something along those lines.
Yes, I think Sema has to check it before indeed. I am not sure it works right with OpenCL rules though for the Builtin functions.  Would it make sense to add a negative test for this then? 


https://reviews.llvm.org/D47154