1// RUN: %libomptarget-compile-generic -fcreate-profile \
2// RUN: -Xarch_device -fprofile-generate \
3// RUN: -Xarch_device -fprofile-update=atomic
4// RUN: env LLVM_PROFILE_FILE=%basename_t.llvm.profraw \
5// RUN: %libomptarget-run-generic 2>&1
6// RUN: llvm-profdata show --all-functions --counts \
7// RUN: %target_triple.%basename_t.llvm.profraw | \
8// RUN: %fcheck-generic --check-prefix="LLVM-PGO"
9
10// RUN: %libomptarget-compile-generic -fcreate-profile \
11// RUN: -Xarch_device -fprofile-instr-generate \
12// RUN: -Xarch_device -fprofile-update=atomic
13// RUN: env LLVM_PROFILE_FILE=%basename_t.clang.profraw \
14// RUN: %libomptarget-run-generic 2>&1
15// RUN: llvm-profdata show --all-functions --counts \
16// RUN: %target_triple.%basename_t.clang.profraw | \
17// RUN: %fcheck-generic --check-prefix="CLANG-PGO"
18
19// REQUIRES: gpu
20// REQUIRES: pgo
21
22int test1(int a) { return a / 2; }
23int test2(int a) { return a * 2; }
24
25int main() {
26 int device_var = 1;
27
28#pragma omp target teams distribute parallel for num_teams(3) \
29 map(tofrom : device_var)
30 for (int i = 1; i <= 30; i++) {
31 device_var *= i;
32 if (i % 2 == 0) {
33 device_var += test1(a: device_var);
34 }
35 if (i % 3 == 0) {
36 device_var += test2(a: device_var);
37 }
38 }
39}
40
41// clang-format off
42// LLVM-PGO-LABEL: __omp_offloading_{{[_0-9a-zA-Z]*}}_main_{{[_0-9a-zA-Z]*}}:
43// LLVM-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
44// LLVM-PGO: Counters: 2
45// LLVM-PGO: Block counts: [0, {{.*}}]
46
47// LLVM-PGO-LABEL: __omp_offloading_{{[_0-9a-zA-Z]*}}_main_{{[_0-9a-zA-Z]*}}_omp_outlined:
48// LLVM-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
49// LLVM-PGO: Counters: 4
50// LLVM-PGO: Block counts: [{{.*}}, 0, {{.*}}, 0]
51
52// LLVM-PGO-LABEL: __omp_offloading_{{[_0-9a-zA-Z]*}}_main_{{[_0-9a-zA-Z]*}}_omp_outlined_omp_outlined:
53// LLVM-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
54// LLVM-PGO: Counters: 4
55// LLVM-PGO: Block counts: [30, 15, 10, {{.*}}]
56
57// LLVM-PGO-LABEL: test1:
58// LLVM-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
59// LLVM-PGO: Counters: 1
60// LLVM-PGO: Block counts: [15]
61
62// LLVM-PGO-LABEL: test2:
63// LLVM-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
64// LLVM-PGO: Counters: 1
65// LLVM-PGO: Block counts: [10]
66
67// LLVM-PGO-LABEL: Instrumentation level:
68// LLVM-PGO-SAME: IR
69
70// CLANG-PGO-LABEL: __omp_offloading_{{[_0-9a-zA-Z]*}}_main_{{[_0-9a-zA-Z]*}}:
71// CLANG-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
72// CLANG-PGO: Counters: 1
73// CLANG-PGO: Function count: {{.*}}
74// CLANG-PGO: Block counts: []
75
76// CLANG-PGO-LABEL: __omp_offloading_{{[_0-9a-zA-Z]*}}_main_{{[_0-9a-zA-Z]*}}_omp_outlined:
77// CLANG-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
78// CLANG-PGO: Counters: 1
79// CLANG-PGO: Function count: {{.*}}
80// CLANG-PGO: Block counts: []
81
82// CLANG-PGO-LABEL: __omp_offloading_{{[_0-9a-zA-Z]*}}_main_{{[_0-9a-zA-Z]*}}_omp_outlined_omp_outlined:
83// CLANG-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
84// CLANG-PGO: Counters: 4
85// CLANG-PGO: Function count: 30
86// CLANG-PGO: Block counts: [{{.*}}, 15, 10]
87
88// CLANG-PGO-LABEL: test1:
89// CLANG-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
90// CLANG-PGO: Counters: 1
91// CLANG-PGO: Function count: 15
92// CLANG-PGO: Block counts: []
93
94// CLANG-PGO-LABEL: test2:
95// CLANG-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
96// CLANG-PGO: Counters: 1
97// CLANG-PGO: Function count: 10
98// CLANG-PGO: Block counts: []
99
100// CLANG-PGO-LABEL: Instrumentation level:
101// CLANG-PGO-SAME: Front-end
102// clang-format on
103

source code of offload/test/offloading/gpupgo/pgo_atomic_teams.c