AProjQ4: imatrix-guided Q4_K dense attention projections

.gitignore

@@ -25,12 +25,18 @@
 /tests/cuda_long_context_smoke
 /tests/test_layer_pack
 /tests/test_metal_session_batch
+/tests/test_metal_indexer_q4
+/tests/test_metal_q4_qb_f16_cache
+/tests/test_metal_q4_prefill_pair
+/tests/test_metal_q4_streams
 /tests/test_mxfp4_cuda
 /tests/test_mxfp4_dot
 /tests/test_mxfp4_metal
 /tests/test_q4k_dot
 /tests/test_sampling
 /tests/test_glm53_kda
+/tests/test_quantizer_indexer_q4
+/tests/test_rocm_q4_dense_pair
 /tests/test_deepseek4_vision_image
 *.o
 *.dSYM/
@@ -41,3 +47,8 @@ __pycache__/
 /misc/
 .*.swp
 .DS_Store
+
+# Local speed benchmark outputs
+/speed-bench/*.csv
+/speed-bench/*_logprobs.json
+/speed-bench/*_logits/

CONTRIBUTING.md

@@ -25,6 +25,14 @@ The C test runner is `ds4_test`. Running it without arguments is equivalent to
 make test
 ```
 
+If a change adds, renames, removes, or changes the parsing/default/effect of an
+environment variable, update `scripts/environment_variables.tsv`, regenerate
+the complete reference, and verify that it is current:
+
+```sh
+make environment-docs
+```
+
 Useful narrower checks:
 
 ```sh

STRIXHALO.md

@@ -18,7 +18,7 @@ sudo apt-get update
 sudo apt-get install -y \
   hipcc rocminfo rocm-smi \
   libamdhip64-dev \
-  libhipblas-dev libhipblaslt-dev \
+  libhipblas-dev \
   librocblas-dev \
   librocwmma-dev \
   libhipcub-dev

cuda/mmq/ds4_ggml_stubs.cu

@@ -7,6 +7,7 @@
 // CUDA uses stream-ordered allocation. HIP retains stream-local allocations
 // because returning live MMQ scratch to its asynchronous pool is unreliable.
 
+#include "ds4_mmq.h"
 #include "common.cuh"   // pulls in ds4_ggml_stubs.h via redirect headers
 
 #if defined(GGML_USE_HIP)
@@ -23,6 +24,20 @@
 #include <mutex>
 #include <vector>
 
+/* Standalone MMQ tests do not link ds4_cuda.cu.  Full ds4 links its strong,
+ * stream-aware registry implementation over this fail-closed weak miss. */
+#if defined(__GNUC__)
+extern "C" __attribute__((weak)) int ds4_cuda_q8_fold_take_q81(
+        const void *src, uint64_t in_dim, cudaStream_t stream,
+        const void **q81) {
+    (void)src;
+    (void)in_dim;
+    (void)stream;
+    if (q81) *q81 = nullptr;
+    return 0;
+}
+#endif
+
 // ----------------------------------------------------------------------------
 // Device info singleton.
 //
@@ -133,10 +148,6 @@ extern "C" void ds4_pool_set_stream(cudaStream_t stream) {
     t_ds4_pool_stream = stream;
 }
 
-extern "C" cudaStream_t ds4_pool_get_stream(void) {
-    return t_ds4_pool_stream;
-}
-
 namespace {
 
 struct ds4_naive_pool : public ggml_cuda_pool {