Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

ENH Use a 2D grid of work items where applicable #98

Merged
merged 1 commit into from
Mar 3, 2023

Conversation

fcharras
Copy link
Collaborator

@fcharras fcharras commented Feb 28, 2023

  • Where applicable, use 2D grid of work items rather than 1D + divisions
    • note: found this piece of information: https://stackoverflow.com/a/15044884 suggesting that:
      • 2D grid really are better for performance because doing // and % ops in the kernel is expensive
      • it indexes work items in a "row major order" meaning that items in the same row belong to the same sub group (/warp)
    • even if the previous information applies to cuda. I assume it's also true for numba-dpex/ sycl
  • also use 2D grids for kmeans kernels

It's WIP because the benchmark show that those changes induce a 30% performance overhead, there seem to be an issue with 2D grid of work items in numba_dpex or SYCL.

Opened an issue IntelPython/numba-dpex#941
TODO: find minimal reproducers.

This PR remains opened as WIP in case those issues are fixed eventually.

Before considering merging the PR, please rebase the branch and check if new kernels on main branch could also benefit from the change.

@fcharras fcharras force-pushed the ENH_use_2D_work_group_shape branch from fd2b014 to dc30df1 Compare March 3, 2023 15:43
@fcharras fcharras marked this pull request as ready for review March 3, 2023 15:44
@fcharras
Copy link
Collaborator Author

fcharras commented Mar 3, 2023

OUT of WIP. Will merge since it's almost identical to what was approved in #88

The performance issue came from assuming the wrong alignment for grid of work items, SYCL uses row-major order, but cuda, numba.cuda and numba_dpex all use column-major order.

The fix consisted in transposing the 2D work group size and inversing the gpex.get_*_id calls.

@fcharras fcharras merged commit 46c451e into main Mar 3, 2023
@fcharras fcharras deleted the ENH_use_2D_work_group_shape branch March 3, 2023 16:08
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

Successfully merging this pull request may close these issues.

1 participant