Merged
Conversation
Contributor
Author
|
It actually reveals few bugs: |
Contributor
|
Is the difference between the two codes are the unoptimized version contains a |
Contributor
Author
Yes
Thank you for the reference! Just added tests |
isVoid
approved these changes
Apr 1, 2025
Contributor
|
This needs the formatting fixing up following the merge of the content of #170. |
ZzEeKkAa
force-pushed
the
yhavrylko/fix/jit_inline
branch
from
9b96a36 to
21718b1
Compare
Contributor
Author
|
@gmarkall it is ready to go now |
gmarkall
added
4 - Waiting on reviewer
Contributor
|
Unfortunately there is user code that uses I think translating it to |
gmarkall
requested changes
Apr 11, 2025
Contributor
gmarkall
left a comment
gmarkall
left a comment
There was a problem hiding this comment.
Other than needing to continue accepting True and False, I think the PR is good.
gmarkall
added
4 - Waiting on author
Contributor
Author
|
Thank you @gmarkall ! I hope it is okay to support bool and show deprecation warning if user uses it. |
gmarkall
added
4 - Waiting on reviewer
gmarkall
reviewed
Apr 30, 2025
gmarkall
requested changes
Apr 30, 2025
Contributor
gmarkall
left a comment
gmarkall
left a comment
There was a problem hiding this comment.
I think this looks good - just one suggestion to link to the docs then this can be merged,
gmarkall
removed
the
4 - Waiting on reviewer
Co-authored-by: Graham Markall <535640+gmarkall@users.noreply.github.com>
gmarkall
approved these changes
May 1, 2025
gmarkall
added
5 - Ready to merge
gmarkall
approved these changes
May 1, 2025
gmarkall
added a commit
to gmarkall/numba-cuda
that referenced
this pull request
May 2, 2025
- Fix Invalid NVVM IR emitted when lowering shfl_sync APIs (NVIDIA#231) - Add Bfloat16 Low++ Bindings (NVIDIA#166) - Fix cuda.jit decorator inline (NVIDIA#181) - Feature: cuda specific make_attribute_wrapper (NVIDIA#193) - return a none tuple if no libdevice path is found (NVIDIA#234)
Closed
gmarkall
added a commit
to gmarkall/numba-cuda
that referenced
this pull request
May 3, 2025
- Local variable debug info deduplication (NVIDIA#222) - Fix package installation for wheels CI (NVIDIA#238) - Fix Invalid NVVM IR emitted when lowering shfl_sync APIs (NVIDIA#231) - Add Bfloat16 Low++ Bindings (NVIDIA#166) - Fix cuda.jit decorator inline (NVIDIA#181) - Feature: cuda specific make_attribute_wrapper (NVIDIA#193) - return a none tuple if no libdevice path is found (NVIDIA#234)
Merged
gmarkall
added a commit
that referenced
this pull request
May 3, 2025
- Local variable debug info deduplication (#222) - Fix package installation for wheels CI (#238) - Fix Invalid NVVM IR emitted when lowering shfl_sync APIs (#231) - Add Bfloat16 Low++ Bindings (#166) - Fix cuda.jit decorator inline (#181) - Feature: cuda specific make_attribute_wrapper (#193) - return a none tuple if no libdevice path is found (#234)
gmarkall
added a commit
to gmarkall/numba-cuda
that referenced
this pull request
May 9, 2025
PR NVIDIA#181 aimed to align the behaviour of the `inline` kwarg with that of upstream Numba, in that it now forces inlining at the Numba IR level. It turns out that this kwarg in Numba-CUDA already had the prior effect of enabling inlining at the NVVM IR level. Because the default value of `inline` is `"never"`, this was interpreted by the `compile_cuda()` function as a `True`ish value and every device function got marked with the `alwaysinline` function attribute. This is a minor problem in that it probably forces a lot of inlining that we don't want, but also a major problem in that it triggers an NVVM bug that was only resolved in CUDA 12.3 that causes a hang in `nvvmCompileProgram()`. To rectify these issues, we add the `forceinline` kwarg to the `@cuda.jit` decorator and the `cuda.compile[_*]()` functions. Now, `compile_cuda()` will only enable inlining at the NVVM IR level for `forceinline` and not `inline`. This is aligned with the behaviour of upstream Numba (see numba/numba#10068). We now document the `inline` and `forceinline` kwargs to clarify the intent and behaviour for users. For clarity: the behaviour is now: - The `inline` kwarg enables inlining only at the Numba IR level. - The `forceinline` kwarg enables inlining only at the NVVM IR level.
gmarkall
added a commit
to gmarkall/numba-cuda
that referenced
this pull request
May 9, 2025
PR NVIDIA#181 aimed to align the behaviour of the `inline` kwarg with that of upstream Numba, in that it now forces inlining at the Numba IR level. It turns out that this kwarg in Numba-CUDA already had the prior effect of enabling inlining at the NVVM IR level. Because the default value of `inline` is `"never"`, this was interpreted by the `compile_cuda()` function as a `True`ish value and every device function got marked with the `alwaysinline` function attribute. This is a minor problem in that it probably forces a lot of inlining that we don't want, but also a major problem in that it triggers an NVVM bug that was only resolved in CUDA 12.3 that causes a hang in `nvvmCompileProgram()`. To rectify these issues, we add the `forceinline` kwarg to the `@cuda.jit` decorator and the `cuda.compile[_*]()` functions. Now, `compile_cuda()` will only enable inlining at the NVVM IR level for `forceinline` and not `inline`. This is aligned with the behaviour of upstream Numba (see numba/numba#10068). We now document the `inline` and `forceinline` kwargs to clarify the intent and behaviour for users. For clarity: the behaviour is now: - The `inline` kwarg enables inlining only at the Numba IR level. - The `forceinline` kwarg enables inlining only at the NVVM IR level.
gmarkall
added a commit
to gmarkall/numba-cuda
that referenced
this pull request
May 9, 2025
PR NVIDIA#181 aimed to align the behaviour of the `inline` kwarg with that of upstream Numba, in that it now forces inlining at the Numba IR level. It turns out that this kwarg in Numba-CUDA already had the prior effect of enabling inlining at the NVVM IR level. Because the default value of `inline` is `"never"`, this was interpreted by the `compile_cuda()` function as a `True`ish value and every device function got marked with the `alwaysinline` function attribute. This is a minor problem in that it probably forces a lot of inlining that we don't want, but also a major problem in that it triggers an NVVM bug that was only resolved in CUDA 12.3 that causes a hang in `nvvmCompileProgram()`. To rectify these issues, we add the `forceinline` kwarg to the `@cuda.jit` decorator and the `cuda.compile[_*]()` functions. Now, `compile_cuda()` will only enable inlining at the NVVM IR level for `forceinline` and not `inline`. This is aligned with the behaviour of upstream Numba (see numba/numba#10068). We now document the `inline` and `forceinline` kwargs to clarify the intent and behaviour for users. For clarity: the behaviour is now: - The `inline` kwarg enables inlining only at the Numba IR level. - The `forceinline` kwarg enables inlining only at the NVVM IR level.
gmarkall
added a commit
that referenced
this pull request
May 9, 2025
PR #181 aimed to align the behaviour of the `inline` kwarg with that of upstream Numba, in that it now forces inlining at the Numba IR level. It turns out that this kwarg in Numba-CUDA already had the prior effect of enabling inlining at the NVVM IR level. Because the default value of `inline` is `"never"`, this was interpreted by the `compile_cuda()` function as a `True`ish value and every device function got marked with the `alwaysinline` function attribute. This is a minor problem in that it probably forces a lot of inlining that we don't want, but also a major problem in that it triggers an NVVM bug that was only resolved in CUDA 12.3 that causes a hang in `nvvmCompileProgram()`. To rectify these issues, we add the `forceinline` kwarg to the `@cuda.jit` decorator and the `cuda.compile[_*]()` functions. Now, `compile_cuda()` will only enable inlining at the NVVM IR level for `forceinline` and not `inline`. This is aligned with the behaviour of upstream Numba (see numba/numba#10068). We now document the `inline` and `forceinline` kwargs to clarify the intent and behaviour for users. For clarity: the behaviour is now: - The `inline` kwarg enables inlining only at the Numba IR level. - The `forceinline` kwarg enables inlining only at the NVVM IR level.
gmarkall
added a commit
that referenced
this pull request
May 9, 2025
PR #181 aimed to align the behaviour of the `inline` kwarg with that of upstream Numba, in that it now forces inlining at the Numba IR level. It turns out that this kwarg in Numba-CUDA already had the prior effect of enabling inlining at the NVVM IR level. Because the default value of `inline` is `"never"`, this was interpreted by the `compile_cuda()` function as a `True`ish value and every device function got marked with the `alwaysinline` function attribute. This is a minor problem in that it probably forces a lot of inlining that we don't want, but also a major problem in that it triggers an NVVM bug that was only resolved in CUDA 12.3 that causes a hang in `nvvmCompileProgram()`. To rectify these issues, we add the `forceinline` kwarg to the `@cuda.jit` decorator and the `cuda.compile[_*]()` functions. Now, `compile_cuda()` will only enable inlining at the NVVM IR level for `forceinline` and not `inline`. This is aligned with the behaviour of upstream Numba (see numba/numba#10068). We now document the `inline` and `forceinline` kwargs to clarify the intent and behaviour for users. For clarity: the behaviour is now: - The `inline` kwarg enables inlining only at the Numba IR level. - The `forceinline` kwarg enables inlining only at the NVVM IR level.
This file contains hidden or 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
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
3 participants
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.
In the
@cuda.jitdecoratorinlineis being ignored and unused. Because it is explicitly declared in the function signature it is not part of thekwsand is not passed further to the dispatcher. This PR fixes it by removing it from function signature.Here is an affected example:
optimized LLVM code before the change:
optimized LLVM code after the change: