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

diff --git a/lib/Sema/SemaDecl.cpp b/lib/Sema/SemaDecl.cpp
index 948c58030af28e17ec193ff749c08cbf268c8d36..cea1c26cf825d23a0955101869124fac20076a02 100644 (file)
--- a/lib/Sema/SemaDecl.cpp
+++ b/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;

diff --git a/test/SemaOpenCL/array-parameters.cl b/test/SemaOpenCL/array-parameters.cl
new file mode 100644 (file)
index 0000000..200acac
--- /dev/null
+++ b/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[]) { }

diff --git a/test/SemaOpenCL/invalid-kernel.cl b/test/SemaOpenCL/invalid-kernel.cl
index df73eddc3b5dee1ce81fb2bac24541a32147b7c1..9a50673ecdfaebdea141f0ab36f78135fa2a4306 100644 (file)
--- a/test/SemaOpenCL/invalid-kernel.cl
+++ b/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;
 }