From a34cd89c7480e3879fd42e49757c3b9c0affb7ed Mon Sep 17 00:00:00 2001 From: Pascal Wachowski Date: Mon, 10 Aug 2026 20:57:31 +0200 Subject: [PATCH 1/4] ci: fork-specific workflow, and include explicitly MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Two things the fork's CI surfaced. fattn-tile.cuh uses std::conditional and std::is_same five times, but neither it nor common.cuh includes — common.cuh only pulls in . HIP/clang provides it transitively, which is why the local build is green; NVCC is stricter. Exactly the class of error that only shows up on hardware I do not have. The upstream workflows assume infrastructure a fork does not have: self-hosted runners for ggml-ci, signing secrets, Apple and Windows matrices. They fail for reasons unrelated to the code, which buries real failures in noise. Disabled the ones that cannot work here (self-hosted, android/apple/cann/riscv/vulkan, release, publishing, AI review) and added turboquant-ci.yml, which checks what matters for this fork: - CPU build plus test-turboquant and test-backend-ops - a CUDA compile check of the ggml-cuda target The CUDA job is compile-only since the runners have no GPU. That is still the most valuable one here: TurboQuant is developed against HIP, and HIP/clang pulls in headers NVCC does not — the type_traits omission above would have been caught by it. Kept active: CI (which carries the ubuntu-cpu and cuda jobs), HIP quality check, sanitize, server and the python linters. Co-Authored-By: Claude Opus 5 (1M context) --- .github/workflows/turboquant-ci.yml | 90 +++++++++++++++++++++++++++++ ggml/src/ggml-cuda/turbo.cuh | 2 + 2 files changed, 92 insertions(+) create mode 100644 .github/workflows/turboquant-ci.yml diff --git a/.github/workflows/turboquant-ci.yml b/.github/workflows/turboquant-ci.yml new file mode 100644 index 000000000..d182cf391 --- /dev/null +++ b/.github/workflows/turboquant-ci.yml @@ -0,0 +1,90 @@ +# CI for this fork. +# +# The upstream workflows assume infrastructure a fork does not have (self-hosted +# runners for ggml-ci, signing secrets, Apple/Windows matrices), so they fail for +# reasons unrelated to the code. This one checks what actually matters here: +# the TurboQuant code builds on CPU and on CUDA, and its tests pass. +# +# The CUDA job is compile-only — the runners have no GPU. That is still the job +# that matters most: TurboQuant is developed against HIP, and HIP/clang pulls in +# headers NVCC does not, so CUDA-only compile errors are the realistic failure +# mode here. + +name: TurboQuant CI + +on: + push: + branches: [master, "feat/**"] + pull_request: + branches: [master] + workflow_dispatch: + +concurrency: + group: ${{ github.workflow }}-${{ github.head_ref || github.run_id }} + cancel-in-progress: true + +jobs: + cpu: + name: CPU build + TurboQuant tests + runs-on: ubuntu-latest + steps: + - uses: actions/checkout@v4 + + - name: Install dependencies + run: | + sudo apt-get update + sudo apt-get install -y build-essential cmake ccache + + - name: Configure + run: | + cmake -B build \ + -DCMAKE_BUILD_TYPE=Release \ + -DLLAMA_BUILD_TESTS=ON \ + -DLLAMA_BUILD_EXAMPLES=OFF \ + -DLLAMA_BUILD_SERVER=OFF \ + -DGGML_NATIVE=OFF + + - name: Build + run: cmake --build build --target test-turboquant test-backend-ops -j $(nproc) + + - name: TurboQuant CPU reference tests + run: ./build/bin/test-turboquant + + # The CPU backend is the reference backend, so test-backend-ops has nothing + # to compare against here. Run it anyway: it still exercises tensor setup + # and quantization, which is where two of the fixed bugs lived. + - name: Backend ops (setup + quantization paths) + run: ./build/bin/test-backend-ops -o FLASH_ATTN_EXT + + cuda: + name: CUDA compile check + runs-on: ubuntu-latest + steps: + - uses: actions/checkout@v4 + + - uses: Jimver/cuda-toolkit@v0.2.19 + id: cuda-toolkit + with: + cuda: "12.6.2" + method: network + sub-packages: '["nvcc", "cudart-dev"]' + + - name: Install dependencies + run: | + sudo apt-get update + sudo apt-get install -y build-essential cmake + + - name: Configure + run: | + cmake -B build \ + -DCMAKE_BUILD_TYPE=Release \ + -DGGML_CUDA=ON \ + -DCMAKE_CUDA_ARCHITECTURES=75 \ + -DLLAMA_BUILD_TESTS=OFF \ + -DLLAMA_BUILD_EXAMPLES=OFF \ + -DLLAMA_BUILD_SERVER=OFF + + # Only the ggml-cuda target: the TurboQuant kernels live there, and a full + # build would take far longer without checking anything additional. + - name: Build ggml-cuda + run: cmake --build build --target ggml-cuda -j $(nproc) diff --git a/ggml/src/ggml-cuda/turbo.cuh b/ggml/src/ggml-cuda/turbo.cuh index ad7c44724..724daca90 100644 --- a/ggml/src/ggml-cuda/turbo.cuh +++ b/ggml/src/ggml-cuda/turbo.cuh @@ -3,6 +3,8 @@ #include "common.cuh" #include "dequantize.cuh" +#include + // ============================================================ // TurboQuant shared building blocks // ============================================================ From 2d82cf115cff1d420e317aae91a67ca196966dcb Mon Sep 17 00:00:00 2001 From: Pascal Wachowski Date: Mon, 10 Aug 2026 21:07:16 +0200 Subject: [PATCH 2/4] fix: two build errors that only -Werror and a CDNA target expose MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Both found by the fork's HIP quality check, which builds with -Werror and GPU_TARGETS=gfx908 — neither of which my local build uses. fattn-tile.cuh: flash_attn_tile_load_tile_turbo_k computes idx / JP and idx % JP. The trailing block instantiates it with J = 0, so JP = 0. The loop never runs, but the compiler evaluates the divisions at compile time and rejects them. The static_assert guarded the condition; the code did not. Now returns early for JP == 0. ops.cpp: ggml_compute_forward_clamp switches over ggml_type without a default, so adding TURBO3_0 and TURBO4_0 to the enum made it incomplete (-Wswitch). Added both to the unsupported branch. Verified by replicating the CI job locally: -Werror, gfx908, ggml-hip target — clean. Neither error was visible in the normal build, which uses gfx1201 and no -Werror. Co-Authored-By: Claude Opus 5 (1M context) --- ggml/src/ggml-cpu/ops.cpp | 2 ++ ggml/src/ggml-cuda/fattn-tile.cuh | 9 +++++++++ 2 files changed, 11 insertions(+) diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp index d950972c8..d7e43f3b5 100644 --- a/ggml/src/ggml-cpu/ops.cpp +++ b/ggml/src/ggml-cpu/ops.cpp @@ -5584,6 +5584,8 @@ void ggml_compute_forward_clamp( case GGML_TYPE_I32: case GGML_TYPE_I64: case GGML_TYPE_F64: + case GGML_TYPE_TURBO3_0: + case GGML_TYPE_TURBO4_0: case GGML_TYPE_COUNT: { GGML_ABORT("fatal error"); diff --git a/ggml/src/ggml-cuda/fattn-tile.cuh b/ggml/src/ggml-cuda/fattn-tile.cuh index d5e9e3e10..e49766f41 100644 --- a/ggml/src/ggml-cuda/fattn-tile.cuh +++ b/ggml/src/ggml-cuda/fattn-tile.cuh @@ -513,6 +513,14 @@ static __device__ __forceinline__ void flash_attn_tile_load_tile_turbo_k( // two threads write the two halves would be a read-modify-write race. constexpr int JP = std::is_same::value ? J/2 : J; + // J == 0 is instantiated for the empty trailing block. Bail out before the + // divisions below — the compiler evaluates them at compile time and rejects + // the division by zero under -Werror even though the loop never executes. + if constexpr (JP == 0) { + GGML_UNUSED_VARS(KV, tile_KV, stride_KV_b, i_sup, elem_off); + return; + } else { + #pragma unroll for (int idx0 = 0; idx0 < I*JP; idx0 += nthreads) { const int idx = idx0 + tid; @@ -547,6 +555,7 @@ static __device__ __forceinline__ void flash_attn_tile_load_tile_turbo_k( tile_KV[i*(J + J_padding) + jp] = fetch(jp); } } + } } // Function that performs a single iteration in for the KQ matrix multiplication: From e463b5e1c239107a48ad15e8920406bd123d9cbe Mon Sep 17 00:00:00 2001 From: Pascal Wachowski Date: Mon, 10 Aug 2026 21:08:21 +0200 Subject: [PATCH 3/4] ci: add cuBLAS to the CUDA compile check The job configured with only nvcc and cudart-dev, so CMake failed at target_link_libraries: ggml-cuda links CUDA::cublas, which was not installed. Configure-time failure, so it never reached the compile step this job exists for. Co-Authored-By: Claude Opus 5 (1M context) --- .github/workflows/turboquant-ci.yml | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/.github/workflows/turboquant-ci.yml b/.github/workflows/turboquant-ci.yml index d182cf391..35d62c1a0 100644 --- a/.github/workflows/turboquant-ci.yml +++ b/.github/workflows/turboquant-ci.yml @@ -67,7 +67,8 @@ jobs: with: cuda: "12.6.2" method: network - sub-packages: '["nvcc", "cudart-dev"]' + # ggml-cuda links against cuBLAS, so nvcc and the runtime alone are not enough. + sub-packages: '["nvcc", "cudart-dev", "cublas", "cublas-dev"]' - name: Install dependencies run: | From 74acf2f452fbfdbff2d0f7a83753baa60c526b2e Mon Sep 17 00:00:00 2001 From: Pascal Wachowski Date: Mon, 10 Aug 2026 21:13:21 +0200 Subject: [PATCH 4/4] ci: install the full CUDA toolkit instead of guessing sub-packages Two attempts at a minimal sub-package list failed: first CMake could not find CUDA::cublas, then apt rejected the package names I added for it. A full install takes longer but does not depend on names I cannot verify without running the job. The CPU job of this workflow passes. Co-Authored-By: Claude Opus 5 (1M context) --- .github/workflows/turboquant-ci.yml | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/.github/workflows/turboquant-ci.yml b/.github/workflows/turboquant-ci.yml index 35d62c1a0..bd01bff84 100644 --- a/.github/workflows/turboquant-ci.yml +++ b/.github/workflows/turboquant-ci.yml @@ -67,8 +67,9 @@ jobs: with: cuda: "12.6.2" method: network - # ggml-cuda links against cuBLAS, so nvcc and the runtime alone are not enough. - sub-packages: '["nvcc", "cudart-dev", "cublas", "cublas-dev"]' + # No sub-packages filter: ggml-cuda needs cuBLAS in addition to nvcc and + # the runtime, and guessing package names cost two failed runs. A full + # install is slower but does not break on naming. - name: Install dependencies run: |