[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:46:41 PDT 2025
https://github.com/EthanLuisMcDonough updated https://github.com/llvm/llvm-project/pull/132262
>From 10cb0f6ac5b02710fd6d78a4f4c3ef039f504e8f 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
---
offload/test/offloading/gpupgo/pgo1.c | 26 ++++++++++++--------------
offload/test/offloading/gpupgo/pgo2.c | 24 ++++++++++++++----------
2 files changed, 26 insertions(+), 24 deletions(-)
diff --git a/offload/test/offloading/gpupgo/pgo1.c b/offload/test/offloading/gpupgo/pgo1.c
index c8011cbae83c0..4fb11bf702aba 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 b75b0beaffdec..820a597bd5108 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