Allow address space qualifiers on OpenCL array parameters

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@206275 91177308-0d34-0410-b5e6-96231b3b80d8

lib/Sema/SemaDecl.cpp

@@ -9469,8 +9469,12 @@ ParmVarDecl *Sema::CheckParameter(DeclContext *DC, SourceLocation StartLoc,
   // Since all parameters have automatic store duration, they can not have
   // an address space.
   if (T.getAddressSpace() != 0) {
-    Diag(NameLoc, diag::err_arg_with_address_space);
-    New->setInvalidDecl();
+    // OpenCL allows function arguments declared to be an array of a type
+    // to be qualified with an address space.
+    if (!(getLangOpts().OpenCL && T->isArrayType())) {
+      Diag(NameLoc, diag::err_arg_with_address_space);
+      New->setInvalidDecl();
+    }
   }   
 
   return New;

test/SemaOpenCL/array-parameters.cl

@@ -0,0 +1,6 @@
+// RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
+// expected-no-diagnostics
+
+kernel void foo(global int a[], local int b[], constant int c[4]) { }
+
+void bar(global int a[], local int b[], constant int c[4], int d[]) { }

test/SemaOpenCL/invalid-kernel.cl

@@ -4,6 +4,8 @@ kernel void no_ptrptr(global int **i) { } // expected-error{{kernel parameter ca
 
 __kernel void no_privateptr(__private int *i) { } // expected-error {{kernel parameter cannot be declared as a pointer to the __private address space}}
 
+__kernel void no_privatearray(__private int i[]) { } // expected-error {{kernel parameter cannot be declared as a pointer to the __private address space}}
+
 kernel int bar()  { // expected-error {{kernel must have void return type}}
   return 6;
 }