[libclc] abb3b98 - libclc: Add work_group_broadcast (#185261)
via cfe-commits
cfe-commits at lists.llvm.org
Sun Mar 8 00:54:34 PST 2026
Author: Matt Arsenault
Date: 2026-03-08T09:54:30+01:00
New Revision: abb3b981a37a04e56cea6439961d6bb3ee058f67
URL: https://github.com/llvm/llvm-project/commit/abb3b981a37a04e56cea6439961d6bb3ee058f67
DIFF: https://github.com/llvm/llvm-project/commit/abb3b981a37a04e56cea6439961d6bb3ee058f67.diff
LOG: libclc: Add work_group_broadcast (#185261)
Added:
libclc/clc/include/clc/collective/clc_work_group_broadcast.h
libclc/clc/include/clc/collective/clc_work_group_broadcast.inc
libclc/clc/lib/generic/collective/clc_work_group_broadcast.cl
libclc/clc/lib/generic/collective/clc_work_group_broadcast.inc
libclc/opencl/lib/generic/collective/work_group_broadcast.cl
libclc/opencl/lib/generic/collective/work_group_broadcast.inc
Modified:
libclc/clc/lib/generic/SOURCES
libclc/opencl/lib/generic/SOURCES
Removed:
################################################################################
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..b8a5a1dff1f31
--- /dev/null
+++ b/libclc/clc/include/clc/collective/clc_work_group_broadcast.h
@@ -0,0 +1,20 @@
+//===----------------------------------------------------------------------===//
+//
+// 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__
+
+#include "clc/internal/clc.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..d54bda65773f8
--- /dev/null
+++ b/libclc/clc/lib/generic/collective/clc_work_group_broadcast.cl
@@ -0,0 +1,23 @@
+//===----------------------------------------------------------------------===//
+//
+// 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/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..7b9695a7181c3
--- /dev/null
+++ b/libclc/opencl/lib/generic/collective/work_group_broadcast.inc
@@ -0,0 +1,29 @@
+//===----------------------------------------------------------------------===//
+//
+// 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_OVERLOAD _CLC_DEF __CLC_GENTYPE work_group_broadcast(__CLC_GENTYPE a,
+ size_t local_id_x) {
+ return __clc_work_group_broadcast(a, local_id_x);
+}
+
+_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE work_group_broadcast(__CLC_GENTYPE a,
+ size_t local_id_x,
+ size_t local_id_y) {
+ return __clc_work_group_broadcast(a, local_id_x, local_id_y);
+}
+
+_CLC_OVERLOAD _CLC_DEF __CLC_GENTYPE work_group_broadcast(__CLC_GENTYPE a,
+ size_t local_id_x,
+ size_t local_id_y,
+ size_t local_id_z) {
+ return __clc_work_group_broadcast(a, local_id_x, local_id_y, local_id_z);
+}
+
+#endif
More information about the cfe-commits
mailing list