[New quant] Q3_PT

ggml/include/ggml-cpu.h

@@ -111,13 +111,14 @@ extern "C" {
     // Internal types and functions exposed for tests and benchmarks
 
     typedef void (*ggml_vec_dot_t)  (int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT x, size_t bx,
-                                       const void * GGML_RESTRICT y, size_t by, int nrc);
+                                       const void * GGML_RESTRICT y, size_t by, int nrc, const void * levels);
 
     struct ggml_type_traits_cpu {
         ggml_from_float_t        from_float;
         ggml_vec_dot_t           vec_dot;
         enum ggml_type           vec_dot_type;
-        int64_t                  nrows; // number of rows to process simultaneously
+        int64_t                  nrows;             // number of rows to process simultaneously
+        size_t                   levels_row_stride;  // bytes to add per row to get next row's quant_levels (0 = per-tensor)
     };
 
     GGML_BACKEND_API const struct ggml_type_traits_cpu * ggml_get_type_traits_cpu(enum ggml_type type);

ggml/include/ggml.h

@@ -429,7 +429,15 @@ extern "C" {
         GGML_TYPE_MXFP4   = 39, // MXFP4 (1 block)
         GGML_TYPE_NVFP4   = 40, // NVFP4 (4 blocks, E4M3 scale)
         GGML_TYPE_Q1_0    = 41,
-        GGML_TYPE_COUNT   = 42,
+        GGML_TYPE_Q3_PT   = 42, // 3.875 bpw per-tensor Lloyd-Max, 16-elem affine sub-blocks
+        GGML_TYPE_Q3_KPT  = 43, // Q3_K with learned per-tensor levels (3.4375 bpw)
+        GGML_TYPE_Q4_DPT  = 44, // IQ4_NL with learned per-tensor int8 levels (4.125 bpw)
+        GGML_TYPE_Q2_DPT  = 45, // 2-bit with learned per-tensor int8 levels (2.5 bpw)
+        GGML_TYPE_Q2_KPT  = 46, // Q2_K with learned per-tensor float levels (2.625 bpw)
+        GGML_TYPE_IQ2_TQ  = 47, // Trellis quantized with RNG codebook (2.5625 bpw)
+        GGML_TYPE_IQ3_TQ  = 48, // 3-bit with per-tensor trained grid table (3.5625 bpw)
+        GGML_TYPE_IQ1_BN  = 49, // 8D vector quantized with per-tensor trained codebook (1.5625 bpw)
+        GGML_TYPE_COUNT   = 50,
     };
 
     // precision
