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=all,safe -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=all,safe -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
bar1(void)9 void bar1(void) {
10 #pragma omp parallel // #0
11 // safe-remark@#0 {{Parallel region is used in unknown ways. Will not attempt to rewrite the state machine. [OMP101]}}
12 {
13 }
14 }
bar2(void)15 void bar2(void) {
16 #pragma omp parallel // #1
17 // safe-remark@#1 {{Parallel region is used in unknown ways. Will not attempt to rewrite the state machine. [OMP101]}}
18 {
19 }
20 }
21
foo1(void)22 void foo1(void) {
23 #pragma omp target teams // #2
24 // all-remark@#2 {{Rewriting generic-mode kernel with a customized state machine. [OMP131]}}
25
26 {
27 baz(); // all-remark {{Value has potential side effects preventing SPMD-mode execution. Add `__attribute__((assume("ompx_spmd_amenable")))` to the called function to override. [OMP121]}}
28 #pragma omp parallel // #3
29 {
30 }
31 bar1();
32 #pragma omp parallel // #4
33 {
34 }
35 }
36 }
37
foo2(void)38 void foo2(void) {
39 #pragma omp target teams // #5
40 // all-remark@#5 {{Rewriting generic-mode kernel with a customized state machine. [OMP131]}}
41 {
42 baz(); // all-remark {{Value has potential side effects preventing SPMD-mode execution. Add `__attribute__((assume("ompx_spmd_amenable")))` to the called function to override. [OMP121]}}
43 #pragma omp parallel // #6
44 {
45 }
46 bar1();
47 bar2();
48 #pragma omp parallel // #7
49 {
50 }
51 bar1();
52 bar2();
53 }
54 }
55
foo3(void)56 void foo3(void) {
57 #pragma omp target teams // #8
58 // all-remark@#8 {{Rewriting generic-mode kernel with a customized state machine. [OMP131]}}
59 {
60 baz(); // all-remark {{Value has potential side effects preventing SPMD-mode execution. Add `__attribute__((assume("ompx_spmd_amenable")))` to the called function to override. [OMP121]}}
61 #pragma omp parallel // #9
62 {
63 }
64 bar1();
65 bar2();
66 #pragma omp parallel // #10
67 {
68 }
69 bar1();
70 bar2();
71 }
72 }
73
spmd(void)74 void spmd(void) {
75 // Verify we do not emit the remarks above for "SPMD" regions.
76 #pragma omp target teams
77 #pragma omp parallel
78 {
79 }
80
81 #pragma omp target teams distribute parallel for
82 for (int i = 0; i < 100; ++i) {
83 }
84 }
85
86 // all-remark@* 9 {{OpenMP runtime call __kmpc_global_thread_num deduplicated. [OMP170]}}
87