[llvm-branch-commits] [libclc] libclc: Add sub_group_broadcast (PR #184846)
Matt Arsenault via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Thu Mar 5 15:06:35 PST 2026
https://github.com/arsenm updated https://github.com/llvm/llvm-project/pull/184846
>From 559951c6c549425eb3806422b0db74957ee3808e Mon Sep 17 00:00:00 2001
From: Matt Arsenault <Matthew.Arsenault at amd.com>
Date: Thu, 5 Mar 2026 18:26:23 +0100
Subject: [PATCH] libclc: Add sub_group_broadcast
---
.../clc/subgroup/sub_group_broadcast.h | 43 +++++++++++++++
libclc/clc/lib/amdgcn/SOURCES | 1 +
.../amdgcn/subgroup/sub_group_broadcast.cl | 55 +++++++++++++++++++
libclc/opencl/lib/generic/SOURCES | 1 +
.../generic/subgroup/sub_group_broadcast.cl | 32 +++++++++++
5 files changed, 132 insertions(+)
create mode 100644 libclc/clc/include/clc/subgroup/sub_group_broadcast.h
create mode 100644 libclc/clc/lib/amdgcn/subgroup/sub_group_broadcast.cl
create mode 100644 libclc/opencl/lib/generic/subgroup/sub_group_broadcast.cl
diff --git a/libclc/clc/include/clc/subgroup/sub_group_broadcast.h b/libclc/clc/include/clc/subgroup/sub_group_broadcast.h
new file mode 100644
index 0000000000000..6698e97182390
--- /dev/null
+++ b/libclc/clc/include/clc/subgroup/sub_group_broadcast.h
@@ -0,0 +1,43 @@
+//===----------------------------------------------------------------------===//
+//
+// 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_BROADCAST_H__
+#define __CLC_SUBGROUP_CLC_SUB_GROUP_BROADCAST_H__
+
+#include <clc/clcfunc.h>
+
+_CLC_DECL _CLC_OVERLOAD _CLC_CONST int
+__clc_sub_group_broadcast(int x, uint sub_group_local_id);
+
+_CLC_DECL _CLC_OVERLOAD _CLC_CONST uint
+__clc_sub_group_broadcast(uint x, uint sub_group_local_id);
+
+_CLC_DECL _CLC_OVERLOAD _CLC_CONST long
+__clc_sub_group_broadcast(long x, uint sub_group_local_id);
+
+_CLC_DECL _CLC_OVERLOAD _CLC_CONST ulong
+__clc_sub_group_broadcast(ulong x, uint sub_group_local_id);
+
+#ifdef cl_khr_fp16
+#pragma OPENCL EXTENSION cl_khr_fp16 : enable
+
+_CLC_DECL _CLC_OVERLOAD _CLC_CONST half
+__clc_sub_group_broadcast(half x, uint sub_group_local_id);
+#endif
+
+_CLC_DECL _CLC_OVERLOAD _CLC_CONST float
+__clc_sub_group_broadcast(float x, uint sub_group_local_id);
+
+#ifdef cl_khr_fp64
+#pragma OPENCL EXTENSION cl_khr_fp64 : enable
+
+_CLC_DECL _CLC_OVERLOAD _CLC_CONST double
+__clc_sub_group_broadcast(double x, uint sub_group_local_id);
+#endif
+
+#endif // __CLC_SUBGROUP_CLC_SUB_GROUP_BROADCAST_H__
diff --git a/libclc/clc/lib/amdgcn/SOURCES b/libclc/clc/lib/amdgcn/SOURCES
index 7006f538d9270..959e4fb48e97a 100644
--- a/libclc/clc/lib/amdgcn/SOURCES
+++ b/libclc/clc/lib/amdgcn/SOURCES
@@ -1,6 +1,7 @@
address_space/qualifier.cl
math/clc_ldexp.cl
mem_fence/clc_mem_fence.cl
+subgroup/sub_group_broadcast.cl
synchronization/clc_work_group_barrier.cl
workitem/clc_get_enqueued_local_size.cl
workitem/clc_get_global_offset.cl
diff --git a/libclc/clc/lib/amdgcn/subgroup/sub_group_broadcast.cl b/libclc/clc/lib/amdgcn/subgroup/sub_group_broadcast.cl
new file mode 100644
index 0000000000000..0a7ea8fc78e96
--- /dev/null
+++ b/libclc/clc/lib/amdgcn/subgroup/sub_group_broadcast.cl
@@ -0,0 +1,55 @@
+//===----------------------------------------------------------------------===//
+//
+// 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/sub_group_broadcast.h"
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint
+__clc_sub_group_broadcast(uint x, uint sub_group_local_id) {
+ uint j = __builtin_amdgcn_readfirstlane(sub_group_local_id);
+ return __builtin_amdgcn_readlane(x, j);
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST ulong
+__clc_sub_group_broadcast(ulong x, uint sub_group_local_id) {
+ uint j = __builtin_amdgcn_readfirstlane(sub_group_local_id);
+ uint2 as_vec = __builtin_astype(x, uint2);
+ as_vec.x = __builtin_amdgcn_readlane(as_vec.x, j);
+ as_vec.y = __builtin_amdgcn_readlane(as_vec.y, j);
+ return __builtin_astype(as_vec, ulong);
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST int
+__clc_sub_group_broadcast(int x, uint sub_group_local_id) {
+ return (int)__clc_sub_group_broadcast((uint)x, sub_group_local_id);
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST float
+__clc_sub_group_broadcast(float x, uint sub_group_local_id) {
+ uint broadcast =
+ __clc_sub_group_broadcast(__builtin_astype(x, uint), sub_group_local_id);
+ return __builtin_astype(broadcast, float);
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST long
+__clc_sub_group_broadcast(long x, uint sub_group_local_id) {
+ return (long)__clc_sub_group_broadcast((ulong)x, sub_group_local_id);
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST double
+__clc_sub_group_broadcast(double x, uint sub_group_local_id) {
+ uint bitcast = __builtin_astype(x, double);
+ ulong broadcast = __clc_sub_group_broadcast(bitcast, sub_group_local_id);
+ return __builtin_astype(broadcast, double);
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST half
+__clc_sub_group_broadcast(half x, uint sub_group_local_id) {
+ ushort bitcast = __builtin_astype(x, ushort);
+ uint broadcast = __clc_sub_group_broadcast((uint)bitcast, sub_group_local_id);
+ return __builtin_astype((ushort)broadcast, half);
+}
diff --git a/libclc/opencl/lib/generic/SOURCES b/libclc/opencl/lib/generic/SOURCES
index a2c1122ca7634..12527ec24fdbe 100644
--- a/libclc/opencl/lib/generic/SOURCES
+++ b/libclc/opencl/lib/generic/SOURCES
@@ -200,6 +200,7 @@ shared/max.cl
shared/min.cl
shared/vload.cl
shared/vstore.cl
+subgroup/sub_group_broadcast.cl
workitem/get_enqueued_local_size.cl
workitem/get_global_id.cl
workitem/get_global_size.cl
diff --git a/libclc/opencl/lib/generic/subgroup/sub_group_broadcast.cl b/libclc/opencl/lib/generic/subgroup/sub_group_broadcast.cl
new file mode 100644
index 0000000000000..e4978c1d2cfc1
--- /dev/null
+++ b/libclc/opencl/lib/generic/subgroup/sub_group_broadcast.cl
@@ -0,0 +1,32 @@
+//===----------------------------------------------------------------------===//
+//
+// 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/sub_group_broadcast.h"
+
+#define CLC_SUB_GROUP_BROADCAST_IMPL(GENTYPE) \
+ _CLC_OVERLOAD _CLC_DEF _CLC_CONST GENTYPE sub_group_broadcast( \
+ GENTYPE x, uint sub_group_local_id) { \
+ return __clc_sub_group_broadcast(x, sub_group_local_id); \
+ }
+
+#define CLC_SUB_GROUP_BROADCAST_TYPES() \
+ CLC_SUB_GROUP_BROADCAST_IMPL(int) \
+ CLC_SUB_GROUP_BROADCAST_IMPL(uint) \
+ CLC_SUB_GROUP_BROADCAST_IMPL(long) \
+ CLC_SUB_GROUP_BROADCAST_IMPL(ulong) \
+ CLC_SUB_GROUP_BROADCAST_IMPL(float)
+
+CLC_SUB_GROUP_BROADCAST_TYPES()
+
+#ifdef cl_khr_fp16
+CLC_SUB_GROUP_BROADCAST_IMPL(half)
+#endif
+
+#ifdef cl_khr_fp64
+CLC_SUB_GROUP_BROADCAST_IMPL(double)
+#endif
More information about the llvm-branch-commits
mailing list