@@ -463,6 +471,7 @@ extern "C" {
         GGML_FTYPE_MOSTLY_IQ2_XXS = 15, // except 1d tensors
         GGML_FTYPE_MOSTLY_IQ2_XS  = 16, // except 1d tensors
         GGML_FTYPE_MOSTLY_IQ3_XXS = 17, // except 1d tensors
+        GGML_FTYPE_MOSTLY_Q3_PT  = 26, // except 1d tensors
         GGML_FTYPE_MOSTLY_IQ1_S   = 18, // except 1d tensors
         GGML_FTYPE_MOSTLY_IQ4_NL  = 19, // except 1d tensors
         GGML_FTYPE_MOSTLY_IQ3_S   = 20, // except 1d tensors
@@ -471,8 +480,11 @@ extern "C" {
         GGML_FTYPE_MOSTLY_IQ1_M   = 23, // except 1d tensors
         GGML_FTYPE_MOSTLY_BF16    = 24, // except 1d tensors
         GGML_FTYPE_MOSTLY_MXFP4   = 25, // except 1d tensors
-        GGML_FTYPE_MOSTLY_NVFP4   = 26, // except 1d tensors
-        GGML_FTYPE_MOSTLY_Q1_0    = 27, // except 1d tensors
+        GGML_FTYPE_MOSTLY_Q3_KPT  = 27, // except 1d tensors
+        GGML_FTYPE_MOSTLY_Q4_DPT  = 28, // except 1d tensors
+        GGML_FTYPE_MOSTLY_Q2_KPT  = 29, // except 1d tensors
+        GGML_FTYPE_MOSTLY_NVFP4   = 30, // except 1d tensors
+        GGML_FTYPE_MOSTLY_Q1_0    = 31, // except 1d tensors
     };
 
     // available tensor operations:
@@ -693,9 +705,8 @@ extern "C" {
 
         char name[GGML_MAX_NAME];
 
-        void * extra; // extra things e.g. for ggml-cuda.cu
-
-        char padding[8];
+        void * extra;        // extra things e.g. for ggml-cuda.cu
+        void * quant_levels; // per-tensor quantization levels (replaces char padding[8]; same size on 64-bit)
     };
 
     static const size_t GGML_TENSOR_SIZE = sizeof(struct ggml_tensor);
@@ -746,6 +757,13 @@ extern "C" {
     GGML_API double ggml_type_sizef(enum ggml_type type), // ggml_type_size()/ggml_blck_size() as float
     "use ggml_row_size() instead");
 
+    // Returns the effective bits per weight for a quantized type.
+    // For standard types this is type_size*8/blck_size.
+    // For types with per-tensor trained codebooks/grids/levels, this includes
+    // the amortized overhead of the auxiliary data stored alongside each tensor.
+    // The overhead is amortized over a reference tensor of 10M elements.
+    GGML_API double ggml_get_bpw(enum ggml_type type);
+
     GGML_API const char * ggml_type_name(enum ggml_type type);
     GGML_API const char * ggml_op_name  (enum ggml_op   op);
     GGML_API const char * ggml_op_symbol(enum ggml_op   op);
@@ -2811,7 +2829,7 @@ extern "C" {
 #        define GGML_RESTRICT restrict
 #    endif
 #endif
-    typedef void (*ggml_to_float_t)  (const void  * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
+    typedef void (*ggml_to_float_t)  (const void  * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k, const void * levels);
     typedef void (*ggml_from_float_t)(const float * GGML_RESTRICT x, void  * GGML_RESTRICT y, int64_t k);
 
     struct ggml_type_traits {
@@ -2822,6 +2840,7 @@ extern "C" {
         bool                     is_quantized;
         ggml_to_float_t          to_float;
         ggml_from_float_t        from_float_ref;
+        size_t                   levels_row_stride;  // bytes to advance quant_levels per row (0 = per-tensor)
     };
 
     GGML_API const struct ggml_type_traits * ggml_get_type_traits(enum ggml_type type);

ggml/src/CMakeLists.txt

@@ -208,6 +208,13 @@ add_library(ggml-base
             ggml-quants.h
             gguf.cpp)
 
+# Enable native SIMD for ggml-quants.c (needed for K-means training in quantization)
+include(CheckCCompilerFlag)
+check_c_compiler_flag("-march=native" GGML_COMPILER_SUPPORTS_MARCH_NATIVE)
+if (GGML_COMPILER_SUPPORTS_MARCH_NATIVE)
+    set_source_files_properties(ggml-quants.c PROPERTIES COMPILE_FLAGS "-march=native")
+endif()
+
 set_target_properties(ggml-base PROPERTIES
     VERSION ${GGML_VERSION}
     SOVERSION ${GGML_VERSION_MAJOR}

ggml/src/ggml-backend-meta.cpp

@@ -424,7 +424,7 @@ struct ggml_backend_meta_buffer_context {
     // FIXME
     // The size of the split state cache is unbounded and can theoretically grow infinitely large.
     // However, it is also expensive to build and clearing it on every rebuild in ggml_backend_meta_graph_compute is too expensive.
-    static constexpr size_t nbtc = GGML_TENSOR_SIZE - sizeof(ggml_tensor::padding);
+    static constexpr size_t nbtc = GGML_TENSOR_SIZE - sizeof(ggml_tensor::quant_levels);
     std::map<std::pair<const ggml_tensor *, bool>, std::pair<ggml_backend_meta_split_state, char[nbtc]>> split_state_cache;
 
     int debug;

ggml/src/ggml-backend.cpp

@@ -2023,6 +2023,7 @@ static struct ggml_tensor * graph_copy_dup_tensor(struct ggml_hash_set hash_set,
     dst->op = src->op;
     dst->flags = src->flags;
     memcpy(dst->op_params, src->op_params, sizeof(dst->op_params));
+    dst->quant_levels = src->quant_levels;
     ggml_set_name(dst, src->name);
 
     // copy src

ggml/src/ggml-blas/ggml-blas.cpp

@@ -1,5 +1,15 @@
 #include "ggml-impl.h"
 #include "ggml-blas.h"
+
+// Helper: compute quant_levels stride for a given row.
+// For Q2_KPT (per-block levels), stride depends on tensor width.
+static inline size_t ggml_quant_levels_stride(ggml_type type, size_t constant_stride, int64_t ne0) {
+    if (type == GGML_TYPE_Q2_KPT) {
+        return (size_t)(ne0 / 256) * 4 * sizeof(float);
+    }
+    return constant_stride;
+}
+
 #include "ggml-backend-impl.h"
 
 #include <future>
@@ -77,10 +87,11 @@ static void ggml_backend_blas_mul_mat(ggml_backend_blas_context * ctx, struct gg
                 const int min_rows_per_thread = std::max((int)(min_cols_per_thread/ne00), 1);
                 const int n_threads = std::max(std::min(ctx->n_threads, (int)(ne01/min_rows_per_thread)), 1);
 
+                const size_t lrs = ggml_quant_levels_stride(src0->type, ggml_get_type_traits(src0->type)->levels_row_stride, src0->ne[0]);
 #ifdef GGML_USE_OPENMP
                 #pragma omp parallel for num_threads(n_threads)
                 for (int64_t i01 = 0; i01 < ne01; i01++) {
-                    to_float((const char *) x + i01*nb01, wplane + i01*ne00, ne00);
+                    to_float((const char *) x + i01*nb01, wplane + i01*ne00, ne00, (const char*)src0->quant_levels + i01*lrs);
                 }
 #else
                 for (int i = 1; i < n_threads; i++) {
@@ -89,7 +100,7 @@ static void ggml_backend_blas_mul_mat(ggml_backend_blas_context * ctx, struct gg
                     if (start < end) {
                         ctx->tasks.push_back(std::async(std::launch::async, [=]() {
                             for (int64_t i01 = start; i01 < end; i01++) {
-                                to_float((const char *) x + i01*nb01, wplane + i01*ne00, ne00);
+                                to_float((const char *) x + i01*nb01, wplane + i01*ne00, ne00, (const char*)src0->quant_levels + i01*lrs);
                             }
                         }));
                     }
@@ -99,7 +110,7 @@ static void ggml_backend_blas_mul_mat(ggml_backend_blas_context * ctx, struct gg
                     const int64_t start = 0;
                     const int64_t end   = ne01/n_threads;
                     for (int64_t i01 = start; i01 < end; i01++) {
-                        to_float((const char *) x + i01*nb01, wplane + i01*ne00, ne00);
+                        to_float((const char *) x + i01*nb01, wplane + i01*ne00, ne00, (const char*)src0->quant_levels + i01*lrs);
                     }
                 }
 #endif

ggml/src/ggml-common.h

@@ -298,6 +298,7 @@ typedef struct {
 } block_q2_K;
 static_assert(sizeof(block_q2_K) == 2*sizeof(ggml_half) + QK_K/16 + QK_K/4, "wrong q2_K block size/padding");
 
+
 // 3-bit quantization
 // weight is represented as x = a * q
 // 16 blocks of 16 elements each
@@ -327,6 +328,12 @@ typedef struct {
 } block_q4_K;
 static_assert(sizeof(block_q4_K) == 2*sizeof(ggml_half) + K_SCALE_SIZE + QK_K/2, "wrong q4_K block size/padding");
 
+// Q3_KPT: Q3_K with learned per-tensor levels
+// Reuses block_q3_K structure but maps 3-bit indices through learned level table
+typedef block_q3_K block_q3_kpt;
+#define Q3KPT_N_LEVELS 8
+
+
 // 5-bit quantization
 // 8 blocks of 32 elements each
 // weight is represented as x = a * q + b
@@ -449,6 +456,115 @@ typedef struct {
 } block_iq4_xs;
 static_assert(sizeof(block_iq4_xs) == sizeof(ggml_half) + sizeof(uint16_t) + QK_K/64 + QK_K/2, "wrong iq4_xs block size/padding");
 
+// 3.875 bpw - per-tensor Lloyd-Max scalar quantization
+// 256 elements = 16 sub-blocks of 16, 8-entry level table trained per tensor
+// Layout: 2 (d) + 2 (dmin) + 24 (scales: 32x6-bit) + 96 (qs: 256x3-bit) = 124 bytes
+typedef struct {
+    ggml_half d;                  //  2 bytes: global scale for 16-elem sub-block ranges
+    ggml_half dmin;               //  2 bytes: global scale for sub-block neg_mins
+    uint8_t scales[3*QK_K/32];   // 24 bytes: 32 x 6-bit (indices 0..15 = ranges, 16..31 = neg_mins)
+    uint8_t qs[3*QK_K/8];        // 96 bytes: 256 x 3-bit Lloyd-Max level index, sequential
+} block_q3_pt;
+static_assert(sizeof(block_q3_pt) == 124, "wrong q3_pt block size");
+
+#define Q3PT_N_LEVELS 8
+
+// Q4_DPT: IQ4_NL with learned per-tensor int8 levels (4.125 bpw)
+// Block format: identical to block_iq4_nl (2 + 16 = 18 bytes per 32 elements)
+typedef block_iq4_nl block_q4_dpt;
+#define Q4DPT_N_LEVELS 16
+
+// Q2_DPT: 2-bit per-tensor Lloyd-Max scalar quantization (2.5 bpw)
+// Block format: 2 bytes (FP16 scale) + 8 bytes (2-bit indices for 32 elements) = 10 bytes per block
+// 4 learned int8 levels per tensor, optimized via Lloyd-Max k-means
+typedef struct {
+    ggml_half d;               // 2 bytes: FP16 scale (delta)
+    uint8_t qs[8];             // 8 bytes: 2-bit indices (4 values per byte, 32 elements total)
+} block_q2_dpt;
+static_assert(sizeof(block_q2_dpt) == sizeof(ggml_half) + 8, "wrong q2_dpt block size/padding");
+
+#define QK2_DPT 32
+#define Q2DPT_N_LEVELS 4
+
+// Q2_KPT: Q2_K with learned per-tensor float levels (2.625 bpw)
+// Reuses block_q2_K structure but maps 2-bit indices through learned level table
+typedef block_q2_K block_q2_kpt;
+#define Q2KPT_N_LEVELS 4
+
+// IQ2_TQ: Trellis Quantized with RNG codebook (2.0625 bpw)
+//
+// Reconstruction: y[i] = d * hash(seed, block_idx, position, trellis_state, qs_idx)
+//   where hash is a deterministic function mapping to [-1, 1]
+//   and trellis_state evolves as: next = (state + idx + 1) & 7
+//
+// Block layout (66 bytes per 256 elements):
+// IQ2_TQ: 2-bit scalar quantization with per-tensor trained asymmetric grid table
+// 32 groups of 8 elements per 256-element super-block
+//   - ggml_half d (2 bytes): super-block scale
+//   - uint8_t scales[16] (16 bytes): 32 × 4-bit grid entry index per group
+//   - uint8_t qs[64] (64 bytes): 256 × 2-bit element index within grid entry
+// recon[j] = d * IQ2TQ_GRID_SCALE * grid[group_idx][elem_idx]
+typedef struct {
+    ggml_half d;                    // Super-block scale (2 bytes)
+    uint8_t scales[QK_K/16];       // 32 × 4-bit grid entry index per group (16 bytes)
+    uint8_t qs[QK_K/4];            // 256 × 2-bit element index (64 bytes)
+} block_iq2_tq;
+static_assert(sizeof(block_iq2_tq) == 82, "wrong iq2_tq block size");
+// 2 + 16 + 64 = 82 bytes per 256 weights = 2.5625 bpw
+
+#define IQ2TQ_GROUP_SIZE 8         // Elements per group
+#define IQ2TQ_N_GROUPS   (QK_K / IQ2TQ_GROUP_SIZE)  // 32 groups per super-block
+#define IQ2TQ_GRID_SCALE 0.125f    // Grid value multiplier: recon = d * GRID_SCALE * grid_int8
+
+// IQ3_TQ: 3-bit scalar quantization with per-tensor trained asymmetric grid table (3.5625 bpw)
+// 32 groups of 8 elements per 256-element super-block
+// Each grid entry has 8 int8 levels (3 bits → 8 values per element)
+// Grid table: 16 entries × 8 int8 = 128 bytes per tensor
+// Block layout:
+//   - ggml_half d (2 bytes): super-block scale
+//   - uint8_t scales[16] (16 bytes): 32 × 4-bit grid entry index per group
+//   - uint8_t qs[96] (96 bytes): 256 × 3-bit element index within grid entry
+// recon[j] = d * IQ3TQ_GRID_SCALE * grid[group_idx][elem_idx]
+typedef struct {
+    ggml_half d;                    // Super-block scale (2 bytes)
+    uint8_t scales[QK_K/16];       // 32 × 4-bit grid entry index per group (16 bytes)
+    uint8_t qs[3*QK_K/8];          // 256 × 3-bit element index (96 bytes)
+} block_iq3_tq;
+static_assert(sizeof(block_iq3_tq) == 114, "wrong iq3_tq block size");
+// 2 + 16 + 96 = 114 bytes per 256 weights = 3.5625 bpw
+
+#define IQ3TQ_GROUP_SIZE 8         // Elements per group
+#define IQ3TQ_N_GROUPS   (QK_K / IQ3TQ_GROUP_SIZE)  // 32 groups per super-block
+#define IQ3TQ_N_LEVELS   8         // 3-bit → 8 levels per grid entry
+#define IQ3TQ_GRID_SCALE 0.125f    // Grid value multiplier
+#define IQ3TQ_GRID_SIZE  128       // 16 entries × 8 int8 = 128 bytes per tensor
+
+// IQ1_BN: 8D vector quantized with per-tensor trained 4096-entry codebook (1.5625 bpw)
+// 32 groups of 8 elements per 256-element super-block
+// Each group selects one of 4096 trained 8D vectors via 12-bit codebook index
+// Codebook: 4096 entries × 8 int8 = 32768 bytes per tensor
+// Block layout:
+//   - ggml_half d (2 bytes): super-block scale
+//   - uint8_t qs[48] (48 bytes): 32 × 12-bit codebook indices packed in pairs
+// 12-bit pair packing (groups 2k, 2k+1 → 3 bytes at qs[3k]):
+//   idx_even = qs[3k] | ((qs[3k+1] & 0x0F) << 8)
+//   idx_odd  = (qs[3k+1] >> 4) | (qs[3k+2] << 4)
+// recon[g*8+k] = d * IQ1BN_GRID_SCALE * codebook[ci][k]
+typedef struct {
+    ggml_half d;                    // Super-block scale (2 bytes)
+    uint8_t qs[3*QK_K/16];         // 32 × 12-bit codebook indices packed in pairs (48 bytes)
+} block_iq1_bn;
+static_assert(sizeof(block_iq1_bn) == 50, "wrong iq1_bn block size");
+// 2 + 48 = 50 bytes per 256 weights = 1.5625 bpw
+
+#define IQ1BN_GROUP_SIZE    8
+#define IQ1BN_N_GROUPS      (QK_K / IQ1BN_GROUP_SIZE)  // 32
+#define IQ1BN_CODEBOOK_K    4096    // number of codebook entries
+#define IQ1BN_CODEBOOK_DIM  8       // vector dimension (= group size)
+#define IQ1BN_GRID_SCALE    0.125f  // Grid value multiplier
+#define IQ1BN_CODEBOOK_SIZE (IQ1BN_CODEBOOK_K * IQ1BN_CODEBOOK_DIM)  // 32768 bytes
+#define IQ1BN_AUX_SIZE      IQ1BN_CODEBOOK_SIZE                      // 32768 bytes
+
 #endif // GGML_COMMON_DECL
 #endif // GGML_COMMON_DECL
 

ggml/src/ggml-cpu/arch-fallback.h

@@ -33,6 +33,9 @@
 #define ggml_vec_dot_iq1_m_q8_K_generic ggml_vec_dot_iq1_m_q8_K
 #define ggml_vec_dot_iq4_nl_q8_0_generic ggml_vec_dot_iq4_nl_q8_0
 #define ggml_vec_dot_iq4_xs_q8_K_generic ggml_vec_dot_iq4_xs_q8_K
+#define ggml_vec_dot_q3_pt_q8_K_generic ggml_vec_dot_q3_pt_q8_K
+#define ggml_vec_dot_q4_dpt_q8_0_generic ggml_vec_dot_q4_dpt_q8_0
+#define ggml_vec_dot_q2_dpt_q8_0_generic ggml_vec_dot_q2_dpt_q8_0
 // repack.cpp
 #define ggml_quantize_mat_q8_0_4x4_generic ggml_quantize_mat_q8_0_4x4
 #define ggml_quantize_mat_q8_0_4x8_generic ggml_quantize_mat_q8_0_4x8
@@ -71,6 +74,9 @@
 #define ggml_gemm_q8_0_4x4_q8_0_generic ggml_gemm_q8_0_4x4_q8_0
 #define ggml_gemm_q8_0_4x8_q8_0_generic ggml_gemm_q8_0_4x8_q8_0
 #elif defined(__aarch64__) || defined(__arm__) || defined(_M_ARM) || defined(_M_ARM64)
+// quants.c
+#define ggml_vec_dot_q2_dpt_q8_0_generic ggml_vec_dot_q2_dpt_q8_0
+#define ggml_vec_dot_q4_dpt_q8_0_generic ggml_vec_dot_q4_dpt_q8_0
 // repack.cpp
 #define ggml_quantize_mat_q8_K_4x4_generic ggml_quantize_mat_q8_K_4x4
 #define ggml_quantize_mat_q8_K_4x8_generic ggml_quantize_mat_q8_K_4x8
@@ -116,6 +122,8 @@
 #define ggml_vec_dot_tq1_0_q8_K_generic ggml_vec_dot_tq1_0_q8_K
 #define ggml_vec_dot_tq2_0_q8_K_generic ggml_vec_dot_tq2_0_q8_K
 #define ggml_vec_dot_iq1_m_q8_K_generic ggml_vec_dot_iq1_m_q8_K
+#define ggml_vec_dot_q2_dpt_q8_0_generic ggml_vec_dot_q2_dpt_q8_0
+#define ggml_vec_dot_q4_dpt_q8_0_generic ggml_vec_dot_q4_dpt_q8_0
 // repack.cpp
 #define ggml_quantize_mat_q8_0_4x4_generic ggml_quantize_mat_q8_0_4x4
 #define ggml_quantize_mat_q8_0_4x8_generic ggml_quantize_mat_q8_0_4x8
@@ -162,6 +170,8 @@
 #define ggml_vec_dot_mxfp4_q8_0_generic ggml_vec_dot_mxfp4_q8_0
 #define ggml_vec_dot_nvfp4_q8_0_generic ggml_vec_dot_nvfp4_q8_0
 #define ggml_vec_dot_q1_0_q8_0_generic ggml_vec_dot_q1_0_q8_0
+#define ggml_vec_dot_q4_dpt_q8_0_generic ggml_vec_dot_q4_dpt_q8_0
+#define ggml_vec_dot_q2_dpt_q8_0_generic ggml_vec_dot_q2_dpt_q8_0
 // repack.cpp
 #define ggml_quantize_mat_q8_0_4x4_generic ggml_quantize_mat_q8_0_4x4
 #define ggml_quantize_mat_q8_0_4x8_generic ggml_quantize_mat_q8_0_4x8
@@ -201,6 +211,7 @@
 #define ggml_gemm_q8_0_4x8_q8_0_generic ggml_gemm_q8_0_4x8_q8_0
 #elif defined(__riscv)
 // quants.c
+#define ggml_vec_dot_q4_dpt_q8_0_generic ggml_vec_dot_q4_dpt_q8_0
 #define ggml_vec_dot_nvfp4_q8_0_generic ggml_vec_dot_nvfp4_q8_0
 // repack.cpp
 #define ggml_quantize_mat_q8_0_4x1_generic ggml_quantize_mat_q8_0_4x1
@@ -303,6 +314,8 @@
 #define ggml_vec_dot_iq1_m_q8_K_generic ggml_vec_dot_iq1_m_q8_K
 #define ggml_vec_dot_iq4_nl_q8_0_generic ggml_vec_dot_iq4_nl_q8_0
 #define ggml_vec_dot_iq4_xs_q8_K_generic ggml_vec_dot_iq4_xs_q8_K
+#define ggml_vec_dot_q4_dpt_q8_0_generic ggml_vec_dot_q4_dpt_q8_0
+#define ggml_vec_dot_q2_dpt_q8_0_generic ggml_vec_dot_q2_dpt_q8_0
 #define ggml_vec_dot_mxfp4_q8_0_generic ggml_vec_dot_mxfp4_q8_0
 #define ggml_vec_dot_nvfp4_q8_0_generic ggml_vec_dot_nvfp4_q8_0
 #define ggml_vec_dot_q1_0_q8_0_generic ggml_vec_dot_q1_0_q8_0