Skip to content

Commit 9d696f5

Browse files
committed
fix docstring for group barrier; and add todo for group barrier overload
1 parent e715ab5 commit 9d696f5

File tree

2 files changed

+6
-5
lines changed

2 files changed

+6
-5
lines changed

numba_dpex/experimental/_kernel_dpcpp_spirv_overloads/_group_barrier_overloads.py

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -120,6 +120,11 @@ def ol_group_barrier(group, fence_scope=MemoryScope.WORK_GROUP):
120120
)
121121

122122
mem_scope = _get_memory_scope(fence_scope)
123+
# TODO: exec_scope needs to be determined based on
124+
# group argument. If group refers to a work_group then,
125+
# exec_scope is MemoryScope.WORK_GROUP.
126+
# If group is sub_group then, exec_scope needs to be
127+
# MemoryScope.SUB_GROUP
123128
exec_scope = get_scope(MemoryScope.WORK_GROUP.value)
124129
spirv_memory_semantics_mask = get_memory_semantics_mask(
125130
MemoryOrder.SEQ_CST.value

numba_dpex/kernel_api/barrier.py

Lines changed: 1 addition & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -15,11 +15,7 @@ def group_barrier(group: Group, fence_scope=MemoryScope.WORK_GROUP):
1515
The function is modeled after the ``sycl::group_barrier`` function. It
1616
synchronizes work within a group of work items. All the work-items
1717
of the group must execute the barrier construct before any work-item
18-
continues execution beyond the barrier. However, unlike
19-
``sycl::group_barrier`` the numba_dpex function implicitly synchronizes at
20-
the level of a work group and does not allow specifying the group as an
21-
argument. The :func:`sub_group_barrier` function should be used if
22-
synchronization has to be performed only across a sub-group.
18+
continues execution beyond the barrier.
2319
2420
The ``group_barrier`` performs mem-fence operations ensuring that memory
2521
accesses issued before the barrier are not re-ordered with those issued

0 commit comments

Comments
 (0)