r206275 - Allow address space qualifiers on OpenCL array parameters

[Fraser Cormack]

Author: fcormack
Date: Tue Apr 15 06:38:29 2014
New Revision: 206275

URL: http://llvm.org/viewvc/llvm-project?rev=206275&view=rev
Log:
Allow address space qualifiers on OpenCL array parameters

Added:
    cfe/trunk/test/SemaOpenCL/array-parameters.cl
Modified:
    cfe/trunk/lib/Sema/SemaDecl.cpp
    cfe/trunk/test/SemaOpenCL/invalid-kernel.cl

Modified: cfe/trunk/lib/Sema/SemaDecl.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDecl.cpp?rev=206275&r1=206274&r2=206275&view=diff
==============================================================================
--- cfe/trunk/lib/Sema/SemaDecl.cpp (original)
+++ cfe/trunk/lib/Sema/SemaDecl.cpp Tue Apr 15 06:38:29 2014
@@ -9469,8 +9469,12 @@ ParmVarDecl *Sema::CheckParameter(DeclCo
   // 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;

Added: cfe/trunk/test/SemaOpenCL/array-parameters.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCL/array-parameters.cl?rev=206275&view=auto
==============================================================================
--- cfe/trunk/test/SemaOpenCL/array-parameters.cl (added)
+++ cfe/trunk/test/SemaOpenCL/array-parameters.cl Tue Apr 15 06:38:29 2014
@@ -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[]) { }

Modified: cfe/trunk/test/SemaOpenCL/invalid-kernel.cl
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCL/invalid-kernel.cl?rev=206275&r1=206274&r2=206275&view=diff
==============================================================================
--- cfe/trunk/test/SemaOpenCL/invalid-kernel.cl (original)
+++ cfe/trunk/test/SemaOpenCL/invalid-kernel.cl Tue Apr 15 06:38:29 2014
@@ -4,6 +4,8 @@ kernel void no_ptrptr(global int **i) {
 
 __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;
 }