Properly defining the semantics of numba_dpex's kernel indexing functions

[diptorupd]

The semantics of numba_dpex.get_global_id, numba_dpex.get_local_id is relatively undocumented and can cause users to make incorrect assumptions about indexing order and run into performance/correctness issues.

Refer: #941, #889

The problem stems from numba-dpex generating an OpenCL or Level Zero interoperability kernel rather than a SYCL kernel. Per the SYCL spec the indexing of SYCL and OpenCL have different orders, refer: https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:opencl:kernel-conventions-sycl.

SYCL kernel query	OpenCL kernel query	Returned Value
With enqueued 3D SYCL global range of range<3> R{r0,r1,r2}
nd_item::get_global_range(0) / item::get_range(0)	get_global_size(2)	r0
nd_item::get_global_range(1) / item::get_range(1)	get_global_size(1)	r1
nd_item::get_global_range(2) / item::get_range(2)	get_global_size(0)	r2
nd_item::get_global_id(0) / item::get_id(0)	get_global_id(2)	Value in range 0..(r0-1)}
nd_item::get_global_id(1) / item::get_id(1)	get_global_id(1)	Value in range 0..(r1-1)}
nd_item::get_global_id(2) / item::get_id(2)	get_global_id(0)	Value in range 0..(r2-1)}
With enqueued 2D SYCL global range of range<2> R{r0,r1}
nd_item::get_global_range(0) / item::get_range(0)	get_global_size(1)	r0
nd_item::get_global_range(1) / item::get_range(1)	get_global_size(0)	r1
nd_item::get_global_id(0) / item::get_id(0)	get_global_id(1)	Value in range 0..(r0-1)}
nd_item::get_global_id(1) / item::get_id(1)	get_global_id(0)	Value in range 0..(r1-1)}
With enqueued 1D SYCL global range of range<1> R{r0}
nd_item::get_global_range(0) / item::get_range(0)	get_global_size(0)	r0
nd_item::get_global_id(0) / item::get_id(0)	get_global_id(0)	Value in range 0..(r0-1)}

Thus, users programming a dpex kernel assuming that the indexing runs in SYCL/C++ order run into issues, as dpex really uses OpenCL's versions of these functions.'

#274 captures the initial aspect of what needs to be done. I am opening the discussion to capture both requirement and design ideas on how the numba_dpex indexing semantics should be fixed and aligned with SYCL irrespective of the backend selected.

[AlexanderKalistratov]

I do not follow. Are you saying that we cannot support SYCL semantics while specifying the range in SYCL indexing order and submitting an OpenCL interop kernel?

I think we cannot. It is not only dimensions values are important, but also dimension id (0, 1, 2). Reverting range we can modify values but not ids.

If I'm right using sycl style get_global_id instead of opencl one should resolve the issue. If it is not we would need to go into details on how exactly igc works.

[fcharras]

Just to add on what I grasp of benefits for performance: there's how the numba_dpex API is designed, and then there's how it actually maps to compute. For GPU computing it's important to know which work items are "contiguous" (depending how the grid of work items is aligned), and it's important to hold the assumption that contiguous work items span threads that are physically contiguous (on the same EU,...).

The same is somewhat true at the work-group level, maybe anticipating in what order the work groups are dispatched can help with cache optimization.