[llvm-branch-commits] [libclc] libclc: Move subgroup functions into clc (PR #185220)
Matt Arsenault via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Sat Mar 7 11:17:15 PST 2026
https://github.com/arsenm created https://github.com/llvm/llvm-project/pull/185220
It turns out there was a generic implementation of the id and sizes.
The practice of splitting every single function into its own file is
kind of a pain here, so introduce a utility header for amdgpu.
>From 2a38eacb6da29bf2aaf2799f2f696a4c5a72dc1c Mon Sep 17 00:00:00 2001
From: Matt Arsenault <Matthew.Arsenault at amd.com>
Date: Sat, 7 Mar 2026 19:14:18 +0100
Subject: [PATCH] libclc: Move subgroup functions into clc
It turns out there was a generic implementation of the id and sizes.
The practice of splitting every single function into its own file is
kind of a pain here, so introduce a utility header for amdgpu.
---
.../clc/include/clc/subgroup/clc_subgroup.h | 23 +++++++
libclc/clc/lib/amdgcn/SOURCES | 4 ++
libclc/clc/lib/amdgcn/subgroup/subgroup.cl | 49 +++++++++++++++
.../workitem/clc_get_max_sub_group_size.cl | 4 +-
.../amdgcn/workitem/clc_get_num_sub_groups.cl | 17 ++++++
.../amdgcn/workitem/clc_get_sub_group_id.cl | 15 +++++
.../amdgcn/workitem/clc_get_sub_group_size.cl | 20 +++++++
libclc/opencl/lib/amdgcn/SOURCES | 2 +-
libclc/opencl/lib/amdgcn/subgroup/subgroup.cl | 60 -------------------
libclc/opencl/lib/generic/SOURCES | 1 +
.../opencl/lib/generic/subgroup/subgroup.cl | 41 +++++++++++++
11 files changed, 174 insertions(+), 62 deletions(-)
create mode 100644 libclc/clc/include/clc/subgroup/clc_subgroup.h
create mode 100644 libclc/clc/lib/amdgcn/subgroup/subgroup.cl
create mode 100644 libclc/clc/lib/amdgcn/workitem/clc_get_num_sub_groups.cl
create mode 100644 libclc/clc/lib/amdgcn/workitem/clc_get_sub_group_id.cl
create mode 100644 libclc/clc/lib/amdgcn/workitem/clc_get_sub_group_size.cl
delete mode 100644 libclc/opencl/lib/amdgcn/subgroup/subgroup.cl
create mode 100644 libclc/opencl/lib/generic/subgroup/subgroup.cl
diff --git a/libclc/clc/include/clc/subgroup/clc_subgroup.h b/libclc/clc/include/clc/subgroup/clc_subgroup.h
new file mode 100644
index 0000000000000..f0a2a11d48445
--- /dev/null
+++ b/libclc/clc/include/clc/subgroup/clc_subgroup.h
@@ -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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef __CLC_SUBGROUP_CLC_SUB_GROUP_SUBGROUP_H__
+#define __CLC_SUBGROUP_CLC_SUB_GROUP_SUBGROUP_H__
+
+#include "clc/internal/clc.h"
+
+_CLC_DECL _CLC_OVERLOAD _CLC_CONST uint __clc_get_sub_group_size(void);
+_CLC_DECL _CLC_OVERLOAD _CLC_CONST uint __clc_get_max_sub_group_size(void);
+_CLC_DECL _CLC_OVERLOAD _CLC_CONST uint __clc_get_num_sub_groups(void);
+_CLC_DECL _CLC_OVERLOAD _CLC_CONST uint __clc_get_enqueued_num_sub_groups(void);
+_CLC_DECL _CLC_OVERLOAD _CLC_CONST uint __clc_get_sub_group_id(void);
+_CLC_DECL _CLC_OVERLOAD _CLC_CONST uint __clc_get_sub_group_local_id(void);
+_CLC_DECL _CLC_OVERLOAD _CLC_CONST int __clc_sub_group_all(int x);
+_CLC_DECL _CLC_OVERLOAD _CLC_CONST int __clc_sub_group_any(int x);
+
+#endif // __CLC_SUBGROUP_CLC_SUB_GROUP_SUBGROUP_H__
diff --git a/libclc/clc/lib/amdgcn/SOURCES b/libclc/clc/lib/amdgcn/SOURCES
index a280461b1664a..a0b6c168b207e 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/subgroup.cl
subgroup/sub_group_broadcast.cl
synchronization/clc_sub_group_barrier.cl
synchronization/clc_work_group_barrier.cl
@@ -12,4 +13,7 @@ workitem/clc_get_local_id.cl
workitem/clc_get_local_size.cl
workitem/clc_get_max_sub_group_size.cl
workitem/clc_get_num_groups.cl
+workitem/clc_get_num_sub_groups.cl
+workitem/clc_get_sub_group_id.cl
+workitem/clc_get_sub_group_size.cl
workitem/clc_get_work_dim.cl
diff --git a/libclc/clc/lib/amdgcn/subgroup/subgroup.cl b/libclc/clc/lib/amdgcn/subgroup/subgroup.cl
new file mode 100644
index 0000000000000..369d74756f986
--- /dev/null
+++ b/libclc/clc/lib/amdgcn/subgroup/subgroup.cl
@@ -0,0 +1,49 @@
+//===----------------------------------------------------------------------===//
+//
+// 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/amdgpu/amdgpu_utils.h"
+#include "clc/integer/clc_mul24.h"
+#include "clc/shared/clc_min.h"
+#include "clc/subgroup/clc_subgroup.h"
+#include "clc/workitem/clc_get_enqueued_local_size.h"
+#include "clc/workitem/clc_get_local_linear_id.h"
+#include "clc/workitem/clc_get_local_size.h"
+#include "clc/workitem/clc_get_max_sub_group_size.h"
+
+static uint enqueued_workgroup_size() {
+ return __clc_mul24((uint)__clc_get_enqueued_local_size(2),
+ __clc_mul24((uint)__clc_get_enqueued_local_size(1),
+ (uint)__clc_get_enqueued_local_size(0)));
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint __clc_get_max_sub_group_size(void) {
+ return __clc_min(__builtin_amdgcn_wavefrontsize(), enqueued_workgroup_size());
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint __clc_get_num_sub_groups(void) {
+ return (__clc_amdgpu_workgroup_size() + __builtin_amdgcn_wavefrontsize() -
+ 1) >>
+ __clc_amdgpu_wavesize_log2();
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint __clc_get_enqueued_num_sub_groups(void) {
+ return (enqueued_workgroup_size() + __builtin_amdgcn_wavefrontsize() - 1) >>
+ __clc_amdgpu_wavesize_log2();
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint __clc_get_sub_group_local_id(void) {
+ return __builtin_amdgcn_mbcnt_hi(~0u, __builtin_amdgcn_mbcnt_lo(~0u, 0u));
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST int __clc_sub_group_all(int x) {
+ return __builtin_amdgcn_ballot_w64(x) == __builtin_amdgcn_read_exec();
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST int __clc_sub_group_any(int x) {
+ return __builtin_amdgcn_ballot_w64(x) != 0;
+}
diff --git a/libclc/clc/lib/amdgcn/workitem/clc_get_max_sub_group_size.cl b/libclc/clc/lib/amdgcn/workitem/clc_get_max_sub_group_size.cl
index cc56f8d9c325d..a947d44e37531 100644
--- a/libclc/clc/lib/amdgcn/workitem/clc_get_max_sub_group_size.cl
+++ b/libclc/clc/lib/amdgcn/workitem/clc_get_max_sub_group_size.cl
@@ -6,8 +6,10 @@
//
//===----------------------------------------------------------------------===//
+#if 0
#include <clc/workitem/clc_get_max_sub_group_size.h>
_CLC_OVERLOAD _CLC_DEF uint __clc_get_max_sub_group_size() {
- return __builtin_amdgcn_wavefrontsize();
+ return __clc_min(__builtin_amdgcn_wavefrontsize(), enqueued_workgroup_size());
}
+#endif
diff --git a/libclc/clc/lib/amdgcn/workitem/clc_get_num_sub_groups.cl b/libclc/clc/lib/amdgcn/workitem/clc_get_num_sub_groups.cl
new file mode 100644
index 0000000000000..2e68adc17ffd4
--- /dev/null
+++ b/libclc/clc/lib/amdgcn/workitem/clc_get_num_sub_groups.cl
@@ -0,0 +1,17 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 0
+#include <clc/workitem/clc_get_local_size.h>
+#include <clc/workitem/clc_get_max_sub_group_size.h>
+#include <clc/workitem/clc_get_num_sub_groups.h>
+
+_CLC_OVERLOAD _CLC_DEF uint __clc_get_num_sub_groups() {
+
+}
+#endif
diff --git a/libclc/clc/lib/amdgcn/workitem/clc_get_sub_group_id.cl b/libclc/clc/lib/amdgcn/workitem/clc_get_sub_group_id.cl
new file mode 100644
index 0000000000000..ba3baf98bda14
--- /dev/null
+++ b/libclc/clc/lib/amdgcn/workitem/clc_get_sub_group_id.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/amdgpu/amdgpu_utils.h"
+#include "clc/workitem/clc_get_local_linear_id.h"
+#include "clc/workitem/clc_get_sub_group_id.h"
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint __clc_get_sub_group_id(void) {
+ return (uint)__clc_get_local_linear_id() >> __clc_amdgpu_wavesize_log2();
+}
diff --git a/libclc/clc/lib/amdgcn/workitem/clc_get_sub_group_size.cl b/libclc/clc/lib/amdgcn/workitem/clc_get_sub_group_size.cl
new file mode 100644
index 0000000000000..585e17ef36f58
--- /dev/null
+++ b/libclc/clc/lib/amdgcn/workitem/clc_get_sub_group_size.cl
@@ -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
+//
+//===----------------------------------------------------------------------===//
+
+#include "clc/amdgpu/amdgpu_utils.h"
+#include "clc/integer/clc_mul24.h"
+#include "clc/shared/clc_min.h"
+#include "clc/workitem/clc_get_local_linear_id.h"
+#include "clc/workitem/clc_get_local_size.h"
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint __clc_get_sub_group_size(void) {
+ uint wavesize = __builtin_amdgcn_wavefrontsize();
+ uint lid = (uint)__clc_get_local_linear_id();
+ return __clc_min(wavesize,
+ __clc_amdgpu_workgroup_size() - (lid & ~(wavesize - 1)));
+}
diff --git a/libclc/opencl/lib/amdgcn/SOURCES b/libclc/opencl/lib/amdgcn/SOURCES
index 524db744339eb..8b137891791fe 100644
--- a/libclc/opencl/lib/amdgcn/SOURCES
+++ b/libclc/opencl/lib/amdgcn/SOURCES
@@ -1 +1 @@
-subgroup/subgroup.cl
+
diff --git a/libclc/opencl/lib/amdgcn/subgroup/subgroup.cl b/libclc/opencl/lib/amdgcn/subgroup/subgroup.cl
deleted file mode 100644
index d67d84e763b4f..0000000000000
--- a/libclc/opencl/lib/amdgcn/subgroup/subgroup.cl
+++ /dev/null
@@ -1,60 +0,0 @@
-//===----------------------------------------------------------------------===//
-//
-// 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/opencl/opencl-base.h>
-
-static uint wavesize_log2() {
- return __builtin_amdgcn_wavefrontsize() == 64 ? 6 : 5;
-}
-
-static uint workgroup_size() {
- return mul24((uint)get_local_size(2),
- mul24((uint)get_local_size(1), (uint)get_local_size(0)));
-}
-
-static uint enqueued_workgroup_size() {
- return mul24((uint)get_enqueued_local_size(2),
- mul24((uint)get_enqueued_local_size(1),
- (uint)get_enqueued_local_size(0)));
-}
-
-_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint get_sub_group_size(void) {
- uint wavesize = __builtin_amdgcn_wavefrontsize();
- uint lid = (uint)get_local_linear_id();
- return min(wavesize, workgroup_size() - (lid & ~(wavesize - 1)));
-}
-
-_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint get_max_sub_group_size(void) {
- return min(__builtin_amdgcn_wavefrontsize(), enqueued_workgroup_size());
-}
-
-_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint get_num_sub_groups(void) {
- return (workgroup_size() + __builtin_amdgcn_wavefrontsize() - 1) >>
- wavesize_log2();
-}
-
-_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint get_enqueued_num_sub_groups(void) {
- return (enqueued_workgroup_size() + __builtin_amdgcn_wavefrontsize() - 1) >>
- wavesize_log2();
-}
-
-_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint get_sub_group_id(void) {
- return (uint)get_local_linear_id() >> wavesize_log2();
-}
-
-_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint get_sub_group_local_id(void) {
- return __builtin_amdgcn_mbcnt_hi(~0u, __builtin_amdgcn_mbcnt_lo(~0u, 0u));
-}
-
-_CLC_DEF _CLC_OVERLOAD _CLC_CONST int sub_group_all(int x) {
- return __builtin_amdgcn_ballot_w64(x) == __builtin_amdgcn_read_exec();
-}
-
-_CLC_DEF _CLC_OVERLOAD _CLC_CONST int sub_group_any(int x) {
- return __builtin_amdgcn_ballot_w64(x) != 0;
-}
diff --git a/libclc/opencl/lib/generic/SOURCES b/libclc/opencl/lib/generic/SOURCES
index 13e8cd97f41c4..ec22a713fe06a 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/subgroup.cl
subgroup/sub_group_broadcast.cl
synchronization/sub_group_barrier.cl
synchronization/work_group_barrier.cl
diff --git a/libclc/opencl/lib/generic/subgroup/subgroup.cl b/libclc/opencl/lib/generic/subgroup/subgroup.cl
new file mode 100644
index 0000000000000..fd552ada4afaf
--- /dev/null
+++ b/libclc/opencl/lib/generic/subgroup/subgroup.cl
@@ -0,0 +1,41 @@
+//===----------------------------------------------------------------------===//
+//
+// 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_subgroup.h"
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint get_sub_group_size(void) {
+ return __clc_get_sub_group_size();
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint get_max_sub_group_size(void) {
+ return __clc_get_max_sub_group_size();
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint get_num_sub_groups(void) {
+ return __clc_get_num_sub_groups();
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint get_enqueued_num_sub_groups(void) {
+ return __clc_get_enqueued_num_sub_groups();
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint get_sub_group_id(void) {
+ return __clc_get_sub_group_id();
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST uint get_sub_group_local_id(void) {
+ return __clc_get_sub_group_local_id();
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST int sub_group_all(int x) {
+ return __clc_sub_group_all(x);
+}
+
+_CLC_DEF _CLC_OVERLOAD _CLC_CONST int sub_group_any(int x) {
+ return __clc_sub_group_any(x);
+}
More information about the llvm-branch-commits
mailing list