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

Enable caching of compiled sycl::kernel_bundle wrapper objects in the experimental KernelDispatcher module #1208

Closed
Tracked by #1120
diptorupd opened this issue Nov 6, 2023 · 1 comment
Assignees

Comments

@diptorupd
Copy link
Contributor

diptorupd commented Nov 6, 2023

When using the kernel decorator numba-dpex generates SYCL interoperability kernels for OpenCL and Level Zero. These interoperability kernels are not cached by the dpcpp runtime and need to be cached at the end of numba-dpex or dpctl to avoid recompilation performance overheads.

The feature was implemented in Python in the numba_dpex.core.kernel_interface.dispatcher.JitKernel class. As the new KernelDispacther class and call_kernel API generate sycl::kernel_bundle compilation using LLVM IR/C the cahing of kernel bundles has to be moved into the C layer, most likely inside dpctl's libsyclinterface.

@diptorupd diptorupd changed the title Enable caching of kernel bundles Enable caching of compiled sycl::kernel_bundle wrapper objects in the experimental KernelDispatcher module Nov 6, 2023
@diptorupd diptorupd self-assigned this Nov 6, 2023
@diptorupd
Copy link
Contributor Author

The feature is now available in numba_dpex.experimental and was implemented as part of #1236

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

No branches or pull requests

1 participant