[clang] fa9dd41 - [opencl] Fix address space deduction on array variables.

[Michael Liao via cfe-commits]

Author: Michael Liao
Date: 2019-12-04T09:37:50-05:00
New Revision: fa9dd410a9a9aa65ce6731cbe1ee12c5941eb3e8

URL: https://github.com/llvm/llvm-project/commit/fa9dd410a9a9aa65ce6731cbe1ee12c5941eb3e8
DIFF: https://github.com/llvm/llvm-project/commit/fa9dd410a9a9aa65ce6731cbe1ee12c5941eb3e8.diff

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

Summary:

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

Reviewers: Anastasia

Subscribers: yaxunl, cfe-commits

Tags: #clang

Differential Revision: https://reviews.llvm.org/D70981

Added: 
    

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

Removed: 
    


################################################################################
diff  --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp
index d35037273106..660be458a698 100644
--- a/clang/lib/Sema/SemaDecl.cpp
+++ b/clang/lib/Sema/SemaDecl.cpp
@@ -6128,7 +6128,26 @@ void Sema::deduceOpenCLAddressSpace(ValueDecl *Decl) {
     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);
   }
 }

diff  --git a/clang/test/SemaOpenCL/address-spaces.cl b/clang/test/SemaOpenCL/address-spaces.cl
index 55a55dc75050..09a6dd0ba53f 100644
--- a/clang/test/SemaOpenCL/address-spaces.cl
+++ b/clang/test/SemaOpenCL/address-spaces.cl
@@ -241,3 +241,10 @@ void func_multiple_addr(void) {
   __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);
+}