metal: add TQ3_4S GPU support

.github/workflows/build-apple.yml

@@ -73,6 +73,39 @@ jobs:
           cd build
           ctest -L main -E "test-llama-archs" --verbose --timeout 900
 
+  macos-latest-arm64-no-metal:
+    runs-on: macos-latest
+
+    steps:
+      - name: Clone
+        id: checkout
+        uses: actions/checkout@v6
+
+      - name: ccache
+        uses: ggml-org/ccache-action@v1.2.21
+        with:
+          key: apple-arm64-no-metal
+          evict-old-files: 1d
+          save: ${{ github.event_name == 'push' && github.ref == 'refs/heads/master' }}
+
+      - name: Build
+        id: cmake_build
+        run: |
+          sysctl -a
+          cmake -B build \
+            -DCMAKE_BUILD_RPATH="@loader_path" \
+            -DLLAMA_FATAL_WARNINGS=ON \
+            -DLLAMA_BUILD_BORINGSSL=ON \
+            -DGGML_METAL=OFF \
+            -DGGML_RPC=ON
+          time cmake --build build --config Release -j $(sysctl -n hw.logicalcpu)
+
+      - name: Test
+        id: cmake_test
+        run: |
+          cd build
+          ctest -L main -E "test-llama-archs" --verbose --timeout 900
+
   macos-latest-x64:
     runs-on: macos-15-intel
 

docs/backend/Metal-TQ3.md

@@ -0,0 +1,83 @@
+# Running TurboQuant `TQ3_4S` on Apple Silicon (Metal)
+
+`TQ3_4S` (ggml type 46) was originally CUDA-only. This documents the Metal GPU
+support and how to run a `TQ3_4S` model on an Apple Silicon Mac.
+
+## Build
+
+```bash
+cmake -S . -B build-metal -DCMAKE_BUILD_TYPE=Release -DGGML_METAL=ON
+cmake --build build-metal --target llama-completion llama-server -j
+```
+
+### Troubleshooting: `unknown type name 'uuid_string_t'`
+
+If the build fails compiling Accelerate/CoreServices users (`ggml-cpu/vec.cpp`,
+`ggml-blas.cpp`) with `error: unknown type name 'uuid_string_t'` in the macOS SDK
+`hfs/hfs_format.h`, a Homebrew package (commonly `util-linux`) has force-linked a
+`uuid/uuid.h` into `/usr/local/include`, which clang searches before the SDK and
+which lacks `uuid_string_t`. Fix by unlinking it:
+
+```bash
+brew unlink util-linux      # keg-only; reversible with: brew link --force util-linux
+# verify it now resolves to the Xcode SDK header:
+echo '#include <uuid/uuid.h>' | clang -E -x c - | grep -m1 uuid/uuid.h
+```
+
+## Run
+
+```bash
+build-metal/bin/llama-completion \
+  -m model-TQ3_4S.gguf \
+  -ngl 99 -c 4096 \
+  -p "your prompt"
+```
+
+Notes:
+- This fork's CLI is `llama-completion` (it rejects `-no-cnv`; use it directly).
+- Some `qwen35`/`qwen36` `TQ3_4S` quants drop the MTP/`nextn` layer while still
+  declaring `nextn_predict_layers > 0`. If loading fails with
+  `missing tensor 'blk.N.attn_norm.weight'`, add:
+  `--override-kv llama.nomtp_trunk_only=bool:true`.
+
+## What is accelerated
+
+| Op | Metal path |
+| --- | --- |
+| `MUL_MAT`, decode (`ne11 == 1`) | coalesced mat-vec (`kernel_mul_mv_tq3_4s_f32`) |
+| `MUL_MAT`, batch (`ne11 > 1`, prefill / spec verify) | `simdgroup_matrix` GEMM (`kernel_mul_mm_tq3_4s_f32`) |
+| `GET_ROWS` | `kernel_get_rows_tq3_4s` |
+
+TQ3_4S dequant applies a per-32-element randomized Hadamard transform (RHT). To
+keep the weight matmul coalesced, the RHT is applied once to the activation in a
+pre-pass (`kernel_tq3_4s_rht_f32`); by RHT orthogonality the weights then need
+only a local codebook lookup. No MoE `_id` kernels yet (dense models only).
+
+Indicative throughput on an M3 Pro (150 GB/s) for a dense 27B `TQ3_4S`:
+decode ~6.4 tok/s, prefill ~32 tok/s (short) up to ~85 tok/s (longer prompts).
+Decode is fundamentally bounded by reading the weights once per token
+(~12.5 GB / 150 GB/s).
+
+## Speculative decoding (draft model)
+
+The GEMM makes batched verification cheap, so draft-model speculation works. The
+draft must share the target's tokenizer (e.g. a small Qwen3.5 model for a Qwen3.5
+target — same 248320 vocab). Tuned example:
+
+```bash
+build-metal/bin/llama-speculative-simple \
+  -m target-TQ3_4S.gguf -ngl 99 \
+  -md draft-Qwen3.5-small.gguf -ngld 99 \
+  --spec-type draft-simple \
+  --spec-draft-n-max 12 --spec-draft-n-min 3 --spec-draft-p-min 0.6 \
+  -fa on -ctk q8_0 -ctv q8_0 \
+  -c 8192 --override-kv llama.nomtp_trunk_only=bool:true \
+  -p "your prompt"
+```
+
+Caveats for hybrid (SSM / gated-delta-net) targets such as `qwen35`: the
+recurrent cache does not support partial sequence removal, so speculation falls
+back to full-state checkpoints. The gain is therefore modest and very sensitive
+to draft acceptance — on an M3 Pro, tuned speculation reaches ~7 tok/s (73%
+acceptance) vs ~6.4 plain. High absolute numbers reported elsewhere (>100 tok/s)
+are high-bandwidth GPUs / Ultra-class Macs, not an M3 Pro.

