Fix nvcc dryrun parsing for CUDA 13.3 - #2722
Conversation
| pub(crate) const CICC_INPUT_SUFFIX: &str = ".cpp1.ii"; | ||
| pub(crate) const PTXAS_INPUT_SUFFIX: &str = ".ptx"; | ||
|
|
||
| pub(crate) fn is_cicc_input(arg: impl AsRef<OsStr>) -> bool { | ||
| arg.as_ref().to_string_lossy().ends_with(CICC_INPUT_SUFFIX) | ||
| } | ||
|
|
||
| pub(crate) fn is_ptxas_input(arg: impl AsRef<OsStr>) -> bool { | ||
| arg.as_ref().to_string_lossy().ends_with(PTXAS_INPUT_SUFFIX) | ||
| } |
There was a problem hiding this comment.
does it have to be pub ?
There was a problem hiding this comment.
They are used in the nvcc module.
Codecov Report❌ Patch coverage is
Additional details and impacted files@@ Coverage Diff @@
## main #2722 +/- ##
==========================================
+ Coverage 74.64% 74.78% +0.14%
==========================================
Files 70 70
Lines 39893 40064 +171
==========================================
+ Hits 29777 29961 +184
+ Misses 10116 10103 -13 ☔ View full report in Codecov by Harness. 🚀 New features to boost your workflow:
|
Merging this PR will not alter performance
Comparing Footnotes
|
be76bf6 to
f90eeef
Compare
f90eeef to
f713f8d
Compare
|
|
||
| [workspace.lints.clippy] | ||
| cloned_instead_of_copied = "warn" | ||
| cloned_ref_to_slice_refs = "warn" |
There was a problem hiding this comment.
please do that a in different PR
nvcc 13.3+ injects a valueless --simt-only flag (Tile compilation) into the cicc/cudafe++ sub-commands of `nvcc --dryrun`, typically right before -o. sccache 0.16.0 locates the cicc/ptxas input by a fixed offset from the end of the argument list and groups the dryrun sub-commands positionally, so the extra token makes it mis-locate the input and skip the device compile. ptxas then produces no .cubin and the build dies with "fatbinary fatal: Could not open input file '*.cubin'". This affects any separable/relocatable (-rdc=true) CUDA target (e.g. gloo_cuda, torch_nvshmem). For CUDA >= 13.3 toolkits, build sccache 0.16.0 from source with the upstream fix backported from mozilla/sccache#2722 (patches/sccache-nvcc-13.3-dryrun-parsing.patch). It locates the cicc/ptxas input by its generated-file suffix (.cpp1.ii / .ptx) instead of by position, in both cicc.rs and nvcc.rs. Older CUDA (and non-CUDA) images keep downloading the prebuilt binary, so their build behavior is unchanged. The gate keys off CUDA_VERSION, so the CUDA 13.4 image picks up the patched sccache automatically. Drop the from-source build once a sccache release including #2722 ships. See #190832. Test Plan: Not buildable on the author's machine (no cargo/docker). To be validated in the CUDA 13.4 image: echo '__global__ void k() {}' > repro.cu sccache nvcc -rdc=true -gencode arch=compute_75,code=sm_75 -c repro.cu -o repro.o # succeeds, no fatbinary error Authored with assistance from Claude. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
nvcc 13.3+ injects a valueless --simt-only flag (Tile compilation) into the cicc/cudafe++ sub-commands of `nvcc --dryrun`, typically right before -o. The prebuilt sccache 0.16.0 locates the cicc/ptxas input by a fixed offset from the end of the argument list and groups the dryrun sub-commands positionally, so the extra token makes it mis-locate the input and skip the device compile. ptxas then produces no .cubin and the build dies with "fatbinary fatal: Could not open input file '*.cubin'". This affects any separable/relocatable (-rdc=true) CUDA target (e.g. gloo_cuda, torch_nvshmem). Revert #189365 (download prebuilt sccache) and build sccache 0.16.0 from source again, applying the upstream fix backported from mozilla/sccache#2722 (patches/sccache-nvcc-13.3-dryrun-parsing.patch). It locates the cicc/ptxas input by its generated-file suffix (.cpp1.ii / .ptx) instead of by position, in both cicc.rs and nvcc.rs. Building from source for all images keeps a single sccache binary across the fleet; the patch is a no-op on older CUDA (it only changes how the input is located and falls back to the old offset). Drop the from-source build and return to the prebuilt binary once a sccache release including #2722 ships. See #190832. Test Plan: Not buildable on the author's machine (no cargo/docker). Validate in a CUDA 13.3+ CI image (e.g. cuda13.4): echo '__global__ void k() {}' > repro.cu sccache nvcc -rdc=true -gencode arch=compute_75,code=sm_75 -c repro.cu -o repro.o # succeeds, no fatbinary error and confirm older-CUDA / non-CUDA images still build and cache normally. Authored with assistance from Claude. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
nvcc 13.3+ injects a valueless --simt-only flag (Tile compilation) into the cicc/cudafe++ sub-commands of `nvcc --dryrun`, typically right before -o. The prebuilt sccache 0.16.0 locates the cicc/ptxas input by a fixed offset from the end of the argument list and groups the dryrun sub-commands positionally, so the extra token makes it mis-locate the input and skip the device compile. ptxas then produces no .cubin and the build dies with "fatbinary fatal: Could not open input file '*.cubin'". This affects any separable/relocatable (-rdc=true) CUDA target (e.g. gloo_cuda, torch_nvshmem). Revert #189365 (download prebuilt sccache) and build sccache 0.16.0 from source again, applying the upstream fix backported from mozilla/sccache#2722 (patches/sccache-nvcc-13.3-dryrun-parsing.patch). It locates the cicc/ptxas input by its generated-file suffix (.cpp1.ii / .ptx) instead of by position, in both cicc.rs and nvcc.rs. Building from source for all images keeps a single sccache binary across the fleet; the patch is a no-op on older CUDA (it only changes how the input is located and falls back to the old offset). Drop the from-source build and return to the prebuilt binary once a sccache release including #2722 ships. See #190832. Test Plan: Not buildable on the author's machine (no cargo/docker). Validate in a CUDA 13.3+ CI image (e.g. cuda13.4): echo '__global__ void k() {}' > repro.cu sccache nvcc -rdc=true -gencode arch=compute_75,code=sm_75 -c repro.cu -o repro.o # succeeds, no fatbinary error and confirm older-CUDA / non-CUDA images still build and cache normally. Authored with assistance from Claude. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
|
Bump @sylvestre, we would very much like to use this change in PyTorch for 13.4 CUDA builds |
nvcc 13.3+ injects a valueless --simt-only flag (Tile compilation) into the cicc/cudafe++ sub-commands of `nvcc --dryrun`, typically right before -o. The prebuilt sccache 0.16.0 locates the cicc/ptxas input by a fixed offset from the end of the argument list and groups the dryrun sub-commands positionally, so the extra token makes it mis-locate the input and skip the device compile. ptxas then produces no .cubin and the build dies with "fatbinary fatal: Could not open input file '*.cubin'". This affects any separable/relocatable (-rdc=true) CUDA target (e.g. gloo_cuda, torch_nvshmem). Revert #189365 (download prebuilt sccache) and build sccache 0.16.0 from source again, applying the upstream fix backported from mozilla/sccache#2722 (patches/sccache-nvcc-13.3-dryrun-parsing.patch). It locates the cicc/ptxas input by its generated-file suffix (.cpp1.ii / .ptx) instead of by position, in both cicc.rs and nvcc.rs. Building from source for all images keeps a single sccache binary across the fleet; the patch is a no-op on older CUDA (it only changes how the input is located and falls back to the old offset). Drop the from-source build and return to the prebuilt binary once a sccache release including #2722 ships. See #190832. Test Plan: Not buildable on the author's machine (no cargo/docker). Validate in a CUDA 13.3+ CI image (e.g. cuda13.4): echo '__global__ void k() {}' > repro.cu sccache nvcc -rdc=true -gencode arch=compute_75,code=sm_75 -c repro.cu -o repro.o # succeeds, no fatbinary error and confirm older-CUDA / non-CUDA images still build and cache normally. Authored with assistance from Claude. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
nvcc 13.3+ injects a valueless --simt-only flag (Tile compilation) into the cicc/cudafe++ sub-commands of `nvcc --dryrun`, typically right before -o. The prebuilt sccache 0.16.0 locates the cicc/ptxas input by a fixed offset from the end of the argument list and groups the dryrun sub-commands positionally, so the extra token makes it mis-locate the input and skip the device compile. ptxas then produces no .cubin and the build dies with "fatbinary fatal: Could not open input file '*.cubin'". This affects any separable/relocatable (-rdc=true) CUDA target (e.g. gloo_cuda, torch_nvshmem). Revert #189365 (download prebuilt sccache) and build sccache 0.16.0 from source again, applying the upstream fix backported from mozilla/sccache#2722 (patches/sccache-nvcc-13.3-dryrun-parsing.patch). It locates the cicc/ptxas input by its generated-file suffix (.cpp1.ii / .ptx) instead of by position, in both cicc.rs and nvcc.rs. Building from source for all images keeps a single sccache binary across the fleet; the patch is a no-op on older CUDA (it only changes how the input is located and falls back to the old offset). Drop the from-source build and return to the prebuilt binary once a sccache release including #2722 ships. See #190832. Test Plan: Not buildable on the author's machine (no cargo/docker). Validate in a CUDA 13.3+ CI image (e.g. cuda13.4): echo '__global__ void k() {}' > repro.cu sccache nvcc -rdc=true -gencode arch=compute_75,code=sm_75 -c repro.cu -o repro.o # succeeds, no fatbinary error and confirm older-CUDA / non-CUDA images still build and cache normally. Authored with assistance from Claude. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…90927) Fixes #190832 nvcc 13.3+ injects a valueless --simt-only flag (Tile compilation) into the cicc sub-command of `nvcc --dryrun`. sccache 0.16.0 locates the cicc input by a fixed offset from the end of the argument list, so the extra token makes it splice out --simt-only as the "input" and skip the device compile. This affects any separable/relocatable (-rdc=true) CUDA target (e.g. gloo_cuda, torch_nvshmem). Building from source in sccache 0.16.0 + the root cause fix, and keeping consistent for all builds. Root cause fix is pending mozilla/sccache#2722, once fixed release ships we will drop this from-source build. Test Plan: To be validated in 13.4 build https://github.com/pytorch/pytorch/actions/runs/29947019889/job/89015281565?pr=190639 Pull Request resolved: #190927 Approved by: https://github.com/Skylion007, https://github.com/atalman Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
nvcc 13.3+ injects a valueless --simt-only flag (Tile compilation) into the cicc/cudafe++ sub-commands of `nvcc --dryrun`, typically right before -o. The prebuilt sccache 0.16.0 locates the cicc/ptxas input by a fixed offset from the end of the argument list and groups the dryrun sub-commands positionally, so the extra token makes it mis-locate the input and skip the device compile. ptxas then produces no .cubin and the build dies with "fatbinary fatal: Could not open input file '*.cubin'". This affects any separable/relocatable (-rdc=true) CUDA target (e.g. gloo_cuda, torch_nvshmem). Revert #189365 (download prebuilt sccache) and build sccache 0.16.0 from source again, applying the upstream fix backported from mozilla/sccache#2722 (patches/sccache-nvcc-13.3-dryrun-parsing.patch). It locates the cicc/ptxas input by its generated-file suffix (.cpp1.ii / .ptx) instead of by position, in both cicc.rs and nvcc.rs. Building from source for all images keeps a single sccache binary across the fleet; the patch is a no-op on older CUDA (it only changes how the input is located and falls back to the old offset). Drop the from-source build and return to the prebuilt binary once a sccache release including #2722 ships. See #190832. Test Plan: Not buildable on the author's machine (no cargo/docker). Validate in a CUDA 13.3+ CI image (e.g. cuda13.4): echo '__global__ void k() {}' > repro.cu sccache nvcc -rdc=true -gencode arch=compute_75,code=sm_75 -c repro.cu -o repro.o # succeeds, no fatbinary error and confirm older-CUDA / non-CUDA images still build and cache normally. Authored with assistance from Claude. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
nvcc 13.3+ injects a valueless --simt-only flag (Tile compilation) into the cicc/cudafe++ sub-commands of `nvcc --dryrun`, typically right before -o. The prebuilt sccache 0.16.0 locates the cicc/ptxas input by a fixed offset from the end of the argument list and groups the dryrun sub-commands positionally, so the extra token makes it mis-locate the input and skip the device compile. ptxas then produces no .cubin and the build dies with "fatbinary fatal: Could not open input file '*.cubin'". This affects any separable/relocatable (-rdc=true) CUDA target (e.g. gloo_cuda, torch_nvshmem). Revert #189365 (download prebuilt sccache) and build sccache 0.16.0 from source again, applying the upstream fix backported from mozilla/sccache#2722 (patches/sccache-nvcc-13.3-dryrun-parsing.patch). It locates the cicc/ptxas input by its generated-file suffix (.cpp1.ii / .ptx) instead of by position, in both cicc.rs and nvcc.rs. Building from source for all images keeps a single sccache binary across the fleet; the patch is a no-op on older CUDA (it only changes how the input is located and falls back to the old offset). Drop the from-source build and return to the prebuilt binary once a sccache release including #2722 ships. See #190832. Test Plan: Not buildable on the author's machine (no cargo/docker). Validate in a CUDA 13.3+ CI image (e.g. cuda13.4): echo '__global__ void k() {}' > repro.cu sccache nvcc -rdc=true -gencode arch=compute_75,code=sm_75 -c repro.cu -o repro.o # succeeds, no fatbinary error and confirm older-CUDA / non-CUDA images still build and cache normally. Authored with assistance from Claude. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
|
When this merging? Waiting for it since CUDA 13.3 |
nvcc 13.3+ injects a valueless --simt-only flag (Tile compilation) into the cicc/cudafe++ sub-commands of `nvcc --dryrun`, typically right before -o. The prebuilt sccache 0.16.0 locates the cicc/ptxas input by a fixed offset from the end of the argument list and groups the dryrun sub-commands positionally, so the extra token makes it mis-locate the input and skip the device compile. ptxas then produces no .cubin and the build dies with "fatbinary fatal: Could not open input file '*.cubin'". This affects any separable/relocatable (-rdc=true) CUDA target (e.g. gloo_cuda, torch_nvshmem). Revert #189365 (prebuilt-binary download) and build sccache 0.16.0 from source again, applying the upstream fix backported from mozilla/sccache#2722 (patches/sccache-nvcc-13.3-dryrun-parsing.patch), which locates the cicc/ptxas input by its generated-file suffix (.cpp1.ii / .ptx) instead of by position, in both cicc.rs and nvcc.rs. Build from source for all image types (riscv64 keeps the prebuilt binary: no nvcc, no sccache-dist, and a from-source build under QEMU emulation would be impractically slow) so the whole fleet shares a single sccache binary; the patch is a no-op on older/non-CUDA compiles (it only changes how the input is located and falls back to the old offset). Source the cargo env from $CARGO_HOME. Since #186302 the builder images pre-install rust at /opt/rust and set CARGO_HOME=/opt/rust, so rustup's env file is /opt/rust/env, not $HOME/.cargo/env. Hardcoding `. "$HOME/.cargo/env"` made every from-source image fail at that line before reaching cargo (this caused the revert of the first landing). Use `. "${CARGO_HOME:-$HOME/.cargo}/env"`. Drop the from-source build and return to the prebuilt binary once a sccache release including #2722 ships. See #190832. Test Plan: Not buildable on the author's machine (no cargo/docker). The patch applies cleanly to sccache v0.16.0 (`git apply --check`). Validate in a CUDA 13.3+ CI image (e.g. cuda13.4): echo '__global__ void k() {}' > repro.cu sccache nvcc -rdc=true -gencode arch=compute_75,code=sm_75 -c repro.cu -o repro.o # succeeds, no fatbinary error and confirm non-CUDA / ROCm / XPU images still build and cache normally. Authored with assistance from Claude. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
nvcc 13.3+ injects a valueless --simt-only flag (Tile compilation) into the cicc/cudafe++ sub-commands of `nvcc --dryrun`, typically right before -o. The prebuilt sccache 0.16.0 locates the cicc/ptxas input by a fixed offset from the end of the argument list and groups the dryrun sub-commands positionally, so the extra token makes it mis-locate the input and skip the device compile. ptxas then produces no .cubin and the build dies with "fatbinary fatal: Could not open input file '*.cubin'". This affects any separable/relocatable (-rdc=true) CUDA target (e.g. gloo_cuda, torch_nvshmem). Revert #189365 (prebuilt-binary download) and build sccache 0.16.0 from source again, applying the upstream fix backported from mozilla/sccache#2722 (patches/sccache-nvcc-13.3-dryrun-parsing.patch), which locates the cicc/ptxas input by its generated-file suffix (.cpp1.ii / .ptx) instead of by position, in both cicc.rs and nvcc.rs. Build from source for all image types (riscv64 keeps the prebuilt binary: no nvcc, no sccache-dist, and a from-source build under QEMU emulation would be impractically slow) so the whole fleet shares a single sccache binary; the patch is a no-op on older/non-CUDA compiles (it only changes how the input is located and falls back to the old offset). Build with the pre-installed rust toolchain rather than reinstalling rustup. Since #186302 the builder images pre-install a pinned rust (1.97.1) at CARGO_HOME=/opt/rust with /opt/rust/bin on PATH. The first landing ran `curl rustup | sh` then `. "$HOME/.cargo/env"`, but with CARGO_HOME=/opt/rust that env file is /opt/rust/env, so the source failed and every from-source image died at that line before reaching cargo -- this caused the revert. Drop the redundant (and pin-diverging) rustup reinstall and just use the cargo already on PATH. Drop the from-source build and return to the prebuilt binary once a sccache release including #2722 ships. See #190832. Test Plan: Not buildable on the author's machine (no cargo/docker). The patch applies cleanly to sccache v0.16.0 (`git apply --check`). Validate in a CUDA 13.3+ CI image (e.g. cuda13.4): echo '__global__ void k() {}' > repro.cu sccache nvcc -rdc=true -gencode arch=compute_75,code=sm_75 -c repro.cu -o repro.o # succeeds, no fatbinary error and confirm non-CUDA / ROCm / XPU images still build and cache normally. Authored with assistance from Claude. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
nvcc 13.3+ injects a valueless --simt-only flag (Tile compilation) into the cicc/cudafe++ sub-commands of `nvcc --dryrun`, typically right before -o. The prebuilt sccache 0.16.0 locates the cicc/ptxas input by a fixed offset from the end of the argument list and groups the dryrun sub-commands positionally, so the extra token makes it mis-locate the input and skip the device compile. ptxas then produces no .cubin and the build dies with "fatbinary fatal: Could not open input file '*.cubin'". This affects any separable/relocatable (-rdc=true) CUDA target (e.g. gloo_cuda, torch_nvshmem). Revert #189365 (prebuilt-binary download) and build sccache 0.16.0 from source again, applying the upstream fix backported from mozilla/sccache#2722 (patches/sccache-nvcc-13.3-dryrun-parsing.patch), which locates the cicc/ptxas input by its generated-file suffix (.cpp1.ii / .ptx) instead of by position, in both cicc.rs and nvcc.rs. Build from source for all image types (riscv64 keeps the prebuilt binary: no nvcc, no sccache-dist, and a from-source build under QEMU emulation would be impractically slow) so the whole fleet shares a single sccache binary; the patch is a no-op on older/non-CUDA compiles (it only changes how the input is located and falls back to the old offset). Build with the pre-installed rust toolchain rather than reinstalling rustup. Since #186302 the builder images pre-install a pinned rust (1.97.1) at CARGO_HOME=/opt/rust with /opt/rust/bin on PATH. The first landing ran `curl rustup | sh` then `. "$HOME/.cargo/env"`, but with CARGO_HOME=/opt/rust that env file is /opt/rust/env, so the source failed and every from-source image died at that line before reaching cargo -- this caused the revert. Drop the redundant (and pin-diverging) rustup reinstall and just use the cargo already on PATH. Drop the from-source build and return to the prebuilt binary once a sccache release including #2722 ships. See #190832. Test Plan: Not buildable on the author's machine (no cargo/docker). The patch applies cleanly to sccache v0.16.0 (`git apply --check`). Validate in a CUDA 13.3+ CI image (e.g. cuda13.4): echo '__global__ void k() {}' > repro.cu sccache nvcc -rdc=true -gencode arch=compute_75,code=sm_75 -c repro.cu -o repro.o # succeeds, no fatbinary error and confirm non-CUDA / ROCm / XPU images still build and cache normally. Authored with assistance from Claude. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
nvcc 13.3+ injects a valueless --simt-only flag (Tile compilation) into the cicc/cudafe++ sub-commands of `nvcc --dryrun`, typically right before -o. The prebuilt sccache 0.16.0 locates the cicc/ptxas input by a fixed offset from the end of the argument list and groups the dryrun sub-commands positionally, so the extra token makes it mis-locate the input and skip the device compile. ptxas then produces no .cubin and the build dies with "fatbinary fatal: Could not open input file '*.cubin'". This affects any separable/relocatable (-rdc=true) CUDA target (e.g. gloo_cuda, torch_nvshmem). Revert #189365 (prebuilt-binary download) and build sccache 0.16.0 from source again, applying the upstream fix backported from mozilla/sccache#2722 (patches/sccache-nvcc-13.3-dryrun-parsing.patch), which locates the cicc/ptxas input by its generated-file suffix (.cpp1.ii / .ptx) instead of by position, in both cicc.rs and nvcc.rs. Build from source for all image types (riscv64 keeps the prebuilt binary: no nvcc, no sccache-dist, and a from-source build under QEMU emulation would be impractically slow) so the whole fleet shares a single sccache binary; the patch is a no-op on older/non-CUDA compiles (it only changes how the input is located and falls back to the old offset). Build with the pre-installed rust toolchain rather than reinstalling rustup. Since #186302 the builder images pre-install a pinned rust (1.97.1) at CARGO_HOME=/opt/rust with /opt/rust/bin on PATH. The first landing ran `curl rustup | sh` then `. "$HOME/.cargo/env"`, but with CARGO_HOME=/opt/rust that env file is /opt/rust/env, so the source failed and every from-source image died at that line before reaching cargo -- this caused the revert. Drop the redundant (and pin-diverging) rustup reinstall and just use the cargo already on PATH. Drop the from-source build and return to the prebuilt binary once a sccache release including #2722 ships. See #190832. Test Plan: Not buildable on the author's machine (no cargo/docker). The patch applies cleanly to sccache v0.16.0 (`git apply --check`). Validate in a CUDA 13.3+ CI image (e.g. cuda13.4): echo '__global__ void k() {}' > repro.cu sccache nvcc -rdc=true -gencode arch=compute_75,code=sm_75 -c repro.cu -o repro.o # succeeds, no fatbinary error and confirm non-CUDA / ROCm / XPU images still build and cache normally. Authored with assistance from Claude. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
nvcc 13.3+ injects a valueless --simt-only flag (Tile compilation) into the cicc/cudafe++ sub-commands of `nvcc --dryrun`, typically right before -o. The prebuilt sccache 0.16.0 locates the cicc/ptxas input by a fixed offset from the end of the argument list and groups the dryrun sub-commands positionally, so the extra token makes it mis-locate the input and skip the device compile. ptxas then produces no .cubin and the build dies with "fatbinary fatal: Could not open input file '*.cubin'". This affects any separable/relocatable (-rdc=true) CUDA target (e.g. gloo_cuda, torch_nvshmem). Revert #189365 (prebuilt-binary download) and build sccache 0.16.0 from source again, applying the upstream fix backported from mozilla/sccache#2722 (patches/sccache-nvcc-13.3-dryrun-parsing.patch), which locates the cicc/ptxas input by its generated-file suffix (.cpp1.ii / .ptx) instead of by position, in both cicc.rs and nvcc.rs. Build from source for all image types (riscv64 keeps the prebuilt binary: no nvcc, no sccache-dist, and a from-source build under QEMU emulation would be impractically slow) so the whole fleet shares a single sccache binary; the patch is a no-op on older/non-CUDA compiles (it only changes how the input is located and falls back to the old offset). Build with the pre-installed rust toolchain rather than reinstalling rustup. Since #186302 the builder images pre-install a pinned rust (1.97.1) at CARGO_HOME=/opt/rust with /opt/rust/bin on PATH. The first landing ran `curl rustup | sh` then `. "$HOME/.cargo/env"`, but with CARGO_HOME=/opt/rust that env file is /opt/rust/env, so the source failed and every from-source image died at that line before reaching cargo -- this caused the revert. Drop the redundant (and pin-diverging) rustup reinstall and just use the cargo already on PATH. Drop the from-source build and return to the prebuilt binary once a sccache release including #2722 ships. See #190832. Test Plan: Not buildable on the author's machine (no cargo/docker). The patch applies cleanly to sccache v0.16.0 (`git apply --check`). Validate in a CUDA 13.3+ CI image (e.g. cuda13.4): echo '__global__ void k() {}' > repro.cu sccache nvcc -rdc=true -gencode arch=compute_75,code=sm_75 -c repro.cu -o repro.o # succeeds, no fatbinary error and confirm non-CUDA / ROCm / XPU images still build and cache normally. Authored with assistance from Claude. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
nvcc 13.3+ injects a valueless --simt-only flag (Tile compilation) into the cicc/cudafe++ sub-commands of `nvcc --dryrun`, typically right before -o. The prebuilt sccache 0.16.0 locates the cicc/ptxas input by a fixed offset from the end of the argument list and groups the dryrun sub-commands positionally, so the extra token makes it mis-locate the input and skip the device compile. ptxas then produces no .cubin and the build dies with "fatbinary fatal: Could not open input file '*.cubin'". This affects any separable/relocatable (-rdc=true) CUDA target (e.g. gloo_cuda, torch_nvshmem). Revert #189365 (prebuilt-binary download) and build sccache 0.16.0 from source again, applying the upstream fix backported from mozilla/sccache#2722 (patches/sccache-nvcc-13.3-dryrun-parsing.patch), which locates the cicc/ptxas input by its generated-file suffix (.cpp1.ii / .ptx) instead of by position, in both cicc.rs and nvcc.rs. Build from source for all image types (riscv64 keeps the prebuilt binary: no nvcc, no sccache-dist, and a from-source build under QEMU emulation would be impractically slow) so the whole fleet shares a single sccache binary; the patch is a no-op on older/non-CUDA compiles (it only changes how the input is located and falls back to the old offset). Build with the pre-installed rust toolchain rather than reinstalling rustup. Since #186302 the builder images pre-install a pinned rust (1.97.1) at CARGO_HOME=/opt/rust with /opt/rust/bin on PATH. The first landing ran `curl rustup | sh` then `. "$HOME/.cargo/env"`, but with CARGO_HOME=/opt/rust that env file is /opt/rust/env, so the source failed and every from-source image died at that line before reaching cargo -- this caused the revert. Drop the redundant (and pin-diverging) rustup reinstall and just use the cargo already on PATH. Drop the from-source build and return to the prebuilt binary once a sccache release including #2722 ships. See #190832. Test Plan: Not buildable on the author's machine (no cargo/docker). The patch applies cleanly to sccache v0.16.0 (`git apply --check`). Validate in a CUDA 13.3+ CI image (e.g. cuda13.4): echo '__global__ void k() {}' > repro.cu sccache nvcc -rdc=true -gencode arch=compute_75,code=sm_75 -c repro.cu -o repro.o # succeeds, no fatbinary error and confirm non-CUDA / ROCm / XPU images still build and cache normally. Authored with assistance from Claude. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…90927) Fixes #190832 nvcc 13.3+ injects a valueless --simt-only flag (Tile compilation) into the cicc sub-command of `nvcc --dryrun`. sccache 0.16.0 locates the cicc input by a fixed offset from the end of the argument list, so the extra token makes it splice out --simt-only as the "input" and skip the device compile. This affects any separable/relocatable (-rdc=true) CUDA target (e.g. gloo_cuda, torch_nvshmem). Building from source in sccache 0.16.0 + the root cause fix, and keeping consistent for all builds. Root cause fix is pending mozilla/sccache#2722, once fixed release ships we will drop this from-source build. Test Plan: To be validated in 13.4 build https://github.com/pytorch/pytorch/actions/runs/29947019889/job/89015281565?pr=190639 Pull Request resolved: #190927 Approved by: https://github.com/Skylion007, https://github.com/atalman Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
|
bumping this, @mbrobbel @sylvestre could we get this merged? |
nvcc 13.3+ injects a valueless --simt-only flag (Tile compilation) into the cicc/cudafe++ sub-commands of `nvcc --dryrun`, typically right before -o. The prebuilt sccache 0.16.0 locates the cicc/ptxas input by a fixed offset from the end of the argument list and groups the dryrun sub-commands positionally, so the extra token makes it mis-locate the input and skip the device compile. ptxas then produces no .cubin and the build dies with "fatbinary fatal: Could not open input file '*.cubin'". This affects any separable/relocatable (-rdc=true) CUDA target (e.g. gloo_cuda, torch_nvshmem). Revert #189365 (prebuilt-binary download) and build sccache 0.16.0 from source again, applying the upstream fix from mozilla/sccache#2722 (patches/sccache-nvcc-13.3-dryrun-parsing.patch). Build from source for all image types (riscv64 keeps the prebuilt binary: no nvcc, no sccache-dist, and a from-source build under QEMU emulation would be impractically slow) so the whole fleet shares a single sccache binary. The backport gates #2722 rather than taking it verbatim. Upstream replaces the positional lookup with a search for the input's generated-file suffix (.cpp1.ii / .ptx) unconditionally, for both cicc and ptxas. Upstream's own dryrun fixture shows --simt-only on the cudafe++ and cicc lines but not on the ptxas line, so ptxas offsets were already correct and did not need changing; running the search there regressed CUDA < 13.3, where torchao's mxfp8 extension fails to build with `Missing "cubin" file output` while stock sccache builds it fine. Here the suffix search runs only when the argument at the historical offset is not a generated input file, so every command line nvcc still emits in the old shape parses exactly as stock 0.16.0 does -- same input, same cache keys -- and 13.3+ still gets the fix on the cicc line that needs it. The "Missing <lang> file output" error now also reports the parsed input and argument lists. Reaching that error means the input was located at the wrong argument and -o was consumed or dropped, and nvcc reshapes these sub-commands between CUDA releases, so the argument vector as sccache saw it is the thing needed to diagnose the next occurrence. Two alternatives were considered. Dropping the ptxas half of #2722 outright is a smaller diff, but it leaves cicc's search unconditional on older CUDA (the same pattern, one compiler over) and gives ptxas no fallback if NVIDIA later shifts its input off the offset the way it did for cicc. Disabling sccache on 13.3+ avoids the parser entirely but gives up CUDA caching, which is the point of having sccache at all. Build with the pre-installed rust toolchain rather than reinstalling rustup. Since #186302 the builder images pre-install a pinned rust (1.97.1) at CARGO_HOME=/opt/rust with /opt/rust/bin on PATH. The first landing ran `curl rustup | sh` then `. "$HOME/.cargo/env"`, but with CARGO_HOME=/opt/rust that env file is /opt/rust/env, so the source failed and every from-source image died at that line before reaching cargo -- this caused the revert. Drop the redundant (and pin-diverging) rustup reinstall and just use the cargo already on PATH. Drop the from-source build and return to the prebuilt binary once a sccache release including #2722 ships. See #190832. Test Plan: Not buildable on the author's machine (no cargo/docker), so the patch is not compiled or unit-tested locally. It applies cleanly to a pristine v0.16.0: ``` git clone --depth 1 --branch v0.16.0 https://github.com/mozilla/sccache /tmp/sccache git -C /tmp/sccache apply --check .ci/docker/common/patches/sccache-nvcc-13.3-dryrun-parsing.patch ``` Validate in a CUDA 13.3+ CI image (e.g. cuda13.4): ``` echo '__global__ void k() {}' > repro.cu sccache nvcc -rdc=true -gencode arch=compute_75,code=sm_75 -c repro.cu -o repro.o # succeeds, no fatbinary error ``` and confirm the CUDA 13.0/13.2 torchao build (which regressed on the ungated backport) plus non-CUDA / ROCm / XPU images still build and cache normally. Authored with assistance from Claude. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
nvcc 13.3+ injects a valueless --simt-only flag (Tile compilation) into the cicc/cudafe++ sub-commands of `nvcc --dryrun`, typically right before -o. The prebuilt sccache 0.16.0 locates the cicc/ptxas input by a fixed offset from the end of the argument list and groups the dryrun sub-commands positionally, so the extra token makes it mis-locate the input and skip the device compile. ptxas then produces no .cubin and the build dies with "fatbinary fatal: Could not open input file '*.cubin'". This affects any separable/relocatable (-rdc=true) CUDA target (e.g. gloo_cuda, torch_nvshmem). Revert #189365 (prebuilt-binary download) and build sccache 0.16.0 from source again, applying the upstream fix from mozilla/sccache#2722 (patches/sccache-nvcc-13.3-dryrun-parsing.patch). Build from source for all image types (riscv64 keeps the prebuilt binary: no nvcc, no sccache-dist, and a from-source build under QEMU emulation would be impractically slow) so the whole fleet shares a single sccache binary. The backport gates #2722 rather than taking it verbatim. Upstream replaces the positional lookup with a search for the input's generated-file suffix (.cpp1.ii / .ptx) unconditionally, for both cicc and ptxas. Upstream's own dryrun fixture shows --simt-only on the cudafe++ and cicc lines but not on the ptxas line, so ptxas offsets were already correct and did not need changing; running the search there regressed CUDA < 13.3, where torchao's mxfp8 extension fails to build with `Missing "cubin" file output` while stock sccache builds it fine. Here the suffix search runs only when the argument at the historical offset is not a generated input file, so every command line nvcc still emits in the old shape parses exactly as stock 0.16.0 does -- same input, same cache keys -- and 13.3+ still gets the fix on the cicc line that needs it. The "Missing <lang> file output" error now also reports the parsed input and argument lists. Reaching that error means the input was located at the wrong argument and -o was consumed or dropped, and nvcc reshapes these sub-commands between CUDA releases, so the argument vector as sccache saw it is the thing needed to diagnose the next occurrence. Two alternatives were considered. Dropping the ptxas half of #2722 outright is a smaller diff, but it leaves cicc's search unconditional on older CUDA (the same pattern, one compiler over) and gives ptxas no fallback if NVIDIA later shifts its input off the offset the way it did for cicc. Disabling sccache on 13.3+ avoids the parser entirely but gives up CUDA caching, which is the point of having sccache at all. Build with the pre-installed rust toolchain rather than reinstalling rustup. Since #186302 the builder images pre-install a pinned rust (1.97.1) at CARGO_HOME=/opt/rust with /opt/rust/bin on PATH. The first landing ran `curl rustup | sh` then `. "$HOME/.cargo/env"`, but with CARGO_HOME=/opt/rust that env file is /opt/rust/env, so the source failed and every from-source image died at that line before reaching cargo -- this caused the revert. Drop the redundant (and pin-diverging) rustup reinstall and just use the cargo already on PATH. Drop the from-source build and return to the prebuilt binary once a sccache release including #2722 ships. See #190832. Test Plan: Not buildable on the author's machine (no cargo/docker), so the patch is not compiled or unit-tested locally. It applies cleanly to a pristine v0.16.0: ``` git clone --depth 1 --branch v0.16.0 https://github.com/mozilla/sccache /tmp/sccache git -C /tmp/sccache apply --check .ci/docker/common/patches/sccache-nvcc-13.3-dryrun-parsing.patch ``` Validate in a CUDA 13.3+ CI image (e.g. cuda13.4): ``` echo '__global__ void k() {}' > repro.cu sccache nvcc -rdc=true -gencode arch=compute_75,code=sm_75 -c repro.cu -o repro.o # succeeds, no fatbinary error ``` and confirm the CUDA 13.0/13.2 torchao build (which regressed on the ungated backport) plus non-CUDA / ROCm / XPU images still build and cache normally. Authored with assistance from Claude. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Reland of #190927 (landed as 2a4b761, auto-reverted in dc90186), with the cause of the revert fixed. Two separate sccache 0.16.0 bugs break CUDA builds; this patches both (patches/sccache-nvcc-13.3-dryrun-parsing.patch) and builds sccache from source until a release carrying them ships. First, nvcc 13.3+ injects a valueless --simt-only flag into the cudafe++ and cicc sub-commands of `nvcc --dryrun`, right before -o. sccache locates the sub-command input by a fixed offset from the end of the argument list, so the extra token makes cicc mis-locate its input and skip the device compile; ptxas then produces no .cubin and the build dies with "fatbinary fatal: Could not open input file '*.cubin'". mozilla/sccache#2722 fixes this by locating the input by its generated-file suffix (.cpp1.ii for cicc, .ptx for ptxas) instead. That fix is taken with one deviation: the suffix search runs only when the argument at the historical offset is not itself a generated input file. Of the three sub-commands sccache parses this way, only cicc is affected -- cudafe++ takes the input as its last argument (offset 1, unaffected by a flag inserted before it) and ptxas never receives --simt-only at all, so its offset was already correct. Gating keeps every command line nvcc still emits in the old shape parsing exactly as stock 0.16.0 does, with identical cache keys, while 13.3+ still gets the fix on the one sub-command that needs it. Second, and this is what the autorevert caught: for a virtual-only gencode such as -gencode arch=compute_120,code=compute_120, nvcc runs ptxas as a PTX verification pass with no -o at all -- ptxas -arch=compute_120 -m64 kernel.compute_120.ptx It syntax-checks the PTX and emits nothing. sccache unconditionally requires an "obj" output, which only -o populates, so it fails the whole compile with `Missing "cubin" file output`. This is not a regression from #2722: stock 0.16.0 fails identically here, since with a three-argument list the offset picks -arch=compute_120 as the input and there is still no -o to find. Any target carrying a +PTX architecture can hit it; here it was torchao's mxfp8 extension, built as a dependency of inductor-build-cuda132. A verification pass has no artifact to cache, so treat the missing output as CannotCache and let the command run directly rather than failing the build. The "Missing <lang> file output" error now also reports the parsed input and argument lists. That instrumentation is what identified the second bug -- the error alone said only that an output was absent, not that the command never had one -- and nvcc reshapes these sub-commands between CUDA releases, so the argument vector as sccache saw it is worth keeping in the message. Build from source for all image types (riscv64 keeps the prebuilt binary: no nvcc, no sccache-dist, and a from-source build under QEMU emulation would be impractically slow) so the whole fleet shares a single sccache binary. Build with the pre-installed rust toolchain rather than reinstalling rustup. Since #186302 the builder images pre-install a pinned rust (1.97.1) at CARGO_HOME=/opt/rust with /opt/rust/bin on PATH. The first landing ran `curl rustup | sh` then `. "$HOME/.cargo/env"`, but with CARGO_HOME=/opt/rust that env file is /opt/rust/env, so the source failed and every from-source image died at that line before reaching cargo. Drop the redundant (and pin-diverging) rustup reinstall and just use the cargo already on PATH. Disabling sccache for the affected CUDA versions was considered and rejected: it gives up CUDA build caching, which is the reason to run sccache at all. Drop the from-source build and return to the prebuilt binary once a sccache release including these fixes ships. See #190832. Test Plan: Not buildable on the author's machine (no cargo/docker), so the patch is not compiled or unit-tested locally. It applies cleanly to a pristine v0.16.0: ``` git clone --depth 1 --branch v0.16.0 https://github.com/mozilla/sccache /tmp/sccache git -C /tmp/sccache apply --check .ci/docker/common/patches/sccache-nvcc-13.3-dryrun-parsing.patch ``` Validate in CI. The torchao mxfp8 extension built by inductor-build-cuda132 is the direct regression test for the second fix, since it carries a code=compute_120 gencode. For the first, in a CUDA 13.3+ image: ``` echo '__global__ void k() {}' > repro.cu sccache nvcc -rdc=true -gencode arch=compute_75,code=sm_75 -c repro.cu -o repro.o # succeeds, no fatbinary error ``` and confirm non-CUDA / ROCm / XPU images still build and cache normally. Worth watching sccache hit rates on CUDA jobs afterwards: CannotCache on the verification pass may suppress caching more of an nvcc invocation than that one sub-command. Authored with assistance from Claude. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Reland of #190927 (landed as 2a4b761, auto-reverted in dc90186), with the cause of the revert fixed. Two separate sccache 0.16.0 bugs break CUDA builds; this patches both (patches/sccache-nvcc-13.3-dryrun-parsing.patch) and builds sccache from source until a release carrying them ships. First, nvcc 13.3+ injects a valueless --simt-only flag into the cudafe++ and cicc sub-commands of `nvcc --dryrun`, right before -o. sccache locates the sub-command input by a fixed offset from the end of the argument list, so the extra token makes cicc mis-locate its input and skip the device compile; ptxas then produces no .cubin and the build dies with "fatbinary fatal: Could not open input file '*.cubin'". mozilla/sccache#2722 fixes this by locating the input by its generated-file suffix (.cpp1.ii for cicc, .ptx for ptxas) instead. That fix is taken with one deviation: the suffix search runs only when the argument at the historical offset is not itself a generated input file. Of the three sub-commands sccache parses this way, only cicc is affected -- cudafe++ takes the input as its last argument (offset 1, unaffected by a flag inserted before it) and ptxas never receives --simt-only at all, so its offset was already correct. Gating keeps every command line nvcc still emits in the old shape parsing exactly as stock 0.16.0 does, with identical cache keys, while 13.3+ still gets the fix on the one sub-command that needs it. Second, and this is what the autorevert caught: for a virtual-only gencode such as -gencode arch=compute_120,code=compute_120, nvcc runs ptxas as a PTX verification pass with no -o at all -- ptxas -arch=compute_120 -m64 kernel.compute_120.ptx It syntax-checks the PTX and emits nothing. sccache marks every ptxas sub-command Cacheable::Yes, and the caching path requires an "obj" output that only -o populates, so the whole compile fails with `Missing "cubin" file output`. This is not a regression from #2722: stock 0.16.0 fails identically here, since with a three-argument list the offset picks -arch=compute_120 as the input and there is still no -o to find. Any target carrying a +PTX architecture can hit it; here it was torchao's mxfp8 extension, built as a dependency of inductor-build-cuda132. Mark that sub-command Cacheable::No instead, which routes it to the path nvcc.rs already uses for non-cacheable sub-commands: build the command and run it. The first attempt at this returned CompilerArguments::CannotCache from cicc::parse_arguments, but the nvcc sub-command dispatch converts CannotCache into an error rather than falling back to running the command, so it only moved the failure ("Cannot cache(no output file)"). Deciding at the grouping stage avoids that path entirely; Cacheable::No with a group is already how the host preprocessor steps are handled. The "Missing <lang> file output" error now also reports the parsed input and argument lists. That instrumentation is what identified the second bug -- the error alone said only that an output was absent, not that the command never had one -- and nvcc reshapes these sub-commands between CUDA releases, so the argument vector as sccache saw it is worth keeping in the message. Build from source for all image types (riscv64 keeps the prebuilt binary: no nvcc, no sccache-dist, and a from-source build under QEMU emulation would be impractically slow) so the whole fleet shares a single sccache binary. Build with the pre-installed rust toolchain rather than reinstalling rustup. Since #186302 the builder images pre-install a pinned rust (1.97.1) at CARGO_HOME=/opt/rust with /opt/rust/bin on PATH. The first landing ran `curl rustup | sh` then `. "$HOME/.cargo/env"`, but with CARGO_HOME=/opt/rust that env file is /opt/rust/env, so the source failed and every from-source image died at that line before reaching cargo. Drop the redundant (and pin-diverging) rustup reinstall and just use the cargo already on PATH. Disabling sccache for the affected CUDA versions was considered and rejected: it gives up CUDA build caching, which is the reason to run sccache at all. Drop the from-source build and return to the prebuilt binary once a sccache release including these fixes ships. See #190832. Test Plan: The patch applies cleanly to a pristine v0.16.0: ``` git clone --depth 1 --branch v0.16.0 https://github.com/mozilla/sccache /tmp/sccache git -C /tmp/sccache apply --check .ci/docker/common/patches/sccache-nvcc-13.3-dryrun-parsing.patch ``` The second bug reproduces without PyTorch or torchao, on any CUDA install, from the virtual-only gencode alone: ``` echo '__global__ void k() {}' > repro.cu cargo build --manifest-path /tmp/sccache/Cargo.toml --release /tmp/sccache/target/release/sccache nvcc -gencode=arch=compute_100,code=sm_100 \ -gencode=arch=compute_120,code=compute_120 -c repro.cu -o repro.o ``` In CI, the torchao mxfp8 extension built by inductor-build-cuda132 is the direct regression test. For the first bug, in a CUDA 13.3+ image: ``` sccache nvcc -rdc=true -gencode arch=compute_75,code=sm_75 -c repro.cu -o repro.o # succeeds, no fatbinary error ``` and confirm non-CUDA / ROCm / XPU images still build and cache normally. Authored with assistance from Claude. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
Reland of #190927 (landed as 2a4b761, auto-reverted in dc90186), with the cause of the revert fixed. Two separate sccache 0.16.0 bugs break CUDA builds; this patches both (patches/sccache-nvcc-13.3-dryrun-parsing.patch) and builds sccache from source until a release carrying them ships. First, nvcc 13.3+ injects a valueless --simt-only flag into the cudafe++ and cicc sub-commands of `nvcc --dryrun`, right before -o. sccache locates the sub-command input by a fixed offset from the end of the argument list, so the extra token makes cicc mis-locate its input and skip the device compile; ptxas then produces no .cubin and the build dies with "fatbinary fatal: Could not open input file '*.cubin'". mozilla/sccache#2722 fixes this by locating the input by its generated-file suffix (.cpp1.ii for cicc, .ptx for ptxas) instead. That fix is taken with one deviation: the suffix search runs only when the argument at the historical offset is not itself a generated input file. Of the three sub-commands sccache parses this way, only cicc is affected -- cudafe++ takes the input as its last argument (offset 1, unaffected by a flag inserted before it) and ptxas never receives --simt-only at all, so its offset was already correct. Gating keeps every command line nvcc still emits in the old shape parsing exactly as stock 0.16.0 does, with identical cache keys, while 13.3+ still gets the fix on the one sub-command that needs it. Second, and this is what the autorevert caught: for a virtual-only gencode such as -gencode arch=compute_120,code=compute_120, nvcc runs ptxas as a PTX verification pass with no -o at all -- ptxas -arch=compute_120 -m64 kernel.compute_120.ptx It syntax-checks the PTX and emits nothing. sccache marks every ptxas sub-command Cacheable::Yes, and the caching path requires an "obj" output that only -o populates, so the whole compile fails with `Missing "cubin" file output`. This is not a regression from #2722: stock 0.16.0 fails identically here, since with a three-argument list the offset picks -arch=compute_120 as the input and there is still no -o to find. Any target carrying a +PTX architecture can hit it; here it was torchao's mxfp8 extension, built as a dependency of inductor-build-cuda132. Mark that sub-command Cacheable::No instead, which routes it to the path nvcc.rs already uses for non-cacheable sub-commands: build the command and run it. The first attempt at this returned CompilerArguments::CannotCache from cicc::parse_arguments, but the nvcc sub-command dispatch converts CannotCache into an error rather than falling back to running the command, so it only moved the failure ("Cannot cache(no output file)"). Deciding at the grouping stage avoids that path entirely; Cacheable::No with a group is already how the host preprocessor steps are handled. The "Missing <lang> file output" error now also reports the parsed input and argument lists. That instrumentation is what identified the second bug -- the error alone said only that an output was absent, not that the command never had one -- and nvcc reshapes these sub-commands between CUDA releases, so the argument vector as sccache saw it is worth keeping in the message. Build from source for all image types (riscv64 keeps the prebuilt binary: no nvcc, no sccache-dist, and a from-source build under QEMU emulation would be impractically slow) so the whole fleet shares a single sccache binary. Build with the pre-installed rust toolchain rather than reinstalling rustup. Since #186302 the builder images pre-install a pinned rust (1.97.1) at CARGO_HOME=/opt/rust with /opt/rust/bin on PATH. The first landing ran `curl rustup | sh` then `. "$HOME/.cargo/env"`, but with CARGO_HOME=/opt/rust that env file is /opt/rust/env, so the source failed and every from-source image died at that line before reaching cargo. Drop the redundant (and pin-diverging) rustup reinstall and just use the cargo already on PATH. Disabling sccache for the affected CUDA versions was considered and rejected: it gives up CUDA build caching, which is the reason to run sccache at all. Drop the from-source build and return to the prebuilt binary once a sccache release including these fixes ships. See #190832. Test Plan: The patch applies cleanly to a pristine v0.16.0: ``` git clone --depth 1 --branch v0.16.0 https://github.com/mozilla/sccache /tmp/sccache git -C /tmp/sccache apply --check .ci/docker/common/patches/sccache-nvcc-13.3-dryrun-parsing.patch ``` The second bug reproduces without PyTorch or torchao, on any CUDA install, from the virtual-only gencode alone: ``` echo '__global__ void k() {}' > repro.cu cargo build --manifest-path /tmp/sccache/Cargo.toml --release /tmp/sccache/target/release/sccache nvcc -gencode=arch=compute_100,code=sm_100 \ -gencode=arch=compute_120,code=compute_120 -c repro.cu -o repro.o ``` In CI, the torchao mxfp8 extension built by inductor-build-cuda132 is the direct regression test. For the first bug, in a CUDA 13.3+ image: ``` sccache nvcc -rdc=true -gencode arch=compute_75,code=sm_75 -c repro.cu -o repro.o # succeeds, no fatbinary error ``` and confirm non-CUDA / ROCm / XPU images still build and cache normally. Authored with assistance from Claude. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
|
Hi @sylvestre @mbrobbel , thanks for merging, we actually caught an error with older CUDA versions (<13.3) with this PR. Error message: Repro on CUDA 13.2/13.0, no PyTorch needed: $ echo '__global__ void k() {}' > repro.cu
$ nvcc --dryrun -gencode=arch=compute_100,code=sm_100 \
-gencode=arch=compute_120,code=compute_120 -c repro.cu -o repro.o 2>&1 | grep ptxas
#$ ptxas -arch=sm_100 -m64 "..._repro.compute_100.ptx" -o "..._repro.compute_100.cubin"
#$ ptxas -arch=compute_120 -m64 "..._repro.compute_120.ptx" # <- no -oFor a virtual-only gencode (compute_120) nvcc runs ptxas just to verify the PTX, with no output. Fix in the let cacheable = if args.iter().any(|arg| arg == "-o") {
Cacheable::Yes
} else {
Cacheable::No
};
(env_vars.clone(), cacheable, group)Our full patch: sccache-nvcc-13.3-dryrun-parsing.patch Could you push another fix for this error? |
…90927) Fixes #190832 nvcc 13.3+ injects a valueless --simt-only flag (Tile compilation) into the cicc sub-command of `nvcc --dryrun`. sccache 0.16.0 locates the cicc input by a fixed offset from the end of the argument list, so the extra token makes it splice out --simt-only as the "input" and skip the device compile. This affects any separable/relocatable (-rdc=true) CUDA target (e.g. gloo_cuda, torch_nvshmem). Building from source in sccache 0.16.0 + the root cause fix, and keeping consistent for all builds. Root cause fix is pending mozilla/sccache#2722, once fixed release ships we will drop this from-source build. Test Plan: To be validated in 13.4 build https://github.com/pytorch/pytorch/actions/runs/29947019889/job/89015281565?pr=190639 Pull Request resolved: #190927 Approved by: https://github.com/Skylion007, https://github.com/atalman Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
"There must be a way" — there is, and it is not a workaround. Wrapping nvcc with released sccache does not miss the cache, it FAILS the build an hour into the media stage: fatbinary fatal : Could not open input file 'activations_impl.compute_80.cubin' ninja: build stopped: subcommand failed. Root cause is upstream and already fixed. sccache decomposes nvcc by parsing `nvcc --dryrun`; with CUDA 13.3.33 the `cicc` dryrun output places `--simt-only` AFTER the input file, so sccache's positional parser took the flag as the input, mis-grouped the cicc/ptxas device steps, and the per-arch .cubin files were never produced. That is mozilla/sccache#2722, whose own description ends "Later, fatbinary fails because the expected .cubin files were never produced" — our exact error. Timing is the whole story: #2722 merged 2026-08-04, v0.17.0 shipped 2026-07-29. No RELEASE can do this yet, which is why my first attempt at the CUDA launcher failed and why SCCACHE_CACHE_MULTIARCH would have been the wrong lever (it disables caching for exactly the TUs we care about; every ONNX CUDA TU builds four -gencode arches). So sccache is now built from that merge commit (SCCACHE_GIT_REV, pinned — not a branch) via cargo in setup-rust-toolchain.ps1, landing in CARGO_BIN which precedes the scoop shims on PATH. Default features are `all`, so the webdav L2 backend survives the source build. The scoop install stays as the version-pinned baseline it overwrites. The assert could NOT be a version check: main still reports package version 0.17.0, so `sccache --version` is identical for the broken and fixed builds. verify-toolchain.ps1 therefore asserts sccache RESOLVES FROM CARGO_BIN — if the source build did not take, the base fails in milliseconds instead of the media stage failing an hour in. CMAKE_CUDA_COMPILER_LAUNCHER is restored on that foundation. 412 tests pass, PSScriptAnalyzer 0/0, ARG consistency and version snapshot green. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
… today's Owner rule: docs move in the same work unit as the change. These two did not, and by this afternoon both were actively wrong — rule 5 still listed the two caching levers as "documented, verified available, and NOT wired", and README still said the Windows lane "carries no --mount=type=cache today". Both were true when written and false a few hours later. Rule 5 now records what actually shipped: the two-tier sccache (SCCACHE_MULTILEVEL_CHAIN=disk,webdav plus a type=cache mount, on all seven compile RUNs including the merge builder, which is not a descendant of `common` so its ENV is repeated), the CUDA launcher on top of it, and the uv/pip wheel cache set inside the RUN rather than as ENV. It also carries the non-obvious constraint that makes the whole thing work: sccache is BUILT FROM SOURCE at SCCACHE_GIT_REV because released builds cannot wrap nvcc on CUDA 13.3 (mozilla/sccache#2722 merged five days after v0.17.0 shipped), and verify-toolchain asserts it resolves from CARGO_BIN since `--version` cannot distinguish the fixed build from the broken one. Source-fetch mounts stay listed as NOT wired, now with the measured reason (clones are shallow, so minutes against hours) so the next reader does not re-derive it. The tier-0 cap note is corrected too: it is 40GB, and it is SHARED by every cache mount plus local sources and git checkouts — the sizes and the cap are one decision. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
ROCm gives AMD cards a real build instead of falling back to Vulkan, so the backend the user picked for OCR now maps 1:1 for every option except SYCL. Adapted from upstream's windows-latest-rocm job; it needs its own job because the toolchain is entirely different (ROCm clang, Unix Makefiles, GPU_TARGETS). sccache mis-parses CUDA 13.3's 'nvcc --dryrun' output, so the per-architecture .ptx intermediates land where fatbinary does not look and the build dies with "Could not open input file 'argsort.compute_75.ptx'". Fixed upstream in mozilla/sccache#2722, merged 2026-08-04, but no release contains it yet (v0.17.0 predates it). Only the CUDA compiler opts out on that job; C/C++ still goes through sccache, and cuda12 keeps caching nvcc since 12.8 parses fine.
## 📌 Description CUDA 13.4 changes `nvcc --dryrun` output in a way that sccache v0.17.0 parses incorrectly. The missing compile steps later surface as `fatbinary` failures because the expected cubins were never produced. - Bypass sccache for cu134 NVCC invocations while keeping host C++ compilation cached. - Accept both CUDA version forms used by the release/nightly (`13.4`) and PR (`134`) build paths. - Trigger the Release dry-run matrix when the shared JIT-cache helper changes, so both cu134 architecture jobs exercise this workaround before merge. - Remove the guard once the pinned sccache release includes the upstream CUDA 13.3+ fix. ## 🔍 Related Issues - Upstream fix: [mozilla/sccache#2722](mozilla/sccache#2722) - Failing nightly: [run 32544126473](https://github.com/flashinfer-ai/flashinfer/actions/runs/32544126473) ## 🚀 Pull Request Checklist Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. ### ✅ Pre-commit Checks - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). ## 🧪 Tests - [x] Tests have been added or updated as needed. - [x] All tests are passing (`unittest`, etc.). Focused validation passed: - `bash -n scripts/jit_cache_build_common.sh` - `shellcheck scripts/jit_cache_build_common.sh` - Mocked launcher checks for CUDA `13.4`, `134`, and `13.0` - Release workflow YAML parse - `pre-commit run --all-files` - `git diff --check` ## Reviewer Notes This intentionally disables only the affected NVCC launcher for cu134. The sccache server and host C++ launcher remain enabled so safe cache hits are preserved. The Release workflow should provide the end-to-end cu134 x86_64 and aarch64 validation. <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **Bug Fixes** * Improved CUDA 13.4 build compatibility by avoiding an incompatible compiler-cache path. * Preserved compiler caching for other supported CUDA versions. * **Chores** * Updated release automation to recognize changes affecting shared build tooling. * Added clearer build logs showing which compiler-cache launchers are enabled. <!-- end of auto-generated comment: release notes by coderabbit.ai -->
ROCm gives AMD cards a real build instead of falling back to Vulkan, so the backend the user picked for OCR now maps 1:1 for every option except SYCL. Adapted from upstream's windows-latest-rocm job; it needs its own job because the toolchain is entirely different (ROCm clang, Unix Makefiles, GPU_TARGETS). sccache mis-parses CUDA 13.3's 'nvcc --dryrun' output, so the per-architecture .ptx intermediates land where fatbinary does not look and the build dies with "Could not open input file 'argsort.compute_75.ptx'". Fixed upstream in mozilla/sccache#2722, merged 2026-08-04, but no release contains it yet (v0.17.0 predates it). Only the CUDA compiler opts out on that job; C/C++ still goes through sccache, and cuda12 keeps caching nvcc since 12.8 parses fine.
## 📌 Description Enable NVCC caching for CUDA 13.4 by building sccache from pinned upstream merge commit `e9b15a35f7240a7edd1b9644583edb388c6cb5f9`, which contains the CUDA 13.3+ dry-run parser fix. - Build the checksum-verified source natively in the existing x86_64 and aarch64 manylinux builders. - Keep CUDA 12.9 and 13.0 on the official sccache v0.17.0 binaries. - Build only the S3-enabled sccache feature set with the locked dependency graph and its supported vendored OpenSSL option. - Remove AWS cache credentials from the Cargo build environment. - Default the Release workflow to the public sccache bucket and region when fork PR variables are unavailable, enabling anonymous read-only caching and exercising the patched build in PR CI. - Re-enable `FLASHINFER_NVCC_LAUNCHER=sccache` for CUDA 13.4. - Build cu134 sccache in a dedicated `Build patched sccache (<arch>)` Actions step in both Release and Nightly Release, preserving elapsed timing even when the step fails. - Pass the resulting native binary into the subsequent wheel-build container through the shared CI cache mount. - Cancel superseded Release dry-runs per PR while keeping real release dispatches independent. - Default the workflow token to contents: read, with actions: read and contents: write granted only to create-release. - Record the sccache install source and exact revision in the existing stats artifact. This is intended as a temporary bridge until an official sccache release contains the upstream parser fix. ## 🔍 Related Issues - mozilla/sccache#2722 ## 🚀 Pull Request Checklist Thank you for contributing to FlashInfer! Before we review your pull request, please make sure the following items are complete. ### ✅ Pre-commit Checks - [x] I have installed `pre-commit` by running `pip install pre-commit` (or used your preferred method). - [x] I have installed the hooks with `pre-commit install`. - [x] I have run the hooks manually with `pre-commit run --all-files` and fixed any reported issues. > If you are unsure about how to set up `pre-commit`, see [the pre-commit documentation](https://pre-commit.com/). ## 🧪 Tests - [x] Tests have been added or updated as needed. - [x] All tests are passing (`unittest`, etc.). Validation completed locally: - `pre-commit run --all-files` - `bash -n scripts/build_patched_sccache.sh scripts/jit_cache_build_common.sh scripts/build_flashinfer_jit_cache_whl.sh scripts/task_test_jit_cache_package_build_import.sh` - `shellcheck scripts/build_patched_sccache.sh scripts/jit_cache_build_common.sh scripts/build_flashinfer_jit_cache_whl.sh` - YAML parsing for both Release workflows - Mocked routing checks for CUDA 13.4 on x86_64 and aarch64 and unchanged release routing for CUDA 12.9/13.0 - Mocked launcher check confirming cu134 enables both C++ and NVCC sccache launchers - Reverified the pinned GitHub source archive SHA256 - Verified the pinned Cargo lockfile contains the vendored OpenSSL dependency - Rebased onto current `main` and reran the changed-file validations Native Linux builds remain for the PR release workflow matrix, which covers both x86_64 and aarch64. ## Reviewer Notes Please compare the dedicated patched-sccache Actions step duration and the resulting NVCC cache statistics for both cu134 architectures. The source build intentionally remains scoped to cu134; other CUDA variants continue using the published v0.17.0 artifacts. <!-- This is an auto-generated comment: release notes by coderabbit.ai --> ## Summary by CodeRabbit * **Bug Fixes** * Improved build-cache reliability for CUDA 13.4 environments. * Ensured the CUDA compiler consistently uses the configured cache. * Added build metadata identifying the cache implementation and revision. * **Chores** * Added validated, reproducible cache builds for CUDA 13.4. * Added support for released or pinned cache binaries. * Added automated verification for patched cache binaries. * Improved release workflow handling and pull request checks. * Added fallback cache storage settings when release configuration is unavailable. <!-- end of auto-generated comment: release notes by coderabbit.ai -->
sccacheassumednvcc-generatedciccandptxasinputs were always at a fixed offset near the end of the command line. With CUDA 13.3.33,ciccdryrun output can place--simt-onlyafter the input file, which makes the old positional parser treat--simt-onlyas the input instead ofkernel.cpp1.ii.That can prevent
sccachefrom grouping theciccandptxasdevice compilation steps correctly. Later,fatbinaryfails because the expected.cubinfiles were never produced.Changes
ciccinputs by.cpp1.iiextension instead of fixed position.ptxasinputs by.ptxextension instead of fixed position.ciccinput before--simt-only.ptxasinput detection by.ptx.nvccdryrun grouping coverage for CUDA 13.3-style--simt-only.