Skip to content

ciflow/trunk/190927

@tinglvv tinglvv tagged this 04 Aug 16:35
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 pytorch#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>
Assets 2
Loading