ggml/src/ggml-metal/ggml-metal-context.m

@@ -706,6 +706,9 @@ void ggml_metal_set_n_cb(ggml_metal_t ctx, int n_cb) {
 
         for (int idx = 0; idx < ggml_metal_op_n_nodes(ctx_op); ++idx) {
             const int res = ggml_metal_op_encode(ctx_op, idx);
+            if (res < 0) {
+                GGML_ABORT("metal op encoding returned invalid fused node count");
+            }
             if (res == 0) {
                 break;
             }

ggml/src/ggml-metal/ggml-metal-device.cpp

@@ -160,6 +160,21 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_get_rows(ggml_me
     return res;
 }
 
+ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_tq3_rht(ggml_metal_library_t lib) {
+    char base[256];
+    char name[256];
+
+    snprintf(base, 256, "kernel_tq3_4s_rht_f32");
+    snprintf(name, 256, "%s", base);
+
+    ggml_metal_pipeline_with_params res = ggml_metal_library_get_pipeline(lib, name);
+    if (!res.pipeline) {
+        res = ggml_metal_library_compile_pipeline(lib, base, name, nullptr);
+    }
+
+    return res;
+}
+
 ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_set_rows(ggml_metal_library_t lib, ggml_type tidx, ggml_type tdst) {
     char base[256];
     char name[256];
@@ -909,6 +924,11 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv(ggml_meta
                 nr0 = N_R0_IQ4_XS;
                 smem = 32*sizeof(float);
             } break;
+        case GGML_TYPE_TQ3_4S:
+            {
+                nsg = N_SG_TQ3_4S;
+                nr0 = N_R0_TQ3_4S;
+            } break;
         default:
             {
                 GGML_LOG_ERROR("Asserting on type %d\n", (int) tsrc0);

ggml/src/ggml-metal/ggml-metal-device.h

@@ -112,6 +112,7 @@ struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_cpy
 struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_pool_1d           (ggml_metal_library_t lib, const struct ggml_tensor * op, enum ggml_op_pool op_pool);
 struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_pool_2d           (ggml_metal_library_t lib, const struct ggml_tensor * op, enum ggml_op_pool op_pool);
 struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_get_rows          (ggml_metal_library_t lib, enum ggml_type tsrc);
+struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_tq3_rht           (ggml_metal_library_t lib);
 struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_set_rows          (ggml_metal_library_t lib, enum ggml_type tidx, enum ggml_type tdst);
 struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_diag              (ggml_metal_library_t lib, const struct ggml_tensor * op);
 struct ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_repeat            (ggml_metal_library_t lib, enum ggml_type tsrc);

ggml/src/ggml-metal/ggml-metal-device.m

@@ -1306,7 +1306,6 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te
                             default:
                                 return false;
                         }
-                    case GGML_TYPE_Q1_0:
                     case GGML_TYPE_Q4_0:
                     case GGML_TYPE_Q4_1:
                     case GGML_TYPE_Q5_0:
@@ -1333,6 +1332,10 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te
                     return false;
                 }
 
+                if (op->src[1]->type != GGML_TYPE_I32 && op->src[1]->type != GGML_TYPE_I64) {
+                    return false;
+                }
+
                 switch (op->type) {
                     case GGML_TYPE_F32:
                     case GGML_TYPE_F16:

ggml/src/ggml-metal/ggml-metal-impl.h

@@ -84,6 +84,9 @@
 #define N_R0_IQ4_XS 2
 #define N_SG_IQ4_XS 2
 
+#define N_R0_TQ3_4S 4
+#define N_SG_TQ3_4S 4
+
 // function constants offsets
 #define FC_FLASH_ATTN_EXT_PAD          100
 #define FC_FLASH_ATTN_EXT_BLK          200
@@ -931,6 +934,18 @@ typedef struct {
     uint64_t nb3;
 } ggml_metal_kargs_get_rows;
 
+// forward randomized Hadamard transform of activations (TQ3_4S mat-mul pre-pass)
+typedef struct {
+    int32_t  nb;    // number of 32-element blocks along K
+    int32_t  ne12;  // src1 dim2 (to unflatten the batch grid index)
+    uint64_t nb11;  // src1 strides (input activation)
+    uint64_t nb12;
+    uint64_t nb13;
+    uint64_t ob11;  // dst strides (contiguous transformed activation)
+    uint64_t ob12;
+    uint64_t ob13;
+} ggml_metal_kargs_tq3_rht;
+
 typedef struct {
     int32_t  nk0;
     int32_t  ne01;