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-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t.out
3 // 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-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o %t.out
4
5 // host-no-diagnostics
6
7 [[omp::assume("omp_no_openmp")]] void baz(void);
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 `[[omp::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 `[[omp::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 `[[omp::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 #pragma omp begin declare target device_type(nohost)
87 struct KernelEnvironmentTy;
88 struct KernelLaunchEnvironmentTy;
89 __attribute__((weak))
__kmpc_target_init(struct KernelEnvironmentTy *,struct KernelLaunchEnvironmentTy *)90 extern "C" int __kmpc_target_init(struct KernelEnvironmentTy *, struct KernelLaunchEnvironmentTy *) { // all-remark {{Could not internalize function. Some optimizations may not be possible. [OMP140]}}
91 return 0;
92 }
93 #pragma omp end declare target
94
95 // all-remark@* 9 {{OpenMP runtime call __kmpc_global_thread_num deduplicated. [OMP170]}}
96