From 7d19104ceb4d42d5d7ebe0758972e1763bea243e Mon Sep 17 00:00:00 2001 From: AntonRydahl Date: Tue, 7 Nov 2023 16:54:53 -0800 Subject: [PATCH 1/2] Added more libomptarget reduction tests --- .../parallel_target_teams_reduction_max.cpp | 49 +++++++++++++++++++ .../parallel_target_teams_reduction_min.cpp | 49 +++++++++++++++++++ 2 files changed, 98 insertions(+) create mode 100644 openmp/libomptarget/test/offloading/parallel_target_teams_reduction_max.cpp create mode 100644 openmp/libomptarget/test/offloading/parallel_target_teams_reduction_min.cpp diff --git a/openmp/libomptarget/test/offloading/parallel_target_teams_reduction_max.cpp b/openmp/libomptarget/test/offloading/parallel_target_teams_reduction_max.cpp new file mode 100644 index 0000000000000..19c05e0bd573c --- /dev/null +++ b/openmp/libomptarget/test/offloading/parallel_target_teams_reduction_max.cpp @@ -0,0 +1,49 @@ +// RUN: %libomptarget-compilexx-and-run-generic +// RUN: %libomptarget-compileoptxx-and-run-generic + +// FIXME: This is a bug in host offload, this should run fine. +// UNSUPPORTED: aarch64-unknown-linux-gnu +// UNSUPPORTED: aarch64-unknown-linux-gnu-LTO +// UNSUPPORTED: x86_64-pc-linux-gnu +// UNSUPPORTED: x86_64-pc-linux-gnu-LTO + +// This test validates that the OpenMP target reductions to find a maximum work +// as indended for a few common data types. + +#include +#include +#include + +template void test_max_reduction() { + const int length = 1000; + const int nmaximas = 8; + std::vector a(length, (Tp)3); + const int step = length / nmaximas; + for (int i = 0; i < nmaximas; i++) { + a[i * step] += (Tp)1; + } + for (int i = nmaximas - 1; i >= 0; i--) { + int idx = 0; + Tp *b = a.data(); +#pragma omp target teams distribute parallel for reduction(max : idx) \ + map(always, to : b[0 : length]) + for (int j = 1; j < length; j++) { + if (b[j] > b[j - 1]) { + idx = std::max(idx, j); + } + } + assert(idx == i * step && + "#pragma omp target teams distribute parallel for " + "reduction(max:) does not work as intended."); + a[idx] -= (Tp)1; + } +} + +int main() { + test_max_reduction(); + test_max_reduction(); + test_max_reduction(); + test_max_reduction(); + test_max_reduction(); + return 0; +} diff --git a/openmp/libomptarget/test/offloading/parallel_target_teams_reduction_min.cpp b/openmp/libomptarget/test/offloading/parallel_target_teams_reduction_min.cpp new file mode 100644 index 0000000000000..e8a8648830b63 --- /dev/null +++ b/openmp/libomptarget/test/offloading/parallel_target_teams_reduction_min.cpp @@ -0,0 +1,49 @@ +// RUN: %libomptarget-compilexx-and-run-generic +// RUN: %libomptarget-compileoptxx-and-run-generic + +// FIXME: This is a bug in host offload, this should run fine. +// UNSUPPORTED: aarch64-unknown-linux-gnu +// UNSUPPORTED: aarch64-unknown-linux-gnu-LTO +// UNSUPPORTED: x86_64-pc-linux-gnu +// UNSUPPORTED: x86_64-pc-linux-gnu-LTO + +// This test validates that the OpenMP target reductions to find a minimum work +// as indended for a few common data types. + +#include +#include +#include + +template void test_min_reduction() { + const int length = 1000; + const int nminimas = 8; + std::vector a(length, (Tp)3); + const int step = length / nminimas; + for (int i = 0; i < nminimas; i++) { + a[i * step] -= (Tp)1; + } + for (int i = 0; i < nminimas; i++) { + int idx = a.size(); + Tp *b = a.data(); +#pragma omp target teams distribute parallel for reduction(min : idx) \ + map(always, to : b[0 : length]) + for (int j = 0; j < length - 1; j++) { + if (b[j] < b[j + 1]) { + idx = std::min(idx, j); + } + } + assert(idx == i * step && + "#pragma omp target teams distribute parallel for " + "reduction(min:) does not work as intended."); + a[idx] += (Tp)1; + } +} + +int main() { + test_min_reduction(); + test_min_reduction(); + test_min_reduction(); + test_min_reduction(); + test_min_reduction(); + return 0; +} From 88e142a99b0d29f52d067ffb4546defa21b76aed Mon Sep 17 00:00:00 2001 From: AntonRydahl Date: Tue, 7 Nov 2023 20:23:22 -0800 Subject: [PATCH 2/2] Added more test cases --- .../parallel_target_teams_reduction_max.cpp | 62 +++++++++++++------ .../parallel_target_teams_reduction_min.cpp | 58 ++++++++++++----- 2 files changed, 84 insertions(+), 36 deletions(-) diff --git a/openmp/libomptarget/test/offloading/parallel_target_teams_reduction_max.cpp b/openmp/libomptarget/test/offloading/parallel_target_teams_reduction_max.cpp index 19c05e0bd573c..8f06802aa772a 100644 --- a/openmp/libomptarget/test/offloading/parallel_target_teams_reduction_max.cpp +++ b/openmp/libomptarget/test/offloading/parallel_target_teams_reduction_max.cpp @@ -12,38 +12,62 @@ #include #include +#include #include -template void test_max_reduction() { - const int length = 1000; - const int nmaximas = 8; - std::vector a(length, (Tp)3); - const int step = length / nmaximas; - for (int i = 0; i < nmaximas; i++) { - a[i * step] += (Tp)1; +template void test_max_idx_reduction() { + const Tp length = 1000; + const Tp nmaximas = 8; + std::vector a(length, 3.0f); + const Tp step = length / nmaximas; + for (Tp i = 0; i < nmaximas; i++) { + a[i * step] += 1.0f; } - for (int i = nmaximas - 1; i >= 0; i--) { - int idx = 0; - Tp *b = a.data(); + for (Tp i = nmaximas; i > 0; i--) { + Tp idx = 0; + float *b = a.data(); #pragma omp target teams distribute parallel for reduction(max : idx) \ map(always, to : b[0 : length]) - for (int j = 1; j < length; j++) { - if (b[j] > b[j - 1]) { + for (Tp j = 0; j < length - 1; j++) { + if (b[j] > b[j + 1]) { idx = std::max(idx, j); } } - assert(idx == i * step && + assert(idx == (i - 1) * step && "#pragma omp target teams distribute parallel for " "reduction(max:) does not work as intended."); - a[idx] -= (Tp)1; + a[idx] -= 1.0f; } } +template void test_max_val_reduction() { + const int length = 1000; + const int half = length / 2; + std::vector a(length, (Tp)3); + a[half] += (Tp)1; + Tp max_val = std::numeric_limits::lowest(); + Tp *b = a.data(); +#pragma omp target teams distribute parallel for reduction(max : max_val) \ + map(always, to : b[0 : length]) + for (int i = 0; i < length; i++) { + max_val = std::max(max_val, b[i]); + } + assert(std::abs(((double)a[half + 1]) - ((double)max_val) + 1.0) < 1e-6 && + "#pragma omp target teams distribute parallel for " + "reduction(max:) does not work as intended."); +} + int main() { - test_max_reduction(); - test_max_reduction(); - test_max_reduction(); - test_max_reduction(); - test_max_reduction(); + // Reducing over indices + test_max_idx_reduction(); + test_max_idx_reduction(); + test_max_idx_reduction(); + + // Reducing over values + test_max_val_reduction(); + test_max_val_reduction(); + test_max_val_reduction(); + test_max_val_reduction(); + test_max_val_reduction(); return 0; } diff --git a/openmp/libomptarget/test/offloading/parallel_target_teams_reduction_min.cpp b/openmp/libomptarget/test/offloading/parallel_target_teams_reduction_min.cpp index e8a8648830b63..66c972f760440 100644 --- a/openmp/libomptarget/test/offloading/parallel_target_teams_reduction_min.cpp +++ b/openmp/libomptarget/test/offloading/parallel_target_teams_reduction_min.cpp @@ -12,22 +12,23 @@ #include #include +#include #include -template void test_min_reduction() { - const int length = 1000; - const int nminimas = 8; - std::vector a(length, (Tp)3); - const int step = length / nminimas; - for (int i = 0; i < nminimas; i++) { - a[i * step] -= (Tp)1; +template void test_min_idx_reduction() { + const Tp length = 1000; + const Tp nminimas = 8; + std::vector a(length, 3.0f); + const Tp step = length / nminimas; + for (Tp i = 0; i < nminimas; i++) { + a[i * step] -= 1.0f; } - for (int i = 0; i < nminimas; i++) { - int idx = a.size(); - Tp *b = a.data(); + for (Tp i = 0; i < nminimas; i++) { + Tp idx = a.size(); + float *b = a.data(); #pragma omp target teams distribute parallel for reduction(min : idx) \ map(always, to : b[0 : length]) - for (int j = 0; j < length - 1; j++) { + for (Tp j = 0; j < length - 1; j++) { if (b[j] < b[j + 1]) { idx = std::min(idx, j); } @@ -35,15 +36,38 @@ template void test_min_reduction() { assert(idx == i * step && "#pragma omp target teams distribute parallel for " "reduction(min:) does not work as intended."); - a[idx] += (Tp)1; + a[idx] += 1.0f; } } +template void test_min_val_reduction() { + const int length = 1000; + const int half = length / 2; + std::vector a(length, (Tp)3); + a[half] -= (Tp)1; + Tp min_val = std::numeric_limits::max(); + Tp *b = a.data(); +#pragma omp target teams distribute parallel for reduction(min : min_val) \ + map(always, to : b[0 : length]) + for (int i = 0; i < length; i++) { + min_val = std::min(min_val, b[i]); + } + assert(std::abs(((double)a[half + 1]) - ((double)min_val) - 1.0) < 1e-6 && + "#pragma omp target teams distribute parallel for " + "reduction(min:) does not work as intended."); +} + int main() { - test_min_reduction(); - test_min_reduction(); - test_min_reduction(); - test_min_reduction(); - test_min_reduction(); + // Reducing over indices + test_min_idx_reduction(); + test_min_idx_reduction(); + test_min_idx_reduction(); + + // Reducing over values + test_min_val_reduction(); + test_min_val_reduction(); + test_min_val_reduction(); + test_min_val_reduction(); + test_min_val_reduction(); return 0; }