1 // RUN: %clang_cc1 -verify=host -Rpass=openmp-opt -Rpass-analysis=openmp-opt -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
2 // RUN: %clang_cc1 -verify -Rpass=openmp-opt -Rpass-analysis=openmp-opt -fopenmp -O2 -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t.out
3 // RUN: %clang_cc1 -fexperimental-new-pass-manager -verify -Rpass=openmp-opt -Rpass-analysis=openmp-opt -fopenmp -O2 -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t.out
4
5 // host-no-diagnostics
6
7 void baz(void) __attribute__((assume("omp_no_openmp")));
8
bar(void)9 void bar(void) {
10 #pragma omp parallel // #1 \
11 // expected-remark@#1 {{Parallel region is used in unknown ways. Will not attempt to rewrite the state machine. [OMP101]}}
12 {
13 }
14 }
15
foo(void)16 void foo(void) {
17 #pragma omp target teams // #2
18 // expected-remark@#2 {{Rewriting generic-mode kernel with a customized state machine. [OMP131]}}
19 {
20 baz(); // expected-remark {{Value has potential side effects preventing SPMD-mode execution. Add `__attribute__((assume("ompx_spmd_amenable")))` to the called function to override. [OMP121]}}
21 #pragma omp parallel
22 {
23 }
24 bar();
25 #pragma omp parallel
26 {
27 }
28 }
29 }
30
spmd(void)31 void spmd(void) {
32 // Verify we do not emit the remarks above for "SPMD" regions.
33 #pragma omp target teams
34 #pragma omp parallel
35 {
36 }
37
38 #pragma omp target teams distribute parallel for
39 for (int i = 0; i < 100; ++i) {
40 }
41 }
42
43 // expected-remark@* {{OpenMP runtime call __kmpc_global_thread_num deduplicated. [OMP170]}}
44