[Openmp-commits] [openmp] [OpenMP ]Added more libomptarget reduction tests (PR #71616)
Anton Rydahl via Openmp-commits
openmp-commits at lists.llvm.org
Tue Nov 7 20:23:58 PST 2023
https://github.com/AntonRydahl updated https://github.com/llvm/llvm-project/pull/71616
>From 7d19104ceb4d42d5d7ebe0758972e1763bea243e Mon Sep 17 00:00:00 2001
From: AntonRydahl <rydahl2610 at gmail.com>
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 000000000000000..19c05e0bd573c20
--- /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 <algorithm>
+#include <cassert>
+#include <vector>
+
+template <class Tp> void test_max_reduction() {
+ const int length = 1000;
+ const int nmaximas = 8;
+ std::vector<Tp> 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:<identifier list>) does not work as intended.");
+ a[idx] -= (Tp)1;
+ }
+}
+
+int main() {
+ test_max_reduction<float>();
+ test_max_reduction<double>();
+ test_max_reduction<int>();
+ test_max_reduction<unsigned int>();
+ test_max_reduction<long>();
+ 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 000000000000000..e8a8648830b6344
--- /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 <algorithm>
+#include <cassert>
+#include <vector>
+
+template <class Tp> void test_min_reduction() {
+ const int length = 1000;
+ const int nminimas = 8;
+ std::vector<Tp> 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:<identifier list>) does not work as intended.");
+ a[idx] += (Tp)1;
+ }
+}
+
+int main() {
+ test_min_reduction<float>();
+ test_min_reduction<double>();
+ test_min_reduction<int>();
+ test_min_reduction<unsigned int>();
+ test_min_reduction<long>();
+ return 0;
+}
>From 88e142a99b0d29f52d067ffb4546defa21b76aed Mon Sep 17 00:00:00 2001
From: AntonRydahl <rydahl2610 at gmail.com>
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 19c05e0bd573c20..8f06802aa772a5c 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 <algorithm>
#include <cassert>
+#include <limits>
#include <vector>
-template <class Tp> void test_max_reduction() {
- const int length = 1000;
- const int nmaximas = 8;
- std::vector<Tp> a(length, (Tp)3);
- const int step = length / nmaximas;
- for (int i = 0; i < nmaximas; i++) {
- a[i * step] += (Tp)1;
+template <class Tp> void test_max_idx_reduction() {
+ const Tp length = 1000;
+ const Tp nmaximas = 8;
+ std::vector<float> 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:<identifier list>) does not work as intended.");
- a[idx] -= (Tp)1;
+ a[idx] -= 1.0f;
}
}
+template <class Tp> void test_max_val_reduction() {
+ const int length = 1000;
+ const int half = length / 2;
+ std::vector<Tp> a(length, (Tp)3);
+ a[half] += (Tp)1;
+ Tp max_val = std::numeric_limits<Tp>::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:<identifier list>) does not work as intended.");
+}
+
int main() {
- test_max_reduction<float>();
- test_max_reduction<double>();
- test_max_reduction<int>();
- test_max_reduction<unsigned int>();
- test_max_reduction<long>();
+ // Reducing over indices
+ test_max_idx_reduction<int>();
+ test_max_idx_reduction<unsigned int>();
+ test_max_idx_reduction<long>();
+
+ // Reducing over values
+ test_max_val_reduction<int>();
+ test_max_val_reduction<unsigned int>();
+ test_max_val_reduction<long>();
+ test_max_val_reduction<float>();
+ test_max_val_reduction<double>();
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 e8a8648830b6344..66c972f76044028 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 <algorithm>
#include <cassert>
+#include <limits>
#include <vector>
-template <class Tp> void test_min_reduction() {
- const int length = 1000;
- const int nminimas = 8;
- std::vector<Tp> a(length, (Tp)3);
- const int step = length / nminimas;
- for (int i = 0; i < nminimas; i++) {
- a[i * step] -= (Tp)1;
+template <class Tp> void test_min_idx_reduction() {
+ const Tp length = 1000;
+ const Tp nminimas = 8;
+ std::vector<float> 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 <class Tp> void test_min_reduction() {
assert(idx == i * step &&
"#pragma omp target teams distribute parallel for "
"reduction(min:<identifier list>) does not work as intended.");
- a[idx] += (Tp)1;
+ a[idx] += 1.0f;
}
}
+template <class Tp> void test_min_val_reduction() {
+ const int length = 1000;
+ const int half = length / 2;
+ std::vector<Tp> a(length, (Tp)3);
+ a[half] -= (Tp)1;
+ Tp min_val = std::numeric_limits<Tp>::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:<identifier list>) does not work as intended.");
+}
+
int main() {
- test_min_reduction<float>();
- test_min_reduction<double>();
- test_min_reduction<int>();
- test_min_reduction<unsigned int>();
- test_min_reduction<long>();
+ // Reducing over indices
+ test_min_idx_reduction<int>();
+ test_min_idx_reduction<unsigned int>();
+ test_min_idx_reduction<long>();
+
+ // Reducing over values
+ test_min_val_reduction<int>();
+ test_min_val_reduction<unsigned int>();
+ test_min_val_reduction<long>();
+ test_min_val_reduction<float>();
+ test_min_val_reduction<double>();
return 0;
}
More information about the Openmp-commits
mailing list