From eeab0f4100ffa2fdcfb7de6a90236a64683b716b Mon Sep 17 00:00:00 2001 From: Fraser Cormack Date: Tue, 15 Apr 2014 11:38:29 +0000 Subject: [PATCH] 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 | 8 ++++++-- test/SemaOpenCL/array-parameters.cl | 6 ++++++ test/SemaOpenCL/invalid-kernel.cl | 2 ++ 3 files changed, 14 insertions(+), 2 deletions(-) create mode 100644 test/SemaOpenCL/array-parameters.cl diff --git a/lib/Sema/SemaDecl.cpp b/lib/Sema/SemaDecl.cpp index 948c58030a..cea1c26cf8 100644 --- 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 index 0000000000..200acacfad --- /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 df73eddc3b..9a50673ecd 100644 --- 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; } -- 2.40.0