Skip to content

CUDA: Initial support for Minor Version Compatibility - #8180

Merged
sklam merged 19 commits into
numba:mainfrom
gmarkall:cuda-mvc
Feb 25, 2023
Merged

sklam merged 19 commits into
numba:mainfrom
gmarkall:cuda-mvc

Conversation

@gmarkall

@gmarkall gmarkall commented Jun 21, 2022 •

Copy link
Copy Markdown
Member

This PR adds support for Minor Version Compatibility on Linux, so that Numba can be used with a CUDA toolkit of a later version than the CUDA version supported by the driver, within the same major version - for example, CUDA toolkit 11.7 can be used with a driver that only supports up to CUDA 11.5.

Minor Version Compatibility is not officially supported by NVIDIA for applications producing PTX, so this PR relies on two external packages to compile PTX and link cubins into a form that can be loaded by the driver. These are the ptxcompiler and cubinlinker packages that I maintain as part of RAPIDS. ptxcompiler is open-source and available in conda-forge, but cubinlinker is binary-only and must be obtained from the RAPIDS channel or NVIDIA PyPI.

There are two small limitations when using MVC at present:

  • It presently does not support linking archives, and
  • therefore linking cudadevrt and the use of Cooperative Groups is not supported.

I plan to address these in a future PR - I just need to find a relatively straightforward and Pythonic way to extract objects from archives to add to the link, but I didn't want to add too much complexity to this initial support PR.

Implementation notes

The implementation works by providing a third choice Linker implementation in numba.cuda.cudadrv.driver. This linker uses the PTX compiler from ptxcompiler and the linker from cubinlinker to compile and link PTX.

Testing adds an additional configuration to test with the CUDA toolkit 11.7 on gpuCI - the gpuCI image only has a driver supporting CUDA 11.5, so this demonstrates correct operation of MVC in Numba.

A small change was made to the interface between the codegen and the linker - instead of the codegen being responsible for taking a copy of the buffer containing the linked executable, the linker now returns a copy that the codegen can own - this keeps things simpler by making the interface between the different linkers and the codegen uniform.

This needs to be enabled using the config variable
`CUDA_ENABLE_MINOR_VERSION_COMPATIBILITY`. Limitations:

- It presently does not support linking archives, and therefore linking
  cudadevrt and Cooperative Groups.
- Performance may be lower than using the driver's linker as relocatable
  code is always generated, even if there is only one PTX input.
@gmarkall gmarkall added 2 - In Progress CUDA CUDA related issue/PR Effort - medium Medium size effort needed labels Jun 21, 2022
@gmarkall gmarkall modified the milestone: Numba 0.57 RC Jun 21, 2022
@stuartarchibald stuartarchibald self-assigned this Jun 28, 2022
@gmarkall gmarkall added this to the Numba 0.57 RC milestone Jun 28, 2022
@stuartarchibald

Copy link
Copy Markdown
Contributor

/AzurePipelines run

@azure-pipelines

Copy link
Copy Markdown
Contributor
Azure Pipelines successfully started running 1 pipeline(s).

@gmarkall

gmarkall commented Jul 5, 2022

Copy link
Copy Markdown
Member Author

This is failing because it needs a not-yet-public package.

@gmarkall

gmarkall commented Sep 1, 2022

Copy link
Copy Markdown
Member Author

gpuci run tests

@gmarkall

Copy link
Copy Markdown
Member Author

gpuci run tests

If they are not available, this should not cause a problem unless the
user requested MVC. If the user did, then an explanatory error message
is more helpful than just the import error.
@gmarkall

Copy link
Copy Markdown
Member Author

gpuci run tests

1 similar comment
@gmarkall

Copy link
Copy Markdown
Member Author

gpuci run tests

@gmarkall

Copy link
Copy Markdown
Member Author

gpuci run tests

@gmarkall
gmarkall marked this pull request as ready for review October 12, 2022 11:28
@gmarkall

gmarkall commented Dec 2, 2022

Copy link
Copy Markdown
Member Author

gpuci run tests

The latest cubinlinker packages have now been copied from the gmarkall
channel to the rapidsai channel.
@gmarkall

gmarkall commented Dec 8, 2022

Copy link
Copy Markdown
Member Author

