perf(ci): cut CUDA build time and archive size - #65
Merged
Conversation
Numbers from the b10269-1.4.0 release run (30937369755): windows cuda-12.4 166 min, cuda-13.3 139 min, and a linux CUDA zip weighing 1339 MB against a 548 MB tar.gz of identical content. Time: - Windows never pinned CMAKE_CUDA_ARCHITECTURES, so it fell back to the ggml default: 8 targets for 12.4, 7 for 13.3 (including 121a-real, ie DGX Spark SASS, on Windows). The CUDA phase alone was 95 of those 139 minutes across 124 translation units. Both lists are now pinned, and no architecture that runs today loses support -- only 90-virtual on 12.4 and 121a-real on 13.3 are dropped, and Hopper still JITs from 80-virtual. - `set /A NINJA_JOBS=%NUMBER_OF_PROCESSORS%-1` left one of the four cores idle for the whole build. - ccache ran with the action's 500M default, which a CUDA build blows through immediately, so the cache thrashed and every run was cold. 2G for the CUDA variants, and eviction moved from 1d to 7d. - Defender scanned every object nvcc emitted; the build tree and the toolchain processes are now excluded, best-effort. - Compress-Archive is single-threaded; 7-Zip from the runner image replaces it where available. Size: - `zip -r` follows symlinks, so libcublas.so -> .so.13 -> .so.13.5.1.27 was stored three times over. `zip -ry` stores the links instead. - -compress-mode=size on the CUDA 13.3 builds (needs CTK >= 12.8, so 12.4 keeps the default). Kernel SASS is unchanged; the fatbin is decompressed once at module load, not per launch. Upstream enables this by default -- this fork disabled it in 02f9af6 as a local GCC 12 workaround, so it is passed from CI rather than by flipping the CMake gate back on. Also adds 80-virtual to the linux x64 13.3 list: cp.async and the Ampere MMA path are gated on __CUDA_ARCH__ >= 800, so A100/H100/B200 were silently JITing Turing-class kernels from the compute_75 fallback.
Putting the arch lists into matrix.cmake_flags renamed the jobs: GitHub builds a matrix job's display name from its matrix values, so the checks started reporting as windows-x64 (cuda-13.3, 13.3, -DGGML_CUDA=ON ... -DCMAKE_CUDA_ARCH...) while master's branch protection requires the old windows-x64 (cuda-13.3, 13.3, -DGGML_CUDA=ON -DGGML_CUDA_CUB_3DOT2=ON) which nothing reports anymore -- the PR hangs on "Expected" forever and cannot be merged. The matrix goes back to its original shape and the flags move into a CUDA_EXTRA variable in the Build step, where they do not touch the job name. Same flags, same build, stable check names.
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
1 participant
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.
Follow-up to #64 (this is the second commit of that branch, which landed after the merge).
Numbers from release run
30937369755: windows cuda-12.4 166 min, cuda-13.3 139 min, and a linux CUDA.zipof 1339 MB against a 548 MB.tar.gzof identical content.Time
CMAKE_CUDA_ARCHITECTURES→ ggml default: 8 targets on 12.4, 7 on 13.3 (including121a-real, i.e. DGX Spark SASS, on Windows). CUDA compile alone was 95 of the 139 minutes across 124 TUs. Both lists pinned; no architecture that runs today loses support — only90-virtualon 12.4 (Hopper still JITs from80-virtual) and121a-realon 13.3 are dropped.set /A NINJA_JOBS=%NUMBER_OF_PROCESSORS%-1idled one of four cores.Compress-Archiveis single-threaded; 7-Zip from the runner image replaces it.Size
zip -rfollows symlinks, solibcublas.so→.so.13→.so.13.5.1.27was stored three times.zip -rystores the links.-compress-mode=sizeon CUDA 13.3 builds. Kernel SASS unchanged; the fatbin decompresses once at module load, not per launch. Upstream enables this by default — this fork disabled it in02f9af612as a local GCC 12 workaround, so it is passed from CI instead of flipping the CMake gate back on. Needs CTK ≥ 12.8, so 12.4 keeps the default.Bonus fix
Adds
80-virtualto the linux x64 13.3 list.CP_ASYNC_AVAILABLEandAMPERE_MMA_AVAILABLEare gated on__CUDA_ARCH__ >= 800(common.cuh:282-292), so A100/H100/B200 were silently JITing Turing-class kernels from thecompute_75fallback — same class of gap as the RDNA/CDNA one, but silent: it ran, just slowly.