[llvm] [PGO][Offload] Make PGO GPU tests atomic (PR #132262)
Ethan Luis McDonough via llvm-commits
llvm-commits at lists.llvm.org
Thu Mar 20 10:43:55 PDT 2025
https://github.com/EthanLuisMcDonough created https://github.com/llvm/llvm-project/pull/132262
This pull request aims to improve the reliability of pgo1.c and pgo2.c
>From 65aae0163974c0f9d7714de966d88c913fb83f94 Mon Sep 17 00:00:00 2001
From: Ethan Luis McDonough <ethanluismcdonough at gmail.com>
Date: Thu, 20 Mar 2025 12:42:10 -0500
Subject: [PATCH] [PGO][Offload] Make PGO GPU tests atomic
---
t.diff | 151 +++++++++++++++++++++++++++++++++++++++++++++++++++++++++
1 file changed, 151 insertions(+)
create mode 100644 t.diff
diff --git a/t.diff b/t.diff
new file mode 100644
index 0000000000000..9b0bfa868e417
--- /dev/null
+++ b/t.diff
@@ -0,0 +1,151 @@
+diff --git a/offload/test/offloading/gpupgo/pgo1.c b/offload/test/offloading/gpupgo/pgo1.c
+index c8011cbae83c..4fb11bf702ab 100644
+--- a/offload/test/offloading/gpupgo/pgo1.c
++++ b/offload/test/offloading/gpupgo/pgo1.c
+@@ -1,5 +1,6 @@
+ // RUN: %libomptarget-compile-generic -fcreate-profile \
+-// RUN: -Xarch_device -fprofile-generate
++// RUN: -Xarch_device -fprofile-generate \
++// RUN: -Xarch_device -fprofile-update=atomic
+ // RUN: env LLVM_PROFILE_FILE=%basename_t.llvm.profraw \
+ // RUN: %libomptarget-run-generic 2>&1
+ // RUN: llvm-profdata show --all-functions --counts \
+@@ -7,7 +8,8 @@
+ // RUN: %fcheck-generic --check-prefix="LLVM-PGO"
+
+ // RUN: %libomptarget-compile-generic -fcreate-profile \
+-// RUN: -Xarch_device -fprofile-instr-generate
++// RUN: -Xarch_device -fprofile-instr-generate \
++// RUN: -Xarch_device -fprofile-update=atomic
+ // RUN: env LLVM_PROFILE_FILE=%basename_t.clang.profraw \
+ // RUN: %libomptarget-run-generic 2>&1
+ // RUN: llvm-profdata show --all-functions --counts \
+@@ -23,10 +25,12 @@ int test2(int a) { return a * 2; }
+ int main() {
+ int m = 2;
+ #pragma omp target
+- for (int i = 0; i < 10; i++) {
+- m = test1(m);
+- for (int j = 0; j < 2; j++) {
+- m = test2(m);
++ {
++ for (int i = 0; i < 10; i++) {
++ m = test1(m);
++ for (int j = 0; j < 2; j++) {
++ m = test2(m);
++ }
+ }
+ }
+ }
+@@ -34,7 +38,7 @@ int main() {
+ // LLVM-PGO-LABEL: __omp_offloading_{{[_0-9a-zA-Z]*}}_main_{{[_0-9a-zA-Z]*}}:
+ // LLVM-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
+ // LLVM-PGO: Counters: 4
+-// LLVM-PGO: Block counts: [20, 10, 2, 1]
++// LLVM-PGO: Block counts: [20, 10, {{.*}}, 1]
+
+ // LLVM-PGO-LABEL: test1:
+ // LLVM-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
+@@ -53,14 +57,10 @@ int main() {
+ // LLVM-PGO-SAME: 3
+ // LLVM-PGO-LABEL: Maximum function count:
+ // LLVM-PGO-SAME: 20
+-// LLVM-PGO-LABEL: Maximum internal block count:
+-// LLVM-PGO-SAME: 10
+
+ // CLANG-PGO-LABEL: __omp_offloading_{{[_0-9a-zA-Z]*}}_main_{{[_0-9a-zA-Z]*}}:
+ // CLANG-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
+-// CLANG-PGO: Counters: 3
+-// CLANG-PGO: Function count: 0
+-// CLANG-PGO: Block counts: [11, 20]
++// CLANG-PGO: Block counts: [10, 20]
+
+ // CLANG-PGO-LABEL: test1:
+ // CLANG-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
+@@ -78,7 +78,5 @@ int main() {
+ // CLANG-PGO-SAME: Front-end
+ // CLANG-PGO-LABEL: Functions shown:
+ // CLANG-PGO-SAME: 3
+-// CLANG-PGO-LABEL: Maximum function count:
+-// CLANG-PGO-SAME: 20
+ // CLANG-PGO-LABEL: Maximum internal block count:
+ // CLANG-PGO-SAME: 20
+diff --git a/offload/test/offloading/gpupgo/pgo2.c b/offload/test/offloading/gpupgo/pgo2.c
+index b75b0beaffde..820a597bd510 100644
+--- a/offload/test/offloading/gpupgo/pgo2.c
++++ b/offload/test/offloading/gpupgo/pgo2.c
+@@ -1,4 +1,5 @@
+-// RUN: %libomptarget-compile-generic -fprofile-generate
++// RUN: %libomptarget-compile-generic -fprofile-generate \
++// RUN: -fprofile-update=atomic
+ // RUN: env LLVM_PROFILE_FILE=%basename_t.llvm.profraw \
+ // RUN: %libomptarget-run-generic 2>&1
+ // RUN: llvm-profdata show --all-functions --counts \
+@@ -8,7 +9,8 @@
+ // RUN: %target_triple.%basename_t.llvm.profraw \
+ // RUN: | %fcheck-generic --check-prefix="LLVM-DEVICE"
+
+-// RUN: %libomptarget-compile-generic -fprofile-instr-generate
++// RUN: %libomptarget-compile-generic -fprofile-instr-generate \
++// RUN: -fprofile-update=atomic
+ // RUN: env LLVM_PROFILE_FILE=%basename_t.clang.profraw \
+ // RUN: %libomptarget-run-generic 2>&1
+ // RUN: llvm-profdata show --all-functions --counts \
+@@ -18,7 +20,8 @@
+ // RUN: %target_triple.%basename_t.clang.profraw | \
+ // RUN: %fcheck-generic --check-prefix="CLANG-DEV"
+
+-// RUN: %libomptarget-compile-generic -Xarch_host -fprofile-generate
++// RUN: %libomptarget-compile-generic -Xarch_host -fprofile-generate \
++// RUN: -fprofile-update=atomic
+ // RUN: env LLVM_PROFILE_FILE=%basename_t.nogpu.profraw \
+ // RUN: %libomptarget-run-generic 2>&1
+ // RUN: llvm-profdata show --all-functions --counts \
+@@ -27,7 +30,7 @@
+ // RUN: not test -e %target_triple.%basename_t.nogpu.profraw
+
+ // RUN: %libomptarget-compile-generic -Xarch_host -fprofile-generate \
+-// RUN: -Xarch_device -fprofile-instr-generate
++// RUN: -Xarch_device -fprofile-instr-generate -fprofile-update=atomic
+ // RUN: env LLVM_PROFILE_FILE=%basename_t.hidf.profraw \
+ // RUN: %libomptarget-run-generic 2>&1
+ // RUN: llvm-profdata show --all-functions --counts \
+@@ -38,7 +41,7 @@
+ // RUN: | %fcheck-generic --check-prefix="CLANG-DEV"
+
+ // RUN: %libomptarget-compile-generic -Xarch_device -fprofile-generate \
+-// RUN: -Xarch_host -fprofile-instr-generate
++// RUN: -Xarch_host -fprofile-instr-generate -fprofile-update=atomic
+ // RUN: env LLVM_PROFILE_FILE=%basename_t.hfdi.profraw \
+ // RUN: %libomptarget-run-generic 2>&1
+ // RUN: llvm-profdata show --all-functions --counts \
+@@ -59,8 +62,10 @@ int main() {
+
+ int device_var = 1;
+ #pragma omp target
+- for (int i = 0; i < 10; i++) {
+- device_var *= i;
++ {
++ for (int i = 0; i < 10; i++) {
++ device_var *= i;
++ }
+ }
+ }
+
+@@ -78,7 +83,7 @@ int main() {
+ // LLVM-DEVICE-LABEL: __omp_offloading_{{[_0-9a-zA-Z]*}}_main_{{[_0-9a-zA-Z]*}}:
+ // LLVM-DEVICE: Hash: {{0[xX][0-9a-fA-F]+}}
+ // LLVM-DEVICE: Counters: 3
+-// LLVM-DEVICE: Block counts: [10, 2, 1]
++// LLVM-DEVICE: Block counts: [10, {{.*}}, 1]
+ // LLVM-DEVICE: Instrumentation level: IR
+
+ // CLANG-HOST-LABEL: main:
+@@ -97,6 +102,5 @@ int main() {
+ // CLANG-DEV-LABEL: __omp_offloading_{{[_0-9a-zA-Z]*}}_main_{{[_0-9a-zA-Z]*}}:
+ // CLANG-DEV: Hash: {{0[xX][0-9a-fA-F]+}}
+ // CLANG-DEV: Counters: 2
+-// CLANG-DEV: Function count: 0
+-// CLANG-DEV: Block counts: [11]
++// CLANG-DEV: Block counts: [10]
+ // CLANG-DEV: Instrumentation level: Front-end
More information about the llvm-commits
mailing list