[llvm] [OpenMP][offload] Add enhanced cross-team reduction test (PR #195940)

[Robert Imschweiler via llvm-commits]

https://github.com/ro-i updated https://github.com/llvm/llvm-project/pull/195940

>From 325463fa8318c972d9612394035d9e1faf77ec36 Mon Sep 17 00:00:00 2001
From: Robert Imschweiler <robert.imschweiler at amd.com>
Date: Tue, 5 May 2026 15:44:49 -0500
Subject: [PATCH] [OpenMP][offload] Add enhanced cross-team reduction test

Tests different patterns of OpenMP cross-team reductions, for multiple
data types.
If run with `LIBOMPTARGET_INFO=16`, shows current register spilling due
to dispatch jump chains (which grow for every reduction in the same
translation unit) for indirect function calls in the reduction runtime.
---
 .../test/offloading/multiple_reductions.cpp   | 129 ++++++++++++++++++
 1 file changed, 129 insertions(+)
 create mode 100644 offload/test/offloading/multiple_reductions.cpp

diff --git a/offload/test/offloading/multiple_reductions.cpp b/offload/test/offloading/multiple_reductions.cpp
new file mode 100644
index 0000000000000..3187076d06f6c
--- /dev/null
+++ b/offload/test/offloading/multiple_reductions.cpp
@@ -0,0 +1,129 @@
+// RUN: %libomptarget-compilexx-run-and-check-generic
+// RUN: %libomptarget-compileoptxx-run-and-check-generic
+// UNSUPPORTED: intelgpu
+
+#include <cassert>
+#include <cstdio>
+#include <cstdlib>
+#include <iostream>
+
+#include "omp.h"
+
+#define N 10000
+
+template <typename T> static void init_data(T *a) {
+  for (int i = 0; i < N; ++i)
+    a[i] = i;
+}
+
+template <typename T> void run_type(void) {
+  T s1, s2;
+
+  T *in1 = static_cast<T *>(malloc(N * sizeof(T)));
+  T *in2 = static_cast<T *>(malloc(N * sizeof(T)));
+  assert(in1 && in2);
+
+  init_data(in1);
+  init_data(in2);
+
+#pragma omp target enter data map(to : in1[0 : N], in2[0 : N])
+
+  // Sum reduction
+  s1 = T(0);
+#pragma omp target teams distribute parallel for reduction(+ : s1)
+  for (int i = 0; i < N; ++i)
+    s1 += in1[i];
+  // CHECK: 49995000
+  std::cout << s1 << '\n';
+
+  s1 = T(0);
+  // Indirect sum reduction
+  auto accumulate = [](T a, T b) { return a + b; };
+#pragma omp target teams distribute parallel for reduction(+ : s1)
+  for (int i = 0; i < N; i++)
+    s1 = accumulate(s1, in1[i]);
+  // CHECK: 49995000
+  std::cout << s1 << '\n';
+
+  // Dot reduction
+  s1 = T(0);
+#pragma omp target teams distribute parallel for reduction(+ : s1)
+  for (int i = 0; i < N; ++i)
+    s1 += in1[i] * in2[i];
+  // CHECK: 2570853208
+  std::cout << s1 << '\n';
+
+  // Combined reduction (sum and max) - in the same loop ...
+  s1 = s2 = T(0);
+#pragma omp target teams distribute parallel for reduction(+ : s1)             \
+    reduction(max : s2)
+  for (int i = 0; i < N; ++i) {
+    s1 += in1[i];
+    s2 = in1[i] > s2 ? in1[i] : s2;
+  }
+  // CHECK: 49995000 : 9999
+  std::cout << s1 << " : " << s2 << '\n';
+
+  // ... and in separate loops
+  s1 = s2 = T(0);
+#pragma omp target map(tofrom : s1, s2)
+#pragma omp teams reduction(+ : s1) reduction(max : s2)
+  {
+#pragma omp distribute parallel for reduction(+ : s1)
+    for (int i = 0; i < N; i++)
+      s1 += in1[i];
+
+#pragma omp distribute parallel for reduction(max : s2)
+    for (int i = 0; i < N; i++)
+      s2 = in1[i] > s2 ? in1[i] : s2;
+  }
+  // CHECK: 49995000 : 9999
+  std::cout << s1 << " : " << s2 << '\n';
+
+  // Reduction in a kernel that is also doing something completely
+  // unrelated to the reduction (pure register work, no memory ops).
+  s1 = T(0);
+#pragma omp target map(tofrom : s1)
+#pragma omp teams reduction(+ : s1)
+  {
+#pragma omp distribute parallel for reduction(+ : s1)
+    for (int i = 0; i < N; i++)
+      s1 += in1[i];
+
+    // Just do something, without actually doing anything
+#pragma omp parallel
+    {
+      int x = omp_get_thread_num();
+      for (int j = 0; j < 100; j++)
+        x = x * 0.9 + j;
+      if (x == -1)
+        s1 += x;
+    }
+  }
+  // CHECK: 49995000
+  std::cout << s1 << '\n';
+
+#pragma omp target exit data map(delete : in1[0 : N], in2[0 : N])
+
+  free(in1);
+  free(in2);
+}
+
+int main(int argc, char **argv) {
+  run_type<double>();
+  run_type<unsigned>();
+  run_type<unsigned long>();
+
+  // Reduction calculating pi
+  double pi = 0.0;
+  // https://en.wikipedia.org/wiki/Leibniz_formula_for_%CF%80
+#pragma omp target teams distribute parallel for reduction(+ : pi)
+  for (int i = 0; i < N; i++) {
+    double term = 1.0 / (2 * i + 1);
+    pi += (i & 0x1) ? -term : term;
+  }
+  // CHECK: 3.141
+  printf("%.3f\n", pi * 4.0);
+
+  return EXIT_SUCCESS;
+}