shiftadd-master-cpu.patch

diff --git a/ggml/include/ggml.h b/ggml/include/ggml.h
index 3c2e6734b..366d04c99 100644
--- a/ggml/include/ggml.h
+++ b/ggml/include/ggml.h
@@ -430,7 +430,23 @@ extern "C" {
         GGML_TYPE_NVFP4   = 40, // NVFP4 (4 blocks, E4M3 scale)
         GGML_TYPE_Q1_0    = 41,
         GGML_TYPE_Q2_0    = 42,
-        GGML_TYPE_COUNT   = 43,
+        GGML_TYPE_SA2     = 43, // ShiftAdd BCQ: 2 bit-planes, alpha group 64
+        GGML_TYPE_SA3     = 44, // ShiftAdd BCQ: 3 bit-planes, alpha group 64
+        GGML_TYPE_SA4     = 45, // ShiftAdd BCQ: 4 bit-planes, alpha group 64
+        // 46 is unused
+        GGML_TYPE_SA2_128 = 47, // ShiftAdd BCQ: 2 bit-planes, alpha group 128
+        GGML_TYPE_SA3_128 = 48, // ShiftAdd BCQ: 3 bit-planes, alpha group 128
+        GGML_TYPE_SA4_128 = 49, // ShiftAdd BCQ: 4 bit-planes, alpha group 128
+        GGML_TYPE_SA2_256 = 50, // ShiftAdd BCQ: 2 bit-planes, alpha group 256
+        GGML_TYPE_SA3_256 = 51, // ShiftAdd BCQ: 3 bit-planes, alpha group 256
+        GGML_TYPE_SA4_256 = 52, // ShiftAdd BCQ: 4 bit-planes, alpha group 256
+        GGML_TYPE_SA5     = 53, // ShiftAdd BCQ: 5 bit-planes, alpha group 64
+        GGML_TYPE_SA5_128 = 54, // ShiftAdd BCQ: 5 bit-planes, alpha group 128
+        GGML_TYPE_SA5_256 = 55, // ShiftAdd BCQ: 5 bit-planes, alpha group 256
+        GGML_TYPE_SA6     = 56, // ShiftAdd BCQ: 6 bit-planes, alpha group 64
+        GGML_TYPE_SA6_128 = 57, // ShiftAdd BCQ: 6 bit-planes, alpha group 128
+        GGML_TYPE_SA6_256 = 58, // ShiftAdd BCQ: 6 bit-planes, alpha group 256
+        GGML_TYPE_COUNT   = 59,
     };
 
     // [TAG_GGML_PREC]
diff --git a/ggml/src/ggml-common.h b/ggml/src/ggml-common.h
index 1dbbe326d..16132fbd2 100644
--- a/ggml/src/ggml-common.h
+++ b/ggml/src/ggml-common.h
@@ -191,6 +191,65 @@ typedef struct {
 } block_q2_0;
 static_assert(sizeof(block_q2_0) == sizeof(ggml_half) + QK2_0 / 4, "wrong q2_0 block size/padding");
 
