Skip to content

Commit 0c48346

Browse files
EthanLuisMcDonoughEthanLuisMcDonough
committed
[PGO][Offload] Update GPU PGO tests
1 parent 10cb0f6 commit 0c48346

File tree

4 files changed

+193
-12
lines changed

4 files changed

+193
-12
lines changed

‎offload/test/offloading/gpupgo/pgo1.c‎

Lines changed: 2 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1,15 +1,13 @@
11
// RUN: %libomptarget-compile-generic -fcreate-profile \
2-
// RUN: -Xarch_device -fprofile-generate \
3-
// RUN: -Xarch_device -fprofile-update=atomic
2+
// RUN: -Xarch_device -fprofile-generate
43
// RUN: env LLVM_PROFILE_FILE=%basename_t.llvm.profraw \
54
// RUN: %libomptarget-run-generic 2>&1
65
// RUN: llvm-profdata show --all-functions --counts \
76
// RUN: %target_triple.%basename_t.llvm.profraw | \
87
// RUN: %fcheck-generic --check-prefix="LLVM-PGO"
98

109
// RUN: %libomptarget-compile-generic -fcreate-profile \
11-
// RUN: -Xarch_device -fprofile-instr-generate \
12-
// RUN: -Xarch_device -fprofile-update=atomic
10+
// RUN: -Xarch_device -fprofile-instr-generate
1311
// RUN: env LLVM_PROFILE_FILE=%basename_t.clang.profraw \
1412
// RUN: %libomptarget-run-generic 2>&1
1513
// RUN: llvm-profdata show --all-functions --counts \

‎offload/test/offloading/gpupgo/pgo2.c‎

Lines changed: 5 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,4 @@
1-
// RUN: %libomptarget-compile-generic -fprofile-generate \
2-
// RUN: -fprofile-update=atomic
1+
// RUN: %libomptarget-compile-generic -fprofile-generate
32
// RUN: env LLVM_PROFILE_FILE=%basename_t.llvm.profraw \
43
// RUN: %libomptarget-run-generic 2>&1
54
// RUN: llvm-profdata show --all-functions --counts \
@@ -9,8 +8,7 @@
98
// RUN: %target_triple.%basename_t.llvm.profraw \
109
// RUN: | %fcheck-generic --check-prefix="LLVM-DEVICE"
1110

12-
// RUN: %libomptarget-compile-generic -fprofile-instr-generate \
13-
// RUN: -fprofile-update=atomic
11+
// RUN: %libomptarget-compile-generic -fprofile-instr-generate
1412
// RUN: env LLVM_PROFILE_FILE=%basename_t.clang.profraw \
1513
// RUN: %libomptarget-run-generic 2>&1
1614
// RUN: llvm-profdata show --all-functions --counts \
@@ -20,8 +18,7 @@
2018
// RUN: %target_triple.%basename_t.clang.profraw | \
2119
// RUN: %fcheck-generic --check-prefix="CLANG-DEV"
2220

23-
// RUN: %libomptarget-compile-generic -Xarch_host -fprofile-generate \
24-
// RUN: -fprofile-update=atomic
21+
// RUN: %libomptarget-compile-generic -Xarch_host -fprofile-generate
2522
// RUN: env LLVM_PROFILE_FILE=%basename_t.nogpu.profraw \
2623
// RUN: %libomptarget-run-generic 2>&1
2724
// RUN: llvm-profdata show --all-functions --counts \
@@ -30,7 +27,7 @@
3027
// RUN: not test -e %target_triple.%basename_t.nogpu.profraw
3128

3229
// RUN: %libomptarget-compile-generic -Xarch_host -fprofile-generate \
33-
// RUN: -Xarch_device -fprofile-instr-generate -fprofile-update=atomic
30+
// RUN: -Xarch_device -fprofile-instr-generate
3431
// RUN: env LLVM_PROFILE_FILE=%basename_t.hidf.profraw \
3532
// RUN: %libomptarget-run-generic 2>&1
3633
// RUN: llvm-profdata show --all-functions --counts \
@@ -41,7 +38,7 @@
4138
// RUN: | %fcheck-generic --check-prefix="CLANG-DEV"
4239

4340
// RUN: %libomptarget-compile-generic -Xarch_device -fprofile-generate \
44-
// RUN: -Xarch_host -fprofile-instr-generate -fprofile-update=atomic
41+
// RUN: -Xarch_host -fprofile-instr-generate
4542
// RUN: env LLVM_PROFILE_FILE=%basename_t.hfdi.profraw \
4643
// RUN: %libomptarget-run-generic 2>&1
4744
// RUN: llvm-profdata show --all-functions --counts \

‎offload/test/offloading/gpupgo/pgo3.c‎

Lines changed: 84 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,84 @@
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+
22+
int test1(int a) { return a / 2; }
23+
24+
int 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(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

‎offload/test/offloading/gpupgo/pgo4.c‎

Lines changed: 102 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,102 @@
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+
22+
int test1(int a) { return a / 2; }
23+
int test2(int a) { return a * 2; }
24+
25+
int 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(device_var);
34+
}
35+
if (i % 3 == 0) {
36+
device_var += test2(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

0 commit comments

Comments
 (0)