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

A sycl::local_accessor-like API for numba-dpex kernel #1331

Merged
merged 12 commits into from
Mar 19, 2024

Conversation

diptorupd
Copy link
Contributor

@diptorupd diptorupd commented Feb 13, 2024

  • Have you provided a meaningful PR description?

The PR adds a sycl::local_accessor like API to numba-dpex's kernel API.

  • Have you added a test, reproducer or referred to an issue with a reproducer?
  • Have you tested your changes locally for CPU and GPU devices?
  • Have you made sure that new changes do not introduce compiler warnings?
  • If this PR is a work in progress, are you filing the PR as a draft?

TODOs

@diptorupd diptorupd marked this pull request as draft February 13, 2024 19:57
@diptorupd diptorupd force-pushed the experimental/local_accessors branch 8 times, most recently from e49d728 to 3fea523 Compare February 21, 2024 19:14
@diptorupd diptorupd force-pushed the experimental/local_accessors branch from 3fea523 to 05678be Compare February 22, 2024 04:34

return self._data[idx_obj]

def __setitem__(self, idx_obj, val):
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Why do you need this?
Any attempt to read/write to local memory on host should result in exception.

Or I misunderstands something?

Copy link
Contributor Author

@diptorupd diptorupd Feb 27, 2024

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Sorry, I missed your comment.

This class is basically a simulator in pure Python for prototyping and testing. Borrowing idea's for numba-mlir's simulator.

The usage is a LocalAccessor will need to be is created in pure Python and passed as a kernel argument. In simulator mode it is just a NumPy array so works as expected. In compiled mode it will be an actual local array. What do you think?

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Since there is no examples I'm not sure how it should work, but I assume, that user need to create instance of this class in order to use local memory inside of kernel. Something like that:

loc_mem = LocalAccessor(...)
kernel(..., loc_mem, ...)

If that's assumption is correct, then user may try to access local memory outside of kernel:

loc_mem = LocalAccessor(...)
loc_mem[0] = 10
kernel(..., loc_mem, ...)

And in this case an exception must be thrown.

Copy link
Contributor

@ZzEeKkAa ZzEeKkAa Feb 28, 2024

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Let's then introduce LocalAccessorMock with those methods implemented and when we call kapi.call_kernel convert LocalAccessor arguments to LocalAccessorMock. This will prevent use of __getitem__ and __setitem__ from use outside of kernel. We still need to raise exception in LocalAccessor. Something like "getting item from outside kernel is prohibited on LocalAccessor".

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I have made these changes. Can you please review again?

@diptorupd diptorupd force-pushed the experimental/local_accessors branch from 05678be to bee7ba3 Compare February 26, 2024 17:47
@diptorupd diptorupd added this to the 0.23 milestone Feb 29, 2024
@diptorupd diptorupd force-pushed the experimental/local_accessors branch 8 times, most recently from 0437184 to f4dc617 Compare March 11, 2024 02:22
@diptorupd diptorupd requested a review from ZzEeKkAa March 11, 2024 04:20
@diptorupd diptorupd marked this pull request as ready for review March 11, 2024 04:20
@diptorupd
Copy link
Contributor Author

@ZzEeKkAa please review once again. I have added all changes and tests. I will add an example and update the kernel programming guide in a separate PR.

@diptorupd diptorupd force-pushed the experimental/local_accessors branch 4 times, most recently from 37566f0 to e44e8ed Compare March 12, 2024 20:34
@diptorupd diptorupd force-pushed the experimental/local_accessors branch 2 times, most recently from 57100c1 to ae0a1b9 Compare March 15, 2024 18:01
Base automatically changed from refactor/kernel_launcher to main March 15, 2024 22:23
@diptorupd diptorupd force-pushed the experimental/local_accessors branch from ae0a1b9 to b12c1a3 Compare March 16, 2024 03:10
@coveralls
Copy link
Collaborator

coveralls commented Mar 16, 2024

Coverage Status

coverage: 83.04% (+0.3%) from 82.76%
when pulling 2200c3c on experimental/local_accessors
into 202f460 on main.

@ZzEeKkAa ZzEeKkAa force-pushed the experimental/local_accessors branch from 102747e to d9d33b3 Compare March 19, 2024 17:10
@ZzEeKkAa ZzEeKkAa force-pushed the experimental/local_accessors branch from d9d33b3 to 52c3a56 Compare March 19, 2024 17:11
ZzEeKkAa
ZzEeKkAa previously approved these changes Mar 19, 2024
Copy link
Contributor

@ZzEeKkAa ZzEeKkAa left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGMT! Thank you!

@ZzEeKkAa ZzEeKkAa force-pushed the experimental/local_accessors branch from 52c3a56 to 2200c3c Compare March 19, 2024 19:43
@ZzEeKkAa ZzEeKkAa merged commit 68b1f39 into main Mar 19, 2024
59 of 65 checks passed
@ZzEeKkAa ZzEeKkAa deleted the experimental/local_accessors branch March 19, 2024 21:37
github-actions bot added a commit that referenced this pull request Mar 19, 2024
A sycl::local_accessor-like API for numba-dpex kernel 68b1f39
github-actions bot added a commit to diptorupd/numba-dpex that referenced this pull request Mar 20, 2024
…al_accessors

A sycl::local_accessor-like API for numba-dpex kernel 68b1f39
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.

4 participants