diff --git a/.github/workflows/turboquant-ci.yml b/.github/workflows/turboquant-ci.yml new file mode 100644 index 000000000..bd01bff84 --- /dev/null +++ b/.github/workflows/turboquant-ci.yml @@ -0,0 +1,92 @@ +# 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 + # 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: | + 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-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: 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 // ============================================================