Skip to content

Commit 1043810

Browse files
EthanLuisMcDonoughEthanLuisMcDonough
authored
[PGO][Offload] Update PGO GPU tests (#132262)
1 parent 847561e commit 1043810

File tree

4 files changed

+200
-17
lines changed

4 files changed

+200
-17
lines changed

‎offload/test/offloading/gpupgo/pgo_atomic_teams.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

‎offload/test/offloading/gpupgo/pgo_atomic_threads.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/pgo2.c renamed to ‎offload/test/offloading/gpupgo/pgo_device_and_host.c

Lines changed: 6 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -59,8 +59,10 @@ int main() {
5959

6060
int device_var = 1;
6161
#pragma omp target
62-
for (int i = 0; i < 10; i++) {
63-
device_var *= i;
62+
{
63+
for (int i = 0; i < 10; i++) {
64+
device_var *= i;
65+
}
6466
}
6567
}
6668

@@ -78,7 +80,7 @@ int main() {
7880
// LLVM-DEVICE-LABEL: __omp_offloading_{{[_0-9a-zA-Z]*}}_main_{{[_0-9a-zA-Z]*}}:
7981
// LLVM-DEVICE: Hash: {{0[xX][0-9a-fA-F]+}}
8082
// LLVM-DEVICE: Counters: 3
81-
// LLVM-DEVICE: Block counts: [10, 2, 1]
83+
// LLVM-DEVICE: Block counts: [10, {{.*}}, 1]
8284
// LLVM-DEVICE: Instrumentation level: IR
8385

8486
// CLANG-HOST-LABEL: main:
@@ -97,6 +99,5 @@ int main() {
9799
// CLANG-DEV-LABEL: __omp_offloading_{{[_0-9a-zA-Z]*}}_main_{{[_0-9a-zA-Z]*}}:
98100
// CLANG-DEV: Hash: {{0[xX][0-9a-fA-F]+}}
99101
// CLANG-DEV: Counters: 2
100-
// CLANG-DEV: Function count: 0
101-
// CLANG-DEV: Block counts: [11]
102+
// CLANG-DEV: Block counts: [10]
102103
// CLANG-DEV: Instrumentation level: Front-end

‎offload/test/offloading/gpupgo/pgo1.c renamed to ‎offload/test/offloading/gpupgo/pgo_device_only.c

Lines changed: 8 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -23,18 +23,20 @@ int test2(int a) { return a * 2; }
2323
int main() {
2424
int m = 2;
2525
#pragma omp target
26-
for (int i = 0; i < 10; i++) {
27-
m = test1(m);
28-
for (int j = 0; j < 2; j++) {
29-
m = test2(m);
26+
{
27+
for (int i = 0; i < 10; i++) {
28+
m = test1(m);
29+
for (int j = 0; j < 2; j++) {
30+
m = test2(m);
31+
}
3032
}
3133
}
3234
}
3335

3436
// LLVM-PGO-LABEL: __omp_offloading_{{[_0-9a-zA-Z]*}}_main_{{[_0-9a-zA-Z]*}}:
3537
// LLVM-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
3638
// LLVM-PGO: Counters: 4
37-
// LLVM-PGO: Block counts: [20, 10, 2, 1]
39+
// LLVM-PGO: Block counts: [20, 10, {{.*}}, 1]
3840

3941
// LLVM-PGO-LABEL: test1:
4042
// LLVM-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
@@ -53,14 +55,10 @@ int main() {
5355
// LLVM-PGO-SAME: 3
5456
// LLVM-PGO-LABEL: Maximum function count:
5557
// LLVM-PGO-SAME: 20
56-
// LLVM-PGO-LABEL: Maximum internal block count:
57-
// LLVM-PGO-SAME: 10
5858

5959
// CLANG-PGO-LABEL: __omp_offloading_{{[_0-9a-zA-Z]*}}_main_{{[_0-9a-zA-Z]*}}:
6060
// CLANG-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
61-
// CLANG-PGO: Counters: 3
62-
// CLANG-PGO: Function count: 0
63-
// CLANG-PGO: Block counts: [11, 20]
61+
// CLANG-PGO: Block counts: [10, 20]
6462

6563
// CLANG-PGO-LABEL: test1:
6664
// CLANG-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
@@ -78,7 +76,5 @@ int main() {
7876
// CLANG-PGO-SAME: Front-end
7977
// CLANG-PGO-LABEL: Functions shown:
8078
// CLANG-PGO-SAME: 3
81-
// CLANG-PGO-LABEL: Maximum function count:
82-
// CLANG-PGO-SAME: 20
8379
// CLANG-PGO-LABEL: Maximum internal block count:
8480
// CLANG-PGO-SAME: 20

0 commit comments

Comments
 (0)