Add ability to link CUDA functions with in-memory PTX. #9470
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
a string. This allows the user to link CUDA functions with dynamically generated PTX code without having to write data to a file.
gmarkall
reviewed
Feb 28, 2024
There was a problem hiding this comment.
Many thanks for the PR, @ed-o-saurus - this is an often-requested feature so I'm happy to see this PR!
There are a few thoughts I need to add on the design and scope of the changes, but I like the idea of using an object to hold the code whilst keeping strings being used for paths, maintaining backwards-compatibility.
I'll post a follow-up once I've had more time to get my thoughts down.
isVoid
pushed a commit
to rapidsai/pynvjitlink
that referenced
this pull request
Mar 19, 2024
Adds support for linking code from memory to Numba's `@cuda.jit` decorator (in addition to the already-supported linking files from disk). New classes are added to Numba's top-level `cuda` module: * `Archive`: An archive of objects * `CUSource`: A CUDA C/C++ source * `Cubin`: A cubin ELF * `Fatbin`: A fatbin ELF * `Object`: An object file * `PTXSource`: PTX assembly source code. These are all used by constructing them with a single argument, the code in memory to use. Once created, they can then be passed to the `link=` kwarg of the `@cuda.jit` decorator. An example showing a use case with a CUDA C/C++ source is added. This implementation is inspired by the approach outlined by @ed-o-saurus in numba/numba#9470. Notes on changes: * Various tests now run on the GPU, so test binaries need to be generated using the relevant compute capability - this change is applied in the `Makefile`, along with some refactoring to tidy it up a little. * Test for new functionality are added. In addition, existing tests that used the test binaries are all modified such that they use a relevant compute capability (usually the one in the test machine, for most of them) because the test binaries are now built for the CC of the current GPU - it's no longer sufficient to hard code CCs like 7.5 or 7.0 for tests. Since fixtures are needed across test files, I've started moving them into `conftest.py`. --------- Co-authored-by: Bradley Dice <bdice@bradleydice.com> Co-authored-by: jakirkham <jakirkham@gmail.com>
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
This commit adds the PTXCode class. It is a simple wrapper around a string. This allows the user to link CUDA functions with dynamically generated PTX code without having to write data to a file.