[libclc] libclc: Add sub_group_reduce_* functions (PR #185294)
Matt Arsenault via cfe-commits
cfe-commits at lists.llvm.org
Sun Mar 8 08:40:56 PDT 2026
https://github.com/arsenm created https://github.com/llvm/llvm-project/pull/185294
None
>From 6c8e690d226ee105d196bf6a7187affbcea8d0dd Mon Sep 17 00:00:00 2001
From: Matt Arsenault <Matthew.Arsenault at amd.com>
Date: Sun, 8 Mar 2026 16:07:09 +0100
Subject: [PATCH] libclc: Add sub_group_reduce_* functions
---
.../clc/subgroup/clc_sub_group_reduce.h | 38 ++++++++
.../clc/subgroup/clc_sub_group_reduce.inc | 13 +++
libclc/clc/lib/amdgcn/SOURCES | 1 +
.../lib/amdgcn/subgroup/sub_group_reduce.cl | 96 +++++++++++++++++++
.../generic/subgroup/clc_sub_group_reduce.cl | 15 +++
libclc/opencl/lib/generic/SOURCES | 1 +
.../lib/generic/subgroup/sub_group_reduce.cl | 15 +++
.../lib/generic/subgroup/sub_group_reduce.inc | 27 ++++++
8 files changed, 206 insertions(+)
create mode 100644 libclc/clc/include/clc/subgroup/clc_sub_group_reduce.h
create mode 100644 libclc/clc/include/clc/subgroup/clc_sub_group_reduce.inc
create mode 100644 libclc/clc/lib/amdgcn/subgroup/sub_group_reduce.cl
create mode 100644 libclc/clc/lib/generic/subgroup/clc_sub_group_reduce.cl
create mode 100644 libclc/opencl/lib/generic/subgroup/sub_group_reduce.cl
create mode 100644 libclc/opencl/lib/generic/subgroup/sub_group_reduce.inc
diff --git a/libclc/clc/include/clc/subgroup/clc_sub_group_reduce.h b/libclc/clc/include/clc/subgroup/clc_sub_group_reduce.h
new file mode 100644
index 0000000000000..2e567929e1351
--- /dev/null
+++ b/libclc/clc/include/clc/subgroup/clc_sub_group_reduce.h
@@ -0,0 +1,38 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CLC_SUBGROUP_CLC_SUB_GROUP_REDUCE_H__
+#define __CLC_SUBGROUP_CLC_SUB_GROUP_REDUCE_H__
+
+#include "clc/internal/clc.h"
+
+#define __CLC_FUNCTION __clc_sub_group_reduce_add
+#define __CLC_BODY <clc/subgroup/clc_sub_group_reduce.inc>
+#include <clc/integer/gentype.inc>
+
+#define __CLC_BODY <clc/subgroup/clc_sub_group_reduce.inc>
+#include <clc/math/gentype.inc>
+#undef __CLC_FUNCTION
+
+#define __CLC_FUNCTION __clc_sub_group_reduce_min
+#define __CLC_BODY <clc/subgroup/clc_sub_group_reduce.inc>
+#include <clc/integer/gentype.inc>
+
+#define __CLC_BODY <clc/subgroup/clc_sub_group_reduce.inc>
+#include <clc/math/gentype.inc>
+#undef __CLC_FUNCTION
+
+#define __CLC_FUNCTION __clc_sub_group_reduce_max
+#define __CLC_BODY <clc/subgroup/clc_sub_group_reduce.inc>
+#include <clc/integer/gentype.inc>
+
+#define __CLC_BODY <clc/subgroup/clc_sub_group_reduce.inc>
+#include <clc/math/gentype.inc>
+#undef __CLC_FUNCTION
+
+#endif // __CLC_SUBGROUP_CLC_SUB_GROUP_REDUCE_H__
diff --git a/libclc/clc/include/clc/subgroup/clc_sub_group_reduce.inc b/libclc/clc/include/clc/subgroup/clc_sub_group_reduce.inc
new file mode 100644
index 0000000000000..73cb5fb46516c
--- /dev/null
+++ b/libclc/clc/include/clc/subgroup/clc_sub_group_reduce.inc
@@ -0,0 +1,13 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#if defined(__CLC_SCALAR) && \
+ (defined(__CLC_FPSIZE) || __CLC_GENSIZE == 32 || __CLC_GENSIZE == 64)
+_CLC_OVERLOAD _CLC_DECL _CLC_CONST __CLC_GENTYPE
+__CLC_FUNCTION(__CLC_GENTYPE x);
+#endif
diff --git a/libclc/clc/lib/amdgcn/SOURCES b/libclc/clc/lib/amdgcn/SOURCES
index a0b6c168b207e..3480554d91e90 100644
--- a/libclc/clc/lib/amdgcn/SOURCES
+++ b/libclc/clc/lib/amdgcn/SOURCES
@@ -3,6 +3,7 @@ math/clc_ldexp.cl
mem_fence/clc_mem_fence.cl
subgroup/subgroup.cl
subgroup/sub_group_broadcast.cl
+subgroup/sub_group_reduce.cl
synchronization/clc_sub_group_barrier.cl
synchronization/clc_work_group_barrier.cl
workitem/clc_get_enqueued_local_size.cl
diff --git a/libclc/clc/lib/amdgcn/subgroup/sub_group_reduce.cl b/libclc/clc/lib/amdgcn/subgroup/sub_group_reduce.cl
new file mode 100644
index 0000000000000..51d877caf3351
--- /dev/null
+++ b/libclc/clc/lib/amdgcn/subgroup/sub_group_reduce.cl
@@ -0,0 +1,96 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "clc/subgroup/clc_sub_group_broadcast.h"
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint __clc_sub_group_reduce_add(uint x) {
+ return __builtin_amdgcn_wave_reduce_add_u32(x, 0);
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST int __clc_sub_group_reduce_add(int x) {
+ return (int)__clc_sub_group_reduce_add((uint)x);
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST ulong __clc_sub_group_reduce_add(ulong x) {
+ return __builtin_amdgcn_wave_reduce_add_u64(x, 0);
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST long __clc_sub_group_reduce_add(long x) {
+ return (long)__clc_sub_group_reduce_add((ulong)x);
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint __clc_sub_group_reduce_min(uint x) {
+ return __builtin_amdgcn_wave_reduce_min_u32(x, 0);
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST int __clc_sub_group_reduce_min(int x) {
+ return __builtin_amdgcn_wave_reduce_min_i32(x, 0);
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST ulong __clc_sub_group_reduce_min(ulong x) {
+ return __builtin_amdgcn_wave_reduce_min_u64(x, 0);
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST long __clc_sub_group_reduce_min(long x) {
+ return __builtin_amdgcn_wave_reduce_min_i64(x, 0);
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint __clc_sub_group_reduce_max(uint x) {
+ return __builtin_amdgcn_wave_reduce_max_u32(x, 0);
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST int __clc_sub_group_reduce_max(int x) {
+ return __builtin_amdgcn_wave_reduce_max_i32(x, 0);
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST ulong __clc_sub_group_reduce_max(ulong x) {
+ return __builtin_amdgcn_wave_reduce_max_u32(x, 0);
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST long __clc_sub_group_reduce_max(long x) {
+ return __builtin_amdgcn_wave_reduce_max_i64(x, 0);
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST float __clc_sub_group_reduce_add(float x) {
+ return __builtin_amdgcn_wave_reduce_fadd_f32(x, 0);
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST double __clc_sub_group_reduce_add(double x) {
+ return __builtin_amdgcn_wave_reduce_fadd_f64(x, 0);
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST float __clc_sub_group_reduce_min(float x) {
+ return __builtin_amdgcn_wave_reduce_fmin_f32(x, 0);
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST double __clc_sub_group_reduce_min(double x) {
+ return __builtin_amdgcn_wave_reduce_fmin_f64(x, 0);
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST float __clc_sub_group_reduce_max(float x) {
+ return __builtin_amdgcn_wave_reduce_fmax_f32(x, 0);
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST double __clc_sub_group_reduce_max(double x) {
+ return __builtin_amdgcn_wave_reduce_fmax_f64(x, 0);
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST half __clc_sub_group_reduce_add(half x) {
+ // FIXME: There should be a direct half builtin available.
+ return (float)__clc_sub_group_reduce_add((float)x);
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST half __clc_sub_group_reduce_min(half x) {
+ // FIXME: There should be a direct half builtin available.
+ return (float)__clc_sub_group_reduce_min((float)x);
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST half __clc_sub_group_reduce_max(half x) {
+ // FIXME: There should be a direct half builtin available.
+ return (float)__clc_sub_group_reduce_max((float)x);
+}
diff --git a/libclc/clc/lib/generic/subgroup/clc_sub_group_reduce.cl b/libclc/clc/lib/generic/subgroup/clc_sub_group_reduce.cl
new file mode 100644
index 0000000000000..d34a30610dc20
--- /dev/null
+++ b/libclc/clc/lib/generic/subgroup/clc_sub_group_reduce.cl
@@ -0,0 +1,15 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "clc/synchronization/clc_sub_group_reduce.h"
+
+#define __CLC_BODY <clc_sub_group_reduce.inc>
+#include <clc/integer/gentype.inc>
+
+#define __CLC_BODY <clc_sub_group_reduce.inc>
+#include <clc/math/gentype.inc>
diff --git a/libclc/opencl/lib/generic/SOURCES b/libclc/opencl/lib/generic/SOURCES
index 0afbb67ebe1e4..415bd466cc58d 100644
--- a/libclc/opencl/lib/generic/SOURCES
+++ b/libclc/opencl/lib/generic/SOURCES
@@ -205,6 +205,7 @@ shared/vload.cl
shared/vstore.cl
subgroup/subgroup.cl
subgroup/sub_group_broadcast.cl
+subgroup/sub_group_reduce.cl
synchronization/sub_group_barrier.cl
synchronization/work_group_barrier.cl
workitem/get_enqueued_local_size.cl
diff --git a/libclc/opencl/lib/generic/subgroup/sub_group_reduce.cl b/libclc/opencl/lib/generic/subgroup/sub_group_reduce.cl
new file mode 100644
index 0000000000000..1614fafd0c6eb
--- /dev/null
+++ b/libclc/opencl/lib/generic/subgroup/sub_group_reduce.cl
@@ -0,0 +1,15 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "clc/subgroup/clc_sub_group_reduce.h"
+
+#define __CLC_BODY <sub_group_reduce.inc>
+#include <clc/integer/gentype.inc>
+
+#define __CLC_BODY <sub_group_reduce.inc>
+#include <clc/math/gentype.inc>
diff --git a/libclc/opencl/lib/generic/subgroup/sub_group_reduce.inc b/libclc/opencl/lib/generic/subgroup/sub_group_reduce.inc
new file mode 100644
index 0000000000000..cd846f8b7b813
--- /dev/null
+++ b/libclc/opencl/lib/generic/subgroup/sub_group_reduce.inc
@@ -0,0 +1,27 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#if defined(__CLC_SCALAR) && \
+ (defined(__CLC_FPSIZE) || __CLC_GENSIZE == 32 || __CLC_GENSIZE == 64)
+
+_CLC_OVERLOAD _CLC_DEF _CLC_CONST __CLC_GENTYPE
+sub_group_reduce_add(__CLC_GENTYPE x) {
+ return __clc_sub_group_reduce_add(x);
+}
+
+_CLC_OVERLOAD _CLC_DEF _CLC_CONST __CLC_GENTYPE
+sub_group_reduce_min(__CLC_GENTYPE x) {
+ return __clc_sub_group_reduce_min(x);
+}
+
+_CLC_OVERLOAD _CLC_DEF _CLC_CONST __CLC_GENTYPE
+sub_group_reduce_max(__CLC_GENTYPE x) {
+ return __clc_sub_group_reduce_max(x);
+}
+
+#endif
More information about the cfe-commits
mailing list