Repository navigation
[CUDA] Fix sanity check PTU compile fail and re-enable unittests - #1161
Merged
Merged
Conversation
Codecov Report✅ All modified and coverable lines are covered by tests. Additional details and impacted files@@ Coverage Diff @@
## main #1161 +/- ##
==========================================
+ Coverage 88.70% 88.80% +0.10%
==========================================
Files 24 24
Lines 6978 6978
==========================================
+ Hits 6190 6197 +7
+ Misses 788 781 -7 see 1 file with indirect coverage changes 🚀 New features to boost your workflow:
|
Contributor
|
clang-tidy review says "All clean, LGTM! 👍" |
1 similar comment
Contributor
|
clang-tidy review says "All clean, LGTM! 👍" |
DetectCudaInstallAndArch ends by creating an interpreter with --cuda-path=/usr/local/cuda-12.3 and --offload-arch=sm_70, a toolkit path that need not exist on the host. createClangInterpreter forwards an explicit path unchecked, clang finds no installation there and selects PTX 4.2, and the NVPTX backend of LLVM 22 reports the pair as a fatal usage error, "PTX version 4.2 does not support target 'sm_70'". The error aborts the unittest binary. LLVM 21 accepted the pair, and on LLVM 22 the test was never reached because the fixture failed first. Create that interpreter from the toolkit and the architecture the test has just detected, so the explicit arguments are exercised on every host with a GPU. The probe of the 12.3 path above it already tolerates an absent directory.
CUDATest::SetUpTestSuite defines test_func and launches it as a top-level statement of the same input. LLVM 22 rejects the launch in device compilation with err_cuda_device_kernel_launch_require_rdc, the Declare fails and all seven tests skip, as they do on the self-hosted CUDA job. On LLVM 21 the input compiles, but the launch runs from the PTU's static initializer before the module constructor has registered the device binary and fails with cudaErrorInvalidResourceHandle, so the four runtime tests skip. A launch from a host function defined with the kernel and called in a later input, the shape a client of the interpreter uses, works on both releases. Define a host function next to the kernel that launches it and call it from the next input. SimpleKernelExecution and CUB_BlockReduceWithDeviceFunction take the same shape, since their own top-level launches fail the same way once the fixture lets them run.
clang-repl builds the device TargetMachine without the target features, so the PTX carries the minimum ISA version of the arch, 5.0 for sm_61. CUB's block reduction uses shfl.sync, which needs 6.0, and the driver JIT rejects the module with cudaErrorInvalidPtx. Skip the test below sm_70, where the minimum is 6.0, until clang-repl passes the features.
aaronj0
force-pushed
the
cudatest-host-launcher
branch
from
October 31, 2026 06:28
fa8610b to
276b239
Compare
Contributor
|
clang-tidy review says "All clean, LGTM! 👍" |
aaronj0
commented
Oct 31, 2026
| // CUB's block reduction needs PTX ISA 6.0, first implied by sm_70. | ||
| std::string arch; | ||
| unsigned sm = 0; | ||
| if (compat::detectNVPTXArch(arch) && |
Collaborator
Author
There was a problem hiding this comment.
Here we are skipping this CUB test based on the arch. This is because the GPU on the self hosted nuc runner is pre-volta which reveals the issue solved by llvm/llvm-project#235754. We can enable it with LLVM 24.
Collaborator
Author
|
@vgvassilev with this PR we are back to running CUDA tests on the self hosted runner, which has been skipping for a while. |
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
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.
The sanity check currently always fails because LLVM 22 rejects the launch in
device compilation: the Declare fails skipping all tests as seen on the self-hosted
CUDA job.
On LLVM 21 the input compiles, but the launch runs from the
PTU's static initializer before the module constructor has registered
the device binary and fails with cudaErrorInvalidResourceHandle.
This is due to llvm/llvm-project#158021 which will be fixed by llvm/llvm-project#234643.
We can workaround this by using a launch from a host function defined with the
kernel and called in a later input works on both releases.
Define a host function next to the kernel that launches it and call it
from the next input. Both the simple kernel and CUB tests take the
same shape, since their own top-level launches fail the same way
once the fixture lets them run. Also drops the hardcoded arch and toolkit path