diff options
| author | Justin Lebar <jlebar@google.com> | 2016-01-20 00:26:57 +0000 |
|---|---|---|
| committer | Justin Lebar <jlebar@google.com> | 2016-01-20 00:26:57 +0000 |
| commit | c66a10652a9404d1d196864fc1b2e9aa413307b8 (patch) | |
| tree | 232fcbb0030c005f33eb22a12210d739dfc22dad /clang/test/SemaCUDA/bad-attributes.cu | |
| parent | 23c4d83aa310903484cb80d2ab8197444a96d2e0 (diff) | |
| download | bcm5719-llvm-c66a10652a9404d1d196864fc1b2e9aa413307b8.tar.gz bcm5719-llvm-c66a10652a9404d1d196864fc1b2e9aa413307b8.zip | |
[CUDA] Only allow __global__ on free functions and static member functions.
Summary:
Warn for NVCC compatibility if you declare a static member function or
inline function as __global__.
Reviewers: tra
Subscribers: jhen, echristo, cfe-commits
Differential Revision: http://reviews.llvm.org/D16261
llvm-svn: 258263
Diffstat (limited to 'clang/test/SemaCUDA/bad-attributes.cu')
| -rw-r--r-- | clang/test/SemaCUDA/bad-attributes.cu | 16 |
1 files changed, 14 insertions, 2 deletions
diff --git a/clang/test/SemaCUDA/bad-attributes.cu b/clang/test/SemaCUDA/bad-attributes.cu index 7e01e141de1..4cb43e240b0 100644 --- a/clang/test/SemaCUDA/bad-attributes.cu +++ b/clang/test/SemaCUDA/bad-attributes.cu @@ -4,8 +4,8 @@ // // You should be able to run this file through nvcc for compatibility testing. // -// RUN: %clang_cc1 -fsyntax-only -verify %s -// RUN: %clang_cc1 -fcuda-is-device -fsyntax-only -verify %s +// RUN: %clang_cc1 -fsyntax-only -Wcuda-compat -verify -DEXPECT_INLINE_WARNING %s +// RUN: %clang_cc1 -fcuda-is-device -fsyntax-only -Wcuda-compat -verify %s #include "Inputs/cuda.h" @@ -47,3 +47,15 @@ __global__ __device__ void z11(); // expected-error {{attributes are not compat // expected-note@-1 {{conflicting attribute is here}} __global__ __host__ void z12(); // expected-error {{attributes are not compatible}} // expected-note@-1 {{conflicting attribute is here}} + +struct S { + __global__ void foo() {}; // expected-error {{must be a free function or static member function}} + __global__ static void bar(); // expected-warning {{kernel function 'bar' is a member function}} + // Although this is implicitly inline, we shouldn't warn. + __global__ static void baz() {}; // expected-warning {{kernel function 'baz' is a member function}} +}; + +__global__ static inline void foobar() {}; +#ifdef EXPECT_INLINE_WARNING +// expected-warning@-2 {{ignored 'inline' attribute on kernel function 'foobar'}} +#endif |
