[opencl] Fix address space deduction on array variables.

[hliao]

Summary

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

[Anastasia]

LGTM! Thanks! Please address the comment in your final commit.

clang/lib/Sema/SemaDecl.cpp

6136

Maybe we should add a comment that we are adding an addr space on an array type first and then propagate to the element type.

clang/lib/Sema/SemaDecl.cpp

void Sema::deduceOpenCLAddressSpace(ValueDecl *Decl) {
   if (VarDecl *Var = dyn_cast<VarDecl>(Decl)) {
     QualType Type = Var->getType();
     if (Type->isSamplerT() || Type->isVoidType())
       return;
     LangAS ImplAS = LangAS::opencl_private;
     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);
   }
 }
 
 static void checkAttributesAfterMerging(Sema &S, NamedDecl &ND) {
   // Ensure that an auto decl is deduced otherwise the checks below might cache
   // the wrong linkage.
   assert(S.ParsingInitForAutoVars.count(&ND) == 0);

clang/test/SemaOpenCL/address-spaces.cl

void func_multiple_addr(void) {
   typedef __private int private_int_t;
   __private __local int var1;   // expected-error {{multiple address spaces specified for type}}
   __private __local int *var2;  // expected-error {{multiple address spaces specified for type}}
   __local private_int_t var3;   // expected-error {{multiple address spaces specified for type}}
   __local private_int_t *var4;  // expected-error {{multiple address spaces specified for type}}
   __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);
+}