+// ShiftAdd / binary-coding quantization (BCQ): w_j = sum_p d*a[g*Q+p] * (bit_{p,j} ? +1 : -1), g = j/G.
+// A mat-vec is Q sign-add passes over the activations plus one scale per (group, plane).
+#define QK_SA 256
+#define GGML_SA_BLOCK(NAME, Q, G)                                   \
+typedef struct {                                                   \
+    uint8_t   qs[(Q)*QK_SA/8]; /* bit-planes, plane-major */        \
+    uint8_t   a[(Q)*(QK_SA/(G))]; /* alpha codes, a[g*Q + p] */     \
+    uint8_t   d[2];            /* fp16 super-block scale, byte-stored so odd sizes stay unpadded */ \
+} NAME;                                                            \
+static_assert(sizeof(NAME) == (Q)*QK_SA/8 + (Q)*(QK_SA/(G)) + 2, "wrong " #NAME " size/padding");
+
+#define QI_SA (QK_SA / 32)
+#define QR_SA 1
+#define SA2_Q 2
+#define SA2_G 64
+#define SA3_Q 3
+#define SA3_G 64
+#define SA4_Q 4
+#define SA4_G 64
+#define SA2_128_Q 2
+#define SA2_128_G 128
+#define SA3_128_Q 3
+#define SA3_128_G 128
+#define SA4_128_Q 4
+#define SA4_128_G 128
+#define SA2_256_Q 2
+#define SA2_256_G 256
+#define SA3_256_Q 3
+#define SA3_256_G 256
+#define SA4_256_Q 4
+#define SA4_256_G 256
+#define SA5_Q 5
+#define SA5_G 64
+#define SA5_128_Q 5
+#define SA5_128_G 128
+#define SA5_256_Q 5
+#define SA5_256_G 256
+#define SA6_Q 6
+#define SA6_G 64
+#define SA6_128_Q 6
+#define SA6_128_G 128
+#define SA6_256_Q 6
+#define SA6_256_G 256
+GGML_SA_BLOCK(block_sa2, SA2_Q, SA2_G)
+GGML_SA_BLOCK(block_sa3, SA3_Q, SA3_G)
+GGML_SA_BLOCK(block_sa4, SA4_Q, SA4_G)
+GGML_SA_BLOCK(block_sa2_128, SA2_128_Q, SA2_128_G)
+GGML_SA_BLOCK(block_sa3_128, SA3_128_Q, SA3_128_G)
+GGML_SA_BLOCK(block_sa4_128, SA4_128_Q, SA4_128_G)
+GGML_SA_BLOCK(block_sa2_256, SA2_256_Q, SA2_256_G)
+GGML_SA_BLOCK(block_sa3_256, SA3_256_Q, SA3_256_G)
+GGML_SA_BLOCK(block_sa4_256, SA4_256_Q, SA4_256_G)
+GGML_SA_BLOCK(block_sa5, SA5_Q, SA5_G)
+GGML_SA_BLOCK(block_sa5_128, SA5_128_Q, SA5_128_G)
+GGML_SA_BLOCK(block_sa5_256, SA5_256_Q, SA5_256_G)
+GGML_SA_BLOCK(block_sa6, SA6_Q, SA6_G)
+GGML_SA_BLOCK(block_sa6_128, SA6_128_Q, SA6_128_G)
+GGML_SA_BLOCK(block_sa6_256, SA6_256_Q, SA6_256_G)
+
 #define QK4_0 32
 typedef struct {
     ggml_half d;           // delta
diff --git a/ggml/src/ggml-cpu/arch-fallback.h b/ggml/src/ggml-cpu/arch-fallback.h
index 4e59d0ae0..e956c3256 100644
--- a/ggml/src/ggml-cpu/arch-fallback.h
+++ b/ggml/src/ggml-cpu/arch-fallback.h
@@ -20,6 +20,21 @@
 #define ggml_vec_dot_q2_0_q8_0_generic ggml_vec_dot_q2_0_q8_0
 #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_sa2_q8_0_generic ggml_vec_dot_sa2_q8_0
+#define ggml_vec_dot_sa3_q8_0_generic ggml_vec_dot_sa3_q8_0
+#define ggml_vec_dot_sa4_q8_0_generic ggml_vec_dot_sa4_q8_0
+#define ggml_vec_dot_sa2_128_q8_0_generic ggml_vec_dot_sa2_128_q8_0
+#define ggml_vec_dot_sa3_128_q8_0_generic ggml_vec_dot_sa3_128_q8_0
+#define ggml_vec_dot_sa4_128_q8_0_generic ggml_vec_dot_sa4_128_q8_0
+#define ggml_vec_dot_sa2_256_q8_0_generic ggml_vec_dot_sa2_256_q8_0
+#define ggml_vec_dot_sa3_256_q8_0_generic ggml_vec_dot_sa3_256_q8_0
+#define ggml_vec_dot_sa4_256_q8_0_generic ggml_vec_dot_sa4_256_q8_0
+#define ggml_vec_dot_sa5_q8_0_generic ggml_vec_dot_sa5_q8_0
+#define ggml_vec_dot_sa5_128_q8_0_generic ggml_vec_dot_sa5_128_q8_0
+#define ggml_vec_dot_sa5_256_q8_0_generic ggml_vec_dot_sa5_256_q8_0
+#define ggml_vec_dot_sa6_q8_0_generic ggml_vec_dot_sa6_q8_0
+#define ggml_vec_dot_sa6_128_q8_0_generic ggml_vec_dot_sa6_128_q8_0
+#define ggml_vec_dot_sa6_256_q8_0_generic ggml_vec_dot_sa6_256_q8_0
 #define ggml_vec_dot_q2_K_q8_K_generic ggml_vec_dot_q2_K_q8_K
 #define ggml_vec_dot_q3_K_q8_K_generic ggml_vec_dot_q3_K_q8_K
 #define ggml_vec_dot_q4_K_q8_K_generic ggml_vec_dot_q4_K_q8_K
@@ -124,6 +139,21 @@
 #define ggml_vec_dot_q2_0_q8_0_generic ggml_vec_dot_q2_0_q8_0
 #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_sa2_q8_0_generic ggml_vec_dot_sa2_q8_0
+#define ggml_vec_dot_sa3_q8_0_generic ggml_vec_dot_sa3_q8_0
+#define ggml_vec_dot_sa4_q8_0_generic ggml_vec_dot_sa4_q8_0
+#define ggml_vec_dot_sa2_128_q8_0_generic ggml_vec_dot_sa2_128_q8_0
+#define ggml_vec_dot_sa3_128_q8_0_generic ggml_vec_dot_sa3_128_q8_0
+#define ggml_vec_dot_sa4_128_q8_0_generic ggml_vec_dot_sa4_128_q8_0
+#define ggml_vec_dot_sa2_256_q8_0_generic ggml_vec_dot_sa2_256_q8_0
+#define ggml_vec_dot_sa3_256_q8_0_generic ggml_vec_dot_sa3_256_q8_0
+#define ggml_vec_dot_sa4_256_q8_0_generic ggml_vec_dot_sa4_256_q8_0
+#define ggml_vec_dot_sa5_q8_0_generic ggml_vec_dot_sa5_q8_0
+#define ggml_vec_dot_sa5_128_q8_0_generic ggml_vec_dot_sa5_128_q8_0
+#define ggml_vec_dot_sa5_256_q8_0_generic ggml_vec_dot_sa5_256_q8_0
+#define ggml_vec_dot_sa6_q8_0_generic ggml_vec_dot_sa6_q8_0
+#define ggml_vec_dot_sa6_128_q8_0_generic ggml_vec_dot_sa6_128_q8_0
+#define ggml_vec_dot_sa6_256_q8_0_generic ggml_vec_dot_sa6_256_q8_0
 #define ggml_vec_dot_iq1_m_q8_K_generic ggml_vec_dot_iq1_m_q8_K
 // repack.cpp
 #define ggml_quantize_mat_q8_0_4x4_generic ggml_quantize_mat_q8_0_4x4
@@ -171,6 +201,21 @@
 #define quantize_row_q8_K_generic quantize_row_q8_K
 #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_sa2_q8_0_generic ggml_vec_dot_sa2_q8_0
+#define ggml_vec_dot_sa3_q8_0_generic ggml_vec_dot_sa3_q8_0
+#define ggml_vec_dot_sa4_q8_0_generic ggml_vec_dot_sa4_q8_0
+#define ggml_vec_dot_sa2_128_q8_0_generic ggml_vec_dot_sa2_128_q8_0
+#define ggml_vec_dot_sa3_128_q8_0_generic ggml_vec_dot_sa3_128_q8_0
+#define ggml_vec_dot_sa4_128_q8_0_generic ggml_vec_dot_sa4_128_q8_0
+#define ggml_vec_dot_sa2_256_q8_0_generic ggml_vec_dot_sa2_256_q8_0
+#define ggml_vec_dot_sa3_256_q8_0_generic ggml_vec_dot_sa3_256_q8_0
+#define ggml_vec_dot_sa4_256_q8_0_generic ggml_vec_dot_sa4_256_q8_0
+#define ggml_vec_dot_sa5_q8_0_generic ggml_vec_dot_sa5_q8_0
+#define ggml_vec_dot_sa5_128_q8_0_generic ggml_vec_dot_sa5_128_q8_0
+#define ggml_vec_dot_sa5_256_q8_0_generic ggml_vec_dot_sa5_256_q8_0
+#define ggml_vec_dot_sa6_q8_0_generic ggml_vec_dot_sa6_q8_0
+#define ggml_vec_dot_sa6_128_q8_0_generic ggml_vec_dot_sa6_128_q8_0
+#define ggml_vec_dot_sa6_256_q8_0_generic ggml_vec_dot_sa6_256_q8_0
 #define ggml_vec_dot_iq1_m_q8_K_generic ggml_vec_dot_iq1_m_q8_K
 #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
@@ -218,6 +263,21 @@
 #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(__riscv)
+#define ggml_vec_dot_sa2_q8_0_generic ggml_vec_dot_sa2_q8_0
+#define ggml_vec_dot_sa3_q8_0_generic ggml_vec_dot_sa3_q8_0
+#define ggml_vec_dot_sa4_q8_0_generic ggml_vec_dot_sa4_q8_0
+#define ggml_vec_dot_sa2_128_q8_0_generic ggml_vec_dot_sa2_128_q8_0
+#define ggml_vec_dot_sa3_128_q8_0_generic ggml_vec_dot_sa3_128_q8_0
+#define ggml_vec_dot_sa4_128_q8_0_generic ggml_vec_dot_sa4_128_q8_0
+#define ggml_vec_dot_sa2_256_q8_0_generic ggml_vec_dot_sa2_256_q8_0
+#define ggml_vec_dot_sa3_256_q8_0_generic ggml_vec_dot_sa3_256_q8_0
+#define ggml_vec_dot_sa4_256_q8_0_generic ggml_vec_dot_sa4_256_q8_0
+#define ggml_vec_dot_sa5_q8_0_generic ggml_vec_dot_sa5_q8_0
+#define ggml_vec_dot_sa5_128_q8_0_generic ggml_vec_dot_sa5_128_q8_0
+#define ggml_vec_dot_sa5_256_q8_0_generic ggml_vec_dot_sa5_256_q8_0
+#define ggml_vec_dot_sa6_q8_0_generic ggml_vec_dot_sa6_q8_0
+#define ggml_vec_dot_sa6_128_q8_0_generic ggml_vec_dot_sa6_128_q8_0
+#define ggml_vec_dot_sa6_256_q8_0_generic ggml_vec_dot_sa6_256_q8_0
 // quants.c
 #define ggml_vec_dot_nvfp4_q8_0_generic ggml_vec_dot_nvfp4_q8_0
 #define ggml_vec_dot_q2_0_q8_0_generic ggml_vec_dot_q2_0_q8_0
@@ -268,6 +328,21 @@
 #define ggml_vec_dot_q2_0_q8_0_generic ggml_vec_dot_q2_0_q8_0
 #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_sa2_q8_0_generic ggml_vec_dot_sa2_q8_0
+#define ggml_vec_dot_sa3_q8_0_generic ggml_vec_dot_sa3_q8_0
+#define ggml_vec_dot_sa4_q8_0_generic ggml_vec_dot_sa4_q8_0
+#define ggml_vec_dot_sa2_128_q8_0_generic ggml_vec_dot_sa2_128_q8_0
+#define ggml_vec_dot_sa3_128_q8_0_generic ggml_vec_dot_sa3_128_q8_0
+#define ggml_vec_dot_sa4_128_q8_0_generic ggml_vec_dot_sa4_128_q8_0
+#define ggml_vec_dot_sa2_256_q8_0_generic ggml_vec_dot_sa2_256_q8_0
+#define ggml_vec_dot_sa3_256_q8_0_generic ggml_vec_dot_sa3_256_q8_0
+#define ggml_vec_dot_sa4_256_q8_0_generic ggml_vec_dot_sa4_256_q8_0
+#define ggml_vec_dot_sa5_q8_0_generic ggml_vec_dot_sa5_q8_0
+#define ggml_vec_dot_sa5_128_q8_0_generic ggml_vec_dot_sa5_128_q8_0
+#define ggml_vec_dot_sa5_256_q8_0_generic ggml_vec_dot_sa5_256_q8_0
+#define ggml_vec_dot_sa6_q8_0_generic ggml_vec_dot_sa6_q8_0
+#define ggml_vec_dot_sa6_128_q8_0_generic ggml_vec_dot_sa6_128_q8_0
+#define ggml_vec_dot_sa6_256_q8_0_generic ggml_vec_dot_sa6_256_q8_0
 #define ggml_vec_dot_q2_K_q8_K_generic ggml_vec_dot_q2_K_q8_K
 #define ggml_vec_dot_iq2_xxs_q8_K_generic ggml_vec_dot_iq2_xxs_q8_K
 #define ggml_vec_dot_iq2_xs_q8_K_generic ggml_vec_dot_iq2_xs_q8_K
@@ -318,6 +393,21 @@
 // quants.c
 #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_sa2_q8_0_generic ggml_vec_dot_sa2_q8_0
+#define ggml_vec_dot_sa3_q8_0_generic ggml_vec_dot_sa3_q8_0
+#define ggml_vec_dot_sa4_q8_0_generic ggml_vec_dot_sa4_q8_0
+#define ggml_vec_dot_sa2_128_q8_0_generic ggml_vec_dot_sa2_128_q8_0
+#define ggml_vec_dot_sa3_128_q8_0_generic ggml_vec_dot_sa3_128_q8_0
+#define ggml_vec_dot_sa4_128_q8_0_generic ggml_vec_dot_sa4_128_q8_0
+#define ggml_vec_dot_sa2_256_q8_0_generic ggml_vec_dot_sa2_256_q8_0
+#define ggml_vec_dot_sa3_256_q8_0_generic ggml_vec_dot_sa3_256_q8_0
+#define ggml_vec_dot_sa4_256_q8_0_generic ggml_vec_dot_sa4_256_q8_0
+#define ggml_vec_dot_sa5_q8_0_generic ggml_vec_dot_sa5_q8_0
+#define ggml_vec_dot_sa5_128_q8_0_generic ggml_vec_dot_sa5_128_q8_0
+#define ggml_vec_dot_sa5_256_q8_0_generic ggml_vec_dot_sa5_256_q8_0
+#define ggml_vec_dot_sa6_q8_0_generic ggml_vec_dot_sa6_q8_0
+#define ggml_vec_dot_sa6_128_q8_0_generic ggml_vec_dot_sa6_128_q8_0
+#define ggml_vec_dot_sa6_256_q8_0_generic ggml_vec_dot_sa6_256_q8_0
 #define ggml_vec_dot_iq2_xxs_q8_K_generic ggml_vec_dot_iq2_xxs_q8_K
 #define ggml_vec_dot_iq2_xs_q8_K_generic ggml_vec_dot_iq2_xs_q8_K
 #define ggml_vec_dot_iq2_s_q8_K_generic ggml_vec_dot_iq2_s_q8_K
diff --git a/ggml/src/ggml-cpu/arch/arm/quants.c b/ggml/src/ggml-cpu/arch/arm/quants.c
index cfbbf6617..dc6938344 100644
--- a/ggml/src/ggml-cpu/arch/arm/quants.c
+++ b/ggml/src/ggml-cpu/arch/arm/quants.c
@@ -1682,6 +1682,84 @@ void ggml_vec_dot_tq2_0_q8_K(int n, float * GGML_RESTRICT s, size_t bs, const vo
 #endif
 }
 
+// ShiftAdd / BCQ: sign-select the q8 activations with each bit-plane, sum, and weight by the plane scale code
+#if defined(__ARM_NEON) && defined(__aarch64__)
+static inline float ggml_vec_dot_sa_q8_0_neon(int n, const uint8_t * GGML_RESTRICT x, const block_q8_0 * GGML_RESTRICT y, int Q, int G) {
+    const int ng = QK_SA / G;
+    const size_t bsz = Q*QK_SA/8 + Q*ng + sizeof(ggml_half);
+    const int nb = n / QK_SA;
+
+    // byte j of a 32-weight chunk takes bit j%8 of sign byte j/8
+    static const uint8_t shuf_lo[16] = {0,0,0,0,0,0,0,0, 1,1,1,1,1,1,1,1};
+    static const uint8_t shuf_hi[16] = {2,2,2,2,2,2,2,2, 3,3,3,3,3,3,3,3};
+    static const uint8_t bits_tab[16] = {1,2,4,8,16,32,64,128, 1,2,4,8,16,32,64,128};
+    const uint8x16_t slo  = vld1q_u8(shuf_lo);
+    const uint8x16_t shi  = vld1q_u8(shuf_hi);
+    const uint8x16_t bits = vld1q_u8(bits_tab);
+
+    float sumf = 0.0f;
+    for (int i = 0; i < nb; ++i) {
+        const uint8_t * qs = x + i*bsz;
+        const uint8_t * a  = qs + Q*QK_SA/8;
+        ggml_half dh; memcpy(&dh, a + Q*ng, sizeof(dh));
+        const float d = GGML_CPU_FP16_TO_FP32(dh);
+
+        float32x4_t acc = vdupq_n_f32(0.0f);
+        for (int c = 0; c < QK_SA/32; ++c) {
+            const block_q8_0 * yb = &y[i*(QK_SA/32) + c];
+            const int8x16_t y0 = vld1q_s8(yb->qs);
+            const int8x16_t y1 = vld1q_s8(yb->qs + 16);
+            const int8x16_t n0 = vnegq_s8(y0);
+            const int8x16_t n1 = vnegq_s8(y1);
+            const int g = c*32/G;
+
+            int32x4_t acci = vdupq_n_s32(0);
+            for (int p = 0; p < Q; ++p) {
+                uint32_t w; memcpy(&w, qs + p*32 + c*4, sizeof(w));
+                const uint8x16_t wv = vreinterpretq_u8_u32(vdupq_n_u32(w));
+                const uint8x16_t m0 = vceqq_u8(vandq_u8(vqtbl1q_u8(wv, slo), bits), bits);
+                const uint8x16_t m1 = vceqq_u8(vandq_u8(vqtbl1q_u8(wv, shi), bits), bits);
+                const int8x16_t  v0 = vbslq_s8(m0, y0, n0);
+                const int8x16_t  v1 = vbslq_s8(m1, y1, n1);
+                const int16x8_t s16 = vaddq_s16(vpaddlq_s8(v0), vpaddlq_s8(v1));
+                acci = vmlaq_n_s32(acci, vpaddlq_s16(s16), (int32_t) a[g*Q + p]);
+            }
+            acc = vmlaq_n_f32(acc, vcvtq_f32_s32(acci), GGML_CPU_FP16_TO_FP32(yb->d));
+        }
+        sumf += d * vaddvq_f32(acc);
+    }
+    return sumf;
+}
+
+#define GGML_SA_VEC_DOT(name, Q, G)                                                                                                                              \
+void ggml_vec_dot_##name##_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) { \
+    assert(nrc == 1);                                                                                                                                            \
+    UNUSED(nrc); UNUSED(bx); UNUSED(by); UNUSED(bs);                                                                                                             \
+    *s = ggml_vec_dot_sa_q8_0_neon(n, (const uint8_t *) vx, (const block_q8_0 *) vy, Q, G);                                                                      \
+}
+#else
+#define GGML_SA_VEC_DOT(name, Q, G)                                                                                                                              \
+void ggml_vec_dot_##name##_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) { \
+    ggml_vec_dot_##name##_q8_0_generic(n, s, bs, vx, bx, vy, by, nrc);                                                                                           \
+}
+#endif
+
+GGML_SA_VEC_DOT(sa2, SA2_Q, SA2_G)
+GGML_SA_VEC_DOT(sa3, SA3_Q, SA3_G)
+GGML_SA_VEC_DOT(sa4, SA4_Q, SA4_G)
+GGML_SA_VEC_DOT(sa2_128, SA2_128_Q, SA2_128_G)
+GGML_SA_VEC_DOT(sa3_128, SA3_128_Q, SA3_128_G)
+GGML_SA_VEC_DOT(sa4_128, SA4_128_Q, SA4_128_G)
+GGML_SA_VEC_DOT(sa2_256, SA2_256_Q, SA2_256_G)
+GGML_SA_VEC_DOT(sa3_256, SA3_256_Q, SA3_256_G)
+GGML_SA_VEC_DOT(sa4_256, SA4_256_Q, SA4_256_G)
+GGML_SA_VEC_DOT(sa5, SA5_Q, SA5_G)
+GGML_SA_VEC_DOT(sa5_128, SA5_128_Q, SA5_128_G)
+GGML_SA_VEC_DOT(sa5_256, SA5_256_Q, SA5_256_G)
+GGML_SA_VEC_DOT(sa6, SA6_Q, SA6_G)
+GGML_SA_VEC_DOT(sa6_128, SA6_128_Q, SA6_128_G)
+GGML_SA_VEC_DOT(sa6_256, SA6_256_Q, SA6_256_G)
+
 void ggml_vec_dot_q2_K_q8_K(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) {
     assert(nrc == 1);
     UNUSED(nrc);
diff --git a/ggml/src/ggml-cpu/arch/x86/quants.c b/ggml/src/ggml-cpu/arch/x86/quants.c
index 21cd6312a..74499f8eb 100644
--- a/ggml/src/ggml-cpu/arch/x86/quants.c
+++ b/ggml/src/ggml-cpu/arch/x86/quants.c
@@ -1635,6 +1635,126 @@ void ggml_vec_dot_tq2_0_q8_K(int n, float * GGML_RESTRICT s, size_t bs, const vo
 #endif
 }
 
+// ShiftAdd / BCQ: per 32-weight chunk and bit-plane, sum_j s_j*y_j = 2*sum_{bit=1} y_j - sum_j y_j, weighted by the plane scale code
+#if defined(__AVX512BW__) && defined(__AVX512VNNI__)
+static inline float ggml_vec_dot_sa_q8_0_avx512(int n, const uint8_t * GGML_RESTRICT x, const block_q8_0 * GGML_RESTRICT y, int Q, int G) {
+    // 64 weights (two q8 chunks) per op: the sign bits form a 64-bit mask that zero-fills the broadcast plane scale code
+    const int ng = QK_SA / G;
+    const size_t bsz = Q*QK_SA/8 + Q*ng + sizeof(ggml_half);
+    const int nb = n / QK_SA;
+    const __m512i ones = _mm512_set1_epi8(1);
+
+    __m256 sumf = _mm256_setzero_ps();
+    for (int i = 0; i < nb; ++i) {
+        const uint8_t * qs = x + i*bsz;
+        const uint8_t * a  = qs + Q*QK_SA/8;
+        ggml_half dh; memcpy(&dh, a + Q*ng, sizeof(dh));
+        const float d = GGML_CPU_FP16_TO_FP32(dh);
+
+        __m256 acc = _mm256_setzero_ps();
+        for (int c = 0; c < QK_SA/32; c += 2) {
+            const block_q8_0 * yb = &y[i*(QK_SA/32) + c];
+            const __m512i yv = _mm512_inserti64x4(_mm512_castsi256_si512(_mm256_loadu_si256((const __m256i *) yb[0].qs)),
+                                                  _mm256_loadu_si256((const __m256i *) yb[1].qs), 1);
+            const uint8_t * ag = a + (c*32/G)*Q;
+
+            __m512i pos = _mm512_setzero_si512();
+            int asum = 0;
+            for (int p = 0; p < Q; ++p) {
+                uint64_t w; memcpy(&w, qs + p*32 + c*4, sizeof(w));
+                pos = _mm512_dpbusd_epi32(pos, _mm512_maskz_set1_epi8(_cvtu64_mask64(w), (char) ag[p]), yv);
+                asum += ag[p];
+            }
+            const __m512i ysum = _mm512_madd_epi16(_mm512_maddubs_epi16(ones, yv), _mm512_set1_epi16(asum));
+            const __m512i acci = _mm512_sub_epi32(_mm512_add_epi32(pos, pos), ysum);
+            acc = _mm256_add_ps(acc, _mm256_mul_ps(_mm256_set1_ps(GGML_CPU_FP16_TO_FP32(yb[0].d)), _mm256_cvtepi32_ps(_mm512_castsi512_si256(acci))));
+            acc = _mm256_add_ps(acc, _mm256_mul_ps(_mm256_set1_ps(GGML_CPU_FP16_TO_FP32(yb[1].d)), _mm256_cvtepi32_ps(_mm512_extracti64x4_epi64(acci, 1))));
+        }
+        sumf = _mm256_add_ps(sumf, _mm256_mul_ps(_mm256_set1_ps(d), acc));
+    }
+    return hsum_float_8(sumf);
+}
+#elif defined(__AVX2__)
+static inline float ggml_vec_dot_sa_q8_0_avx2(int n, const uint8_t * GGML_RESTRICT x, const block_q8_0 * GGML_RESTRICT y, int Q, int G) {
+    const int ng = QK_SA / G;
+    const size_t bsz = Q*QK_SA/8 + Q*ng + sizeof(ggml_half);
+    const int nb = n / QK_SA;
+    const __m256i ones = _mm256_set1_epi8(1);
+
+    __m256 sumf = _mm256_setzero_ps();
+    for (int i = 0; i < nb; ++i) {
+        const uint8_t * qs = x + i*bsz;
+        const uint8_t * a  = qs + Q*QK_SA/8;
+        ggml_half dh; memcpy(&dh, a + Q*ng, sizeof(dh));
+        const float d = GGML_CPU_FP16_TO_FP32(dh);
+
+        __m256 acc = _mm256_setzero_ps();
+        for (int c = 0; c < QK_SA/32; ++c) {
+            const block_q8_0 * yb = &y[i*(QK_SA/32) + c];
+            const __m256i yv = _mm256_loadu_si256((const __m256i *) yb->qs);
+            const uint8_t * ag = a + (c*32/G)*Q;
+
+            __m256i pos = _mm256_setzero_si256();
+            int asum = 0;
+            for (int p = 0; p < Q; ++p) {
+                const __m256i m = bytes_from_bits_32(qs + p*32 + c*4);
+                asum += ag[p];
+#if defined(__AVX512VNNI__) && defined(__AVX512VL__)
+                pos = _mm256_dpbusd_epi32(pos, _mm256_and_si256(m, _mm256_set1_epi8((char) ag[p])), yv);
+#elif defined(__AVXVNNI__)
+                pos = _mm256_dpbusd_avx_epi32(pos, _mm256_and_si256(m, _mm256_set1_epi8((char) ag[p])), yv);
+#else
+                const __m256i sel = _mm256_maddubs_epi16(_mm256_and_si256(m, ones), yv);
+                pos = _mm256_add_epi32(pos, _mm256_madd_epi16(sel, _mm256_set1_epi16(ag[p])));
+#endif
+            }
+            const __m256i ysum = _mm256_madd_epi16(_mm256_maddubs_epi16(ones, yv), _mm256_set1_epi16(asum));
+            const __m256i acci = _mm256_sub_epi32(_mm256_add_epi32(pos, pos), ysum);
+            acc = _mm256_add_ps(acc, _mm256_mul_ps(_mm256_set1_ps(GGML_CPU_FP16_TO_FP32(yb->d)), _mm256_cvtepi32_ps(acci)));
+        }
+        sumf = _mm256_add_ps(sumf, _mm256_mul_ps(_mm256_set1_ps(d), acc));
+    }
+    return hsum_float_8(sumf);
+}
+#endif
+
+#if defined(__AVX512BW__) && defined(__AVX512VNNI__)
+#define GGML_SA_VEC_DOT(name, Q, G)                                                                                                                              \
+void ggml_vec_dot_##name##_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) { \
+    assert(nrc == 1);                                                                                                                                            \
+    UNUSED(nrc); UNUSED(bx); UNUSED(by); UNUSED(bs);                                                                                                             \
+    *s = ggml_vec_dot_sa_q8_0_avx512(n, (const uint8_t *) vx, (const block_q8_0 *) vy, Q, G);                                                                    \
+}
+#elif defined(__AVX2__)
+#define GGML_SA_VEC_DOT(name, Q, G)                                                                                                                              \
+void ggml_vec_dot_##name##_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) { \
+    assert(nrc == 1);                                                                                                                                            \
+    UNUSED(nrc); UNUSED(bx); UNUSED(by); UNUSED(bs);                                                                                                             \
+    *s = ggml_vec_dot_sa_q8_0_avx2(n, (const uint8_t *) vx, (const block_q8_0 *) vy, Q, G);                                                                    \
+}
+#else
+#define GGML_SA_VEC_DOT(name, Q, G)                                                                                                                              \
+void ggml_vec_dot_##name##_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) { \
+    ggml_vec_dot_##name##_q8_0_generic(n, s, bs, vx, bx, vy, by, nrc);                                                                                           \
+}
+#endif
+
+GGML_SA_VEC_DOT(sa2, SA2_Q, SA2_G)
+GGML_SA_VEC_DOT(sa3, SA3_Q, SA3_G)
+GGML_SA_VEC_DOT(sa4, SA4_Q, SA4_G)
+GGML_SA_VEC_DOT(sa2_128, SA2_128_Q, SA2_128_G)
+GGML_SA_VEC_DOT(sa3_128, SA3_128_Q, SA3_128_G)
+GGML_SA_VEC_DOT(sa4_128, SA4_128_Q, SA4_128_G)
+GGML_SA_VEC_DOT(sa2_256, SA2_256_Q, SA2_256_G)
+GGML_SA_VEC_DOT(sa3_256, SA3_256_Q, SA3_256_G)
+GGML_SA_VEC_DOT(sa4_256, SA4_256_Q, SA4_256_G)
+GGML_SA_VEC_DOT(sa5, SA5_Q, SA5_G)
+GGML_SA_VEC_DOT(sa5_128, SA5_128_Q, SA5_128_G)
+GGML_SA_VEC_DOT(sa5_256, SA5_256_Q, SA5_256_G)
+GGML_SA_VEC_DOT(sa6, SA6_Q, SA6_G)
+GGML_SA_VEC_DOT(sa6_128, SA6_128_Q, SA6_128_G)
+GGML_SA_VEC_DOT(sa6_256, SA6_256_Q, SA6_256_G)
+
 void ggml_vec_dot_q2_K_q8_K(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) {
     assert(nrc == 1);
     UNUSED(nrc);
diff --git a/ggml/src/ggml-cpu/ggml-cpu.c b/ggml/src/ggml-cpu/ggml-cpu.c
index 4c33369ee..d8430a9a4 100644
--- a/ggml/src/ggml-cpu/ggml-cpu.c
+++ b/ggml/src/ggml-cpu/ggml-cpu.c
@@ -237,6 +237,96 @@ static const struct ggml_type_traits_cpu type_traits_cpu[GGML_TYPE_COUNT] = {
         .vec_dot_type             = GGML_TYPE_Q8_0,
         .nrows                    = 1,
     },
+    [GGML_TYPE_SA2] = {
+        .from_float               = quantize_row_sa2,
+        .vec_dot                  = ggml_vec_dot_sa2_q8_0,
+        .vec_dot_type             = GGML_TYPE_Q8_0,
+        .nrows                    = 1,
+    },
+    [GGML_TYPE_SA3] = {
+        .from_float               = quantize_row_sa3,
+        .vec_dot                  = ggml_vec_dot_sa3_q8_0,
+        .vec_dot_type             = GGML_TYPE_Q8_0,
+        .nrows                    = 1,
+    },
+    [GGML_TYPE_SA4] = {
+        .from_float               = quantize_row_sa4,
+        .vec_dot                  = ggml_vec_dot_sa4_q8_0,
+        .vec_dot_type             = GGML_TYPE_Q8_0,
+        .nrows                    = 1,
+    },
+    [GGML_TYPE_SA2_128] = {
+        .from_float               = quantize_row_sa2_128,
+        .vec_dot                  = ggml_vec_dot_sa2_128_q8_0,
+        .vec_dot_type             = GGML_TYPE_Q8_0,
+        .nrows                    = 1,
+    },
+    [GGML_TYPE_SA3_128] = {
+        .from_float               = quantize_row_sa3_128,
+        .vec_dot                  = ggml_vec_dot_sa3_128_q8_0,
+        .vec_dot_type             = GGML_TYPE_Q8_0,
+        .nrows                    = 1,
+    },
+    [GGML_TYPE_SA4_128] = {
+        .from_float               = quantize_row_sa4_128,
+        .vec_dot                  = ggml_vec_dot_sa4_128_q8_0,
+        .vec_dot_type             = GGML_TYPE_Q8_0,
+        .nrows                    = 1,
+    },
+    [GGML_TYPE_SA2_256] = {
+        .from_float               = quantize_row_sa2_256,
+        .vec_dot                  = ggml_vec_dot_sa2_256_q8_0,
+        .vec_dot_type             = GGML_TYPE_Q8_0,
+        .nrows                    = 1,
+    },
+    [GGML_TYPE_SA3_256] = {
+        .from_float               = quantize_row_sa3_256,
+        .vec_dot                  = ggml_vec_dot_sa3_256_q8_0,
+        .vec_dot_type             = GGML_TYPE_Q8_0,
+        .nrows                    = 1,
+    },
+    [GGML_TYPE_SA4_256] = {
+        .from_float               = quantize_row_sa4_256,
+        .vec_dot                  = ggml_vec_dot_sa4_256_q8_0,
+        .vec_dot_type             = GGML_TYPE_Q8_0,
+        .nrows                    = 1,
+    },
+    [GGML_TYPE_SA5] = {
+        .from_float               = quantize_row_sa5,
+        .vec_dot                  = ggml_vec_dot_sa5_q8_0,
+        .vec_dot_type             = GGML_TYPE_Q8_0,
+        .nrows                    = 1,
+    },
+    [GGML_TYPE_SA5_128] = {
+        .from_float               = quantize_row_sa5_128,
+        .vec_dot                  = ggml_vec_dot_sa5_128_q8_0,
+        .vec_dot_type             = GGML_TYPE_Q8_0,
+        .nrows                    = 1,
+    },
+    [GGML_TYPE_SA5_256] = {
+        .from_float               = quantize_row_sa5_256,
+        .vec_dot                  = ggml_vec_dot_sa5_256_q8_0,
+        .vec_dot_type             = GGML_TYPE_Q8_0,
+        .nrows                    = 1,
+    },
+    [GGML_TYPE_SA6] = {
+        .from_float               = quantize_row_sa6,
+        .vec_dot                  = ggml_vec_dot_sa6_q8_0,
+        .vec_dot_type             = GGML_TYPE_Q8_0,
+        .nrows                    = 1,
+    },
+    [GGML_TYPE_SA6_128] = {
+        .from_float               = quantize_row_sa6_128,
+        .vec_dot                  = ggml_vec_dot_sa6_128_q8_0,
+        .vec_dot_type             = GGML_TYPE_Q8_0,
+        .nrows                    = 1,
+    },
+    [GGML_TYPE_SA6_256] = {
+        .from_float               = quantize_row_sa6_256,
+        .vec_dot                  = ggml_vec_dot_sa6_256_q8_0,
+        .vec_dot_type             = GGML_TYPE_Q8_0,
+        .nrows                    = 1,
+    },
     [GGML_TYPE_Q4_0] = {
         .from_float               = quantize_row_q4_0,
         .vec_dot                  = ggml_vec_dot_q4_0_q8_0,
diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp
index 02d5efffa..783a7119e 100644
--- a/ggml/src/ggml-cpu/ops.cpp
+++ b/ggml/src/ggml-cpu/ops.cpp
@@ -666,6 +666,21 @@ void ggml_compute_forward_add(
             } break;
         case GGML_TYPE_Q1_0:
         case GGML_TYPE_Q2_0:
+        case GGML_TYPE_SA2:
+        case GGML_TYPE_SA3:
+        case GGML_TYPE_SA4:
+        case GGML_TYPE_SA2_128:
+        case GGML_TYPE_SA3_128:
+        case GGML_TYPE_SA4_128:
+        case GGML_TYPE_SA2_256:
+        case GGML_TYPE_SA3_256:
+        case GGML_TYPE_SA4_256:
+        case GGML_TYPE_SA5:
+        case GGML_TYPE_SA5_128:
+        case GGML_TYPE_SA5_256:
+        case GGML_TYPE_SA6:
+        case GGML_TYPE_SA6_128:
+        case GGML_TYPE_SA6_256:
         case GGML_TYPE_Q4_0:
         case GGML_TYPE_Q4_1:
         case GGML_TYPE_Q5_0:
@@ -1117,6 +1132,21 @@ void ggml_compute_forward_add1(
             } break;
         case GGML_TYPE_Q1_0:
         case GGML_TYPE_Q2_0:
+        case GGML_TYPE_SA2:
+        case GGML_TYPE_SA3:
+        case GGML_TYPE_SA4:
+        case GGML_TYPE_SA2_128:
+        case GGML_TYPE_SA3_128:
+        case GGML_TYPE_SA4_128:
+        case GGML_TYPE_SA2_256:
+        case GGML_TYPE_SA3_256:
+        case GGML_TYPE_SA4_256:
+        case GGML_TYPE_SA5:
+        case GGML_TYPE_SA5_128:
+        case GGML_TYPE_SA5_256:
+        case GGML_TYPE_SA6:
+        case GGML_TYPE_SA6_128:
+        case GGML_TYPE_SA6_256:
         case GGML_TYPE_Q4_0:
         case GGML_TYPE_Q4_1:
         case GGML_TYPE_Q5_0:
@@ -1248,6 +1278,21 @@ void ggml_compute_forward_acc(
         case GGML_TYPE_BF16:
         case GGML_TYPE_Q1_0:
         case GGML_TYPE_Q2_0:
+        case GGML_TYPE_SA2:
+        case GGML_TYPE_SA3:
+        case GGML_TYPE_SA4:
+        case GGML_TYPE_SA2_128:
+        case GGML_TYPE_SA3_128:
+        case GGML_TYPE_SA4_128:
+        case GGML_TYPE_SA2_256:
+        case GGML_TYPE_SA3_256:
+        case GGML_TYPE_SA4_256:
+        case GGML_TYPE_SA5:
+        case GGML_TYPE_SA5_128:
+        case GGML_TYPE_SA5_256:
+        case GGML_TYPE_SA6:
+        case GGML_TYPE_SA6_128:
+        case GGML_TYPE_SA6_256:
         case GGML_TYPE_Q4_0:
         case GGML_TYPE_Q4_1:
         case GGML_TYPE_Q5_0:
@@ -4918,6 +4963,21 @@ void ggml_compute_forward_out_prod(
     switch (src0->type) {
         case GGML_TYPE_Q1_0:
         case GGML_TYPE_Q2_0:
+        case GGML_TYPE_SA2:
+        case GGML_TYPE_SA3:
+        case GGML_TYPE_SA4:
+        case GGML_TYPE_SA2_128:
+        case GGML_TYPE_SA3_128:
+        case GGML_TYPE_SA4_128:
+        case GGML_TYPE_SA2_256:
+        case GGML_TYPE_SA3_256:
+        case GGML_TYPE_SA4_256:
+        case GGML_TYPE_SA5:
+        case GGML_TYPE_SA5_128:
+        case GGML_TYPE_SA5_256:
+        case GGML_TYPE_SA6:
+        case GGML_TYPE_SA6_128:
+        case GGML_TYPE_SA6_256:
         case GGML_TYPE_Q4_0:
         case GGML_TYPE_Q4_1:
         case GGML_TYPE_Q5_0:
@@ -5236,6 +5296,21 @@ void ggml_compute_forward_set(
         case GGML_TYPE_BF16:
         case GGML_TYPE_Q1_0:
         case GGML_TYPE_Q2_0:
+        case GGML_TYPE_SA2:
+        case GGML_TYPE_SA3:
+        case GGML_TYPE_SA4:
+        case GGML_TYPE_SA2_128:
+        case GGML_TYPE_SA3_128:
+        case GGML_TYPE_SA4_128:
+        case GGML_TYPE_SA2_256:
+        case GGML_TYPE_SA3_256:
+        case GGML_TYPE_SA4_256:
+        case GGML_TYPE_SA5:
+        case GGML_TYPE_SA5_128:
+        case GGML_TYPE_SA5_256:
+        case GGML_TYPE_SA6:
+        case GGML_TYPE_SA6_128:
+        case GGML_TYPE_SA6_256:
         case GGML_TYPE_Q4_0:
         case GGML_TYPE_Q4_1:
         case GGML_TYPE_Q5_0:
@@ -5461,6 +5536,21 @@ void ggml_compute_forward_get_rows(
     switch (src0->type) {
         case GGML_TYPE_Q1_0:
         case GGML_TYPE_Q2_0:
+        case GGML_TYPE_SA2:
+        case GGML_TYPE_SA3:
+        case GGML_TYPE_SA4:
+        case GGML_TYPE_SA2_128:
+        case GGML_TYPE_SA3_128:
+        case GGML_TYPE_SA4_128:
+        case GGML_TYPE_SA2_256:
+        case GGML_TYPE_SA3_256:
+        case GGML_TYPE_SA4_256:
+        case GGML_TYPE_SA5:
+        case GGML_TYPE_SA5_128:
+        case GGML_TYPE_SA5_256:
+        case GGML_TYPE_SA6:
+        case GGML_TYPE_SA6_128:
+        case GGML_TYPE_SA6_256:
         case GGML_TYPE_Q4_0:
         case GGML_TYPE_Q4_1:
         case GGML_TYPE_Q5_0:
@@ -6221,6 +6311,21 @@ void ggml_compute_forward_clamp(
         case GGML_TYPE_BF16:
         case GGML_TYPE_Q1_0:
         case GGML_TYPE_Q2_0:
+        case GGML_TYPE_SA2:
+        case GGML_TYPE_SA3:
+        case GGML_TYPE_SA4:
+        case GGML_TYPE_SA2_128:
+        case GGML_TYPE_SA3_128:
+        case GGML_TYPE_SA4_128:
+        case GGML_TYPE_SA2_256:
+        case GGML_TYPE_SA3_256:
+        case GGML_TYPE_SA4_256:
+        case GGML_TYPE_SA5:
+        case GGML_TYPE_SA5_128:
+        case GGML_TYPE_SA5_256:
+        case GGML_TYPE_SA6:
+        case GGML_TYPE_SA6_128:
+        case GGML_TYPE_SA6_256:
         case GGML_TYPE_Q4_0:
         case GGML_TYPE_Q4_1:
         case GGML_TYPE_Q5_0:
diff --git a/ggml/src/ggml-cpu/quants.c b/ggml/src/ggml-cpu/quants.c
index 5e36459f8..507e96c4a 100644
--- a/ggml/src/ggml-cpu/quants.c
+++ b/ggml/src/ggml-cpu/quants.c
@@ -174,6 +174,67 @@ void ggml_vec_dot_q1_0_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, c
     *s = sumf;
 }
 
+// ShiftAdd / BCQ: per 32-element chunk and bit-plane, sum_j s_j*y_j = 2*sum_{bit=1} y_j - sum_j y_j
+static inline float ggml_vec_dot_sa_q8_0_impl(int n, const uint8_t * GGML_RESTRICT x, const block_q8_0 * GGML_RESTRICT y, int Q, int G) {
+    assert(n % QK_SA == 0);
+    const int ng = QK_SA / G;
+    const size_t bsz = Q*QK_SA/8 + Q*ng + sizeof(ggml_half);
+    const int nb = n / QK_SA;
+    float sumf = 0.0f;
+    for (int i = 0; i < nb; ++i) {
+        const uint8_t * qs = x + i*bsz;
+        const uint8_t * a  = qs + Q*QK_SA/8;
+        ggml_half dh; memcpy(&dh, a + Q*ng, sizeof(dh));
+        const float d = GGML_CPU_FP16_TO_FP32(dh);
+        float acc = 0.0f;
+        for (int c = 0; c < QK_SA/32; ++c) {
+            const block_q8_0 * yb = &y[i*(QK_SA/32) + c];
+            const float d8 = GGML_CPU_FP16_TO_FP32(yb->d);
+            int ysum = 0;
+            for (int j = 0; j < 32; ++j) ysum += yb->qs[j];
+            const int g = c*32/G;
+            for (int p = 0; p < Q; ++p) {
+                const uint8_t * bits = qs + p*32 + c*4;
+                int pos = 0;
+                for (int j = 0; j < 32; ++j) {
+                    if ((bits[j/8] >> (j%8)) & 1) pos += yb->qs[j];
+                }
+                acc += (float) a[g*Q + p] * d8 * (float)(2*pos - ysum);
+            }
+        }
+        sumf += d * acc;
+    }
+    return sumf;
+}
+
+#define GGML_SA_CPU_IMPL(name, Q, G)                                                                                    \
+void quantize_row_##name(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) {                           \
+    quantize_row_##name##_ref(x, (block_##name *) y, k);                                                                \
+}                                                                                                                       \
+void ggml_vec_dot_##name##_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, \
+                                const void * GGML_RESTRICT vy, size_t by, int nrc) {                                   \
+    assert(nrc == 1);                                                                                                   \
+    UNUSED(nrc); UNUSED(bx); UNUSED(by); UNUSED(bs);                                                                    \
+    *s = ggml_vec_dot_sa_q8_0_impl(n, (const uint8_t *) vx, (const block_q8_0 *) vy, Q, G);                            \
+}
+
+GGML_SA_CPU_IMPL(sa2, SA2_Q, SA2_G)
+GGML_SA_CPU_IMPL(sa3, SA3_Q, SA3_G)
+GGML_SA_CPU_IMPL(sa4, SA4_Q, SA4_G)
+GGML_SA_CPU_IMPL(sa2_128, SA2_128_Q, SA2_128_G)
+GGML_SA_CPU_IMPL(sa3_128, SA3_128_Q, SA3_128_G)
+GGML_SA_CPU_IMPL(sa4_128, SA4_128_Q, SA4_128_G)
+GGML_SA_CPU_IMPL(sa2_256, SA2_256_Q, SA2_256_G)
+GGML_SA_CPU_IMPL(sa3_256, SA3_256_Q, SA3_256_G)
+GGML_SA_CPU_IMPL(sa4_256, SA4_256_Q, SA4_256_G)
+GGML_SA_CPU_IMPL(sa5, SA5_Q, SA5_G)
+GGML_SA_CPU_IMPL(sa5_128, SA5_128_Q, SA5_128_G)
+GGML_SA_CPU_IMPL(sa5_256, SA5_256_Q, SA5_256_G)
+GGML_SA_CPU_IMPL(sa6, SA6_Q, SA6_G)
+GGML_SA_CPU_IMPL(sa6_128, SA6_128_Q, SA6_128_G)
+GGML_SA_CPU_IMPL(sa6_256, SA6_256_Q, SA6_256_G)
+
+
 void ggml_vec_dot_q2_0_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) {
     const int qk = QK2_0;
     const int nb = n / qk;
diff --git a/ggml/src/ggml-cpu/quants.h b/ggml/src/ggml-cpu/quants.h
index 93ea7eeff..0897bc694 100644
--- a/ggml/src/ggml-cpu/quants.h
+++ b/ggml/src/ggml-cpu/quants.h
@@ -40,6 +40,36 @@ void quantize_row_iq4_xs (const float * GGML_RESTRICT x, void * GGML_RESTRICT y,
 // Dot product
 void ggml_vec_dot_q1_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
 void ggml_vec_dot_q2_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
+void quantize_row_sa2(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
+void ggml_vec_dot_sa2_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
+void quantize_row_sa3(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
+void ggml_vec_dot_sa3_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
+void quantize_row_sa4(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
+void quantize_row_sa2_128(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
+void quantize_row_sa3_128(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
+void quantize_row_sa4_128(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
+void quantize_row_sa2_256(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
+void quantize_row_sa3_256(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
+void quantize_row_sa4_256(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
+void quantize_row_sa5(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
+void quantize_row_sa5_128(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
+void quantize_row_sa5_256(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
+void quantize_row_sa6(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
+void quantize_row_sa6_128(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
+void quantize_row_sa6_256(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
+void ggml_vec_dot_sa4_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
+void ggml_vec_dot_sa2_128_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
+void ggml_vec_dot_sa3_128_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
+void ggml_vec_dot_sa4_128_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
+void ggml_vec_dot_sa2_256_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
+void ggml_vec_dot_sa3_256_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
+void ggml_vec_dot_sa4_256_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
+void ggml_vec_dot_sa5_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
+void ggml_vec_dot_sa5_128_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
+void ggml_vec_dot_sa5_256_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
+void ggml_vec_dot_sa6_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
+void ggml_vec_dot_sa6_128_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
+void ggml_vec_dot_sa6_256_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
 void ggml_vec_dot_q4_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
 void ggml_vec_dot_q4_1_q8_1(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
 void ggml_vec_dot_q5_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
@@ -85,6 +115,21 @@ void ggml_vec_dot_nvfp4_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs,
 
 void ggml_vec_dot_tq1_0_q8_K_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
 void ggml_vec_dot_tq2_0_q8_K_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
+void ggml_vec_dot_sa2_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
+void ggml_vec_dot_sa3_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
+void ggml_vec_dot_sa4_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
+void ggml_vec_dot_sa2_128_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
+void ggml_vec_dot_sa3_128_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
+void ggml_vec_dot_sa4_128_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
+void ggml_vec_dot_sa2_256_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
+void ggml_vec_dot_sa3_256_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
+void ggml_vec_dot_sa4_256_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
+void ggml_vec_dot_sa5_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
+void ggml_vec_dot_sa5_128_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
+void ggml_vec_dot_sa5_256_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
+void ggml_vec_dot_sa6_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
+void ggml_vec_dot_sa6_128_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
+void ggml_vec_dot_sa6_256_q8_0_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
 
 void ggml_vec_dot_q2_K_q8_K_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
 void ggml_vec_dot_q3_K_q8_K_generic(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
diff --git a/ggml/src/ggml-quants.c b/ggml/src/ggml-quants.c
index 7750a72ce..6e9a66e78 100644
--- a/ggml/src/ggml-quants.c
+++ b/ggml/src/ggml-quants.c
@@ -2111,6 +2111,158 @@ size_t quantize_q1_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst,
     return nrow * row_size;
 }
 
+// ====================== ShiftAdd / BCQ (sa2 .. sa6, group 64 / 128 / 256) ======================
+// Generic over Q (bit-planes) and G (alpha group size); block layout: qs[Q*32] | a[Q*256/G] | d.
+
+static void sa_dequant_row(const uint8_t * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k, int Q, int G) {
+    assert(k % QK_SA == 0);
+    const int ng = QK_SA / G;
+    const size_t bs = Q*QK_SA/8 + Q*ng + sizeof(ggml_half);
+    const int64_t nb = k / QK_SA;
+    for (int64_t i = 0; i < nb; ++i) {
+        const uint8_t * qs = x + i*bs;
+        const uint8_t * a  = qs + Q*QK_SA/8;
+        ggml_half dh; memcpy(&dh, a + Q*ng, sizeof(dh));
+        const float d = GGML_FP16_TO_FP32(dh);
+        for (int j = 0; j < QK_SA; ++j) {
+            const int g = j / G;
+            float v = 0.0f;
+            for (int p = 0; p < Q; ++p) {
+                const float al = d * a[g*Q + p];
+                v += ((qs[p*32 + j/8] >> (j%8)) & 1) ? al : -al;
+            }
+            y[i*QK_SA + j] = v;
+        }
+    }
+}
+
+// nearest of the 2^Q levels sum_p +-alpha_p; returns code with bit p set for +alpha_p
+static inline int sa_nearest(float w, const float * alpha, int Q) {
+    int best = 0; float bestd = INFINITY;
+    for (int c = 0; c < (1 << Q); ++c) {
+        float v = 0.0f;
+        for (int p = 0; p < Q; ++p) v += ((c >> p) & 1) ? alpha[p] : -alpha[p];
+        const float dd = fabsf(w - v);
+        if (dd < bestd) { bestd = dd; best = c; }
+    }
+    return best;
+}
+
+// greedy BCQ init + alternating LS refinement for one group of n weights (optional importance)
+static void sa_fit_group(const float * w, const float * imp, int n, int Q, float * alpha, float * r) {
+    float sw = 0.0f;
+    for (int j = 0; j < n; ++j) { r[j] = w[j]; sw += imp ? imp[j] : 1.0f; }
+    if (sw <= 0.0f) sw = 1.0f;
+    for (int p = 0; p < Q; ++p) {
+        float s = 0.0f;
+        for (int j = 0; j < n; ++j) s += (imp ? imp[j] : 1.0f) * fabsf(r[j]);
+        alpha[p] = s / sw;
+        for (int j = 0; j < n; ++j) r[j] -= r[j] >= 0 ? alpha[p] : -alpha[p];
+    }
+    for (int it = 0; it < 4; ++it) {
+        float A[8][9] = {{0}};
+        for (int j = 0; j < n; ++j) {
+            const int c = sa_nearest(w[j], alpha, Q);
+            const float h = imp ? imp[j] : 1.0f;
+            float b[8];
+            for (int p = 0; p < Q; ++p) b[p] = ((c >> p) & 1) ? 1.0f : -1.0f;
+            for (int p = 0; p < Q; ++p) {
+                for (int q = 0; q < Q; ++q) A[p][q] += h*b[p]*b[q];
+                A[p][Q] += h*b[p]*w[j];
+            }
+        }
+        // Gaussian elimination with partial pivoting
+        bool ok = true;
+        for (int p = 0; p < Q && ok; ++p) {
+            int piv = p;
+            for (int q = p + 1; q < Q; ++q) if (fabsf(A[q][p]) > fabsf(A[piv][p])) piv = q;
+            if (fabsf(A[piv][p]) < 1e-20f) { ok = false; break; }
+            if (piv != p) for (int c = 0; c <= Q; ++c) { float t = A[p][c]; A[p][c] = A[piv][c]; A[piv][c] = t; }
+            for (int q = 0; q < Q; ++q) {
+                if (q == p) continue;
+                const float f = A[q][p] / A[p][p];
+                for (int c = p; c <= Q; ++c) A[q][c] -= f*A[p][c];
+            }
+        }
+        if (!ok) break;
+        for (int p = 0; p < Q; ++p) alpha[p] = fabsf(A[p][Q] / A[p][p]);
+        // sort descending (sign folded into the bit-planes)
+        for (int p = 0; p < Q; ++p) for (int q = p + 1; q < Q; ++q)
+            if (alpha[q] > alpha[p]) { float t = alpha[p]; alpha[p] = alpha[q]; alpha[q] = t; }
+    }
+}
+
+static void sa_quant_row(const float * GGML_RESTRICT x, uint8_t * GGML_RESTRICT y, int64_t k, int Q, int G, const float * imp) {
+    assert(k % QK_SA == 0);
+    const int ng = QK_SA / G;
+    const size_t bs = Q*QK_SA/8 + Q*ng + sizeof(ggml_half);
+    const int64_t nb = k / QK_SA;
+    float alpha[8*8];
+    float r[QK_SA];
+    for (int64_t i = 0; i < nb; ++i) {
+        const float * xb = x + i*QK_SA;
+        uint8_t * qs = y + i*bs;
+        uint8_t * a  = qs + Q*QK_SA/8;
+        float amax = 0.0f;
+        for (int g = 0; g < ng; ++g) {
+            sa_fit_group(xb + g*G, imp ? imp + g*G : NULL, G, Q, alpha + g*Q, r);
+            for (int p = 0; p < Q; ++p) amax = MAX(amax, alpha[g*Q + p]);
+        }
+        const float d = amax / 255.0f;
+        const ggml_half dh = GGML_FP32_TO_FP16(d);
+        memcpy(a + Q*ng, &dh, sizeof(dh));
+        const float dq = GGML_FP16_TO_FP32(dh);
+        const float id = dq > 0 ? 1.0f/dq : 0.0f;
+        memset(qs, 0, Q*QK_SA/8);
+        for (int g = 0; g < ng; ++g) {
+            float aq[8];
+            for (int p = 0; p < Q; ++p) {
+                const int c = MIN(255, MAX(0, (int) nearbyintf(alpha[g*Q + p]*id)));
+                a[g*Q + p] = (uint8_t) c;
+                aq[p] = dq*c;
+            }
+            for (int j = g*G; j < (g + 1)*G; ++j) {
+                const int c = sa_nearest(xb[j], aq, Q);
+                for (int p = 0; p < Q; ++p) if ((c >> p) & 1) qs[p*32 + j/8] |= (uint8_t)(1 << (j%8));
+            }
+        }
+    }
+}
+
+#define GGML_SA_IMPL(name, TYPE, Q, G)                                                                                  \
+void quantize_row_##name##_ref(const float * GGML_RESTRICT x, block_##name * GGML_RESTRICT y, int64_t k) {             \
+    sa_quant_row(x, (uint8_t *) y, k, Q, G, NULL);                                                                      \
+}                                                                                                                       \
+void dequantize_row_##name(const block_##name * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) {                 \
+    sa_dequant_row((const uint8_t *) x, y, k, Q, G);                                                                    \
+}                                                                                                                       \
+size_t quantize_##name(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row,     \
+                       const float * quant_weights) {                                                                   \
+    const size_t row_size = ggml_row_size(TYPE, n_per_row);                                                             \
+    char * qrow = (char *) dst;                                                                                         \
+    for (int64_t row = 0; row < nrow; ++row) {                                                                          \
+        sa_quant_row(src + row*n_per_row, (uint8_t *) qrow, n_per_row, Q, G, quant_weights);                            \
+        qrow += row_size;                                                                                               \
+    }                                                                                                                   \
+    return nrow * row_size;                                                                                             \
+}
+
+GGML_SA_IMPL(sa2, GGML_TYPE_SA2, SA2_Q, SA2_G)
+GGML_SA_IMPL(sa3, GGML_TYPE_SA3, SA3_Q, SA3_G)
+GGML_SA_IMPL(sa4, GGML_TYPE_SA4, SA4_Q, SA4_G)
+GGML_SA_IMPL(sa2_128, GGML_TYPE_SA2_128, SA2_128_Q, SA2_128_G)
+GGML_SA_IMPL(sa3_128, GGML_TYPE_SA3_128, SA3_128_Q, SA3_128_G)
+GGML_SA_IMPL(sa4_128, GGML_TYPE_SA4_128, SA4_128_Q, SA4_128_G)
+GGML_SA_IMPL(sa2_256, GGML_TYPE_SA2_256, SA2_256_Q, SA2_256_G)
+GGML_SA_IMPL(sa3_256, GGML_TYPE_SA3_256, SA3_256_Q, SA3_256_G)
+GGML_SA_IMPL(sa4_256, GGML_TYPE_SA4_256, SA4_256_Q, SA4_256_G)
+GGML_SA_IMPL(sa5, GGML_TYPE_SA5, SA5_Q, SA5_G)
+GGML_SA_IMPL(sa5_128, GGML_TYPE_SA5_128, SA5_128_Q, SA5_128_G)
+GGML_SA_IMPL(sa5_256, GGML_TYPE_SA5_256, SA5_256_Q, SA5_256_G)
+GGML_SA_IMPL(sa6, GGML_TYPE_SA6, SA6_Q, SA6_G)
+GGML_SA_IMPL(sa6_128, GGML_TYPE_SA6_128, SA6_128_Q, SA6_128_G)
+GGML_SA_IMPL(sa6_256, GGML_TYPE_SA6_256, SA6_256_Q, SA6_256_G)
+
 size_t quantize_q2_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) {
     if (!quant_weights) {
         quantize_row_q2_0_ref(src, dst, (int64_t)nrow*n_per_row);
@@ -5509,6 +5661,156 @@ bool ggml_validate_row_data(enum ggml_type type, const void * data, size_t nbyte
             {
                 VALIDATE_ROW_DATA_D_F16_IMPL(block_q2_0, data, nb);
             } break;
+        case GGML_TYPE_SA2:
+            {
+                const block_sa2 * q = (const block_sa2 *) data;
+                for (size_t i = 0; i < nb; ++i) {
+                    ggml_half d; memcpy(&d, q[i].d, sizeof(d));
+                    if (!validate_fp16(d, i)) {
+                        return false;
+                    }
+                }
+            } break;
+        case GGML_TYPE_SA3:
+            {
+                const block_sa3 * q = (const block_sa3 *) data;
+                for (size_t i = 0; i < nb; ++i) {
+                    ggml_half d; memcpy(&d, q[i].d, sizeof(d));
+                    if (!validate_fp16(d, i)) {
+                        return false;
+                    }
+                }
+            } break;
+        case GGML_TYPE_SA4:
+            {
+                const block_sa4 * q = (const block_sa4 *) data;
+                for (size_t i = 0; i < nb; ++i) {
+                    ggml_half d; memcpy(&d, q[i].d, sizeof(d));
+                    if (!validate_fp16(d, i)) {
+                        return false;
+                    }
+                }
+            } break;
+        case GGML_TYPE_SA2_128:
+            {
+                const block_sa2_128 * q = (const block_sa2_128 *) data;
+                for (size_t i = 0; i < nb; ++i) {
+                    ggml_half d; memcpy(&d, q[i].d, sizeof(d));
+                    if (!validate_fp16(d, i)) {
+                        return false;
+                    }
+                }
+            } break;
+        case GGML_TYPE_SA3_128:
+            {
+                const block_sa3_128 * q = (const block_sa3_128 *) data;
+                for (size_t i = 0; i < nb; ++i) {
+                    ggml_half d; memcpy(&d, q[i].d, sizeof(d));
+                    if (!validate_fp16(d, i)) {
+                        return false;
+                    }
+                }
+            } break;
+        case GGML_TYPE_SA4_128:
+            {
+                const block_sa4_128 * q = (const block_sa4_128 *) data;
+                for (size_t i = 0; i < nb; ++i) {
+                    ggml_half d; memcpy(&d, q[i].d, sizeof(d));
+                    if (!validate_fp16(d, i)) {
+                        return false;
+                    }
+                }
+            } break;
+        case GGML_TYPE_SA2_256:
+            {
+                const block_sa2_256 * q = (const block_sa2_256 *) data;
+                for (size_t i = 0; i < nb; ++i) {
+                    ggml_half d; memcpy(&d, q[i].d, sizeof(d));
+                    if (!validate_fp16(d, i)) {
+                        return false;
+                    }
+                }
+            } break;
+        case GGML_TYPE_SA3_256:
+            {
+                const block_sa3_256 * q = (const block_sa3_256 *) data;
+                for (size_t i = 0; i < nb; ++i) {
+                    ggml_half d; memcpy(&d, q[i].d, sizeof(d));
+                    if (!validate_fp16(d, i)) {
+                        return false;
+                    }
+                }
+            } break;
+        case GGML_TYPE_SA4_256:
+            {
+                const block_sa4_256 * q = (const block_sa4_256 *) data;
+                for (size_t i = 0; i < nb; ++i) {
+                    ggml_half d; memcpy(&d, q[i].d, sizeof(d));
+                    if (!validate_fp16(d, i)) {
+                        return false;
+                    }
+                }
+            } break;
+        case GGML_TYPE_SA5:
+            {
+                const block_sa5 * q = (const block_sa5 *) data;
+                for (size_t i = 0; i < nb; ++i) {
+                    ggml_half d; memcpy(&d, q[i].d, sizeof(d));
+                    if (!validate_fp16(d, i)) {
+                        return false;
+                    }
+                }
+            } break;
+        case GGML_TYPE_SA5_128:
+            {
+                const block_sa5_128 * q = (const block_sa5_128 *) data;
+                for (size_t i = 0; i < nb; ++i) {
+                    ggml_half d; memcpy(&d, q[i].d, sizeof(d));
+                    if (!validate_fp16(d, i)) {
+                        return false;
+                    }
+                }
+            } break;
+        case GGML_TYPE_SA5_256:
+            {
+                const block_sa5_256 * q = (const block_sa5_256 *) data;
+                for (size_t i = 0; i < nb; ++i) {
+                    ggml_half d; memcpy(&d, q[i].d, sizeof(d));
+                    if (!validate_fp16(d, i)) {
+                        return false;
+                    }
+                }
+            } break;
+        case GGML_TYPE_SA6:
+            {
+                const block_sa6 * q = (const block_sa6 *) data;
+                for (size_t i = 0; i < nb; ++i) {
+                    ggml_half d; memcpy(&d, q[i].d, sizeof(d));
+                    if (!validate_fp16(d, i)) {
+                        return false;
+                    }
+                }
+            } break;
+        case GGML_TYPE_SA6_128:
+            {
+                const block_sa6_128 * q = (const block_sa6_128 *) data;
+                for (size_t i = 0; i < nb; ++i) {
+                    ggml_half d; memcpy(&d, q[i].d, sizeof(d));
+                    if (!validate_fp16(d, i)) {
+                        return false;
+                    }
+                }
+            } break;
+        case GGML_TYPE_SA6_256:
+            {
+                const block_sa6_256 * q = (const block_sa6_256 *) data;
+                for (size_t i = 0; i < nb; ++i) {
+                    ggml_half d; memcpy(&d, q[i].d, sizeof(d));
+                    if (!validate_fp16(d, i)) {
+                        return false;
+                    }
+                }
+            } break;
         case GGML_TYPE_Q4_0:
             {
                 VALIDATE_ROW_DATA_D_F16_IMPL(block_q4_0, data, nb);
diff --git a/ggml/src/ggml-quants.h b/ggml/src/ggml-quants.h
index 75188f1af..3be990ca3 100644
--- a/ggml/src/ggml-quants.h
+++ b/ggml/src/ggml-quants.h
@@ -96,6 +96,51 @@ GGML_API size_t quantize_q5_K(const float * GGML_RESTRICT src, void * GGML_RESTR
 GGML_API size_t quantize_q6_K(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
 GGML_API size_t quantize_q1_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
 GGML_API size_t quantize_q2_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
+GGML_API void quantize_row_sa2_ref(const float * GGML_RESTRICT x, block_sa2 * GGML_RESTRICT y, int64_t k);
+GGML_API void dequantize_row_sa2(const block_sa2 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
+GGML_API size_t quantize_sa2(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
+GGML_API void quantize_row_sa3_ref(const float * GGML_RESTRICT x, block_sa3 * GGML_RESTRICT y, int64_t k);
+GGML_API void dequantize_row_sa3(const block_sa3 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
+GGML_API size_t quantize_sa3(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
+GGML_API void quantize_row_sa4_ref(const float * GGML_RESTRICT x, block_sa4 * GGML_RESTRICT y, int64_t k);
+GGML_API void quantize_row_sa2_128_ref(const float * GGML_RESTRICT x, block_sa2_128 * GGML_RESTRICT y, int64_t k);
+GGML_API void quantize_row_sa3_128_ref(const float * GGML_RESTRICT x, block_sa3_128 * GGML_RESTRICT y, int64_t k);
+GGML_API void quantize_row_sa4_128_ref(const float * GGML_RESTRICT x, block_sa4_128 * GGML_RESTRICT y, int64_t k);
+GGML_API void quantize_row_sa2_256_ref(const float * GGML_RESTRICT x, block_sa2_256 * GGML_RESTRICT y, int64_t k);
+GGML_API void quantize_row_sa3_256_ref(const float * GGML_RESTRICT x, block_sa3_256 * GGML_RESTRICT y, int64_t k);
+GGML_API void quantize_row_sa4_256_ref(const float * GGML_RESTRICT x, block_sa4_256 * GGML_RESTRICT y, int64_t k);
+GGML_API void quantize_row_sa5_ref(const float * GGML_RESTRICT x, block_sa5 * GGML_RESTRICT y, int64_t k);
+GGML_API void quantize_row_sa5_128_ref(const float * GGML_RESTRICT x, block_sa5_128 * GGML_RESTRICT y, int64_t k);
+GGML_API void quantize_row_sa5_256_ref(const float * GGML_RESTRICT x, block_sa5_256 * GGML_RESTRICT y, int64_t k);
+GGML_API void quantize_row_sa6_ref(const float * GGML_RESTRICT x, block_sa6 * GGML_RESTRICT y, int64_t k);
+GGML_API void quantize_row_sa6_128_ref(const float * GGML_RESTRICT x, block_sa6_128 * GGML_RESTRICT y, int64_t k);
+GGML_API void quantize_row_sa6_256_ref(const float * GGML_RESTRICT x, block_sa6_256 * GGML_RESTRICT y, int64_t k);
+GGML_API void dequantize_row_sa4(const block_sa4 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
+GGML_API void dequantize_row_sa2_128(const block_sa2_128 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
+GGML_API void dequantize_row_sa3_128(const block_sa3_128 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
+GGML_API void dequantize_row_sa4_128(const block_sa4_128 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
+GGML_API void dequantize_row_sa2_256(const block_sa2_256 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
+GGML_API void dequantize_row_sa3_256(const block_sa3_256 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
+GGML_API void dequantize_row_sa4_256(const block_sa4_256 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
+GGML_API void dequantize_row_sa5(const block_sa5 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
+GGML_API void dequantize_row_sa5_128(const block_sa5_128 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
+GGML_API void dequantize_row_sa5_256(const block_sa5_256 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
+GGML_API void dequantize_row_sa6(const block_sa6 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
+GGML_API void dequantize_row_sa6_128(const block_sa6_128 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
+GGML_API void dequantize_row_sa6_256(const block_sa6_256 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
+GGML_API size_t quantize_sa4(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
+GGML_API size_t quantize_sa2_128(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
+GGML_API size_t quantize_sa3_128(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
+GGML_API size_t quantize_sa4_128(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
+GGML_API size_t quantize_sa2_256(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
+GGML_API size_t quantize_sa3_256(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
+GGML_API size_t quantize_sa4_256(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
+GGML_API size_t quantize_sa5(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
+GGML_API size_t quantize_sa5_128(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
+GGML_API size_t quantize_sa5_256(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
+GGML_API size_t quantize_sa6(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
+GGML_API size_t quantize_sa6_128(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
+GGML_API size_t quantize_sa6_256(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
 GGML_API size_t quantize_q4_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
 GGML_API size_t quantize_q4_1(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
 GGML_API size_t quantize_q5_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
diff --git a/ggml/src/ggml.c b/ggml/src/ggml.c
index 0b2adf6e4..b5dc05369 100644
--- a/ggml/src/ggml.c
+++ b/ggml/src/ggml.c
@@ -698,6 +698,126 @@ static const struct ggml_type_traits type_traits[GGML_TYPE_COUNT] = {
         .to_float                 = (ggml_to_float_t) dequantize_row_q2_0,
         .from_float_ref           = (ggml_from_float_t) quantize_row_q2_0_ref,
     },
+    [GGML_TYPE_SA2] = {
+        .type_name                = "sa2",
+        .blck_size                = QK_SA,
+        .type_size                = sizeof(block_sa2),
+        .is_quantized             = true,
+        .to_float                 = (ggml_to_float_t) dequantize_row_sa2,
+        .from_float_ref           = (ggml_from_float_t) quantize_row_sa2_ref,
+    },
+    [GGML_TYPE_SA3] = {
+        .type_name                = "sa3",
+        .blck_size                = QK_SA,
+        .type_size                = sizeof(block_sa3),
+        .is_quantized             = true,
+        .to_float                 = (ggml_to_float_t) dequantize_row_sa3,
+        .from_float_ref           = (ggml_from_float_t) quantize_row_sa3_ref,
+    },
+    [GGML_TYPE_SA4] = {
+        .type_name                = "sa4",
+        .blck_size                = QK_SA,
+        .type_size                = sizeof(block_sa4),
+        .is_quantized             = true,
+        .to_float                 = (ggml_to_float_t) dequantize_row_sa4,
+        .from_float_ref           = (ggml_from_float_t) quantize_row_sa4_ref,
+    },
+    [GGML_TYPE_SA2_128] = {
+        .type_name                = "sa2_128",
+        .blck_size                = QK_SA,
+        .type_size                = sizeof(block_sa2_128),
+        .is_quantized             = true,
+        .to_float                 = (ggml_to_float_t) dequantize_row_sa2_128,
+        .from_float_ref           = (ggml_from_float_t) quantize_row_sa2_128_ref,
+    },
+    [GGML_TYPE_SA3_128] = {
+        .type_name                = "sa3_128",
+        .blck_size                = QK_SA,
+        .type_size                = sizeof(block_sa3_128),
+        .is_quantized             = true,
+        .to_float                 = (ggml_to_float_t) dequantize_row_sa3_128,
+        .from_float_ref           = (ggml_from_float_t) quantize_row_sa3_128_ref,
+    },
+    [GGML_TYPE_SA4_128] = {
+        .type_name                = "sa4_128",
+        .blck_size                = QK_SA,
+        .type_size                = sizeof(block_sa4_128),
+        .is_quantized             = true,
+        .to_float                 = (ggml_to_float_t) dequantize_row_sa4_128,
+        .from_float_ref           = (ggml_from_float_t) quantize_row_sa4_128_ref,
+    },
+    [GGML_TYPE_SA2_256] = {
+        .type_name                = "sa2_256",
+        .blck_size                = QK_SA,
+        .type_size                = sizeof(block_sa2_256),
+        .is_quantized             = true,
+        .to_float                 = (ggml_to_float_t) dequantize_row_sa2_256,
+        .from_float_ref           = (ggml_from_float_t) quantize_row_sa2_256_ref,
+    },
+    [GGML_TYPE_SA3_256] = {
+        .type_name                = "sa3_256",
+        .blck_size                = QK_SA,
+        .type_size                = sizeof(block_sa3_256),
+        .is_quantized             = true,
+        .to_float                 = (ggml_to_float_t) dequantize_row_sa3_256,
+        .from_float_ref           = (ggml_from_float_t) quantize_row_sa3_256_ref,
+    },
+    [GGML_TYPE_SA4_256] = {
+        .type_name                = "sa4_256",
+        .blck_size                = QK_SA,
+        .type_size                = sizeof(block_sa4_256),
+        .is_quantized             = true,
+        .to_float                 = (ggml_to_float_t) dequantize_row_sa4_256,
+        .from_float_ref           = (ggml_from_float_t) quantize_row_sa4_256_ref,
+    },
+    [GGML_TYPE_SA5] = {
+        .type_name                = "sa5",
+        .blck_size                = QK_SA,
+        .type_size                = sizeof(block_sa5),
+        .is_quantized             = true,
+        .to_float                 = (ggml_to_float_t) dequantize_row_sa5,
+        .from_float_ref           = (ggml_from_float_t) quantize_row_sa5_ref,
+    },
+    [GGML_TYPE_SA5_128] = {
+        .type_name                = "sa5_128",
+        .blck_size                = QK_SA,
+        .type_size                = sizeof(block_sa5_128),
+        .is_quantized             = true,
+        .to_float                 = (ggml_to_float_t) dequantize_row_sa5_128,
+        .from_float_ref           = (ggml_from_float_t) quantize_row_sa5_128_ref,
+    },
+    [GGML_TYPE_SA5_256] = {
+        .type_name                = "sa5_256",
+        .blck_size                = QK_SA,
+        .type_size                = sizeof(block_sa5_256),
+        .is_quantized             = true,
+        .to_float                 = (ggml_to_float_t) dequantize_row_sa5_256,
+        .from_float_ref           = (ggml_from_float_t) quantize_row_sa5_256_ref,
+    },
+    [GGML_TYPE_SA6] = {
+        .type_name                = "sa6",
+        .blck_size                = QK_SA,
+        .type_size                = sizeof(block_sa6),
+        .is_quantized             = true,
+        .to_float                 = (ggml_to_float_t) dequantize_row_sa6,
+        .from_float_ref           = (ggml_from_float_t) quantize_row_sa6_ref,
+    },
+    [GGML_TYPE_SA6_128] = {
+        .type_name                = "sa6_128",
+        .blck_size                = QK_SA,
+        .type_size                = sizeof(block_sa6_128),
+        .is_quantized             = true,
+        .to_float                 = (ggml_to_float_t) dequantize_row_sa6_128,
+        .from_float_ref           = (ggml_from_float_t) quantize_row_sa6_128_ref,
+    },
+    [GGML_TYPE_SA6_256] = {
+        .type_name                = "sa6_256",
+        .blck_size                = QK_SA,
+        .type_size                = sizeof(block_sa6_256),
+        .is_quantized             = true,
+        .to_float                 = (ggml_to_float_t) dequantize_row_sa6_256,
+        .from_float_ref           = (ggml_from_float_t) quantize_row_sa6_256_ref,
+    },
     [GGML_TYPE_Q4_0] = {
         .type_name                = "q4_0",
         .blck_size                = QK4_0,
@@ -8145,6 +8265,21 @@ size_t ggml_quantize_chunk(
     switch (type) {
         case GGML_TYPE_Q1_0:    result = quantize_q1_0   (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
         case GGML_TYPE_Q2_0:    result = quantize_q2_0   (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
+        case GGML_TYPE_SA2:     result = quantize_sa2    (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
+        case GGML_TYPE_SA3:     result = quantize_sa3    (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
+        case GGML_TYPE_SA4:     result = quantize_sa4    (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
+        case GGML_TYPE_SA2_128:     result = quantize_sa2_128    (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
+        case GGML_TYPE_SA3_128:     result = quantize_sa3_128    (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
+        case GGML_TYPE_SA4_128:     result = quantize_sa4_128    (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
+        case GGML_TYPE_SA2_256:     result = quantize_sa2_256    (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
+        case GGML_TYPE_SA3_256:     result = quantize_sa3_256    (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
+        case GGML_TYPE_SA4_256:     result = quantize_sa4_256    (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
+        case GGML_TYPE_SA5:     result = quantize_sa5    (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
+        case GGML_TYPE_SA5_128:     result = quantize_sa5_128    (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
+        case GGML_TYPE_SA5_256:     result = quantize_sa5_256    (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
+        case GGML_TYPE_SA6:     result = quantize_sa6    (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
+        case GGML_TYPE_SA6_128:     result = quantize_sa6_128    (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
+        case GGML_TYPE_SA6_256:     result = quantize_sa6_256    (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
         case GGML_TYPE_Q4_0:    result = quantize_q4_0   (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
         case GGML_TYPE_Q4_1:    result = quantize_q4_1   (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
         case GGML_TYPE_Q5_0:    result = quantize_q5_0   (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
diff --git a/gguf-py/gguf/constants.py b/gguf-py/gguf/constants.py
index b8cfa3a85..3d38851d4 100644
--- a/gguf-py/gguf/constants.py
+++ b/gguf-py/gguf/constants.py
@@ -6283,6 +6283,21 @@ class GGMLQuantizationType(IntEnum):
     NVFP4   = 40
     Q1_0    = 41
     Q2_0    = 42
+    SA2     = 43
+    SA3     = 44
+    SA4     = 45
+    SA2_128 = 47
+    SA3_128 = 48
+    SA4_128 = 49
+    SA2_256 = 50
+    SA3_256 = 51
+    SA4_256 = 52
+    SA5     = 53
+    SA5_128 = 54
+    SA5_256 = 55
+    SA6     = 56
+    SA6_128 = 57
+    SA6_256 = 58
 
 
 class ExpertGatingFuncType(IntEnum):
@@ -6502,6 +6517,21 @@ GGML_QUANT_SIZES: dict[GGMLQuantizationType, tuple[int, int]] = {
     GGMLQuantizationType.NVFP4:   (64, 4 + 32),
     GGMLQuantizationType.Q1_0:    (128, 2 + 16),
     GGMLQuantizationType.Q2_0:    (64, 2 + 16),
+    GGMLQuantizationType.SA2:     (256, 2 * 32 + 2 * 4 + 2),
+    GGMLQuantizationType.SA3:     (256, 3 * 32 + 3 * 4 + 2),
+    GGMLQuantizationType.SA4:     (256, 4 * 32 + 4 * 4 + 2),
+    GGMLQuantizationType.SA2_128: (256, 2 * 32 + 2 * 2 + 2),
+    GGMLQuantizationType.SA3_128: (256, 3 * 32 + 3 * 2 + 2),
+    GGMLQuantizationType.SA4_128: (256, 4 * 32 + 4 * 2 + 2),
+    GGMLQuantizationType.SA2_256: (256, 2 * 32 + 2 * 1 + 2),
+    GGMLQuantizationType.SA3_256: (256, 3 * 32 + 3 * 1 + 2),
+    GGMLQuantizationType.SA4_256: (256, 4 * 32 + 4 * 1 + 2),
+    GGMLQuantizationType.SA5: (256, 5 * 32 + 5 * 4 + 2),
+    GGMLQuantizationType.SA5_128: (256, 5 * 32 + 5 * 2 + 2),
+    GGMLQuantizationType.SA5_256: (256, 5 * 32 + 5 * 1 + 2),
+    GGMLQuantizationType.SA6: (256, 6 * 32 + 6 * 4 + 2),
+    GGMLQuantizationType.SA6_128: (256, 6 * 32 + 6 * 2 + 2),
+    GGMLQuantizationType.SA6_256: (256, 6 * 32 + 6 * 1 + 2),
 }
 
 
diff --git a/gguf-py/gguf/quants.py b/gguf-py/gguf/quants.py
index cc0772913..20fe2ec51 100644
--- a/gguf-py/gguf/quants.py
+++ b/gguf-py/gguf/quants.py
@@ -1429,3 +1429,130 @@ class IQ4_XS(__Quant, qtype=GGMLQuantizationType.IQ4_XS):
         qs = np.take_along_axis(kvalues, qs, axis=-1).astype(np.float32).reshape((n_blocks, -1, 32))
 
         return (dl * qs).reshape((n_blocks, -1))
+
+
+class _SABase:
+    """ShiftAdd / BCQ super-block of 256: qs[Q*32] bit-planes | a[Q*256/G] uint8 alpha codes | d fp16.
+    w_j = sum_p d*a[g*Q+p] * (+1 if bit_{p,j} else -1),  g = j // G."""
+    Q = 3
+    G = 64
+
+    @classmethod
+    def pack(cls, bits: np.ndarray, codes: np.ndarray, d: np.ndarray) -> np.ndarray:
+        """bits [n, Q, 256] in {0,1}; codes [n, 256//G, Q] uint8; d [n] float -> raw blocks [n, type_size]"""
+        n = bits.shape[0]
+        qs = np.packbits(bits.astype(np.uint8).reshape(n, cls.Q, 32, 8), axis=-1, bitorder="little").reshape(n, -1)
+        a = codes.astype(np.uint8).reshape(n, -1)
+        dd = d.astype(np.float16).reshape(n, 1).view(np.uint8)
+        return np.concatenate([qs, a, dd], axis=-1)
+
+    @classmethod
+    def quantize_blocks(cls, blocks: np.ndarray) -> np.ndarray:
+        # greedy BCQ per group (reference only; real models are exported from the calibrated quantizer)
+        n = blocks.shape[0]
+        Q, G = cls.Q, cls.G
+        w = blocks.reshape(n, 256 // G, G).astype(np.float32)
+        r = w.copy()
+        alphas = []
+        for _ in range(Q):
+            a = np.abs(r).mean(-1, keepdims=True)
+            r = r - a * np.where(r >= 0, 1.0, -1.0)
+            alphas.append(a[..., 0])
+        alpha = np.stack(alphas, -1)  # [n, ng, Q]
+        d = alpha.reshape(n, -1).max(-1) / 255.0
+        d = d.astype(np.float16).astype(np.float32)
+        with np.errstate(divide="ignore", invalid="ignore"):
+            codes = np.clip(np.round(np.where(d[:, None, None] > 0, alpha / d[:, None, None], 0)), 0, 255)
+        aq = codes * d[:, None, None]
+        signs = np.array([[1 if (c >> p) & 1 else -1 for p in range(Q)] for c in range(2 ** Q)], dtype=np.float32)
+        levels = aq @ signs.T  # [n, ng, 2^Q]
+        idx = np.abs(w[..., None] - levels[:, :, None, :]).argmin(-1)  # [n, ng, G]
+        bits = ((idx.reshape(n, 256)[:, None, :] >> np.arange(Q)[None, :, None]) & 1)
+        return cls.pack(bits, codes, d)
+
+    @classmethod
+    def dequantize_blocks(cls, blocks: np.ndarray) -> np.ndarray:
+        n = blocks.shape[0]
+        Q, G = cls.Q, cls.G
+        qs, a, d = np.hsplit(blocks, [Q * 32, Q * 32 + Q * (256 // G)])
+        d = d.view(np.float16).astype(np.float32)  # [n, 1]
+        bits = np.unpackbits(qs.reshape(n, Q, 32), axis=-1, bitorder="little").astype(np.float32)  # [n, Q, 256]
+        s = bits * 2 - 1
+        alpha = a.reshape(n, 256 // G, Q).astype(np.float32) * d[:, :, None]
+        alpha = np.repeat(alpha, G, axis=1)  # [n, 256, Q]
+        return (s.transpose(0, 2, 1) * alpha).sum(-1)
+
+
+class SA2(_SABase, __Quant, qtype=GGMLQuantizationType.SA2):
+    Q = 2
+    G = 64
+
+
+class SA3(_SABase, __Quant, qtype=GGMLQuantizationType.SA3):
+    Q = 3
+    G = 64
+
+
+class SA4(_SABase, __Quant, qtype=GGMLQuantizationType.SA4):
+    Q = 4
+    G = 64
+
+
+class SA2_128(_SABase, __Quant, qtype=GGMLQuantizationType.SA2_128):
+    Q = 2
+    G = 128
+
+
+class SA3_128(_SABase, __Quant, qtype=GGMLQuantizationType.SA3_128):
+    Q = 3
+    G = 128
+
+
+class SA4_128(_SABase, __Quant, qtype=GGMLQuantizationType.SA4_128):
+    Q = 4
+    G = 128
+
+
+class SA2_256(_SABase, __Quant, qtype=GGMLQuantizationType.SA2_256):
+    Q = 2
+    G = 256
+
+
+class SA3_256(_SABase, __Quant, qtype=GGMLQuantizationType.SA3_256):
+    Q = 3
+    G = 256
+
+
+class SA4_256(_SABase, __Quant, qtype=GGMLQuantizationType.SA4_256):
+    Q = 4
+    G = 256
+
+
+class SA5(_SABase, __Quant, qtype=GGMLQuantizationType.SA5):
+    Q = 5
+    G = 64
+
+
+class SA5_128(_SABase, __Quant, qtype=GGMLQuantizationType.SA5_128):
+    Q = 5
+    G = 128
+
+
+class SA5_256(_SABase, __Quant, qtype=GGMLQuantizationType.SA5_256):
+    Q = 5
+    G = 256
+
+
+class SA6(_SABase, __Quant, qtype=GGMLQuantizationType.SA6):
+    Q = 6
+    G = 64
+
+
+class SA6_128(_SABase, __Quant, qtype=GGMLQuantizationType.SA6_128):
+    Q = 6
+    G = 128
+
+
+class SA6_256(_SABase, __Quant, qtype=GGMLQuantizationType.SA6_256):
+    Q = 6
+    G = 256
diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp
index 0fc5c0b5f..f35c13196 100644
--- a/tests/test-backend-ops.cpp
+++ b/tests/test-backend-ops.cpp
@@ -9910,6 +9910,11 @@ static const ggml_type all_types[] = {
     GGML_TYPE_IQ2_XXS, GGML_TYPE_IQ2_XS, GGML_TYPE_IQ2_S,
     GGML_TYPE_IQ3_XXS, GGML_TYPE_IQ1_S, GGML_TYPE_IQ1_M,
     GGML_TYPE_IQ4_NL, GGML_TYPE_IQ3_S, GGML_TYPE_IQ4_XS,
+    GGML_TYPE_SA2, GGML_TYPE_SA2_128, GGML_TYPE_SA2_256,
+    GGML_TYPE_SA3, GGML_TYPE_SA3_128, GGML_TYPE_SA3_256,
+    GGML_TYPE_SA4, GGML_TYPE_SA4_128, GGML_TYPE_SA4_256,
+    GGML_TYPE_SA5, GGML_TYPE_SA5_128, GGML_TYPE_SA5_256,
+    GGML_TYPE_SA6, GGML_TYPE_SA6_128, GGML_TYPE_SA6_256,
 };
 
 static const ggml_type base_types[] = {
@@ -9939,6 +9944,11 @@ static const ggml_type other_types[] = {
     GGML_TYPE_IQ3_XXS, GGML_TYPE_IQ1_S, GGML_TYPE_IQ1_M,
     GGML_TYPE_IQ4_NL, GGML_TYPE_IQ3_S, GGML_TYPE_IQ4_XS,
     GGML_TYPE_BF16,
+    GGML_TYPE_SA2, GGML_TYPE_SA2_128, GGML_TYPE_SA2_256,
+    GGML_TYPE_SA3, GGML_TYPE_SA3_128, GGML_TYPE_SA3_256,
+    GGML_TYPE_SA4, GGML_TYPE_SA4_128, GGML_TYPE_SA4_256,
+    GGML_TYPE_SA5, GGML_TYPE_SA5_128, GGML_TYPE_SA5_256,
+    GGML_TYPE_SA6, GGML_TYPE_SA6_128, GGML_TYPE_SA6_256,
 };
 
 #ifdef _MSC_VER
@@ -11104,7 +11114,10 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
                                  GGML_TYPE_Q8_0, GGML_TYPE_Q1_0, GGML_TYPE_Q2_0, GGML_TYPE_MXFP4, GGML_TYPE_Q2_K, GGML_TYPE_Q3_K,
                                  GGML_TYPE_Q4_K, GGML_TYPE_Q5_K, GGML_TYPE_Q6_K, GGML_TYPE_TQ2_0,
                                  GGML_TYPE_IQ2_XXS, GGML_TYPE_IQ2_XS, GGML_TYPE_IQ2_S, GGML_TYPE_IQ3_XXS, GGML_TYPE_IQ3_S,
-                                 GGML_TYPE_IQ1_S, GGML_TYPE_IQ1_M, GGML_TYPE_IQ4_NL, GGML_TYPE_IQ4_XS}) {
+                                 GGML_TYPE_IQ1_S, GGML_TYPE_IQ1_M, GGML_TYPE_IQ4_NL, GGML_TYPE_IQ4_XS,
+                                 GGML_TYPE_SA2, GGML_TYPE_SA2_128, GGML_TYPE_SA2_256, GGML_TYPE_SA3, GGML_TYPE_SA3_128, GGML_TYPE_SA3_256,
+                                 GGML_TYPE_SA4, GGML_TYPE_SA4_128, GGML_TYPE_SA4_256, GGML_TYPE_SA5, GGML_TYPE_SA5_128, GGML_TYPE_SA5_256,
+                                 GGML_TYPE_SA6, GGML_TYPE_SA6_128, GGML_TYPE_SA6_256}) {
             test_cases.emplace_back(new test_mul_mat(type_a, GGML_TYPE_F32, 48,  n, 2560, {1, 1}, {1, 1}));
             test_cases.emplace_back(new test_mul_mat(type_a, GGML_TYPE_F32, 1000, n, 1024, {1, 1}, {1, 1}));
             test_cases.emplace_back(new test_mul_mat(type_a, GGML_TYPE_F32, 3000, n, 512,  {1, 1}, {1, 1}));
diff --git a/tests/test-quantize-fns.cpp b/tests/test-quantize-fns.cpp
index 1930a09b2..778c1360a 100644
--- a/tests/test-quantize-fns.cpp
+++ b/tests/test-quantize-fns.cpp
@@ -29,6 +29,19 @@ constexpr float MAX_DOT_PRODUCT_ERROR_LOWBIT = 0.04f;
 constexpr float MAX_DOT_PRODUCT_ERROR_FP4 = 0.03f;
 constexpr float MAX_DOT_PRODUCT_ERROR_BINARY = 0.40f;
 constexpr float MAX_DOT_PRODUCT_ERROR_TERNARY = 0.15f;
+constexpr float MAX_DOT_PRODUCT_ERROR_2PLANES = 0.06f; // ShiftAdd with 2 bit-planes: 4 levels per group, like a 2-bit grid without minimums
+
+// ShiftAdd / BCQ types: number of bit-planes (0 for other types)
+static int sa_bit_planes(ggml_type type) {
+    switch (type) {
+        case GGML_TYPE_SA2: case GGML_TYPE_SA2_128: case GGML_TYPE_SA2_256: return 2;
+        case GGML_TYPE_SA3: case GGML_TYPE_SA3_128: case GGML_TYPE_SA3_256: return 3;
+        case GGML_TYPE_SA4: case GGML_TYPE_SA4_128: case GGML_TYPE_SA4_256: return 4;
+        case GGML_TYPE_SA5: case GGML_TYPE_SA5_128: case GGML_TYPE_SA5_256: return 5;
+        case GGML_TYPE_SA6: case GGML_TYPE_SA6_128: case GGML_TYPE_SA6_256: return 6;
+        default: return 0;
+    }
+}
 
 static const char* RESULT_STR[] = {"ok", "FAILED"};
 
@@ -197,7 +210,9 @@ static int test_vec_dot_q(bool verbose) {
                 type == GGML_TYPE_Q3_K    ? MAX_QUANTIZATION_TOTAL_ERROR_3BITS :
                 type == GGML_TYPE_IQ3_S   ? MAX_QUANTIZATION_TOTAL_ERROR_3BITS :
                 type == GGML_TYPE_IQ3_XXS ? MAX_QUANTIZATION_TOTAL_ERROR_3BITS_XXS :
-                type == GGML_TYPE_NVFP4   ? MAX_QUANTIZATION_TOTAL_ERROR_FP4 : MAX_QUANTIZATION_TOTAL_ERROR;
+                type == GGML_TYPE_NVFP4   ? MAX_QUANTIZATION_TOTAL_ERROR_FP4 :
+                sa_bit_planes(type) == 2  ? MAX_QUANTIZATION_TOTAL_ERROR_2BITS :
+                sa_bit_planes(type) == 3  ? MAX_QUANTIZATION_TOTAL_ERROR_3BITS : MAX_QUANTIZATION_TOTAL_ERROR;
             bool failed = !(total_error < max_quantization_error);
             num_failed += failed;
             if (failed || verbose) {
@@ -213,8 +228,10 @@ static int test_vec_dot_q(bool verbose) {
 
             const float vec_dot_error = dot_product_error(qfns_cpu, type, test_size, test_data.data(), test_data2.data(), nullptr, nullptr, 1);
             const float max_allowed_error = type == GGML_TYPE_Q2_K || type == GGML_TYPE_IQ2_XS || type == GGML_TYPE_IQ2_XXS ||
-                type == GGML_TYPE_IQ3_XXS || type == GGML_TYPE_IQ3_S || type == GGML_TYPE_IQ2_S
+                type == GGML_TYPE_IQ3_XXS || type == GGML_TYPE_IQ3_S || type == GGML_TYPE_IQ2_S || sa_bit_planes(type) == 3
                 ? MAX_DOT_PRODUCT_ERROR_LOWBIT
+                : sa_bit_planes(type) == 2
+                ? MAX_DOT_PRODUCT_ERROR_2PLANES
                 : type == GGML_TYPE_Q1_0
                 ? MAX_DOT_PRODUCT_ERROR_BINARY
                 : type == GGML_TYPE_TQ1_0 || type == GGML_TYPE_TQ2_0 || type == GGML_TYPE_Q2_0