Step B: Gyro's trellis MMQ prompt kernels on an upstream ggml, so the packaged engine reads Gyro prompts as fast as rc2

cmake/strata_ggml_tq_overlay.cmake writes, at configure time, a copy of the pinned ggml-cuda headers (and ggml.h, quantize.cu) into
the build folder with TQ2_T / TQK6 / TQK7 routed the way IQ2_XS is: same Q8_0_16 tile layout, dot kernels and tile configs (the
fork's own configs for them are copies of IQ2_XS's), three template instances, the tile loader and tq.cuh from third_party/ggml-tq.
Only the strata_mmq target sees the copy; every other type is compiled from the untouched files. Anchors are checked: another
llama.cpp turns the overlay off with a message (FP16 prompt path). moe_mmq.cu knows the trellis block geometry itself.

The tile loader and the trellis tables are agentionai/llama.cpp's (MIT, ggml authors; Laurent Zuijdwijk), ported from
agentionai/llama.cpp 93a24e5 ggml/src/ggml-cuda/{tq.cuh,mmq-load-tiles.cuh,tq-lut-i8.cuh}.

Checked: Gyro-S and Gyro-M exact_pp 4096/20000 residual dumps IDENTICAL to rc2 (R9700 gfx1201); Q2_0 and IQ3_XXS IDENTICAL to f-split
on gfx1201 and sm_86; ctest ggml_tq_overlay.

CMakeLists.txt

@@ -76,6 +76,7 @@ option(STRATA_PREFILL_MMQ "Build the optional GGML MMQ prefill path for HIP" OFF
 # CUDA_MODULE_LOADING=EAGER, and these took ~3.8 MB of VRAM - 3 IQ3_XXS expert slots on a P100 - from every model); on
 # HIP a code object loads on first use, so they cost other models nothing there (R9700: the same slots).  Used at run
 # time with STRATA_MMQ_KQ=1 (opt-in: it changes the prompt's rounding against the FP16 path).
+option(STRATA_MMQ_TQ_OVERLAY "Gyro trellis MMQ prompt kernels on an upstream llama.cpp: an overlay of its ggml-cuda headers written at configure time (cmake/strata_ggml_tq_overlay.cmake)" ON)
 option(STRATA_MMQ_KQUANTS "Build the MMQ prompt kernels for Q4_K / Q5_K / Q5_1 / Q6_K: the experts of Unsloth's K-quant files, and on HIP the dense GGUF projections (STRATA_DENSE_MMQ)" ON)
 # Q6_K gate/up in the grouped expert kernels (Unsloth's UD-Q6_K_XL, and layer 2 of UD-Q5_K_XL; docs/QUANTS.md).  On by
 # default since MODEL-UPDATE Phase 1: one more instance per grouped kernel in iq_kernels.cu, +19 KB of sm_60 code,
@@ -1449,6 +1450,28 @@ if(STRATA_NATIVE_EXPERTS)
   if(STRATA_GGML_DIR AND EXISTS "${STRATA_GGML_DIR}/ggml/src/ggml-cuda/template-instances/mmq-instance-tqk6.cu")
     list(APPEND _strata_mmq_types tq2_t tqk6 tqk7)
     set(_strata_mmq_tq ON)
+  elseif(STRATA_MMQ_TQ_OVERLAY AND (STRATA_ENABLE_CUDA OR STRATA_ENABLE_HIP))
+    # an upstream llama.cpp: the trellis MMQ instances come from an overlay of the pinned ggml-cuda headers written at
+    # configure time (cmake/strata_ggml_tq_overlay.cmake); off with a message if the llama.cpp tree is not the pinned one
+    if(STRATA_GGML_DIR)
+      set(_strata_tq_llama "${STRATA_GGML_DIR}")
+    else()
+      set(_strata_tq_llama "${strata_llamacpp_SOURCE_DIR}")
+    endif()
+    include(${CMAKE_CURRENT_SOURCE_DIR}/cmake/strata_ggml_tq_overlay.cmake)
+    set(_strata_tq_overlay_dir "${CMAKE_BINARY_DIR}/ggml-tq-overlay")
+    strata_ggml_tq_overlay("${_strata_tq_llama}" "${_strata_tq_overlay_dir}")
+    if(STRATA_TQ_OVERLAY_OK)
+      configure_file(third_party/ggml/ggml-common-tq.h "${_strata_tq_overlay_dir}/ggml-common-tq.h" COPYONLY)
+      configure_file(src/kernels/cuda/tq_lut_i8.cuh "${_strata_tq_overlay_dir}/tq_lut_i8.cuh" COPYONLY)
+      set(_strata_mmq_tq ON)
+      set(_strata_mmq_tq_overlay ON)
+      message(STATUS "trellis MMQ (Gyro prompts): overlay on the pinned ggml-cuda in ${_strata_tq_overlay_dir}")
+      if(STRATA_BUILD_TESTS)
+        add_test(NAME ggml_tq_overlay COMMAND ${CMAKE_COMMAND} -DLLAMA=${_strata_tq_llama}
+                 -DOUT=${CMAKE_BINARY_DIR}/ggml-tq-overlay-test -P ${CMAKE_CURRENT_SOURCE_DIR}/cmake/ggml_tq_overlay_test.cmake)
+      endif()
+    endif()
   endif()
   # #296: OrcaRouter Q4_K_S's Q5_0 experts (Q4_K and Q5_1 come with STRATA_MMQ_KQUANTS on CUDA, S5's branch)
   if(STRATA_ORCA_Q4KS_MMQ)
@@ -1464,11 +1487,15 @@ if(STRATA_NATIVE_EXPERTS)
       set(_strata_gg ${strata_llamacpp_SOURCE_DIR}/ggml)
     endif()
     set(_strata_mmq ${_strata_gg}/src/ggml-cuda)
+    set(_strata_quantize_cu ${_strata_mmq}/quantize.cu)   # the overlay's copy when the trellis overlay is on: it includes mmq.cuh
+    if(_strata_mmq_tq_overlay)
+      set(_strata_quantize_cu ${_strata_tq_overlay_dir}/quantize.cu)
+    endif()
     # #136: the Q2_0 pack's prompt experts on Strata's own int8 kernels (moe_fused.cu, opt-in STRATA_PF_FUSED=1; CUDA
     # only: the HIP branch below keeps MMQ, and an sm_75 device or an sm_75-only build reports them unavailable)
     add_library(strata_mmq STATIC src/prefill/moe_mmq.cu src/prefill/moe_fused.cu src/prefill/moe_fused_iq.cu
                 src/prefill/mmq_nvfp4_w4a8.cu
-                src/prefill/ggml_cuda_host.cu ${_strata_mmq}/quantize.cu)
+                src/prefill/ggml_cuda_host.cu ${_strata_quantize_cu})
     # q8_0: the drafter's matrices (E-9); STRATA_ORCA_Q4KS_MMQ adds q5_0 (#296, in _strata_mmq_types), and
     # STRATA_MMQ_KQUANTS q4_k q5_k q5_1 q6_k, the experts of Unsloth's UD-Q4/Q5/Q6_K_XL (docs/QUANTS.md; the dispatch
     # follows eddoursul/Strata 7706c1b).  MODEL-UPDATE Phase 1: on CUDA they are a library of their own, strata_kq
