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

[Michael Liao via Phabricator via cfe-commits]

This revision was automatically updated to reflect the committed changes.
Closed by commit rGfa9dd410a9a9: [opencl] Fix address space deduction on array variables. (authored by hliao).

Changed prior to commit:
  https://reviews.llvm.org/D70981?vs=231970&id=232123#toc

Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D70981/new/

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,26 @@
     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()) {
+        // Add the address space to the original array type and then propagate
+        // that to the element type through `getAsArrayType`.
+        OrigTy = Context.getAddrSpaceQualType(OrigTy, ImplAS);
+        OrigTy = QualType(Context.getAsArrayType(OrigTy), 0);
+        // Re-generate the decayed type.
+        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.232123.patch
Type: text/x-patch
Size: 2064 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20191204/39b5a3f8/attachment-0001.bin>