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

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

Closed
wants to merge 12 commits into from

Conversation

gmarkall
Copy link
Member

Testing on CI as I work on the tests.

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.
Presently fails because the codegen can't be pickled due to ctypes
pointer.
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 added 2 - In Progress CUDA CUDA related issue/PR labels May 19, 2022
@gmarkall gmarkall added this to the Numba 0.56 RC milestone May 19, 2022
@gmarkall
Copy link
Member Author

I started with the wrong branch here and public CI is busy, so I'll close this now.

@gmarkall gmarkall closed this May 19, 2022
@gmarkall gmarkall added abandoned PR is abandoned (no reason required) and removed 2 - In Progress labels May 19, 2022
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
abandoned PR is abandoned (no reason required) CUDA CUDA related issue/PR
Projects
None yet
Development

Successfully merging this pull request may close these issues.

None yet

1 participant