[PATCH] D70981: [opencl] Fix address space deduction on array variables.

[Michael Liao via Phabricator via cfe-commits]

hliao created this revision.
hliao added a reviewer: Anastasia.
Herald added subscribers: cfe-commits, yaxunl.
Herald added a project: clang.

- The deduced address space needs applying to its element type as well.


Repository:
  rG LLVM Github Monorepo

https://reviews.llvm.org/D70981

Files:
  clang/lib/Sema/SemaDecl.cpp
  clang/test/SemaOpenCL/address-spaces.cl


Index: clang/test/SemaOpenCL/address-spaces.cl
===================================================================
--- clang/test/SemaOpenCL/address-spaces.cl
+++ clang/test/SemaOpenCL/address-spaces.cl
@@ -241,3 +241,10 @@
   __private private_int_t var5; // expected-warning {{multiple identical address spaces specified for type}}
   __private private_int_t *var6;// expected-warning {{multiple identical address spaces specified for type}}
 }
+
+void func_with_array_param(const unsigned data[16]);
+
+__kernel void k() {
+  unsigned data[16];
+  func_with_array_param(data);
+}
Index: clang/lib/Sema/SemaDecl.cpp
===================================================================
--- clang/lib/Sema/SemaDecl.cpp
+++ clang/lib/Sema/SemaDecl.cpp
@@ -6128,7 +6128,23 @@
     if ((getLangOpts().OpenCLCPlusPlus || getLangOpts().OpenCLVersion >= 200) &&
         Var->hasGlobalStorage())
       ImplAS = LangAS::opencl_global;
+    // If the original type from a decayed type is an array type and that array
+    // type has no address space yet, deduce it now.
+    if (auto DT = dyn_cast<DecayedType>(Type)) {
+      auto OrigTy = DT->getOriginalType();
+      if (!OrigTy.getQualifiers().hasAddressSpace() && OrigTy->isArrayType()) {
+        OrigTy = Context.getAddrSpaceQualType(OrigTy, ImplAS);
+        OrigTy = QualType(Context.getAsArrayType(OrigTy), 0);
+        Type = Context.getDecayedType(OrigTy);
+      }
+    }
     Type = Context.getAddrSpaceQualType(Type, ImplAS);
+    // Apply any qualifiers (including address space) from the array type to
+    // the element type. This implements C99 6.7.3p8: "If the specification of
+    // an array type includes any type qualifiers, the element type is so
+    // qualified, not the array type."
+    if (Type->isArrayType())
+      Type = QualType(Context.getAsArrayType(Type), 0);
     Decl->setType(Type);
   }
 }


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D70981.231970.patch
Type: text/x-patch
Size: 1879 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20191203/aa702ef9/attachment-0001.bin>