rapidsai/pynvjitlink

 pynvjitlink

The RAPIDS pynvjitlink library provides a Python binding for the nvJitLink library.

Installation with pip

pip install --extra-index-url https://pypi.nvidia.com pynvjitlink-cu12

Installation from source

Install with either:

python -m pip install .

or

python -m pip install -e .

for an editable install.

Installation with Conda

conda install -c rapidsai pynvjitlink

Contributing Guide

Review the CONTRIBUTING.md file for information on how to contribute code and issues to the project.

pynvjitlink's Issues

[BUG] The Numba patch breaks some functionality for versions 0.58 and above

The Numba test suite, when run with the pynvjitlink patch, fails some tests. To run the test suite with the patch, the following script can be used:

from pynvjitlink import patch
patch.patch_numba_linker()

if __name__ == '__main__':
    from numba.testing._runtests import _main
    import sys

    sys.exit(0 if _main(sys.argv) else 1)

invoked with

python test_patched_numba.py numba.cuda.tests -v -m

ends with

Ran 1462 tests in 134.056s

FAILED (errors=28, skipped=33, expected failures=11)

The errors are of this form:

======================================================================
ERROR: test_ex_linking_cu (numba.cuda.tests.doc_examples.test_ffi.TestFFI.test_ex_linking_cu)
----------------------------------------------------------------------
Traceback (most recent call last):
  File "/home/gmarkall/numbadev/numba/numba/cuda/tests/doc_examples/test_ffi.py", line 42, in test_ex_linking_cu
    multiply_vectors[1, 32](r, x, y)
  File "/home/gmarkall/numbadev/numba/numba/cuda/dispatcher.py", line 539, in __call__
    return self.dispatcher.call(args, self.griddim, self.blockdim,
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/gmarkall/numbadev/numba/numba/cuda/dispatcher.py", line 673, in call
    kernel = _dispatcher.Dispatcher._cuda_call(self, *args)
             ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/gmarkall/numbadev/numba/numba/cuda/dispatcher.py", line 681, in _compile_for_args
    return self.compile(tuple(argtypes))
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/gmarkall/numbadev/numba/numba/cuda/dispatcher.py", line 926, in compile
    kernel.bind()
  File "/home/gmarkall/numbadev/numba/numba/cuda/dispatcher.py", line 197, in bind
    self._codelibrary.get_cufunc()
  File "/home/gmarkall/numbadev/numba/numba/cuda/codegen.py", line 195, in get_cufunc
    cubin = self.get_cubin(cc=device.compute_capability)
            ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  File "/home/gmarkall/numbadev/numba/numba/cuda/codegen.py", line 172, in get_cubin
    linker.add_file_guess_ext(path)
  File "/home/gmarkall/numbadev/numba/numba/cuda/cudadrv/driver.py", line 2656, in add_file_guess_ext
    self.add_cu_file(path)
  File "/home/gmarkall/numbadev/numba/numba/cuda/cudadrv/driver.py", line 2648, in add_cu_file
    self.add_cu(cu, os.path.basename(path))
  File "/home/gmarkall/numbadev/nvjitlink/pynvjitlink/patch.py", line 129, in add_cu
    self.add_ptx(program.ptx.rstrip(b"\x00"), ptx_name)
                 ^^^^^^^^^^^
AttributeError: 'NvrtcProgram' object has no attribute 'ptx'

----------------------------------------------------------------------

[FEA] Add an `update-version.sh`

Like other RAPIDS repos, it would be nice to have a script that bumps the patch version in the three places it's hardcoded:

https://github.com/rapidsai/pynvjitlink/blob/main/CMakeLists.txt#L5
https://github.com/rapidsai/pynvjitlink/blob/main/pynvjitlink/VERSION#L1
https://github.com/rapidsai/pynvjitlink/blob/main/pyproject.toml#L14

Concretely update_version.sh should += the patch version in these three places and possibly remind the user the commit the result.

Embed build-time nvjitlink version and make accessible

pynvjitlink should embed the CUDA version used at build-time as a compile-time constant accessible via a Python API. That would be very helpful for us for debugging in the future. Since we're statically linking our only recourse for debugging purposes right now is to check when a particular pynvjitlink package was built and then see what version of CUDA the build scripts used at that time.

Once we provide this constant, we could also automatically check at import-time whether the CUDA runtime/driver on the system could lead to version mismatches that are unsupported by the current version of pynvjitlink.

Originally posted by @vyasr in #33 (comment)

[DOC] Broken link in readme

Report incorrect documentation

https://docs.nvidia.com/cuda/nvJitLink/index.html is broken

Location of incorrect documentation

https://github.com/rapidsai/pynvjitlink/blob/main/README.md?plain=1#L4

Describe the problems or issues found in the documentation
A clear and concise description of what you found to be incorrect.

Steps taken to verify documentation is incorrect
List any steps you have taken:

Suggested fix for documentation
Detail proposed changes to fix the documentation if you have any.

---

Report needed documentation

Report needed documentation
A clear and concise description of what documentation you believe it is needed and why.

Describe the documentation you'd like
A clear and concise description of what you want to happen.

Steps taken to search for needed documentation
List any steps you have taken:

[BUG] Remove xfails on error checking test cases when possible

PR #3 skips some tests that corrupt memory due to an underlying c++ bug. We should remove these when a future CUDA version is released that fixes said bug in libnvjitlink.

build_wheel.sh needs to stop on error

https://github.com/rapidsai/pynvjitlink/actions/runs/7102603110/job/19333372146#step:9:343

ci/build_wheel.sh: line 12: conda: command not found

We need to add set -euo pipefail to the build and test CI scripts.

Fix `__version__` to load `VERSION` file

Currently __version__ is hard-coded to 0.0.0

pynvjitlink/pynvjitlink/__init__.py

Line 6 in 1d4bcbd

__version__ = "0.0.0"

However that is not the actual version. Instead we should replace this by loading the VERSION file

pynvjitlink/pynvjitlink/VERSION

Line 1 in 1d4bcbd

0.2.0

Compatibility of `pynvjitlink` across versions

What kind of versioning scheme do we intend to use with pynvjitlink?

Do we intended to hold compatibility with major versions? So 1.x.y would be compatible with x increments for new features and y for bug fixes

Or do we intend to keep compatibility only within minor versions? So 1.2.y would be compatible for all y, but 1.3 would break compatibility?

Or do we have a different model of compatibility in mind?

cc @brandon-b-miller @gmarkall

[FEA] Expand supported python versions

Is your feature request related to a problem? Please describe.
Currently wheels are only built for python 3.9 and 3.10 because we rely on the rapids ci images, this limits non rapids consumers of pynvjitlink.

Describe the solution you'd like
If we want to support other versions of python we'll need either our own images or new rapids images for this one purpose, or something else.

Describe alternatives you've considered
Initially we tried installing python manually via conda within each CI job allowing us to build with any python version, however this was messy and error prone

Additional context
N/A

[FEA] Add more pre-commit hooks

We should add a few more pre-commit hooks, such as those that check the year of our copyright headers, among some other potential things.

[BUG] Patching breaks cooperative groups

Patching numba seems to disable launching kernels that involve cooperative groups:

from numba import cuda
from pynvjitlink import patch

patch.patch_numba_linker()

@cuda.jit
def cg_kernel():
    grid = cuda.cg.this_grid()

cg_kernel[1,1]()

this produces the following error from here

numba.cuda.cudadrv.driver.LinkerError: Don't know how to link archives

Using numba 0.59.0.

Is `.ycm_extra_conf.py` needed?

Am seeing this file .ycm_extra_conf.py and am wondering if it is needed

It appears to use CUDA 12.0 headers (not sure if that is expected)

pynvjitlink/.ycm_extra_conf.py

Line 14 in 72a76ee

CUDA_INCLUDE_DIR = "/usr/local/cuda-12.0/include"

Also pins to Python 3.10 headers

pynvjitlink/.ycm_extra_conf.py

Line 9 in 72a76ee

PYTHON_INCLUDE_NAME = "python3.10" # Not ideal

If this file is still needed, maybe it could benefit from some cleanup

If not, maybe we can just drop it

Note: It appears the file is not used in the repo