Skip to content

Commit 65aae01

Browse files
EthanLuisMcDonoughEthanLuisMcDonough
committed
[PGO][Offload] Make PGO GPU tests atomic
1 parent 0489447 commit 65aae01

File tree

1 file changed

+151
-0
lines changed

1 file changed

+151
-0
lines changed

‎t.diff

Lines changed: 151 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,151 @@
1+
diff --git a/offload/test/offloading/gpupgo/pgo1.c b/offload/test/offloading/gpupgo/pgo1.c
2+
index c8011cbae83c..4fb11bf702ab 100644
3+
--- a/offload/test/offloading/gpupgo/pgo1.c
4+
+++ b/offload/test/offloading/gpupgo/pgo1.c
5+
@@ -1,5 +1,6 @@
6+
// RUN: %libomptarget-compile-generic -fcreate-profile \
7+
-// RUN: -Xarch_device -fprofile-generate
8+
+// RUN: -Xarch_device -fprofile-generate \
9+
+// RUN: -Xarch_device -fprofile-update=atomic
10+
// RUN: env LLVM_PROFILE_FILE=%basename_t.llvm.profraw \
11+
// RUN: %libomptarget-run-generic 2>&1
12+
// RUN: llvm-profdata show --all-functions --counts \
13+
@@ -7,7 +8,8 @@
14+
// RUN: %fcheck-generic --check-prefix="LLVM-PGO"
15+
16+
// RUN: %libomptarget-compile-generic -fcreate-profile \
17+
-// RUN: -Xarch_device -fprofile-instr-generate
18+
+// RUN: -Xarch_device -fprofile-instr-generate \
19+
+// RUN: -Xarch_device -fprofile-update=atomic
20+
// RUN: env LLVM_PROFILE_FILE=%basename_t.clang.profraw \
21+
// RUN: %libomptarget-run-generic 2>&1
22+
// RUN: llvm-profdata show --all-functions --counts \
23+
@@ -23,10 +25,12 @@ int test2(int a) { return a * 2; }
24+
int main() {
25+
int m = 2;
26+
#pragma omp target
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+
+ {
32+
+ for (int i = 0; i < 10; i++) {
33+
+ m = test1(m);
34+
+ for (int j = 0; j < 2; j++) {
35+
+ m = test2(m);
36+
+ }
37+
}
38+
}
39+
}
40+
@@ -34,7 +38,7 @@ int main() {
41+
// LLVM-PGO-LABEL: __omp_offloading_{{[_0-9a-zA-Z]*}}_main_{{[_0-9a-zA-Z]*}}:
42+
// LLVM-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
43+
// LLVM-PGO: Counters: 4
44+
-// LLVM-PGO: Block counts: [20, 10, 2, 1]
45+
+// LLVM-PGO: Block counts: [20, 10, {{.*}}, 1]
46+
47+
// LLVM-PGO-LABEL: test1:
48+
// LLVM-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
49+
@@ -53,14 +57,10 @@ int main() {
50+
// LLVM-PGO-SAME: 3
51+
// LLVM-PGO-LABEL: Maximum function count:
52+
// LLVM-PGO-SAME: 20
53+
-// LLVM-PGO-LABEL: Maximum internal block count:
54+
-// LLVM-PGO-SAME: 10
55+
56+
// CLANG-PGO-LABEL: __omp_offloading_{{[_0-9a-zA-Z]*}}_main_{{[_0-9a-zA-Z]*}}:
57+
// CLANG-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
58+
-// CLANG-PGO: Counters: 3
59+
-// CLANG-PGO: Function count: 0
60+
-// CLANG-PGO: Block counts: [11, 20]
61+
+// CLANG-PGO: Block counts: [10, 20]
62+
63+
// CLANG-PGO-LABEL: test1:
64+
// CLANG-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
65+
@@ -78,7 +78,5 @@ int main() {
66+
// CLANG-PGO-SAME: Front-end
67+
// CLANG-PGO-LABEL: Functions shown:
68+
// CLANG-PGO-SAME: 3
69+
-// CLANG-PGO-LABEL: Maximum function count:
70+
-// CLANG-PGO-SAME: 20
71+
// CLANG-PGO-LABEL: Maximum internal block count:
72+
// CLANG-PGO-SAME: 20
73+
diff --git a/offload/test/offloading/gpupgo/pgo2.c b/offload/test/offloading/gpupgo/pgo2.c
74+
index b75b0beaffde..820a597bd510 100644
75+
--- a/offload/test/offloading/gpupgo/pgo2.c
76+
+++ b/offload/test/offloading/gpupgo/pgo2.c
77+
@@ -1,4 +1,5 @@
78+
-// RUN: %libomptarget-compile-generic -fprofile-generate
79+
+// RUN: %libomptarget-compile-generic -fprofile-generate \
80+
+// RUN: -fprofile-update=atomic
81+
// RUN: env LLVM_PROFILE_FILE=%basename_t.llvm.profraw \
82+
// RUN: %libomptarget-run-generic 2>&1
83+
// RUN: llvm-profdata show --all-functions --counts \
84+
@@ -8,7 +9,8 @@
85+
// RUN: %target_triple.%basename_t.llvm.profraw \
86+
// RUN: | %fcheck-generic --check-prefix="LLVM-DEVICE"
87+
88+
-// RUN: %libomptarget-compile-generic -fprofile-instr-generate
89+
+// RUN: %libomptarget-compile-generic -fprofile-instr-generate \
90+
+// RUN: -fprofile-update=atomic
91+
// RUN: env LLVM_PROFILE_FILE=%basename_t.clang.profraw \
92+
// RUN: %libomptarget-run-generic 2>&1
93+
// RUN: llvm-profdata show --all-functions --counts \
94+
@@ -18,7 +20,8 @@
95+
// RUN: %target_triple.%basename_t.clang.profraw | \
96+
// RUN: %fcheck-generic --check-prefix="CLANG-DEV"
97+
98+
-// RUN: %libomptarget-compile-generic -Xarch_host -fprofile-generate
99+
+// RUN: %libomptarget-compile-generic -Xarch_host -fprofile-generate \
100+
+// RUN: -fprofile-update=atomic
101+
// RUN: env LLVM_PROFILE_FILE=%basename_t.nogpu.profraw \
102+
// RUN: %libomptarget-run-generic 2>&1
103+
// RUN: llvm-profdata show --all-functions --counts \
104+
@@ -27,7 +30,7 @@
105+
// RUN: not test -e %target_triple.%basename_t.nogpu.profraw
106+
107+
// RUN: %libomptarget-compile-generic -Xarch_host -fprofile-generate \
108+
-// RUN: -Xarch_device -fprofile-instr-generate
109+
+// RUN: -Xarch_device -fprofile-instr-generate -fprofile-update=atomic
110+
// RUN: env LLVM_PROFILE_FILE=%basename_t.hidf.profraw \
111+
// RUN: %libomptarget-run-generic 2>&1
112+
// RUN: llvm-profdata show --all-functions --counts \
113+
@@ -38,7 +41,7 @@
114+
// RUN: | %fcheck-generic --check-prefix="CLANG-DEV"
115+
116+
// RUN: %libomptarget-compile-generic -Xarch_device -fprofile-generate \
117+
-// RUN: -Xarch_host -fprofile-instr-generate
118+
+// RUN: -Xarch_host -fprofile-instr-generate -fprofile-update=atomic
119+
// RUN: env LLVM_PROFILE_FILE=%basename_t.hfdi.profraw \
120+
// RUN: %libomptarget-run-generic 2>&1
121+
// RUN: llvm-profdata show --all-functions --counts \
122+
@@ -59,8 +62,10 @@ int main() {
123+
124+
int device_var = 1;
125+
#pragma omp target
126+
- for (int i = 0; i < 10; i++) {
127+
- device_var *= i;
128+
+ {
129+
+ for (int i = 0; i < 10; i++) {
130+
+ device_var *= i;
131+
+ }
132+
}
133+
}
134+
135+
@@ -78,7 +83,7 @@ int main() {
136+
// LLVM-DEVICE-LABEL: __omp_offloading_{{[_0-9a-zA-Z]*}}_main_{{[_0-9a-zA-Z]*}}:
137+
// LLVM-DEVICE: Hash: {{0[xX][0-9a-fA-F]+}}
138+
// LLVM-DEVICE: Counters: 3
139+
-// LLVM-DEVICE: Block counts: [10, 2, 1]
140+
+// LLVM-DEVICE: Block counts: [10, {{.*}}, 1]
141+
// LLVM-DEVICE: Instrumentation level: IR
142+
143+
// CLANG-HOST-LABEL: main:
144+
@@ -97,6 +102,5 @@ int main() {
145+
// CLANG-DEV-LABEL: __omp_offloading_{{[_0-9a-zA-Z]*}}_main_{{[_0-9a-zA-Z]*}}:
146+
// CLANG-DEV: Hash: {{0[xX][0-9a-fA-F]+}}
147+
// CLANG-DEV: Counters: 2
148+
-// CLANG-DEV: Function count: 0
149+
-// CLANG-DEV: Block counts: [11]
150+
+// CLANG-DEV: Block counts: [10]
151+
// CLANG-DEV: Instrumentation level: Front-end

0 commit comments

Comments
 (0)