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.

Sign up for GitHub

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

Jump to bottom

[WIP] on-disk caching of CUDA kernels #8079

Closed
wants to merge 12 commits into from
Closed

[WIP] on-disk caching of CUDA kernels #8079

wants to merge 12 commits into from

Commits on May 19, 2022

  1. CUDA caching: enough hacks to save empty function 

    The following:

```python
from numba import cuda

@cuda.jit(cache=True)
def f():
    pass

f[1, 1]()
```

runs and saves cache files to the disk. The subsequent run loading from
the cache still fails.
    @gmarkall
    gmarkall committed May 19, 2022
    Configuration menu
    Copy the full SHA
    ea4fc08 View commit details
    Browse the repository at this point in the history

  2. [WIP] Attempt to make the load work 

    Presently fails because the codegen can't be pickled due to ctypes
pointer.
    @gmarkall
    gmarkall committed May 19, 2022
    Configuration menu
    Copy the full SHA
    b0fcf79 View commit details
    Browse the repository at this point in the history

  3. Successful pickle and unpickle of empty kernel 

    This incorporates fixes to the PTX pickling from @c200chromebook:

https://gist.github.com/c200chromebook/1ee304161a39b247e8a029ff9d2b3cd7

For the simple test:

```python
from numba import cuda

@cuda.jit(cache=True)
def f():
    pass

f[1, 1]()
```

The following CUDA Driver API calls are seen on the first invocation:

```
== CUDA [427] DEBUG -- call driver api: cuLinkCreate_v2
== CUDA [428] DEBUG -- call driver api: cuLinkAddData_v2
== CUDA [428] DEBUG -- call driver api: cuLinkComplete
== CUDA [428] DEBUG -- call driver api: cuLinkDestroy
```

but are absent from a subsequent run.
    @gmarkall
    gmarkall committed May 19, 2022
    Configuration menu
    Copy the full SHA
    cdede56 View commit details
    Browse the repository at this point in the history

  4. Add some temporary debug prints to make it easier to see when compile… 

    … / link happens
    @gmarkall
    gmarkall committed May 19, 2022
    Configuration menu
    Copy the full SHA
    a4c85fe View commit details
    Browse the repository at this point in the history

  5. Pickle dispatcher overloads too - issue with pickling device function… 

    …s though
    @gmarkall
    gmarkall committed May 19, 2022
    Configuration menu
    Copy the full SHA
    89362c9 View commit details
    Browse the repository at this point in the history

  6. Revert "Pickle dispatcher overloads too - issue with pickling device … 

    …functions though"

This reverts commit 65b9615.
    @gmarkall
    gmarkall committed May 19, 2022
    Configuration menu
    Copy the full SHA
    62a15d0 View commit details
    Browse the repository at this point in the history

  7. Fix damaged line for CPU caching

    @gmarkall
    gmarkall committed May 19, 2022
    Configuration menu
    Copy the full SHA
    49c8fcf View commit details
    Browse the repository at this point in the history

  8. Also serialize the LLVM module in CUDACodeLibrary

    @gmarkall
    gmarkall committed May 19, 2022
    Configuration menu
    Copy the full SHA
    88423a2 View commit details
    Browse the repository at this point in the history

  9. Use compute capability in cache key

    @gmarkall
    gmarkall committed May 19, 2022
    Configuration menu
    Copy the full SHA
    926627e View commit details
    Browse the repository at this point in the history

  10. Use base dispatcher serialization methods

    @gmarkall
    gmarkall committed May 19, 2022
    Configuration menu
    Copy the full SHA
    f8dc8ee View commit details
    Browse the repository at this point in the history

  11. Remove debug messages

    @gmarkall
    gmarkall committed May 19, 2022
    Configuration menu
    Copy the full SHA
    0d4e447 View commit details
    Browse the repository at this point in the history

  12. Fixup to match current main

    @gmarkall
    gmarkall committed May 19, 2022
    Configuration menu
    Copy the full SHA
    4864bfa View commit details
    Browse the repository at this point in the history