[Openmp-commits] [openmp] 446e11a - [OpenMP ]Adding more libomptarget reduction tests (#71616)

[via Openmp-commits]

Author: Anton Rydahl
Date: 2023-11-07T20:39:08-08:00
New Revision: 446e11aceffbee027d010b7ccdfbe36117ce602e

URL: https://github.com/llvm/llvm-project/commit/446e11aceffbee027d010b7ccdfbe36117ce602e
DIFF: https://github.com/llvm/llvm-project/commit/446e11aceffbee027d010b7ccdfbe36117ce602e.diff

LOG: [OpenMP ]Adding more libomptarget reduction tests (#71616)

Based on https://github.com/llvm/llvm-project/pull/70766 I think it
would be good to have a few more offloading reduction tests, so we do
not accidentally break minimum and maximum reductions another time.

Added: 
    openmp/libomptarget/test/offloading/parallel_target_teams_reduction_max.cpp
    openmp/libomptarget/test/offloading/parallel_target_teams_reduction_min.cpp

Modified: 
    

Removed: 
    


################################################################################
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..8f06802aa772a5c
--- /dev/null
+++ b/openmp/libomptarget/test/offloading/parallel_target_teams_reduction_max.cpp
@@ -0,0 +1,73 @@
+// 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 <limits>
+#include <vector>
+
+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 (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 (Tp j = 0; j < length - 1; j++) {
+      if (b[j] > b[j + 1]) {
+        idx = std::max(idx, j);
+      }
+    }
+    assert(idx == (i - 1) * step &&
+           "#pragma omp target teams distribute parallel for "
+           "reduction(max:<identifier list>) does not work as intended.");
+    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() {
+  // 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
new file mode 100644
index 000000000000000..66c972f76044028
--- /dev/null
+++ b/openmp/libomptarget/test/offloading/parallel_target_teams_reduction_min.cpp
@@ -0,0 +1,73 @@
+// 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 <limits>
+#include <vector>
+
+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 (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 (Tp 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] += 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() {
+  // 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;
+}