gpuci run tests

@stuartarchibald stuartarchibald left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks for the patch. This in general looks good, there's a few minor thing to resolve and this could do with a manual test to confirm the new code paths, otherwise should be ok to merge once comments are resolved. Thanks again.

Comment thread docs/source/cuda/minor_version_compatibility.rst Outdated
Comment thread numba/cuda/cudadrv/driver.py Outdated
Comment thread docs/source/user/installing.rst
Comment thread numba/cuda/cudadrv/driver.py Outdated
Comment thread numba/cuda/cudadrv/driver.py Outdated
Comment thread numba/cuda/cudadrv/driver.py
Comment thread numba/cuda/cudadrv/driver.py
Comment thread numba/cuda/cudadrv/driver.py
Comment thread numba/cuda/cudadrv/driver.py
@stuartarchibald stuartarchibald added 4 - Waiting on author Waiting for author to respond to review and removed 3 - Ready for Review labels Feb 17, 2023
- Change title to "CUDA Minor Version Compatibility" to make it clearer
  it only pertains to CUDA.
- Explicitly state that Windows is unsupported.
- Modify grammar describing the use of external packages.
- Edit wording and correct spelling in exception message given when
  required packages are not available.
- Clarify reference to cc kwarg in exception message.
- Minor edit to generation of arch string.
@gmarkall

Copy link
Copy Markdown
Member Author

gpuci run tests

@stuartarchibald

Copy link
Copy Markdown
Contributor

I've reviewed through to cb163aa and approve the code changes, need to do a manual check that this is working as expected.

@stuartarchibald stuartarchibald added 4 - Waiting on reviewer Waiting for reviewer to respond to author and removed 4 - Waiting on author Waiting for author to respond to review labels Feb 23, 2023
@stuartarchibald

Copy link
Copy Markdown
Contributor

I've reviewed through to cb163aa and approve the code changes, need to do a manual check that this is working as expected.

As noted in an OOB discussion, gpuCI tests are doing an automated check for MVC so there's no need to do a manual check. I'll approve this patch!

@stuartarchibald stuartarchibald left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks for the patch and fixes!

@stuartarchibald stuartarchibald added Pending BuildFarm For PRs that have been reviewed but pending a push through our buildfarm 4 - Waiting on CI Review etc done, waiting for CI to finish and removed 4 - Waiting on reviewer Waiting for reviewer to respond to author labels Feb 24, 2023
@stuartarchibald

Copy link
Copy Markdown
Contributor

Buildfarm ID: numba_smoketest_cuda_yaml_184.

@stuartarchibald

Copy link
Copy Markdown
Contributor

Buildfarm ID: numba_smoketest_cuda_yaml_184.

Passed.

@stuartarchibald stuartarchibald added 5 - Ready to merge Review and testing done, is ready to merge BuildFarm Passed For PRs that have been through the buildfarm and passed and removed 4 - Waiting on CI Review etc done, waiting for CI to finish Pending BuildFarm For PRs that have been reviewed but pending a push through our buildfarm labels Feb 24, 2023
@sklam
sklam merged commit 9992d56 into numba:main Feb 25, 2023
esc added a commit to esc/numba that referenced this pull request Mar 2, 2023
* main: (2743 commits)
  Upgrade stale GitHub action
  Changed wording of the escape error
  Correct documentation for CAS
  Revert "Merge pull request numba#8751 from apmasell/mulitprocess_recycling"
  CUDA driver.py: PR numba#8180 feedback
  CUDA Docs: Refine MVC docs based on PR numba#8180 feedback
  revert build_ext change
  add empty line back
  revert some changes
  fix compiler arg, revert build script
  add build_doc back, fix compiler arg
  Correct grammar in comment
  fix import
  Stale bot: Use `abandoned - stale` label for closed PRs
  refactor build_ext
  BLD: remove distutils
  CUDA tidy-up: remove some unneeded methods
  Update numba/tests/support.py
  Update .github/workflows/stale.yml
  Fix line length error from flake8
  ...
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

5 - Ready to merge Review and testing done, is ready to merge BuildFarm Passed For PRs that have been through the buildfarm and passed CUDA CUDA related issue/PR Effort - medium Medium size effort needed highpriority

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants