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

[Matt Arsenault via Phabricator via cfe-commits]

arsenm 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")
----------------
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


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


================
Comment at: lib/CodeGen/CGBuiltin.cpp:3500
+        if (auto *PtrTy = dyn_cast<llvm::PointerType>(PTy)) {
+          if (PtrTy->getAddressSpace() !=
+              ArgValue->getType()->getPointerAddressSpace()) {
----------------
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.


================
Comment at: test/CodeGenOpenCL/numbered-address-space.cl:36
+#if 0
+// XXX: Should this compile?
+void test_generic_as_to_builtin_parameter_explicit_cast_numeric(__attribute__((address_space(3))) int *local_ptr, float src) {
----------------
Anastasia wrote:
> `__attribute__((address_space(N)))` is not an OpenCL feature and I think it's not specified in C either? But I think generally non matching address spaces don't compile in Clang. So it might be useful to disallow this?
I'm pretty sure it's a C extension. The way things seem to work now is address spaces are accepted anywhere and everywhere.


https://reviews.llvm.org/D47154