From 9d696f52723a7fc7cd2c25d61faaed8dc1f1fb8b Mon Sep 17 00:00:00 2001 From: Adarsh Yoga Date: Fri, 16 Feb 2024 13:25:24 -0600 Subject: [PATCH] fix docstring for group barrier; and add todo for group barrier overload --- .../_group_barrier_overloads.py | 5 +++++ numba_dpex/kernel_api/barrier.py | 6 +----- 2 files changed, 6 insertions(+), 5 deletions(-) diff --git a/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_group_barrier_overloads.py b/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_group_barrier_overloads.py index 4a21d98ecb..85485127f2 100644 --- a/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_group_barrier_overloads.py +++ b/numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_group_barrier_overloads.py @@ -120,6 +120,11 @@ def ol_group_barrier(group, fence_scope=MemoryScope.WORK_GROUP): ) mem_scope = _get_memory_scope(fence_scope) + # TODO: exec_scope needs to be determined based on + # group argument. If group refers to a work_group then, + # exec_scope is MemoryScope.WORK_GROUP. + # If group is sub_group then, exec_scope needs to be + # MemoryScope.SUB_GROUP exec_scope = get_scope(MemoryScope.WORK_GROUP.value) spirv_memory_semantics_mask = get_memory_semantics_mask( MemoryOrder.SEQ_CST.value diff --git a/numba_dpex/kernel_api/barrier.py b/numba_dpex/kernel_api/barrier.py index 4b0b8ae199..744262403d 100644 --- a/numba_dpex/kernel_api/barrier.py +++ b/numba_dpex/kernel_api/barrier.py @@ -15,11 +15,7 @@ def group_barrier(group: Group, fence_scope=MemoryScope.WORK_GROUP): The function is modeled after the ``sycl::group_barrier`` function. It synchronizes work within a group of work items. All the work-items of the group must execute the barrier construct before any work-item - continues execution beyond the barrier. However, unlike - ``sycl::group_barrier`` the numba_dpex function implicitly synchronizes at - the level of a work group and does not allow specifying the group as an - argument. The :func:`sub_group_barrier` function should be used if - synchronization has to be performed only across a sub-group. + continues execution beyond the barrier. The ``group_barrier`` performs mem-fence operations ensuring that memory accesses issued before the barrier are not re-ordered with those issued