[llvm] [OpenMP][offload] Add cross-team reduction performance test (PR #195940)
via llvm-commits
llvm-commits at lists.llvm.org
Tue May 5 14:01:51 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-offload
Author: Robert Imschweiler (ro-i)
<details>
<summary>Changes</summary>
Tests different patterns of OpenMP cross-team reductions. Should not be included in regular `ninja check-offload` runs, thus excluded.
Note: not sure if this the right directory or if this rather should live somewhere else?
See also the discussion in https://github.com/llvm/llvm-project/pull/195102.
---
Patch is 20.96 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/195940.diff
6 Files Affected:
- (modified) offload/test/lit.cfg (+1-1)
- (added) offload/test/offloading/perf/.gitignore (+4)
- (added) offload/test/offloading/perf/Makefile (+27)
- (added) offload/test/offloading/perf/common.h (+183)
- (added) offload/test/offloading/perf/reduction.cpp (+306)
- (added) offload/test/offloading/perf/reduction.h (+60)
``````````diff
diff --git a/offload/test/lit.cfg b/offload/test/lit.cfg
index 14475c59fc20d..e13d14efa7d26 100644
--- a/offload/test/lit.cfg
+++ b/offload/test/lit.cfg
@@ -86,7 +86,7 @@ config.name = 'libomptarget :: ' + config.libomptarget_current_target
config.suffixes = ['.c', '.cpp', '.cc', '.f90', '.cu', '.td']
# excludes: A list of directories to exclude from the testuites.
-config.excludes = ['Inputs', 'unit']
+config.excludes = ['Inputs', 'unit', 'perf']
# test_source_root: The root path where tests are located.
config.test_source_root = os.path.dirname(__file__)
diff --git a/offload/test/offloading/perf/.gitignore b/offload/test/offloading/perf/.gitignore
new file mode 100644
index 0000000000000..ea6c07ca0e86d
--- /dev/null
+++ b/offload/test/offloading/perf/.gitignore
@@ -0,0 +1,4 @@
+.clangd
+compile_commands.json
+reduction_*
+out_reduction_*
diff --git a/offload/test/offloading/perf/Makefile b/offload/test/offloading/perf/Makefile
new file mode 100644
index 0000000000000..369c623e68a59
--- /dev/null
+++ b/offload/test/offloading/perf/Makefile
@@ -0,0 +1,27 @@
+.PHONY: all reduction clean
+
+OFFLOAD_ARCH ?= gfx90a
+COMMON_FLAGS = -O2 -fopenmp --offload-arch=$(OFFLOAD_ARCH) -std=c++20 -save-temps=obj
+
+PROJECTS = reduction
+COMMON_HEADERS = common.h
+
+RED_SRC = reduction.cpp
+RED_HEADERS = reduction.h
+# List of num_team values to build for.
+RED_NUM_TEAMS = 208 10400
+
+all: $(PROJECTS)
+
+reduction: $(foreach TEAM,$(RED_NUM_TEAMS),reduction_$(TEAM))
+
+reduction_%: $(RED_SRC) $(RED_HEADERS) $(COMMON_HEADERS)
+ @echo "Building reduction for $* teams ..."
+ rm -rf out_reduction_$*
+ mkdir -p out_reduction_$*
+ cd out_reduction_$* && $(CXX) -DXTEAM_NUM_TEAMS=$* $(COMMON_FLAGS) -o $@ $(addprefix ../,$(RED_SRC)) && cp $@ ..
+ cd out_reduction_$* && $(dir $(CXX))llvm-dis *.bc
+
+clean:
+ rm -rf $(foreach TEAM,$(RED_NUM_TEAMS),reduction_$(TEAM))
+ rm -rf $(foreach TEAM,$(RED_NUM_TEAMS),out_reduction_$(TEAM))
diff --git a/offload/test/offloading/perf/common.h b/offload/test/offloading/perf/common.h
new file mode 100644
index 0000000000000..fdc7dbb8b59fa
--- /dev/null
+++ b/offload/test/offloading/perf/common.h
@@ -0,0 +1,183 @@
+#pragma once
+
+#include <algorithm>
+#include <chrono>
+#include <cstdint>
+#include <cstdlib>
+#include <format>
+#include <iostream>
+#include <limits>
+#include <numeric>
+#include <optional>
+#include <vector>
+
+#include "omp.h"
+
+#ifndef XTEAM_NUM_THREADS
+#define XTEAM_NUM_THREADS 512
+#endif
+#ifndef XTEAM_NUM_TEAMS
+#define XTEAM_NUM_TEAMS 208 // gfx90a number of CUs: 104
+#endif
+// If true, let codegen for reduction/scan determine num_teams and num_threads.
+#ifndef CODEGEN_AUTODETECTION
+#define CODEGEN_AUTODETECTION 0
+#endif
+#if CODEGEN_AUTODETECTION
+#define TEAMS
+#define THREADS
+#else // CODEGEN_AUTODETECTION
+#define TEAMS num_teams(XTEAM_NUM_TEAMS)
+#define THREADS num_threads(XTEAM_NUM_THREADS)
+#endif // CODEGEN_AUTODETECTION
+#define TEAMS_THREADS TEAMS THREADS
+
+// Represents the total of threads in the Grid
+#define XTEAM_TOTAL_NUM_THREADS (XTEAM_NUM_TEAMS * XTEAM_NUM_THREADS)
+
+// Benchmark warmup iterations
+#define WARMUP_ITERS 2
+// Benchmark minimum number of measured iterations (as a lower bound in case the
+// test is so slow that it wouldn't get enough iterations in the auto-scale
+// timeframe)
+#define BENCH_MIN_ITERS 10
+// Auto-scale timeframe in seconds. Benchmarks will be repeated until they reach
+// at least this amount of seconds.
+#define AUTO_SCALE_TIME 1.0
+
+// Floating point absolute and relative tolerance for comparison.
+#define FP_ABS_TOL 1e-12
+#define FP_REL_TOL 1e-6
+
+// default alignment for aligned_alloc
+#define ALIGNMENT 128
+
+#define duration_cast(x) \
+ std::chrono::duration_cast<std::chrono::duration<double>>(x)
+using Clock = std::chrono::steady_clock;
+
+struct TimingResult {
+ double MinS, MaxS, AvgS;
+ double BestMbps, AvgMbps;
+};
+
+// =========================================================================
+// Utility functions
+// =========================================================================
+
+template <typename T> T *alloc(uint64_t N) {
+ if (N > std::numeric_limits<size_t>::max() / sizeof(T)) {
+ std::cerr << std::format("alloc size overflow n={} sizeof(T)={}\n", N,
+ sizeof(T));
+ exit(EXIT_FAILURE);
+ }
+ size_t Bytes = sizeof(T) * N;
+ Bytes = ((Bytes + ALIGNMENT - 1) / ALIGNMENT) * ALIGNMENT;
+
+ T *Ret = static_cast<T *>(aligned_alloc(ALIGNMENT, Bytes));
+ if (!Ret) {
+ std::cerr << std::format("aligned_alloc failed bytes={}\n", Bytes);
+ exit(EXIT_FAILURE);
+ }
+ return Ret;
+}
+
+template <typename T> T *targetAlloc(uint64_t N, int Devid) {
+ if (N > std::numeric_limits<size_t>::max() / sizeof(T)) {
+ std::cerr << std::format("target_alloc size overflow n={} sizeof(T)={}\n",
+ N, sizeof(T));
+ exit(EXIT_FAILURE);
+ }
+ T *Ret = static_cast<T *>(omp_target_alloc(sizeof(T) * N, Devid));
+ if (!Ret) {
+ std::cerr << std::format("omp_target_alloc failed n={} devid={}\n", N,
+ Devid);
+ exit(EXIT_FAILURE);
+ }
+ return Ret;
+}
+
+// The values are deterministic for reproducibility.
+template <typename T, bool is_fp> void initData(T *Arr1, T *Arr2, uint64_t N) {
+ srand(42);
+ for (uint64_t I = 0; I < N; I++) {
+ if constexpr (is_fp) {
+ Arr1[I] = T((rand() % 100) / 100.0); // NOLINT(misc-predictable-rand)
+ Arr2[I] = T((rand() % 100) / 100.0); // NOLINT(misc-predictable-rand)
+ } else {
+ Arr1[I] = T(rand() % 1000); // NOLINT(misc-predictable-rand)
+ Arr2[I] = T(rand() % 1000); // NOLINT(misc-predictable-rand)
+ }
+ }
+}
+
+// =========================================================================
+// Benchmark utilities
+// =========================================================================
+
+template <typename T, bool is_fp>
+static bool check(T Result, T Gold, std::string_view Label) {
+ if constexpr (!is_fp) {
+ if (Result == Gold)
+ return true;
+ std::cerr << std::format("FAIL {}: got {}, expected {}\n", Label, Result,
+ Gold);
+ return false;
+ }
+ double G = (double)Gold, C = (double)Result;
+ double AbsErr = std::abs(C - G);
+ double Scale = std::max({1.0, std::abs(G), std::abs(C)});
+ double RelErr = AbsErr / Scale;
+ if (AbsErr <= FP_ABS_TOL || RelErr <= FP_REL_TOL)
+ return true;
+ std::cerr << std::format("FAIL {}: got {}, expected {} (abs={}, rel={})\n",
+ Label, Result, Gold, AbsErr, RelErr);
+ return false;
+}
+
+static TimingResult createTimingResult(const std::vector<double> &Times,
+ uint64_t DataBytes) {
+ if (Times.empty()) {
+ std::cerr << "internal error: no timing samples collected\n";
+ return TimingResult{0.0, 0.0, 0.0, 0.0, 0.0};
+ }
+ auto [Mn, Mx] = std::minmax_element(Times.begin(), Times.end());
+ // NOLINTNEXTLINE(llvm-use-ranges)
+ double Avg = std::accumulate(Times.begin(), Times.end(), 0.0) /
+ static_cast<double>(Times.size());
+ double BestMbps = (*Mn > 0.0) ? (1e-6 * DataBytes / *Mn) : 0.0;
+ double AvgMbps = (Avg > 0.0) ? (1e-6 * DataBytes / Avg) : 0.0;
+ return TimingResult{*Mn, *Mx, Avg, BestMbps, AvgMbps};
+}
+
+// Add locale-independent thousand separators to make visual number parsing
+// easier
+static std::string fmtNumSep(std::string S) {
+ for (int Pos = S.length() - 3; Pos > 0; Pos -= 3)
+ S.insert(Pos, ",");
+ return S;
+}
+
+static void printResult(std::string_view Test, std::string_view Type,
+ uint64_t N, const std::optional<TimingResult> &R) {
+ if (!R) {
+ std::cerr << std::format("{:<24} {:<8} {:>15} FAIL\n", Test, Type,
+ fmtNumSep(std::format("{}", N)));
+ return;
+ }
+ std::cout << std::format("{:<24} {:<8} {:>15} {:>10.6f} {:>10.6f} "
+ "{:>10.6f} {:>12} {:>12}\n",
+ Test, Type, fmtNumSep(std::format("{}", N)), R->MinS,
+ R->MaxS, R->AvgS,
+ fmtNumSep(std::format("{:.0f}", R->BestMbps)),
+ fmtNumSep(std::format("{:.0f}", R->AvgMbps)));
+}
+
+static void printHeader() {
+ std::cout << std::format(
+ "{:>24} {:>8} {:>15} {:>10} {:>10} {:>10} {:>12} {:>12}\n", "test",
+ "type", "N", "min(s)", "max(s)", "avg(s)", "best MB/s", "avg MB/s");
+ std::cout << std::format(
+ "{:->24} {:->8} {:->15} {:->10} {:->10} {:->10} {:->12} {:->12}\n",
+ "", "", "", "", "", "", "", "");
+}
diff --git a/offload/test/offloading/perf/reduction.cpp b/offload/test/offloading/perf/reduction.cpp
new file mode 100644
index 0000000000000..7f8fa05186730
--- /dev/null
+++ b/offload/test/offloading/perf/reduction.cpp
@@ -0,0 +1,306 @@
+// OpenMP cross-team reduction performance & correctness benchmark
+// Have a look at common.h for configuration options.
+
+#include <numbers>
+#include <unistd.h>
+
+#include "omp.h"
+
+#include "common.h"
+#include "reduction.h"
+
+static Config Conf;
+
+template <typename T> static T redSum(const T *__restrict In, uint64_t N) {
+ T S = redIdentity<T, RedOp::Sum>();
+#pragma omp target teams distribute parallel for TEAMS_THREADS reduction(+ : S)
+ for (uint64_t I = 0; I < N; I++)
+ S += In[I];
+ return S;
+}
+
+template <typename T> static T redMax(const T *__restrict In, uint64_t N) {
+ T M = redIdentity<T, RedOp::Max>();
+#pragma omp target teams distribute parallel for TEAMS_THREADS reduction( \
+ max : M)
+ for (uint64_t I = 0; I < N; I++)
+ M = std::max(M, In[I]);
+ return M;
+}
+
+template <typename T> static T redMult(const T *__restrict In, uint64_t N) {
+ T M = redIdentity<T, RedOp::Mult>();
+#pragma omp target teams distribute parallel for TEAMS_THREADS reduction(* : M)
+ for (uint64_t I = 0; I < N; I++)
+ M *= In[I];
+ return M;
+}
+
+template <typename T>
+static T redDot(const T *__restrict A, const T *__restrict B, uint64_t N) {
+ T S = redIdentity<T, RedOp::Sum>();
+#pragma omp target teams distribute parallel for TEAMS_THREADS reduction(+ : S)
+ for (uint64_t I = 0; I < N; I++)
+ S += A[I] * B[I];
+ return S;
+}
+
+// Reduction through an indirect call, checking for potential frontend pattern
+// matching.
+template <typename T> static T redIndirect(const T *__restrict In, uint64_t N) {
+ T S = redIdentity<T, RedOp::Sum>();
+ auto Accumulate = [](T A, T B) { return A + B; };
+#pragma omp target teams distribute parallel for TEAMS_THREADS reduction(+ : S)
+ for (uint64_t I = 0; I < N; I++)
+ S = Accumulate(S, In[I]);
+ return S;
+}
+
+// Combined reduction (sum and max) in the same loop.
+template <typename T> static T redCombined(const T *__restrict In, uint64_t N) {
+ T S = redIdentity<T, RedOp::Sum>();
+ T M = redIdentity<T, RedOp::Max>();
+#pragma omp target teams distribute parallel for TEAMS_THREADS reduction( \
+ + : S) reduction(max : M)
+ for (uint64_t I = 0; I < N; I++) {
+ S += In[I];
+ M = std::max(M, In[I]);
+ }
+ return (S / 2) + (M / 2);
+}
+
+// Combined reduction (sum and max) in separate loops.
+template <typename T>
+static T redCombinedSeparate(const T *__restrict In, uint64_t N) {
+ T S = redIdentity<T, RedOp::Sum>();
+ T M = redIdentity<T, RedOp::Max>();
+#pragma omp target map(tofrom : S, M)
+#pragma omp teams TEAMS reduction(+ : S) reduction(max : M)
+ {
+#pragma omp distribute parallel for THREADS reduction(+ : S)
+ for (uint64_t I = 0; I < N; I++)
+ S += In[I];
+
+#pragma omp distribute parallel for THREADS reduction(max : M)
+ for (uint64_t I = 0; I < N; I++)
+ M = std::max(M, In[I]);
+ }
+ return (S / 2) + (M / 2);
+}
+
+// Have a reduction in a kernel that is also doing something completely
+// unrelated to the reduction (pure register work, no memory ops).
+template <typename T>
+static T redKernelPart(const T *__restrict In, uint64_t N) {
+ T S = redIdentity<T, RedOp::Sum>();
+
+#pragma omp target map(tofrom : S)
+#pragma omp teams TEAMS reduction(+ : S)
+ {
+#pragma omp distribute parallel for THREADS reduction(+ : S)
+ for (uint64_t I = 0; I < N; I++)
+ S += In[I];
+
+ // Just do something, without actually doing anything
+#pragma omp parallel THREADS
+ {
+ int TID = omp_get_thread_num();
+ T X = static_cast<T>(TID);
+ for (int J = 0; J < 100; J++)
+ X = X * static_cast<T>(0.9) + static_cast<T>(J);
+ if (X == static_cast<T>(-1))
+ S += X;
+ }
+ }
+
+ return S;
+}
+
+// Reduction without any memory access.
+static double redPi(uint64_t N) {
+ double Pi = 0.0;
+
+ // https://en.wikipedia.org/wiki/Leibniz_formula_for_%CF%80
+#pragma omp target teams distribute parallel for TEAMS_THREADS reduction(+ : Pi)
+ for (uint64_t I = 0; I < N; I++) {
+ double Term = 1.0 / (2 * I + 1);
+ Pi += (I & 0x1) ? -Term : Term;
+ }
+
+ return Pi * 4.0;
+}
+
+// =========================================================================
+// Benchmark utilities
+// =========================================================================
+
+template <typename T, bool is_fp, typename KernelFunc, typename... InputArgs>
+static std::optional<TimingResult>
+runBenchRed(KernelFunc Kernel, T Gold, uint64_t N, std::string_view Label,
+ InputArgs... Inputs) {
+ std::vector<double> Times;
+ double TotalTime = 0.0;
+ int WarmUpIters = WARMUP_ITERS;
+
+ while (TotalTime < AUTO_SCALE_TIME || Times.size() < BENCH_MIN_ITERS) {
+ auto T1 = Clock::now();
+ T Result = Kernel(Inputs..., N);
+ auto T2 = Clock::now();
+ if (!check<T, is_fp>(Result, Gold, Label))
+ return std::nullopt;
+
+ if (WarmUpIters > 0) {
+ WarmUpIters--;
+ continue;
+ }
+ double D = duration_cast(T2 - T1).count();
+ Times.push_back(D);
+ TotalTime += D;
+ }
+
+ return createTimingResult(Times, sizeof(T) * N * sizeof...(Inputs));
+}
+
+// Run the benchmarks for the given data type (is_fp: true for floating point
+// types).
+template <typename T, bool is_fp>
+static void runType(std::string_view TypeName) {
+ std::optional<TimingResult> R;
+
+ for (uint64_t N : Conf.ArraySizes) {
+ T *In1 = alloc<T>(N);
+ T *In2 = alloc<T>(N);
+ initData<T, is_fp>(In1, In2, N);
+
+#pragma omp target enter data map(to : In1[0 : N], In2[0 : N])
+
+ T GoldSum = goldRed<T, RedOp::Sum>(In1, N);
+ T GoldMax = goldRed<T, RedOp::Max>(In1, N);
+ T GoldMult = goldRed<T, RedOp::Mult>(In1, N);
+ T GoldDot = goldRedDot<T>(In1, In2, N);
+
+ // ================================================================
+ // dot reduction
+ // ================================================================
+ R = runBenchRed<T, is_fp>(redDot<T>, GoldDot, N, "red_dot", In1, In2);
+ printResult("red_dot", TypeName, N, R);
+
+ // ================================================================
+ // max reduction
+ // ================================================================
+ R = runBenchRed<T, is_fp>(redMax<T>, GoldMax, N, "red_max", In1);
+ printResult("red_max", TypeName, N, R);
+
+ // ================================================================
+ // sum reduction
+ // ================================================================
+ R = runBenchRed<T, is_fp>(redSum<T>, GoldSum, N, "red_sum", In1);
+ printResult("red_sum", TypeName, N, R);
+
+ if (!Conf.QuickRun || std::is_same_v<T, double>) {
+ // ================================================================
+ // mult reduction
+ // ================================================================
+ R = runBenchRed<T, is_fp>(redMult<T>, GoldMult, N, "red_mult", In1);
+ printResult("red_mult", TypeName, N, R);
+
+ // ================================================================
+ // indirect reduction (sum)
+ // ================================================================
+ R = runBenchRed<T, is_fp>(redIndirect<T>, GoldSum, N, "red_indirect",
+ In1);
+ printResult("red_indirect", TypeName, N, R);
+
+ // ================================================================
+ // combined (sum and max) reduction - in the same loop ...
+ // ================================================================
+ T GoldCombined = (GoldSum / 2) + (GoldMax / 2);
+ R = runBenchRed<T, is_fp>(redCombined<T>, GoldCombined, N, "red_combined",
+ In1);
+ printResult("red_combined", TypeName, N, R);
+
+ // ================================================================
+ // ... and in separate loops
+ // ================================================================
+ R = runBenchRed<T, is_fp>(redCombinedSeparate<T>, GoldCombined, N,
+ "red_combined_separate", In1);
+ printResult("red_combined_separate", TypeName, N, R);
+
+ // ================================================================
+ // reduction (sum) in a kernel that is also doing something completely
+ // unrelated to the reduction.
+ // ================================================================
+ R = runBenchRed<T, is_fp>(redKernelPart<T>, GoldSum, N, "red_kernel_part",
+ In1);
+ printResult("red_kernel_part", TypeName, N, R);
+ }
+
+#pragma omp target exit data map(delete : In1[0 : N], In2[0 : N])
+
+ free(In1);
+ free(In2);
+ }
+
+ if (std::is_same_v<T, double>) {
+ // ================================================================
+ // reduction computing Pi
+ // ================================================================
+ double GoldPi = std::numbers::pi;
+ uint64_t N = 5000000000;
+ R = runBenchRed<double, true>(redPi, GoldPi, N, "red_pi");
+ printResult("red_pi", TypeName, N, R);
+ }
+}
+
+static void usage(std::string_view Argv0) {
+ std::cout << "Usage: " << Argv0 << " [-v] [-h]\n"
+ << " -v: Verbose run (test all array sizes)\n"
+ << " -h: Show this help message\n";
+}
+
+int main(int argc, char *const *argv) {
+ int Opt;
+
+ while ((Opt = getopt(argc, argv, "vh")) != -1) {
+ switch (Opt) {
+ case 'v':
+ Conf.QuickRun = false;
+ break;
+ case 'h':
+ usage(argv[0]);
+ return EXIT_SUCCESS;
+ default:
+ usage(argv[0]);
+ return EXIT_FAILURE;
+ }
+ }
+
+ if (Conf.QuickRun)
+ Conf.ArraySizes.assign(ArraySizesQuick.begin(), ArraySizesQuick.end());
+ else
+ Conf.ArraySizes.assign(ArraySizes.begin(), ArraySizes.end());
+
+ std::cout << std::format(
+ "xteam reduction benchmark (quick run: {}) — {} warmup iterations "
+ "{} teams, {} threads, codegen autodetection: {}\n",
+ Conf.QuickRun ? "true" : "false", WARMUP_ITERS, XTEAM_NUM_TEAMS,
+ XTEAM_NUM_THREADS, CODEGEN_AUTODETECTION ? "true" : "false");
+
+ std::cout << "Array sizes: ";
+ for (uint64_t SZ : Conf.ArraySizes)
+ std::cout << " " << fmtNumSep(std::format("{}", SZ));
+ std::cout << "\n\n";
+
+ printHeader();
+
+ std::cout << "\n--- double ---\n";
+ runType<double, true>("double");
+
+ std::cout << "\n--- uint ---\n";
+ runType<unsigned, false>("uint");
+
+ std::cout << "\n--- ulong ---\n";
+ runType<unsigned long, false>("ulong");
+
+ return EXIT_SUCCESS;
+}
diff --git a/offload/test/offloading/perf/reduction.h b/offload/test/offloading/perf/reduction.h
new file mode 100644
index 0000000000000..83c141f430f73
--- /dev/null
+++ b/offload/test/offloading/perf/reduction.h
@@ -0,0 +1,60 @@
+#pragma once
+
+#include <algorithm>
+#include <array>
+#include <cstdint>
+#include <limits>
+#include <type_traits>
+#include <vector>
+
+static const std::array<uint64_t, 1> ArraySizesQuick{177777777};
+static const std::array<uint64_t, 14> ArraySizes{
+ 1, 100, 1024, 2048, 4096, 8192, 10000,
+ 81920, 1000000, 4194304, 23445657, 41943040, 100000000, 177777777};
+
+struct Config {
+ bool QuickRun = true;
+ std::vector<uint64_t> ArraySizes;
+};
+
+enum class RedOp { Sum, Max, Min, Mult };
+
+template <typename T, RedOp Op> constexpr T redIdentity() {
+ if constexpr (Op == RedOp::Sum)
+ return T(0);
+ else if constexpr (Op == RedOp::Max)
+ return std::numeric_limits<T>::lowest();
+ else if constexpr (Op == R...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/195940
More information about the llvm-commits
mailing list