@@ -1486,6 +1513,13 @@ if(STRATA_NATIVE_EXPERTS)
     foreach(_t IN LISTS _strata_mmq_cuda_types)
       target_sources(strata_mmq PRIVATE ${_strata_mmq}/template-instances/mmq-instance-${_t}.cu)
     endforeach()
+    if(_strata_mmq_tq_overlay)
+      foreach(_t tq2_t tqk6 tqk7)
+        target_sources(strata_mmq PRIVATE ${_strata_tq_overlay_dir}/template-instances/mmq-instance-${_t}.cu)
+      endforeach()
+      # first on the include path: the patched mmq.cuh / ggml.h shadow the pinned ones for this target only
+      target_include_directories(strata_mmq BEFORE PRIVATE ${_strata_tq_overlay_dir} ${_strata_tq_overlay_dir}/include)
+    endif()
     # NVFP4: W4A8 (default) in mmq_nvfp4_w4a8.cu with the plain architectures; the opt-in W4A4 needs RTX 50's FP4 MMA,
     # which exists only in the arch-specific target, so its one unit alone is built with 12x -> 12xa (the default
     # binary is unchanged and runs on sm_121; a card without that image falls back to W4A8 at run time)
@@ -1639,6 +1673,10 @@ if(STRATA_NATIVE_EXPERTS)
       set(_strata_gg "${strata_llamacpp_SOURCE_DIR}/ggml")
     endif()
     set(_strata_mmq "${_strata_gg}/src/ggml-cuda")
+    set(_strata_quantize_cu "${_strata_mmq}/quantize.cu")   # the overlay's copy when the trellis overlay is on: it includes mmq.cuh
+    if(_strata_mmq_tq_overlay)
+      set(_strata_quantize_cu "${_strata_tq_overlay_dir}/quantize.cu")
+    endif()
     if(NOT EXISTS "${_strata_mmq}/common.cuh" OR NOT EXISTS "${_strata_mmq}/quantize.cu")
       message(FATAL_ERROR "The selected GGML source is missing the ggml-cuda MMQ sources: ${_strata_mmq}")
     endif()
@@ -1648,7 +1686,7 @@ if(STRATA_NATIVE_EXPERTS)
       src/prefill/moe_fused.cu
       src/prefill/moe_fused_iq.cu
       src/prefill/ggml_cuda_host.cu
-      "${_strata_mmq}/quantize.cu")
+      "${_strata_quantize_cu}")
     if(STRATA_MMQ_KQUANTS)
       # the same K-quant instances as the CUDA branch: on HIP they serve the dense GGUF projections
       # of Gemm::native (STRATA_DENSE_MMQ=1), the mixed-quant packs' attention/PLE tensors
@@ -1657,6 +1695,11 @@ if(STRATA_NATIVE_EXPERTS)
     foreach(_t IN LISTS _strata_mmq_types)
       list(APPEND _strata_mmq_sources "${_strata_mmq}/template-instances/mmq-instance-${_t}.cu")
     endforeach()
