[llvm-branch-commits] [libclc] libclc: Add work_group_broadcast (PR #185261)
Matt Arsenault via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Sun Mar 8 00:14:38 PST 2026
https://github.com/arsenm created https://github.com/llvm/llvm-project/pull/185261
None
>From 8ba783764fa69e5fea9c716f67850378c3521dbc Mon Sep 17 00:00:00 2001
From: Matt Arsenault <Matthew.Arsenault at amd.com>
Date: Sat, 7 Mar 2026 21:00:20 +0100
Subject: [PATCH] libclc: Add work_group_broadcast
---
.../clc/collective/clc_work_group_broadcast.h | 18 +++++++
.../collective/clc_work_group_broadcast.inc | 16 ++++++
libclc/clc/lib/generic/SOURCES | 1 +
.../collective/clc_work_group_broadcast.cl | 24 +++++++++
.../collective/clc_work_group_broadcast.inc | 51 +++++++++++++++++++
libclc/opencl/lib/generic/SOURCES | 1 +
.../collective/work_group_broadcast.cl | 15 ++++++
.../collective/work_group_broadcast.inc | 28 ++++++++++
8 files changed, 154 insertions(+)
create mode 100644 libclc/clc/include/clc/collective/clc_work_group_broadcast.h
create mode 100644 libclc/clc/include/clc/collective/clc_work_group_broadcast.inc
create mode 100644 libclc/clc/lib/generic/collective/clc_work_group_broadcast.cl
create mode 100644 libclc/clc/lib/generic/collective/clc_work_group_broadcast.inc
create mode 100644 libclc/opencl/lib/generic/collective/work_group_broadcast.cl
create mode 100644 libclc/opencl/lib/generic/collective/work_group_broadcast.inc
diff --git a/libclc/clc/include/clc/collective/clc_work_group_broadcast.h b/libclc/clc/include/clc/collective/clc_work_group_broadcast.h
new file mode 100644
index 0000000000000..e942edf6bbd64
--- /dev/null
+++ b/libclc/clc/include/clc/collective/clc_work_group_broadcast.h
@@ -0,0 +1,18 @@
+//===----------------------------------------------------------------------===//
+//
+// 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_COLLECTIVE_CLC_WORK_GROUP_BROADCAST_H__
+#define __CLC_COLLECTIVE_CLC_WORK_GROUP_BROADCAST_H__
+
+#define __CLC_BODY <clc/collective/clc_work_group_broadcast.inc>
+#include <clc/integer/gentype.inc>
+
+#define __CLC_BODY <clc/collective/clc_work_group_broadcast.inc>
+#include <clc/math/gentype.inc>
+
+#endif // __CLC_COLLECTIVE_CLC_WORK_GROUP_BROADCAST_H__
diff --git a/libclc/clc/include/clc/collective/clc_work_group_broadcast.inc b/libclc/clc/include/clc/collective/clc_work_group_broadcast.inc
new file mode 100644
index 0000000000000..64633fa06a047
--- /dev/null
+++ b/libclc/clc/include/clc/collective/clc_work_group_broadcast.inc
@@ -0,0 +1,16 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE
+__clc_work_group_broadcast(__CLC_GENTYPE a, size_t local_id_x);
+
+_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE __clc_work_group_broadcast(
+ __CLC_GENTYPE a, size_t local_id_x, size_t local_id_y);
+
+_CLC_OVERLOAD _CLC_DECL __CLC_GENTYPE __clc_work_group_broadcast(
+ __CLC_GENTYPE a, size_t local_id_x, size_t local_id_y, size_t local_id_z);
diff --git a/libclc/clc/lib/generic/SOURCES b/libclc/clc/lib/generic/SOURCES
index 0e9166a2b249e..05ea37869bbf8 100644
--- a/libclc/clc/lib/generic/SOURCES
+++ b/libclc/clc/lib/generic/SOURCES
@@ -15,6 +15,7 @@ atomic/clc_atomic_inc.cl
atomic/clc_atomic_load.cl
atomic/clc_atomic_store.cl
collective/clc_work_group_any_all.cl
+collective/clc_work_group_broadcast.cl
common/clc_degrees.cl
common/clc_radians.cl
common/clc_sign.cl
diff --git a/libclc/clc/lib/generic/collective/clc_work_group_broadcast.cl b/libclc/clc/lib/generic/collective/clc_work_group_broadcast.cl
new file mode 100644
index 0000000000000..4db8df99ce234
--- /dev/null
+++ b/libclc/clc/lib/generic/collective/clc_work_group_broadcast.cl
@@ -0,0 +1,24 @@
+//===----------------------------------------------------------------------===//
+//
+// 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/atomic/clc_atomic_load.h"
+#include "clc/atomic/clc_atomic_store.h"
+#include "clc/collective/clc_work_group_broadcast.h"
+#include "clc/internal/clc.h"
+#include "clc/subgroup/clc_sub_group_broadcast.h"
+#include "clc/subgroup/clc_subgroup.h"
+#include "clc/synchronization/clc_work_group_barrier.h"
+#include "clc/workitem/clc_get_local_id.h"
+
+#pragma OPENCL EXTENSION __cl_clang_function_scope_local_variables : enable
+
+#define __CLC_BODY <clc_work_group_broadcast.inc>
+#include <clc/integer/gentype.inc>
+
+#define __CLC_BODY <clc_work_group_broadcast.inc>
+#include <clc/math/gentype.inc>
diff --git a/libclc/clc/lib/generic/collective/clc_work_group_broadcast.inc b/libclc/clc/lib/generic/collective/clc_work_group_broadcast.inc
new file mode 100644
index 0000000000000..a431560cf10d0
--- /dev/null
+++ b/libclc/clc/lib/generic/collective/clc_work_group_broadcast.inc
@@ -0,0 +1,51 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifdef __CLC_SCALAR
+
+_CLC_OVERLOAD
+static __CLC_GENTYPE __clc_work_group_broadcast_impl(__CLC_GENTYPE a,
+ bool is_leader) {
+ __local __CLC_GENTYPE scratch;
+ if (is_leader) {
+ __scoped_atomic_store_n(&scratch, a, __ATOMIC_RELAXED,
+ __MEMORY_SCOPE_WRKGRP);
+ }
+
+ __clc_work_group_barrier(__MEMORY_SCOPE_WRKGRP, __CLC_MEMORY_LOCAL);
+ __CLC_GENTYPE result =
+ __scoped_atomic_load_n(&scratch, __ATOMIC_RELAXED, __MEMORY_SCOPE_WRKGRP);
+ __clc_work_group_barrier(__MEMORY_SCOPE_WRKGRP, __CLC_MEMORY_LOCAL);
+ return result;
+}
+
+_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE
+__clc_work_group_broadcast(__CLC_GENTYPE a, size_t local_id_x) {
+ if (__clc_get_num_sub_groups() == 1)
+ return __clc_sub_group_broadcast(a, local_id_x);
+
+ bool is_leader = __clc_get_local_id(0) == local_id_x;
+ return __clc_work_group_broadcast_impl(a, is_leader);
+}
+
+_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE __clc_work_group_broadcast(
+ __CLC_GENTYPE a, size_t local_id_x, size_t local_id_y) {
+ bool is_leader = __clc_get_local_id(0) == local_id_x &&
+ __clc_get_local_id(1) == local_id_y;
+ return __clc_work_group_broadcast_impl(a, is_leader);
+}
+
+_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE __clc_work_group_broadcast(
+ __CLC_GENTYPE a, size_t local_id_x, size_t local_id_y, size_t local_id_z) {
+ bool is_leader = __clc_get_local_id(0) == local_id_x &&
+ __clc_get_local_id(1) == local_id_y &&
+ __clc_get_local_id(2) == local_id_z;
+ return __clc_work_group_broadcast_impl(a, is_leader);
+}
+
+#endif // __CLC_SCALAR
diff --git a/libclc/opencl/lib/generic/SOURCES b/libclc/opencl/lib/generic/SOURCES
index 5dc92c27f9d74..0afbb67ebe1e4 100644
--- a/libclc/opencl/lib/generic/SOURCES
+++ b/libclc/opencl/lib/generic/SOURCES
@@ -43,6 +43,7 @@ atomic/atom_xchg.cl
atomic/atom_xor.cl
atomic/atomic_work_item_fence.cl
collective/work_group_any_all.cl
+collective/work_group_broadcast.cl
common/degrees.cl
common/mix.cl
common/radians.cl
diff --git a/libclc/opencl/lib/generic/collective/work_group_broadcast.cl b/libclc/opencl/lib/generic/collective/work_group_broadcast.cl
new file mode 100644
index 0000000000000..fa4e8fbc26bb3
--- /dev/null
+++ b/libclc/opencl/lib/generic/collective/work_group_broadcast.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/collective/clc_work_group_broadcast.h"
+
+#define __CLC_BODY <work_group_broadcast.inc>
+#include <clc/integer/gentype.inc>
+
+#define __CLC_BODY <work_group_broadcast.inc>
+#include <clc/math/gentype.inc>
diff --git a/libclc/opencl/lib/generic/collective/work_group_broadcast.inc b/libclc/opencl/lib/generic/collective/work_group_broadcast.inc
new file mode 100644
index 0000000000000..5c1d2098b8861
--- /dev/null
+++ b/libclc/opencl/lib/generic/collective/work_group_broadcast.inc
@@ -0,0 +1,28 @@
+//===----------------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifdef __CLC_SCALAR
+
+_CLC_OVERLOAD _CLC_DEF void work_group_broadcast(__CLC_GENTYPE a,
+ size_t local_id_x) {
+ __clc_work_group_broadcast(a, local_id_x);
+}
+
+_CLC_OVERLOAD _CLC_DEF void
+work_group_broadcast(__CLC_GENTYPE a, size_t local_id_x, size_t local_id_y) {
+ __clc_work_group_broadcast(a, local_id_x, local_id_y);
+}
+
+_CLC_OVERLOAD _CLC_DEF void work_group_broadcast(__CLC_GENTYPE a,
+ size_t local_id_x,
+ size_t local_id_y,
+ size_t local_id_z) {
+ __clc_work_group_broadcast(a, local_id_x, local_id_y, local_id_z);
+}
+
+#endif
More information about the llvm-branch-commits
mailing list