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