+    if(_strata_mmq_tq_overlay)
+      foreach(_t tq2_t tqk6 tqk7)
+        list(APPEND _strata_mmq_sources "${_strata_tq_overlay_dir}/template-instances/mmq-instance-${_t}.cu")
+      endforeach()
+    endif()
     set_source_files_properties(${_strata_mmq_sources} PROPERTIES LANGUAGE HIP)
     # moe_fused.cu is Strata's own CUDA-shaped source: it takes the HIP compat headers, ggml's MMQ sources do not
     set_source_files_properties(src/prefill/moe_fused.cu src/prefill/moe_fused_iq.cu
@@ -1665,6 +1708,9 @@ if(STRATA_NATIVE_EXPERTS)
     target_include_directories(strata_mmq
       PRIVATE "${_strata_gg}/include" "${_strata_gg}/src" "${_strata_mmq}"
       PUBLIC ${CMAKE_CURRENT_SOURCE_DIR}/include)
+    if(_strata_mmq_tq_overlay)
+      target_include_directories(strata_mmq BEFORE PRIVATE ${_strata_tq_overlay_dir} ${_strata_tq_overlay_dir}/include)
+    endif()
     target_compile_definitions(strata_mmq PRIVATE GGML_USE_HIP=1)
     if(STRATA_MMQ_KQUANTS)
       target_compile_definitions(strata_mmq PRIVATE STRATA_MMQ_KQUANTS=1)

cmake/ggml_tq_overlay_test.cmake

@@ -0,0 +1,49 @@
+# cmake -DLLAMA=<llama.cpp dir> -DOUT=<scratch dir> -P cmake/ggml_tq_overlay_test.cmake
+# The trellis MMQ overlay on a pinned llama.cpp: it generates, every copied header is the pinned one except the patched
+# ones, and the patched ones only gain lines (TQ routing, three externs, config copies of IQ2_XS's).
+get_filename_component(CMAKE_SOURCE_DIR "${CMAKE_CURRENT_LIST_DIR}/.." ABSOLUTE)
+include(${CMAKE_CURRENT_LIST_DIR}/strata_ggml_tq_overlay.cmake)
+file(REMOVE_RECURSE "${OUT}")
+strata_ggml_tq_overlay("${LLAMA}" "${OUT}")
+if(NOT STRATA_TQ_OVERLAY_OK)
+  message(FATAL_ERROR "overlay did not generate on ${LLAMA}")
+endif()
+set(CU "${LLAMA}/ggml/src/ggml-cuda")
+file(GLOB _cuh "${CU}/*.cuh" "${CU}/*.h")
+set(_patched mmq.cuh mmq-load-tiles.cuh)
+foreach(_f IN LISTS _cuh)
+  get_filename_component(_n "${_f}" NAME)
+  if(_n MATCHES "^mmq-config-" OR _n IN_LIST _patched)
+    continue()
+  endif()
+  file(READ "${_f}" _a)
+  file(READ "${OUT}/${_n}" _b)
+  if(NOT _a STREQUAL _b)
+    message(FATAL_ERROR "${_n} differs from the pinned file but is not meant to")
+  endif()
+endforeach()
+file(READ "${OUT}/mmq.cuh" _m)
+foreach(_pat "case GGML_TYPE_TQ2_T:   return MMQ_DP4A_TXS_Q8_0_16;" "ggml_cuda_mmq_load_tiles_tq<type, J, fallback>,\n                    ggml_cuda_mmq_vec_dot_q8_0_16_q8_1_dp4a"
+             "extern DECL_MMQ_CASE(GGML_TYPE_TQK7);")
+  string(FIND "${_m}" "${_pat}" _p)
+  if(_p LESS 0)
+    message(FATAL_ERROR "mmq.cuh lacks: ${_pat}")
+  endif()
+endforeach()
+file(READ "${OUT}/include/ggml.h" _g)
+string(FIND "${_g}" "GGML_TYPE_TQK7    = 146," _p)
+if(_p LESS 0)
+  message(FATAL_ERROR "ggml.h lacks the trellis enumerators")
+endif()
+foreach(_c IN ITEMS ampere rdna4)
+  file(READ "${CU}/mmq-config-${_c}.cuh" _o)
+  file(READ "${OUT}/mmq-config-${_c}.cuh" _n2)
+  string(REGEX MATCHALL "CASE[(]GGML_TYPE_IQ2_XS," _x "${_o}")
+  string(REGEX MATCHALL "CASE[(]GGML_TYPE_TQK6," _y "${_n2}")
+  list(LENGTH _x _lx)
+  list(LENGTH _y _ly)
+  if(NOT _lx EQUAL _ly OR _lx EQUAL 0)
+    message(FATAL_ERROR "mmq-config-${_c}.cuh: ${_lx} IQ2_XS cases, ${_ly} TQK6 cases")
+  endif()
+endforeach()
+message(STATUS "ggml_tq_overlay_test: ok")

cmake/strata_ggml_tq_overlay.cmake

@@ -0,0 +1,137 @@
+# cmake/strata_ggml_tq_overlay.cmake - MMQ prompt kernels for the Gyro trellis types (TQ2_T, TQK6, TQK7) on an UPSTREAM llama.cpp.
+#
+# The agentionai llama.cpp fork has these types in its ggml (ggml.h enum, ggml-cuda/tq.cuh, a tile loader, MMQ configs and
+# routing, three template instances).  Upstream does not.  This script writes, at configure time, a small OVERLAY of the
+# pinned ggml-cuda headers into the build tree: copies of ggml.h, mmq.cuh, mmq-load-tiles.cuh and the mmq-config-*.cuh with
+# the trellis types added the way IQ2_XS is handled (same Q8_0_16 tile layout, same dot kernels, same configs - the fork's
+# own numbers for them are copies of IQ2_XS's), plus tq.cuh / tq-mmq-tiles.cuh from third_party/ggml-tq and three template
+# instances.  Only the strata_mmq target sees the overlay (it comes first on its include path); every other type's
+# instances are compiled from the untouched pinned files, and the additions are additive (identity gate: Q2_0, IQ3_XXS).
+#
+# Every edit is anchored on pinned text; if an anchor is missing (another llama.cpp), the script says so and the trellis
+# MMQ stays off (STRATA_TQ_OVERLAY_OK=OFF): the engine then reads trellis prompts through the FP16 path, as before.
+#
+#   strata_ggml_tq_overlay(<llama.cpp dir> <out dir>)  ->  STRATA_TQ_OVERLAY_OK (parent scope)
+
+function(strata_ggml_tq_overlay LLAMA OUT)
+  set(STRATA_TQ_OVERLAY_OK OFF PARENT_SCOPE)
+  set(_cu "${LLAMA}/ggml/src/ggml-cuda")
+  set(_inc "${LLAMA}/ggml/include")
+  set(_tq "${CMAKE_SOURCE_DIR}/third_party/ggml-tq")
+  foreach(_f "${_inc}/ggml.h" "${_cu}/mmq.cuh" "${_cu}/mmq-load-tiles.cuh" "${_cu}/common.cuh")
+    if(NOT EXISTS "${_f}")
+      message(STATUS "trellis MMQ overlay: ${_f} is missing: off")
+      return()
+    endif()
+  endforeach()
+
+  # ---- helpers: read a file with LF line ends; insert text before / after the single occurrence of an anchor ----------
+  macro(_tq_read VAR PATH)
+    file(READ "${PATH}" ${VAR})
+    string(REPLACE "\r\n" "\n" ${VAR} "${${VAR}}")
+    set_property(DIRECTORY APPEND PROPERTY CMAKE_CONFIGURE_DEPENDS "${PATH}")
+  endmacro()
+  macro(_tq_edit VAR WHERE ANCHOR TEXT)   # WHERE = BEFORE | AFTER; the anchor must occur exactly once
+    string(FIND "${${VAR}}" "${ANCHOR}" _p)
+    if(_p LESS 0)
+      set(_tq_fail "anchor not found: ${ANCHOR}")
+    else()
+      string(LENGTH "${ANCHOR}" _alen)
+      math(EXPR _after "${_p} + ${_alen}")
+      string(SUBSTRING "${${VAR}}" ${_after} -1 _tail)
+      string(FIND "${_tail}" "${ANCHOR}" _p2)
+      if(NOT _p2 LESS 0)
+        set(_tq_fail "anchor is not unique: ${ANCHOR}")
+      else()
+        string(SUBSTRING "${${VAR}}" 0 ${_p} _head)
+        if("${WHERE}" STREQUAL "AFTER")
+          set(${VAR} "${_head}${ANCHOR}${TEXT}${_tail}")
+        else()
+          set(${VAR} "${_head}${TEXT}${ANCHOR}${_tail}")
+        endif()
+      endif()
+    endif()
+  endmacro()
+  set(_tq_fail "")
+
+  file(MAKE_DIRECTORY "${OUT}/include" "${OUT}/template-instances")
+
+  # ---- the whole header set of ggml-cuda and ggml/include is copied: quote includes resolve next to the including file,
+  # so a header left in the pinned directory would pull the pinned mmq.cuh / ggml.h back in beside the patched ones -------
+  file(GLOB _all_cu "${_cu}/*.cuh" "${_cu}/*.h")
+  foreach(_f IN LISTS _all_cu)
+    get_filename_component(_n "${_f}" NAME)
+    configure_file("${_f}" "${OUT}/${_n}" COPYONLY)
+  endforeach()
+  configure_file("${_cu}/quantize.cu" "${OUT}/quantize.cu" COPYONLY)   # compiled from here: it includes mmq.cuh
+  if(IS_DIRECTORY "${_cu}/vendors")
+    file(COPY "${_cu}/vendors" DESTINATION "${OUT}")
+  endif()
+  file(GLOB _all_inc "${_inc}/*.h")
+  foreach(_f IN LISTS _all_inc)
+    get_filename_component(_n "${_f}" NAME)
+    configure_file("${_f}" "${OUT}/include/${_n}" COPYONLY)
+  endforeach()
+
+  # ---- ggml.h: the three enumerators (values above GGML_TYPE_COUNT: ggml's own tables never see them) ---------------------
+  _tq_read(_h "${_inc}/ggml.h")
+  _tq_edit(_h BEFORE "        GGML_TYPE_COUNT   = 43,"
+    "        GGML_TYPE_TQ2_T   = 144, // Agention trellis types (Strata overlay; the fork's ids)
+        GGML_TYPE_TQK6    = 145,
+        GGML_TYPE_TQK7    = 146,
+")
+  file(WRITE "${OUT}/include/ggml.h" "${_h}")
+
+  # ---- mmq.cuh: routing of the trellis types as IQ2_XS ----------------------------------------------------------------
+  _tq_read(_m "${_cu}/mmq.cuh")
+  _tq_edit(_m AFTER "        case GGML_TYPE_IQ2_XXS:\n        case GGML_TYPE_IQ2_XS:\n"
+    "        case GGML_TYPE_TQ2_T:\n        case GGML_TYPE_TQK6:\n        case GGML_TYPE_TQK7:\n")
+  _tq_edit(_m AFTER "        case GGML_TYPE_IQ2_XS:  return MMQ_DP4A_TXS_Q8_0_16;\n"
+    "        case GGML_TYPE_TQ2_T:   return MMQ_DP4A_TXS_Q8_0_16;\n        case GGML_TYPE_TQK6:    return MMQ_DP4A_TXS_Q8_0_16;\n        case GGML_TYPE_TQK7:    return MMQ_DP4A_TXS_Q8_0_16;\n")
+  _tq_edit(_m BEFORE "            case GGML_TYPE_IQ3_XXS:\n                return ggml_cuda_mmq_util_funcs(\n                    VDR_IQ3_XXS_Q8_1_MMQ,"
+    "            case GGML_TYPE_TQ2_T:\n            case GGML_TYPE_TQK6:\n            case GGML_TYPE_TQK7:\n                return ggml_cuda_mmq_util_funcs(\n                    VDR_IQ2_XS_Q8_1_MMQ,\n                    ggml_cuda_mmq_load_tiles_tq<type, J, fallback>,\n                    ggml_cuda_mmq_vec_dot_q8_0_16_q8_1_dp4a<type, J, fallback>,\n                    ggml_cuda_mmq_write_back_dp4a<type, J, fallback>);\n")
+  _tq_edit(_m BEFORE "        case GGML_TYPE_IQ3_XXS:\n            return ggml_cuda_mmq_util_funcs(\n                -1,\n                ggml_cuda_mmq_load_tiles_iq3_xxs<type, J, fallback>,"
+    "        case GGML_TYPE_TQ2_T:\n        case GGML_TYPE_TQK6:\n        case GGML_TYPE_TQK7:\n            return ggml_cuda_mmq_util_funcs(\n                -1,\n                ggml_cuda_mmq_load_tiles_tq<type, J, fallback>,\n                ggml_cuda_mmq_vec_dot_q8_0_16_q8_1_mma<type, J, fallback>,\n                ggml_cuda_mmq_write_back_mma<type, J, fallback>);\n")
+  _tq_edit(_m AFTER "extern DECL_MMQ_CASE(GGML_TYPE_IQ4_XS);\n"
+    "extern DECL_MMQ_CASE(GGML_TYPE_TQ2_T);\nextern DECL_MMQ_CASE(GGML_TYPE_TQK6);\nextern DECL_MMQ_CASE(GGML_TYPE_TQK7);\n")
+  file(WRITE "${OUT}/mmq.cuh" "${_m}")
+
+  # ---- mmq-load-tiles.cuh: the loader, at the end -------------------------------------------------------------------
+  _tq_read(_l "${_cu}/mmq-load-tiles.cuh")
+  file(WRITE "${OUT}/mmq-load-tiles.cuh" "${_l}\n#include \"tq-mmq-tiles.cuh\"\n")
+  configure_file("${_tq}/tq.cuh" "${OUT}/tq.cuh" COPYONLY)
+  configure_file("${_tq}/tq-mmq-tiles.cuh" "${OUT}/tq-mmq-tiles.cuh" COPYONLY)
+
+  # ---- mmq-config-*.cuh: every IQ2_XS CASE line gets one copy per trellis type ---------------------------------------
+  file(GLOB _cfgs "${_cu}/mmq-config-*.cuh")
+  foreach(_c ${_cfgs})
+    get_filename_component(_cn "${_c}" NAME)
+    _tq_read(_t "${_c}")
+    string(REPLACE ";" "<SEMI>" _t "${_t}")   # the CASE lines hold semicolons; a CMake list would split them
+    string(REGEX MATCHALL "[^\n]*CASE\\(GGML_TYPE_IQ2_XS,[^\n]*\n" _lines "${_t}")
+    list(REMOVE_DUPLICATES _lines)
+    foreach(_ln IN LISTS _lines)
+      set(_add "")
+      foreach(_ty TQ2_T TQK6 TQK7)
+        string(REPLACE "GGML_TYPE_IQ2_XS," "GGML_TYPE_${_ty}," _x "${_ln}")
+        string(APPEND _add "${_x}")
+      endforeach()
+      _tq_edit(_t AFTER "${_ln}" "${_add}")
+    endforeach()
+    string(REPLACE "<SEMI>" ";" _t "${_t}")
+    file(WRITE "${OUT}/${_cn}" "${_t}")
+  endforeach()
+
+  # ---- template instances -------------------------------------------------------------------------------------------
+  foreach(_ty tq2_t tqk6 tqk7)
+    string(TOUPPER "${_ty}" _TY)
+    file(WRITE "${OUT}/template-instances/mmq-instance-${_ty}.cu"
+         "// Strata trellis overlay (cmake/strata_ggml_tq_overlay.cmake)\n\n#include \"../mmq.cuh\"\n\nDECL_MMQ_CASE(GGML_TYPE_${_TY});\n")
+  endforeach()
+
+  if(_tq_fail)
+    message(STATUS "trellis MMQ overlay off: ${_tq_fail}")
+    return()
+  endif()
+  set(STRATA_TQ_OVERLAY_OK ON PARENT_SCOPE)
+endfunction()

docs/GYRO.md

@@ -40,16 +40,14 @@ ready-made engine that is older is replaced by one compiled from this source.
 gfx1201 with identical output to rc2). **Not supported:** Intel Arc (SYCL) and Apple Metal: the engines there have no
 trellis kernels; the Metal engine refuses the file at load with the reason. Images are off for these files.
 
-**Prompt speed and the ggml build.** The trellis prompt kernels (MMQ) come from agentionai's llama.cpp fork
-(MIT, `STRATA_GGML_DIR`). A build against upstream ggml works and gives the same answers, but reads prompts through the
-FP16 path, which is slower. For the fast prompt path compile with the fork's ggml:
-
-```sh
-git clone https://github.com/agentionai/llama.cpp ../agention-llama.cpp
-cmake -S . -B build -G Ninja -DCMAKE_BUILD_TYPE=Release -DSTRATA_ENABLE_CUDA=ON \
-  -DCMAKE_CUDA_ARCHITECTURES=86 -DSTRATA_GGML_DIR=$PWD/../agention-llama.cpp
-# AMD: -DSTRATA_ENABLE_HIP=ON -DSTRATA_PREFILL_MMQ=ON -DCMAKE_HIP_ARCHITECTURES=gfx1201 -DSTRATA_GGML_DIR=...
-```
+**Prompt speed and the ggml build.** The trellis prompt kernels (MMQ) are agentionai's, from their llama.cpp fork (MIT).
+Strata builds against upstream ggml, which does not have them, so it adds them itself at configure time:
+`cmake/strata_ggml_tq_overlay.cmake` writes a copy of the pinned ggml-cuda headers into the build folder with the three
+trellis types routed like IQ2_XS (the tile loader is the fork's, `third_party/ggml-tq/`); only Strata's own MMQ library sees
+the copy, every other type is compiled from the untouched files (identity gate below). The ready-made engines read Gyro
+prompts through it as fast as rc2 does, and bit for bit the same. If the pinned ggml changes and an anchor of the script
+no longer matches, the overlay switches itself off with a message and prompts take the slower FP16 path. A build against the
+fork (`-DSTRATA_GGML_DIR=<agentionai/llama.cpp>`, as rc2 does) uses the fork's own instances instead.
 
 ## Credit and licenses
 
@@ -147,21 +145,24 @@ Gyro release (2026-10-07). A = agentionai/Strata rc2 9ff02c9, B = this port.
 | --- | --- | ---: | --- | ---: | ---: |
 | agentionai fork | Gyro-S | 85.4 / 89.7 (1.051) | 5/5 | 5418 / 5313 (0.98) | 14358 / 13859 (0.97) |
 | agentionai fork | Gyro-M | 72.7 / 82.7 (1.136) | 5/5 | 5183 / 5520 (1.07) | 13421 / 14164 (1.06) |
-| upstream | Gyro-S | 85.3 / 89.6 (1.050) | 5/5 | 5425 / 6939 (1.28) | 14364 / 21677 (1.51) |
+| upstream + overlay (the ready-made engine) | Gyro-S | 85.3 / 89.6 (1.050) | 5/5 | 5413 / 5298 (0.98) | 14372 / 13854 (0.96) |
+| upstream + overlay | Gyro-M | 72.6 / 82.6 (1.138) | 5/5 | 5682 / 5524 (0.97) | 15005 / 14160 (0.94) |
+| upstream, no overlay (FP16 prompt path) | Gyro-S | 85.3 / 89.6 (1.050) | 5/5 | 5425 / 6939 (1.28) | 14364 / 21677 (1.51) |
 
 Decode is faster than rc2's because the port sits on the newer Strata (0.1.45); the same engine change makes IQ3_XXS 7% faster
 (rc2 86.9, port 93.1 tok/s, 2 pairs). Prompts of the ISTA files gained 11-16% from it too (IQ3_XXS 4K 6184 to 5480 ms, 20K
-16587 to 13864 ms); Gyro-M's prompts did not and are 5-7% slower than rc2's, Gyro-S's are 2-3% faster. The cause for Gyro-M is
-not isolated (`STRATA_GFX12_DEFAULTS=0` does not give it back). With **upstream ggml** (the ready-made engine) the trellis
+16587 to 13864 ms). Gyro-M's first run against rc2 showed its prompts 5-7% slower; in the second rc2's own Gyro-M prompts
+took 15.0 s instead of 13.4 s (the 92 GB file's page-cache state), the port's stayed at 14.2 s, and Gyro-M's prompts are 3-6%
+faster than rc2's: that difference was rc2's baseline, not a regression. Without the overlay (plain upstream ggml) the trellis
 prompt kernels are missing and prompts go through the FP16 path: 28% (4K) to 51% (20K) slower than rc2, and a different
-rounding (the answer after a 4K prompt can differ). Decode is unaffected. That is why a source build against the fork's ggml is
-the fast configuration.
+rounding. Decode is unaffected either way.
 
 Identity (the fixed gate flags, 10 prompts, exact_pp 4096 and 20000 tokens, same card for both arms):
 
 - Q2_0 and IQ3_XXS, f-split against the port (upstream ggml): 10/10 identical, residual dumps and outputs IDENTICAL.
-- Gyro-S and Gyro-M, rc2 against the port (fork ggml): 10/10 identical; exact_pp 4096 / 20000 residual dumps and outputs
-  IDENTICAL when run alone (a run beside other engines on the box once showed a different dump: not reproducible alone).
+- Gyro-S and Gyro-M, rc2 against the port with the fork's ggml and with upstream + overlay: 10/10 identical; exact_pp 4096 /
+  20000 residual dumps and outputs IDENTICAL (run alone; a run beside other engines on the box once showed a different dump).
+- Q2_0 and IQ3_XXS, f-split against the port with the overlay on, on AMD gfx1201 and on NVIDIA sm_86: IDENTICAL.
 - The pack step (`tools/iq_pack.py --compat-bf16`) writes byte-identical files to rc2's tools for Gyro-S.
 
 ## Measured (the authors', rc2)

setup.py

@@ -6601,10 +6601,6 @@ def main() -> int:
     ok(f"size: {model}")
     if MODELS[model].get("license"):                   # the ready-made packs: each size has its own license
         say(f"  Its license: {MODELS[model]['license']}")
-    if family == "gyro":                               # a tip, nothing changes (recommend, never force)
-        say("  Tip: Gyro's prompt kernels come from agentionai's llama.cpp fork. The engine setup installs reads prompts "
-            "through a slower path (measured 20-30% slower on an AMD R9700; answers speed is the same); compile with "
-            "-DSTRATA_GGML_DIR=<agentionai/llama.cpp> for the fast one (docs/GYRO.md).")
     tag = fam["tag"] + model                           # names of the pack, config and start script
     small = min(x["vram_gb"] for x in chosen)         # each card keeps its layers' KV of the whole context
     rec_ctx = 32768 if small < 14 else 65536 if small < 20 else 131072

src/prefill/moe_mmq.cu

@@ -18,6 +18,16 @@
 #include <cstring>
 
 namespace strata::prefill::mmq {
+namespace {
+// The Gyro trellis types (ids 144-146) are outside the type table of an upstream ggml (indexed by the id): their block
+// geometry is known here (128 weights; 34 / 26 / 30 bytes), everything else asks ggml.
+inline bool is_tq(int t) { return t >= 144 && t <= 146; }
+inline int64_t type_blck(int t) { return is_tq(t) ? 128 : (int64_t) ggml_blck_size((ggml_type) t); }
+inline size_t type_bytes(int t) { return t == 144 ? 34 : t == 145 ? 26 : t == 146 ? 30 : ggml_type_size((ggml_type) t); }
+inline const char* type_label(int t) {
+    return t == 144 ? "tq2_t" : t == 145 ? "tqk6" : t == 146 ? "tqk7" : ggml_type_name((ggml_type) t);
+}
+}  // namespace
 #if defined(STRATA_KQ_DSO)
 namespace kq {
 int abi_args_size() { return (int) sizeof(mmq_args); }
@@ -301,7 +311,7 @@ bool fits(int t, int64_t w_rows) {
                 said[t] = true;
                 std::fprintf(stderr, "strata: prompt kernels: llama.cpp's MMQ has no tile for %s (%lld rows) on GPU %d "
                                      "(cc %d, %zu bytes of shared memory per block): that product takes the non-MMQ path "
-                                     "(#420)\n", ggml_type_name((ggml_type) t), (long long) w_rows, id, cc, smpbo);
+                                     "(#420)\n", type_label((int) t), (long long) w_rows, id, cc, smpbo);
             }
             return false;
         }
@@ -310,7 +320,7 @@ bool fits(int t, int64_t w_rows) {
 }
 
 size_t matrix_bytes(int t, int64_t rows, int64_t cols) {
-    return (size_t) rows * (size_t) (cols / ggml_blck_size((ggml_type) t)) * ggml_type_size((ggml_type) t);
+    return (size_t) rows * (size_t) (cols / type_blck((int) t)) * type_bytes((int) t);
 }
 
 size_t q8_bytes(int64_t rows, int64_t cols) {
@@ -394,10 +404,10 @@ Context::~Context() { delete (ggml_backend_cuda_context*) ctx_; }
 void Context::run(const Product& p, void* stream) {
     if (p.n <= 0 || p.max_rows <= 0) return;
     const ggml_type t = (ggml_type) p.type;
-    const int64_t qk = ggml_blck_size(t), bpr = p.w_cols / qk;
+    const int64_t qk = type_blck((int) t), bpr = p.w_cols / qk;
     const mmq_args a = {(const char*) p.w, t, (const int*) p.xq, p.ids, p.bounds, p.dst, p.y_scale,
                         p.w_cols, p.w_rows, p.total_rows, bpr, p.total_rows, p.ld_dst,
-                        p.n, p.n, p.tab ? 0 : (int64_t) (p.expert_bytes / ggml_type_size(t)), 0, 0,
+                        p.n, p.n, p.tab ? 0 : (int64_t) (p.expert_bytes / type_bytes((int) t)), 0, 0,
                         1, 1, 0, 0, 0,
                         p.max_rows, p.max_rows};
     auto& ctx = *(ggml_backend_cuda_context*) ctx_;
@@ -411,7 +421,7 @@ void Context::run(const Product& p, void* stream) {
         mmq_x_tab tab{};
         for (int z = 0; z < p.n; ++z) tab.p[z] = (const char*) p.tab->p[z];
         if (!run_tab(&ctx, &a, &tab, s)) {
-            std::fprintf(stderr, "prefill mmq: no pointer-table MMQ for %s with %lld rows\n", ggml_type_name(t),
+            std::fprintf(stderr, "prefill mmq: no pointer-table MMQ for %s with %lld rows\n", type_label((int) t),
                          (long long) p.w_rows);
             std::exit(1);
         }
@@ -485,7 +495,7 @@ size_t overread_bytes(int t, int64_t cols) {
     // the K tile's over-read for each (fallback=false) config on any visible GPU, the largest of them: the over-read
     // is not monotonic in the tile length (20 blocks: 12 more with a 16-block tile, 4 with a 24-block one), so each
     // config's own; MMQ_ITER_K's is the floor (it is the tile when no config is known)
-    const int64_t qk = ggml_blck_size((ggml_type) t), bpn = cols / qk;
+    const int64_t qk = type_blck((int) t), bpn = cols / qk;
     auto over = [&](int64_t k_vram) -> int64_t {
         const int64_t bpi = k_vram / qk;
         return bpi > 0 ? (bpi - bpn % bpi) % bpi : 0;
@@ -497,7 +507,7 @@ size_t overread_bytes(int t, int64_t cols) {
             const ggml_cuda_mmq_config c = ggml_cuda_mmq_get_config((ggml_type) t, J, false, info.devices[id].cc);
             if (c.type != GGML_TYPE_COUNT) blocks = std::max<int64_t>(blocks, over(c.K_vram));
         }
-    return (size_t) blocks * ggml_type_size((ggml_type) t);
+    return (size_t) blocks * type_bytes((int) t);
 }
 
 bool whole_tiles(int t, int64_t w_rows, int n_experts) {

third_party/ggml-tq/tq-mmq-tiles.cuh

@@ -0,0 +1,67 @@
+// third_party/ggml-tq/tq-mmq-tiles.cuh - the MMQ tile loader of the trellis types (TQ2_T, TQK6, TQK7), appended to ggml-cuda's
+// mmq-load-tiles.cuh by cmake/strata_ggml_tq_overlay.cmake when ggml comes from upstream llama.cpp (the agentionai fork has
+// it natively).  The function is agentionai/llama.cpp's (ggml/src/ggml-cuda/mmq-load-tiles.cuh, MIT, the ggml authors,
+// written by Laurent Zuijdwijk), unchanged; the other lines of this patch only route the trellis types to it the way
+// IQ2_XS is routed (same Q8_0_16 tile layout, same dot kernels, same configs).
+#pragma once
+#include "tq.cuh"
+
+// Trellis types (TQ2_T, TQK6, TQK7): a 128-weight block is TQ_LANES lanes of 16 weights. Each lane
+// decodes its 4 trellis steps and requantizes them to int8 with one scale per 16 weights, filling
+// the Q8_0_16 tile layout (as IQ2_XS). One MMQ_ITER_K slice of a row is 2 blocks = 16 lanes.
+template <ggml_type type, int J, bool fallback> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_tq(
+        const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) {
+    // ggml_cuda_mmq_get_util_funcs names this loader for every type; only trellis types reach it.
+    constexpr ggml_type tq_type = ggml_cuda_type_is_tq(type) ? type : GGML_TYPE_TQ2_T;
+    typedef typename tq_block_info<tq_type>::block_t block_t;
+    constexpr int warp_size   = ggml_cuda_get_physical_warp_size();
+    constexpr int nwarps      = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size;
+    constexpr int I           = ggml_cuda_mmq_get_I(type, J, fallback);
+    constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback);
+
+#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
+    int   * x_qs = (int   *)  x_tile;
+    float * x_df = (float *) (x_qs + MMQ_TILE_NE_K*2);
+#else
+    constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(type, I);
+    int   * x_qs = (int   *)  x_tile;
+    float * x_df = (float *) (x_qs + txs.qs);
+#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
+
+    static_assert(MMQ_ITER_K % QK_TQK == 0, "bad MMQ_ITER_K");
+    constexpr int threads_per_row = (MMQ_ITER_K / QK_TQK) * TQ_LANES; // 16 lanes x 16 weights = 256
+    static_assert(warp_size % threads_per_row == 0, "bad warp size");
+    constexpr int nrows = warp_size / threads_per_row;
+    const int kqsx = threadIdx.x % threads_per_row;
+    const int kbx  = kqsx / TQ_LANES;
+    const int lane = kqsx % TQ_LANES;
+
+#pragma unroll
+    for (int i0 = 0; i0 < I; i0 += nwarps * nrows) {
+        int i = i0 + threadIdx.y*nrows + threadIdx.x/threads_per_row;
+
+        if (fallback) {
+            i = min(i, i_max);
+        }
+
+        const block_t * bxi = (const block_t *) x + kbx0 + i*stride + kbx;
+
+        int q[TQ_STEPS_LANE];
+        const float dq = tq_lane_to_q8<tq_type>(bxi, lane, q);
+
+#pragma unroll
+        for (int k = 0; k < TQ_STEPS_LANE; ++k) {
+#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
+            x_qs[i*sram_stride + 4*kqsx + k] = q[k];
+#else
+            x_qs[i*(2*MMQ_TILE_NE_K + 1) + 4*kqsx + k] = q[k];
+#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
+        }
+
+#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
+        x_df[i*sram_stride                             + kqsx] = dq;
+#else
+        x_df[i*(2*MMQ_TILE_NE_K*2/QI8_0) + i/(QI8_0/4) + kqsx] = dq;
+#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
+    }
+}

third_party/ggml-tq/tq.cuh

@@ -0,0 +1,108 @@
+// third_party/ggml-tq/tq.cuh - the trellis expert types (TQ2_T, TQK6, TQK7) for ggml-cuda's MMQ tile loader.
+// Derived from agentionai/llama.cpp ggml/src/ggml-cuda/tq.cuh (MIT, the ggml authors; Laurent Zuijdwijk): only what the
+// MMQ path needs (no mat-vec kernels), reading Strata's copies of the codebook tables (third_party/ggml/ggml-common-tq.h,
+// src/kernels/cuda/tq_lut_i8.cuh) instead of the fork's ggml-common.h.
+//
+// A 128-weight block is 32 trellis steps of 4 weights.  Step t has a 16-bit state s (the codebook index uses its low
+// TQ_STATE_BITS bits) and its weights are d * lut[x >> 21].{lo,hi}, d * lut[(x >> 10) & 2047].{lo,hi},
+// x = (s & TQ_STATE_MASK) * 0x9e3779b1.  A block is split into 8 lanes of 4 consecutive steps (16 weights).
+#pragma once
+
+#include "common.cuh"
+#include "ggml-common-tq.h"
+#include "tq_lut_i8.cuh"
+
+#define TQ_LUT_POINTS   2048
+#define TQ_LANES        8   // lanes per 128-weight block
+#define TQ_STEPS_LANE   4   // trellis steps per lane
+
+#define QR_TQ 4
+#define QI_TQ (QK_TQK/QR_TQ)
+static_assert(QK_TQ2_T == QK_TQK, "trellis block sizes differ");
+
+template<>
+struct ggml_cuda_type_traits<GGML_TYPE_TQ2_T> {
+    static constexpr int qk = QK_TQ2_T;
+    static constexpr int qr = QR_TQ;
+    static constexpr int qi = QI_TQ;
+    static constexpr int bs = sizeof(block_tq2_t);
+};
+
+template<>
+struct ggml_cuda_type_traits<GGML_TYPE_TQK6> {
+    static constexpr int qk = QK_TQK;
+    static constexpr int qr = QR_TQ;
+    static constexpr int qi = QI_TQ;
+    static constexpr int bs = sizeof(block_tqk6);
+};
+
+template<>
+struct ggml_cuda_type_traits<GGML_TYPE_TQK7> {
+    static constexpr int qk = QK_TQK;
+    static constexpr int qr = QR_TQ;
+    static constexpr int qi = QI_TQ;
+    static constexpr int bs = sizeof(block_tqk7);
+};
+
+static constexpr __host__ __device__ bool ggml_cuda_type_is_tq(ggml_type type) {
+    return type == GGML_TYPE_TQ2_T || type == GGML_TYPE_TQK6 || type == GGML_TYPE_TQK7;
+}
+
+template <ggml_type type> struct tq_block_info;
+template <> struct tq_block_info<GGML_TYPE_TQ2_T> { typedef block_tq2_t block_t; static constexpr int K = 8; };
+template <> struct tq_block_info<GGML_TYPE_TQK6>  { typedef block_tqk6  block_t; static constexpr int K = 6; };
+template <> struct tq_block_info<GGML_TYPE_TQK7>  { typedef block_tqk7  block_t; static constexpr int K = 7; };
+
+// States of steps 4l..4l+3 of one block. qs is the block's qs array, 2-byte aligned
+// (every trellis block has an even size and qs sits right after the fp16 scale).
+template <ggml_type type>
+static __device__ __forceinline__ void tq_states4(const uint8_t * qs, const int l, uint32_t s[TQ_STEPS_LANE]) {
+    const uint16_t * qs16 = (const uint16_t *) qs;
+    if constexpr (type == GGML_TYPE_TQ2_T) {
+        // s_t = qs[(t + 31) % 32] << 8 | qs[t]
+        const uint32_t w01  = qs16[2*l + 0];
+        const uint32_t w23  = qs16[2*l + 1];
+        const uint32_t prev = qs[(4*l + 31) & 31];
+        const uint32_t b0 = w01 & 0xFF, b1 = w01 >> 8, b2 = w23 & 0xFF, b3 = w23 >> 8;
+        s[0] = (prev << 8) | b0;
+        s[1] = (b0   << 8) | b1;
+        s[2] = (b1   << 8) | b2;
+        s[3] = (b2   << 8) | b3;
+    } else {
+        // Step t's state is the 16 stream bits from bit (31 - t)*K (circular over 32*K bits, i.e. 2*K u16 words).
+        // Steps 4l..4l+3 all lie in the four words from wb = off3 >> 4, off3 = (28 - 4l)*K.
+        constexpr int K      = tq_block_info<type>::K;
+        constexpr int nwords = 2*K;
+        const int off3 = (28 - 4*l)*K;
+        const int wb   = off3 >> 4;
+        const int r    = off3 & 15;
+        const int wi1  = wb + 1 >= nwords ? wb + 1 - nwords : wb + 1;
+        const int wi2  = wb + 2 >= nwords ? wb + 2 - nwords : wb + 2;
+        const int wi3  = wb + 3 >= nwords ? wb + 3 - nwords : wb + 3;
+        const uint64_t v = (uint64_t) qs16[wb] | ((uint64_t) qs16[wi1] << 16) |
+                           ((uint64_t) qs16[wi2] << 32) | ((uint64_t) qs16[wi3] << 48);
+#pragma unroll
+        for (int i = 0; i < TQ_STEPS_LANE; ++i) {
+            s[i] = (uint32_t) (v >> (r + (3 - i)*K)) & 0xFFFFu;
+        }
+    }
+}
+
+// MMQ (batched) path: lane l of block x as 16 int8 values (4 packed ints, weight order) with one float scale per
+// lane, i.e. weights ~= scale * q.  The int8 codebook (one global scale) is already int8: no decode to float.
+template <ggml_type type>
+static __device__ __forceinline__ float tq_lane_to_q8(const void * x, const int l, int q[TQ_STEPS_LANE]) {
+    typedef typename tq_block_info<type>::block_t block_t;
+    const block_t * b = (const block_t *) x;
+
+    uint32_t s[TQ_STEPS_LANE];
+    tq_states4<type>(b->qs, l, s);
+
+#pragma unroll
+    for (int i = 0; i < TQ_STEPS_LANE; ++i) {
+        const uint32_t xh = (s[i] & TQ_STATE_MASK) * 0x9e3779b1u;
+        q[i] = (int) ((uint32_t) tq_lut_i8[xh >> 21] | ((uint32_t) tq_lut_i8[(xh >> 10) & 2047u] << 16));
+    }
+    // Strata's block structs hold the fp16 scale as uint16_t (the fork's as ggml_half): reinterpret, do not convert
+    return __half2float(__ushort_as_half(b->d)) * TQ_LUT_I8_SCALE;
+}