diff --git a/offload/test/offloading/gpupgo/pgo_atomic_teams.c b/offload/test/offloading/gpupgo/pgo_atomic_teams.c new file mode 100644 index 0000000000000..7bf3b1c11f28b --- /dev/null +++ b/offload/test/offloading/gpupgo/pgo_atomic_teams.c @@ -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 diff --git a/offload/test/offloading/gpupgo/pgo_atomic_threads.c b/offload/test/offloading/gpupgo/pgo_atomic_threads.c new file mode 100644 index 0000000000000..f0e7111f7a64b --- /dev/null +++ b/offload/test/offloading/gpupgo/pgo_atomic_threads.c @@ -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 diff --git a/offload/test/offloading/gpupgo/pgo2.c b/offload/test/offloading/gpupgo/pgo_device_and_host.c similarity index 95% rename from offload/test/offloading/gpupgo/pgo2.c rename to offload/test/offloading/gpupgo/pgo_device_and_host.c index af3ad9e4a6c19..3e95791ce9a50 100644 --- a/offload/test/offloading/gpupgo/pgo2.c +++ b/offload/test/offloading/gpupgo/pgo_device_and_host.c @@ -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; + } } } @@ -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: @@ -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 diff --git a/offload/test/offloading/gpupgo/pgo1.c b/offload/test/offloading/gpupgo/pgo_device_only.c similarity index 85% rename from offload/test/offloading/gpupgo/pgo1.c rename to offload/test/offloading/gpupgo/pgo_device_only.c index 1159858c51218..2939af613b6dd 100644 --- a/offload/test/offloading/gpupgo/pgo1.c +++ b/offload/test/offloading/gpupgo/pgo_device_only.c @@ -23,10 +23,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 +36,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 +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]+}} @@ -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