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; }
23
24int main() {
25 int device_var = 1;
26#pragma omp target map(tofrom : device_var)
27 {
28#pragma omp parallel for
29 for (int i = 1; i <= 10; i++) {
30 device_var *= i;
31 if (i % 2 == 0) {
32 device_var += test1(a: device_var);
33 }
34 }
35 }
36}
37
38// clang-format off
39// LLVM-PGO-LABEL: __omp_offloading_{{[_0-9a-zA-Z]*}}_main_{{[_0-9a-zA-Z]*}}:
40// LLVM-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
41// LLVM-PGO: Counters: 2
42// LLVM-PGO: Block counts: [0, {{.*}}]
43
44// LLVM-PGO-LABEL: __omp_offloading_{{[_0-9a-zA-Z]*}}_main_{{[_0-9a-zA-Z]*}}_omp_outlined:
45// LLVM-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
46// LLVM-PGO: Counters: 5
47// LLVM-PGO: Block counts: [10, 5, {{.*}}, 10, {{.*}}]
48
49// LLVM-PGO-LABEL: test1:
50// LLVM-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
51// LLVM-PGO: Counters: 1
52// LLVM-PGO: Block counts: [5]
53
54// LLVM-PGO-LABEL: Instrumentation level:
55// LLVM-PGO-SAME: IR
56// LLVM-PGO-SAME: entry_first = 0
57// LLVM-PGO-LABEL: Functions shown:
58// LLVM-PGO-SAME: 3
59// LLVM-PGO-LABEL: Maximum function count:
60// LLVM-PGO-SAME: 10
61
62// CLANG-PGO-LABEL: __omp_offloading_{{[_0-9a-zA-Z]*}}_main_{{[_0-9a-zA-Z]*}}:
63// CLANG-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
64// CLANG-PGO: Counters: 1
65// CLANG-PGO: Function count: {{.*}}
66// CLANG-PGO: Block counts: []
67
68// CLANG-PGO-LABEL: __omp_offloading_{{[_0-9a-zA-Z]*}}_main_{{[_0-9a-zA-Z]*}}_omp_outlined:
69// CLANG-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
70// CLANG-PGO: Counters: 3
71// CLANG-PGO: Function count: {{.*}}
72// CLANG-PGO: Block counts: [{{.*}}, 5]
73
74// CLANG-PGO-LABEL: test1:
75// CLANG-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
76// CLANG-PGO: Counters: 1
77// CLANG-PGO: Function count: 5
78// CLANG-PGO: Block counts: []
79
80// CLANG-PGO-LABEL: Instrumentation level:
81// CLANG-PGO-SAME: Front-end
82// CLANG-PGO-LABEL: Functions shown:
83// CLANG-PGO-SAME: 3
84// clang-format on
85

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