https://github.com/arsenm created https://github.com/llvm/llvm-project/pull/185208
None >From 4dc91679dd9d17855c8376853e1a1bad809dcc9d Mon Sep 17 00:00:00 2001 From: Matt Arsenault <[email protected]> Date: Sat, 7 Mar 2026 18:03:47 +0100 Subject: [PATCH] libclc: Move sub_group_barrier to clc --- .../synchronization/clc_sub_group_barrier.h | 21 +++++++++++++++++++ libclc/clc/lib/amdgcn/SOURCES | 1 + .../synchronization/clc_sub_group_barrier.cl | 18 ++++++++++++++++ libclc/clc/lib/generic/SOURCES | 1 + .../lib/generic/subgroup/sub_group_barrier.cl | 14 +++++++++++++ libclc/opencl/lib/amdgcn/SOURCES | 1 - libclc/opencl/lib/generic/SOURCES | 1 + .../synchronization/sub_group_barrier.cl | 14 +++++++------ 8 files changed, 64 insertions(+), 7 deletions(-) create mode 100644 libclc/clc/include/clc/synchronization/clc_sub_group_barrier.h create mode 100644 libclc/clc/lib/amdgcn/synchronization/clc_sub_group_barrier.cl create mode 100644 libclc/clc/lib/generic/subgroup/sub_group_barrier.cl rename libclc/opencl/lib/{amdgcn => generic}/synchronization/sub_group_barrier.cl (56%) diff --git a/libclc/clc/include/clc/synchronization/clc_sub_group_barrier.h b/libclc/clc/include/clc/synchronization/clc_sub_group_barrier.h new file mode 100644 index 0000000000000..58ca3d5b5a028 --- /dev/null +++ b/libclc/clc/include/clc/synchronization/clc_sub_group_barrier.h @@ -0,0 +1,21 @@ +//===----------------------------------------------------------------------===// +// +// 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_BARRIER_H__ +#define __CLC_SUBGROUP_CLC_SUB_GROUP_BARRIER_H__ + +#include "clc/internal/clc.h" +#include "clc/mem_fence/clc_mem_semantic.h" + +_CLC_DECL _CLC_OVERLOAD void +__clc_sub_group_barrier(__CLC_MemorySemantics memory_semantics, + int memory_scope); +_CLC_DECL _CLC_OVERLOAD void +__clc_sub_group_barrier(__CLC_MemorySemantics memory_semantics); + +#endif // __CLC_SUBGROUP_CLC_SUB_GROUP_BARRIER_H__ diff --git a/libclc/clc/lib/amdgcn/SOURCES b/libclc/clc/lib/amdgcn/SOURCES index 28ce5f1943825..a280461b1664a 100644 --- a/libclc/clc/lib/amdgcn/SOURCES +++ b/libclc/clc/lib/amdgcn/SOURCES @@ -2,6 +2,7 @@ address_space/qualifier.cl math/clc_ldexp.cl mem_fence/clc_mem_fence.cl subgroup/sub_group_broadcast.cl +synchronization/clc_sub_group_barrier.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/synchronization/clc_sub_group_barrier.cl b/libclc/clc/lib/amdgcn/synchronization/clc_sub_group_barrier.cl new file mode 100644 index 0000000000000..ae26b1125f5c0 --- /dev/null +++ b/libclc/clc/lib/amdgcn/synchronization/clc_sub_group_barrier.cl @@ -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 +// +//===----------------------------------------------------------------------===// + +#include "clc/mem_fence/clc_mem_fence.h" +#include "clc/synchronization/clc_sub_group_barrier.h" + +_CLC_DEF _CLC_OVERLOAD void +__clc_sub_group_barrier(__CLC_MemorySemantics memory_semantics, int scope) { + __builtin_amdgcn_wave_barrier(); + + if (memory_semantics) + __clc_mem_fence(scope, __ATOMIC_ACQ_REL, memory_semantics); +} diff --git a/libclc/clc/lib/generic/SOURCES b/libclc/clc/lib/generic/SOURCES index b352382fbe6d9..fb991d1ac624e 100644 --- a/libclc/clc/lib/generic/SOURCES +++ b/libclc/clc/lib/generic/SOURCES @@ -176,6 +176,7 @@ shared/clc_min.cl shared/clc_qualifier.cl shared/clc_vload.cl shared/clc_vstore.cl +subgroup/sub_group_barrier.cl workitem/clc_get_global_linear_id.cl workitem/clc_get_local_linear_id.cl workitem/clc_get_num_sub_groups.cl diff --git a/libclc/clc/lib/generic/subgroup/sub_group_barrier.cl b/libclc/clc/lib/generic/subgroup/sub_group_barrier.cl new file mode 100644 index 0000000000000..7a50acf2c54d2 --- /dev/null +++ b/libclc/clc/lib/generic/subgroup/sub_group_barrier.cl @@ -0,0 +1,14 @@ +//===----------------------------------------------------------------------===// +// +// 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_barrier.h" + +_CLC_DEF _CLC_OVERLOAD void +__clc_sub_group_barrier(__CLC_MemorySemantics memory_semantics) { + __clc_sub_group_barrier(memory_semantics, __MEMORY_SCOPE_WRKGRP); +} diff --git a/libclc/opencl/lib/amdgcn/SOURCES b/libclc/opencl/lib/amdgcn/SOURCES index a21f6f0ad44b0..524db744339eb 100644 --- a/libclc/opencl/lib/amdgcn/SOURCES +++ b/libclc/opencl/lib/amdgcn/SOURCES @@ -1,2 +1 @@ subgroup/subgroup.cl -synchronization/sub_group_barrier.cl diff --git a/libclc/opencl/lib/generic/SOURCES b/libclc/opencl/lib/generic/SOURCES index 398e326ffe482..13e8cd97f41c4 100644 --- a/libclc/opencl/lib/generic/SOURCES +++ b/libclc/opencl/lib/generic/SOURCES @@ -201,6 +201,7 @@ shared/min.cl shared/vload.cl shared/vstore.cl subgroup/sub_group_broadcast.cl +synchronization/sub_group_barrier.cl synchronization/work_group_barrier.cl workitem/get_enqueued_local_size.cl workitem/get_global_id.cl diff --git a/libclc/opencl/lib/amdgcn/synchronization/sub_group_barrier.cl b/libclc/opencl/lib/generic/synchronization/sub_group_barrier.cl similarity index 56% rename from libclc/opencl/lib/amdgcn/synchronization/sub_group_barrier.cl rename to libclc/opencl/lib/generic/synchronization/sub_group_barrier.cl index 2b57d86294ef5..e03eb01cd767f 100644 --- a/libclc/opencl/lib/amdgcn/synchronization/sub_group_barrier.cl +++ b/libclc/opencl/lib/generic/synchronization/sub_group_barrier.cl @@ -6,16 +6,18 @@ // //===----------------------------------------------------------------------===// -#include <clc/opencl/opencl-base.h> +#include "clc/opencl/synchronization/utils.h" +#include "clc/opencl/utils.h" +#include "clc/synchronization/clc_sub_group_barrier.h" _CLC_DEF _CLC_OVERLOAD void sub_group_barrier(cl_mem_fence_flags flags, memory_scope scope) { - __builtin_amdgcn_wave_barrier(); - - if (flags) - atomic_work_item_fence(flags, memory_order_acq_rel, scope); + __CLC_MemorySemantics memory_semantics = __opencl_get_memory_semantics(flags); + __clc_sub_group_barrier(memory_semantics, + __opencl_get_clang_memory_scope(scope)); } _CLC_DEF _CLC_OVERLOAD void sub_group_barrier(cl_mem_fence_flags flags) { - sub_group_barrier(flags, memory_scope_sub_group); + __CLC_MemorySemantics memory_semantics = __opencl_get_memory_semantics(flags); + __clc_sub_group_barrier(memory_semantics); } _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
