Skip to content

[PGO][Offload] Update PGO GPU tests #132262

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 5 commits into from
May 14, 2025
Merged

[PGO][Offload] Update PGO GPU tests #132262

Show file tree
Hide file tree
Changes from 2 commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
10cb0f6
[PGO][Offload] Make PGO GPU tests atomic
EthanLuisMcDonough Mar 20, 2025
0c48346
[PGO][Offload] Update GPU PGO tests
EthanLuisMcDonough Apr 30, 2025
db194f2
Merge branch 'main' into gpu_pgo_atomic_tests2
EthanLuisMcDonough May 14, 2025
c200700
Merge branch 'main' into gpu_pgo_atomic_tests2
EthanLuisMcDonough May 14, 2025
5c5979c
[PGO][Offload] Change GPU PGO test names
EthanLuisMcDonough May 14, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
20 changes: 8 additions & 12 deletions offload/test/offloading/gpupgo/pgo1.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -23,18 +23,20 @@ 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);
}
}
}
}

// 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]+}}
Expand All @@ -53,14 +55,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]+}}
Expand All @@ -78,7 +76,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
11 changes: 6 additions & 5 deletions offload/test/offloading/gpupgo/pgo2.c
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -59,8 +59,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;
}
}
}

Expand All @@ -78,7 +80,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:
Expand All @@ -97,6 +99,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
84 changes: 84 additions & 0 deletions offload/test/offloading/gpupgo/pgo3.c
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,84 @@
// RUN: %libomptarget-compile-generic -fcreate-profile \
// 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 \
// RUN: %target_triple.%basename_t.llvm.profraw | \
// RUN: %fcheck-generic --check-prefix="LLVM-PGO"

// RUN: %libomptarget-compile-generic -fcreate-profile \
// 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 \
// RUN: %target_triple.%basename_t.clang.profraw | \
// RUN: %fcheck-generic --check-prefix="CLANG-PGO"

// REQUIRES: gpu
// REQUIRES: pgo

int test1(int a) { return a / 2; }

int main() {
int device_var = 1;
#pragma omp target map(tofrom : device_var)
{
#pragma omp parallel for
for (int i = 1; i <= 10; i++) {
device_var *= i;
if (i % 2 == 0) {
device_var += test1(device_var);
}
}
}
}

// clang-format off
// 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: 2
// LLVM-PGO: Block counts: [0, {{.*}}]

// LLVM-PGO-LABEL: __omp_offloading_{{[_0-9a-zA-Z]*}}_main_{{[_0-9a-zA-Z]*}}_omp_outlined:
// LLVM-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
// LLVM-PGO: Counters: 5
// LLVM-PGO: Block counts: [10, 5, {{.*}}, 10, {{.*}}]

// LLVM-PGO-LABEL: test1:
// LLVM-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
// LLVM-PGO: Counters: 1
// LLVM-PGO: Block counts: [5]

// LLVM-PGO-LABEL: Instrumentation level:
// LLVM-PGO-SAME: IR
// LLVM-PGO-SAME: entry_first = 0
// LLVM-PGO-LABEL: Functions shown:
// LLVM-PGO-SAME: 3
// LLVM-PGO-LABEL: Maximum function 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: 1
// CLANG-PGO: Function count: {{.*}}
// CLANG-PGO: Block counts: []

// CLANG-PGO-LABEL: __omp_offloading_{{[_0-9a-zA-Z]*}}_main_{{[_0-9a-zA-Z]*}}_omp_outlined:
// CLANG-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
// CLANG-PGO: Counters: 3
// CLANG-PGO: Function count: {{.*}}
// CLANG-PGO: Block counts: [{{.*}}, 5]

// CLANG-PGO-LABEL: test1:
// CLANG-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
// CLANG-PGO: Counters: 1
// CLANG-PGO: Function count: 5
// CLANG-PGO: Block counts: []

