1*08a22076SJohannes Doerfert // RUN: %clang_cc1 -verify=expected -fopenmp -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized
2*08a22076SJohannes Doerfert
bad()3*08a22076SJohannes Doerfert void bad() {
4*08a22076SJohannes Doerfert #pragma omp target data ompx_attribute() // expected-error {{unexpected OpenMP clause 'ompx_attribute' in directive '#pragma omp target data'}}
5*08a22076SJohannes Doerfert #pragma omp target data ompx_attribute(__attribute__((launch_bounds(1, 2)))) // expected-error {{unexpected OpenMP clause 'ompx_attribute' in directive '#pragma omp target data'}} expected-error {{expected at least one 'map', 'use_device_ptr', or 'use_device_addr' clause for '#pragma omp target data'}}
6*08a22076SJohannes Doerfert
7*08a22076SJohannes Doerfert #pragma omp target ompx_attribute()
8*08a22076SJohannes Doerfert {}
9*08a22076SJohannes Doerfert #pragma omp target ompx_attribute(__attribute__(()))
10*08a22076SJohannes Doerfert {}
11*08a22076SJohannes Doerfert #pragma omp target ompx_attribute(__attribute__((pure))) // expected-warning {{'ompx_attribute' clause only allows 'amdgpu_flat_work_group_size', 'amdgpu_waves_per_eu', and 'launch_bounds'; 'pure' is ignored}}
12*08a22076SJohannes Doerfert {}
13*08a22076SJohannes Doerfert #pragma omp target ompx_attribute(__attribute__((pure,amdgpu_waves_per_eu(1, 2), const))) // expected-warning {{'ompx_attribute' clause only allows 'amdgpu_flat_work_group_size', 'amdgpu_waves_per_eu', and 'launch_bounds'; 'pure' is ignored}} expected-warning {{'ompx_attribute' clause only allows 'amdgpu_flat_work_group_size', 'amdgpu_waves_per_eu', and 'launch_bounds'; 'const' is ignored}}
14*08a22076SJohannes Doerfert {}
15*08a22076SJohannes Doerfert #pragma omp target ompx_attribute(__attribute__((amdgpu_waves_per_eu()))) // expected-error {{'amdgpu_waves_per_eu' attribute takes at least 1 argument}}
16*08a22076SJohannes Doerfert {}
17*08a22076SJohannes Doerfert #pragma omp target ompx_attribute(__attribute__((amdgpu_waves_per_eu(1, 2, 3)))) // expected-error {{'amdgpu_waves_per_eu' attribute takes no more than 2 arguments}}
18*08a22076SJohannes Doerfert {}
19*08a22076SJohannes Doerfert #pragma omp target ompx_attribute(__attribute__((amdgpu_flat_work_group_size(1)))) // expected-error {{'amdgpu_flat_work_group_size' attribute requires exactly 2 arguments}}
20*08a22076SJohannes Doerfert {}
21*08a22076SJohannes Doerfert #pragma omp target ompx_attribute(__attribute__((amdgpu_flat_work_group_size(1, 2, 3,)))) // expected-error {{expected expression}}
22*08a22076SJohannes Doerfert {}
23*08a22076SJohannes Doerfert #pragma omp target ompx_attribute([[clang::amdgpu_waves_per_eu(1, 2, 3)]]) // expected-error {{'amdgpu_waves_per_eu' attribute takes no more than 2 arguments}}
24*08a22076SJohannes Doerfert {}
25*08a22076SJohannes Doerfert #pragma omp target ompx_attribute([[clang::unknown]]) // expected-warning {{'ompx_attribute' clause only allows 'amdgpu_flat_work_group_size', 'amdgpu_waves_per_eu', and 'launch_bounds'; 'unknown' is ignored}}
26*08a22076SJohannes Doerfert {}
27*08a22076SJohannes Doerfert #pragma omp target ompx_attribute(baz) // expected-error {{expected ')'}} expected-note {{to match this '('}}
28*08a22076SJohannes Doerfert {}
29*08a22076SJohannes Doerfert #pragma omp target ompx_attribute(__attribute__((launch_bounds(1))))
30*08a22076SJohannes Doerfert {}
31*08a22076SJohannes Doerfert #pragma omp target ompx_attribute(__attribute__((launch_bounds(bad)))) // expected-error {{'launch_bounds' attribute requires parameter 0 to be an integer constant}}
32*08a22076SJohannes Doerfert {}
33*08a22076SJohannes Doerfert #pragma omp target ompx_attribute(__attribute__((launch_bounds(1, // expected-error {{expected expression}} expected-error {{expected ')'}} expected-error {{expected ')'}} expected-error {{expected ')'}} expected-note {{to match this '('}}
34*08a22076SJohannes Doerfert {}
35*08a22076SJohannes Doerfert #pragma omp target ompx_attribute(__attribute__((launch_bounds(1, 2 // expected-error {{expected ')'}} expected-error {{expected ')'}} expected-error {{expected ')'}} expected-error {{expected ')'}} expected-note {{to match this '('}}
36*08a22076SJohannes Doerfert {}
37*08a22076SJohannes Doerfert #pragma omp target ompx_attribute(__attribute__((launch_bounds(1, 2) // expected-error {{expected ')'}} expected-error {{expected ')'}} expected-error {{expected ')'}} expected-note {{to match this '('}}
38*08a22076SJohannes Doerfert {}
39*08a22076SJohannes Doerfert #pragma omp target ompx_attribute(__attribute__((launch_bounds(1, 2)) // expected-error {{expected ')'}} expected-error {{expected ')'}} expected-note {{to match this '('}}
40*08a22076SJohannes Doerfert {}
41*08a22076SJohannes Doerfert #pragma omp target ompx_attribute(__attribute__((launch_bounds(1, 2))) // expected-error {{expected ')'}} expected-note {{to match this '('}}
42*08a22076SJohannes Doerfert {}
43*08a22076SJohannes Doerfert #pragma omp target ompx_attribute(__attribute__((launch_bounds(1, -3)))) // expected-warning {{'launch_bounds' attribute parameter 1 is negative and will be ignored}}
44*08a22076SJohannes Doerfert {}
45*08a22076SJohannes Doerfert #pragma omp target ompx_attribute(__attribute__((amdgpu_waves_per_eu(10, 1)))) // expected-error {{'amdgpu_waves_per_eu' attribute argument is invalid: min must not be greater than max}}
46*08a22076SJohannes Doerfert {}
47*08a22076SJohannes Doerfert }
48