// CLANG-PGO-LABEL: Instrumentation level:
// CLANG-PGO-SAME: Front-end
// CLANG-PGO-LABEL: Functions shown:
// CLANG-PGO-SAME: 3
// clang-format on
102 changes: 102 additions & 0 deletions offload/test/offloading/gpupgo/pgo4.c
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,102 @@
// RUN: %libomptarget-compile-generic -fcreate-profile \
// 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 \
// RUN: %target_triple.%basename_t.llvm.profraw | \
// RUN: %fcheck-generic --check-prefix="LLVM-PGO"

// RUN: %libomptarget-compile-generic -fcreate-profile \
// 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 \
// RUN: %target_triple.%basename_t.clang.profraw | \
// RUN: %fcheck-generic --check-prefix="CLANG-PGO"

// REQUIRES: gpu
// REQUIRES: pgo

int test1(int a) { return a / 2; }
int test2(int a) { return a * 2; }

int main() {
int device_var = 1;

#pragma omp target teams distribute parallel for num_teams(3) \
map(tofrom : device_var)
for (int i = 1; i <= 30; i++) {
device_var *= i;
if (i % 2 == 0) {
device_var += test1(device_var);
}
if (i % 3 == 0) {
device_var += test2(device_var);
}
}
}

// clang-format off
// 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: 2
// LLVM-PGO: Block counts: [0, {{.*}}]

// LLVM-PGO-LABEL: __omp_offloading_{{[_0-9a-zA-Z]*}}_main_{{[_0-9a-zA-Z]*}}_omp_outlined:
// LLVM-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
// LLVM-PGO: Counters: 4
// LLVM-PGO: Block counts: [{{.*}}, 0, {{.*}}, 0]

// LLVM-PGO-LABEL: __omp_offloading_{{[_0-9a-zA-Z]*}}_main_{{[_0-9a-zA-Z]*}}_omp_outlined_omp_outlined:
// LLVM-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
// LLVM-PGO: Counters: 4
// LLVM-PGO: Block counts: [30, 15, 10, {{.*}}]

// LLVM-PGO-LABEL: test1:
// LLVM-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
// LLVM-PGO: Counters: 1
// LLVM-PGO: Block counts: [15]

// LLVM-PGO-LABEL: test2:
// LLVM-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
// LLVM-PGO: Counters: 1
// LLVM-PGO: Block counts: [10]

// LLVM-PGO-LABEL: Instrumentation level:
// LLVM-PGO-SAME: IR

// 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: 1
// CLANG-PGO: Function count: {{.*}}
// CLANG-PGO: Block counts: []

// CLANG-PGO-LABEL: __omp_offloading_{{[_0-9a-zA-Z]*}}_main_{{[_0-9a-zA-Z]*}}_omp_outlined:
// CLANG-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
// CLANG-PGO: Counters: 1
// CLANG-PGO: Function count: {{.*}}
// CLANG-PGO: Block counts: []

// CLANG-PGO-LABEL: __omp_offloading_{{[_0-9a-zA-Z]*}}_main_{{[_0-9a-zA-Z]*}}_omp_outlined_omp_outlined:
// CLANG-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
// CLANG-PGO: Counters: 4
// CLANG-PGO: Function count: 30
// CLANG-PGO: Block counts: [{{.*}}, 15, 10]

// CLANG-PGO-LABEL: test1:
// CLANG-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
// CLANG-PGO: Counters: 1
// CLANG-PGO: Function count: 15
// CLANG-PGO: Block counts: []

// CLANG-PGO-LABEL: test2:
// CLANG-PGO: Hash: {{0[xX][0-9a-fA-F]+}}
// CLANG-PGO: Counters: 1
// CLANG-PGO: Function count: 10
// CLANG-PGO: Block counts: []

// CLANG-PGO-LABEL: Instrumentation level:
// CLANG-PGO-SAME: Front-end
// clang-format on
Loading