sync : ggml

* metal: fix memory leak if model is freed without any GPU operations

* metal: run dummy work only if residency sets are used

* metal: wrap function in #if defined

* metal: measure system-wide wired memory in test

* metal: always build regression test

Co-authored-by: YiChen Lv <63285796+forforever73@users.noreply.github.com>

---------

Co-authored-by: YiChen Lv <63285796+forforever73@users.noreply.github.com>

ggml/CMakeLists.txt

@@ -4,7 +4,7 @@ project("ggml" C CXX ASM)
 
 ### GGML Version
 set(GGML_VERSION_MAJOR 0)
-set(GGML_VERSION_MINOR 17)
+set(GGML_VERSION_MINOR 18)
 set(GGML_VERSION_PATCH 0)
 set(GGML_VERSION_BASE "${GGML_VERSION_MAJOR}.${GGML_VERSION_MINOR}.${GGML_VERSION_PATCH}")
 

ggml/src/ggml-cuda/common.cuh

@@ -977,6 +977,13 @@ struct ggml_cuda_type_traits<GGML_TYPE_Q1_0> {
     static constexpr int qi = QI1_0;
 };
 
+template<>
+struct ggml_cuda_type_traits<GGML_TYPE_Q2_0> {
+    static constexpr int qk = QK2_0;
+    static constexpr int qr = QR2_0;
+    static constexpr int qi = QI2_0;
+};
+
 template<>
 struct ggml_cuda_type_traits<GGML_TYPE_Q4_0> {
     static constexpr int qk = QK4_0;

ggml/src/ggml-cuda/convert.cu

@@ -459,6 +459,8 @@ to_bf16_cuda_t ggml_get_to_bf16_cuda(ggml_type type) {
     switch (type) {
         case GGML_TYPE_Q1_0:
             return dequantize_block_cont_cuda<QK1_0, QR1_0, dequantize_q1_0>;
+        case GGML_TYPE_Q2_0:
+            return dequantize_block_cont_cuda<QK2_0, QR2_0, dequantize_q2_0>;
         case GGML_TYPE_Q4_0:
             return dequantize_row_q4_0_cuda;
         case GGML_TYPE_Q4_1:
@@ -514,6 +516,8 @@ to_fp16_cuda_t ggml_get_to_fp16_cuda(ggml_type type) {
     switch (type) {
         case GGML_TYPE_Q1_0:
             return dequantize_block_cont_cuda<QK1_0, QR1_0, dequantize_q1_0>;
+        case GGML_TYPE_Q2_0:
+            return dequantize_block_cont_cuda<QK2_0, QR2_0, dequantize_q2_0>;
         case GGML_TYPE_Q4_0:
             return dequantize_row_q4_0_cuda;
         case GGML_TYPE_Q4_1:
@@ -572,6 +576,8 @@ to_fp32_cuda_t ggml_get_to_fp32_cuda(ggml_type type) {
     switch (type) {
         case GGML_TYPE_Q1_0:
             return dequantize_block_cont_cuda<QK1_0, QR1_0, dequantize_q1_0>;
+        case GGML_TYPE_Q2_0:
+            return dequantize_block_cont_cuda<QK2_0, QR2_0, dequantize_q2_0>;
         case GGML_TYPE_Q4_0:
             return dequantize_row_q4_0_cuda;
         case GGML_TYPE_Q4_1:
@@ -629,6 +635,8 @@ to_fp16_nc_cuda_t ggml_get_to_fp16_nc_cuda(ggml_type type) {
             return convert_unary_cuda<float>;
         case GGML_TYPE_Q1_0:
             return dequantize_block_cuda<QK1_0, QR1_0, dequantize_q1_0>;
+        case GGML_TYPE_Q2_0:
+            return dequantize_block_cuda<QK2_0, QR2_0, dequantize_q2_0>;
         case GGML_TYPE_Q4_0:
             return dequantize_block_cuda<QK4_0, QR4_0, dequantize_q4_0>;
         case GGML_TYPE_Q4_1:
@@ -652,6 +660,8 @@ to_bf16_nc_cuda_t ggml_get_to_bf16_nc_cuda(ggml_type type) {
             return convert_unary_cuda<float, nv_bfloat16>;
         case GGML_TYPE_Q1_0:
             return dequantize_block_cuda<QK1_0, QR1_0, dequantize_q1_0>;
+        case GGML_TYPE_Q2_0:
+            return dequantize_block_cuda<QK2_0, QR2_0, dequantize_q2_0>;
         case GGML_TYPE_Q4_0:
             return dequantize_block_cuda<QK4_0, QR4_0, dequantize_q4_0>;
         case GGML_TYPE_Q4_1:
@@ -675,6 +685,8 @@ to_fp32_nc_cuda_t ggml_get_to_fp32_nc_cuda(ggml_type type) {
             return convert_unary_cuda<half, float>;
         case GGML_TYPE_Q1_0:
             return dequantize_block_cuda<QK1_0, QR1_0, dequantize_q1_0>;
+        case GGML_TYPE_Q2_0:
+            return dequantize_block_cuda<QK2_0, QR2_0, dequantize_q2_0>;
         case GGML_TYPE_Q4_0:
             return dequantize_block_cuda<QK4_0, QR4_0, dequantize_q4_0>;
         case GGML_TYPE_Q4_1:

ggml/src/ggml-cuda/dequantize.cuh

@@ -23,6 +23,26 @@ static __device__ __forceinline__ void dequantize_q1_0(const void * vx, const in
     v.y = (2*bit_1 - 1) * d;
 }
 
+static __device__ __forceinline__ void dequantize_q2_0(const void * vx, const int64_t ib, const int iqs, float2 & v){
+    const block_q2_0 * x = (const block_q2_0 *) vx;
+
+    const float d = x[ib].d;
+
+    // Q2_0: 2 bits per element, 4 elements per byte.
+    // Stored code c in {0,1,2,3} maps to symbol s = c - 1 in {-1, 0, +1, +2}.
+    const int byte_index_0 = iqs / 4;
+    const int bit_offset_0 = (iqs % 4) * 2;
+
+    const int byte_index_1 = (iqs + 1) / 4;
+    const int bit_offset_1 = ((iqs + 1) % 4) * 2;
+
+    const int c0 = (x[ib].qs[byte_index_0] >> bit_offset_0) & 0x3;
+    const int c1 = (x[ib].qs[byte_index_1] >> bit_offset_1) & 0x3;
+
+    v.x = (c0 - 1) * d;
+    v.y = (c1 - 1) * d;
+}
+
 static __device__ __forceinline__ void dequantize_q4_0(const void * vx, const int64_t ib, const int iqs, float2 & v){
     const block_q4_0 * x = (const block_q4_0 *) vx;
 

ggml/src/ggml-cuda/getrows.cu

@@ -320,6 +320,10 @@ static void ggml_cuda_get_rows_switch_src0_type(
             get_rows_cuda_q<QK1_0, QR1_0, dequantize_q1_0>(src0_d, src1_d, dst_d,
                 ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream);
             break;
+        case GGML_TYPE_Q2_0:
+            get_rows_cuda_q<QK2_0, QR2_0, dequantize_q2_0>(src0_d, src1_d, dst_d,
+                ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream);
+            break;
         case GGML_TYPE_Q4_0:
             get_rows_cuda_q<QK4_0, QR4_0, dequantize_q4_0>(src0_d, src1_d, dst_d,
                 ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream);

ggml/src/ggml-cuda/ggml-cuda.cu

@@ -4802,6 +4802,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
                     case GGML_TYPE_F32:
                     case GGML_TYPE_F16:
                     case GGML_TYPE_Q1_0:
+                    case GGML_TYPE_Q2_0:
                     case GGML_TYPE_Q4_0:
                     case GGML_TYPE_Q4_1:
                     case GGML_TYPE_Q5_0:
@@ -4840,6 +4841,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
                     case GGML_TYPE_BF16:
                     case GGML_TYPE_I32:
                     case GGML_TYPE_Q1_0:
+                    case GGML_TYPE_Q2_0:
                     case GGML_TYPE_Q4_0:
                     case GGML_TYPE_Q4_1:
                     case GGML_TYPE_Q5_0:

ggml/src/ggml-cuda/mmq-config-ampere.cuh

@@ -16,6 +16,23 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf
     CASE(GGML_TYPE_Q1_0, 256, 1, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
     CASE(GGML_TYPE_Q1_0, 256, 1, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
 
+    CASE(GGML_TYPE_Q2_0, 256, 1, 128,   8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);
+    CASE(GGML_TYPE_Q2_0, 256, 1, 128,  16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);
+    CASE(GGML_TYPE_Q2_0, 256, 1, 128,  32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);
+    CASE(GGML_TYPE_Q2_0, 256, 1, 128,  64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);
+    CASE(GGML_TYPE_Q2_0, 256, 1, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);
+    CASE(GGML_TYPE_Q2_0, 256, 1, 128,   8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
+    CASE(GGML_TYPE_Q2_0, 256, 1, 128,  16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
+    CASE(GGML_TYPE_Q2_0, 256, 1, 128,  24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
+    CASE(GGML_TYPE_Q2_0, 256, 1, 128,  32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
+    CASE(GGML_TYPE_Q2_0, 256, 1, 128,  40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
+    CASE(GGML_TYPE_Q2_0, 256, 1, 128,  48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
+    CASE(GGML_TYPE_Q2_0, 256, 1, 128,  64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
+    CASE(GGML_TYPE_Q2_0, 256, 1, 128,  80, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
+    CASE(GGML_TYPE_Q2_0, 256, 1, 128,  96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
+    CASE(GGML_TYPE_Q2_0, 256, 1, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
+    CASE(GGML_TYPE_Q2_0, 256, 1, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
+
     CASE(GGML_TYPE_Q4_0, 256, 1, 128,   8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);
     CASE(GGML_TYPE_Q4_0, 256, 1, 128,  16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);
     CASE(GGML_TYPE_Q4_0, 256, 1, 128,  32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);

ggml/src/ggml-cuda/mmq-config-cdna.cuh

@@ -7,6 +7,14 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf
     CASE(GGML_TYPE_Q1_0, 512, 1, 128,  48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
     CASE(GGML_TYPE_Q1_0, 512, 1, 128,  64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
 
+    CASE(GGML_TYPE_Q2_0, 512, 1, 128,  16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);
+    CASE(GGML_TYPE_Q2_0, 512, 1, 128,  32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);
+    CASE(GGML_TYPE_Q2_0, 512, 1, 128,  64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);
+    CASE(GGML_TYPE_Q2_0, 512, 1, 128,  16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
+    CASE(GGML_TYPE_Q2_0, 512, 1, 128,  32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
+    CASE(GGML_TYPE_Q2_0, 512, 1, 128,  48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
+    CASE(GGML_TYPE_Q2_0, 512, 1, 128,  64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, false);
+
     CASE(GGML_TYPE_Q4_0, 512, 1, 128,  16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);
     CASE(GGML_TYPE_Q4_0, 512, 1, 128,  32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);
     CASE(GGML_TYPE_Q4_0, 512, 1, 128,  64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, true, true);

ggml/src/ggml-cuda/mmq-config-pascal.cuh

@@ -11,6 +11,18 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf
     CASE(GGML_TYPE_Q1_0, 256, 2, 64,  48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
     CASE(GGML_TYPE_Q1_0, 256, 2, 64,  64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
 
+    CASE(GGML_TYPE_Q2_0, 256, 2, 64,   8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 64,  16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 64,  32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 64,  64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 64,   8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 64,  16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 64,  24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 64,  32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 64,  40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 64,  48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 64,  64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+
     CASE(GGML_TYPE_Q4_0, 256, 2, 64,   8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
     CASE(GGML_TYPE_Q4_0, 256, 2, 64,  16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
     CASE(GGML_TYPE_Q4_0, 256, 2, 64,  32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);

ggml/src/ggml-cuda/mmq-config-rdna2.cuh

@@ -11,6 +11,18 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf
     CASE(GGML_TYPE_Q1_0, 256, 2, 128,  48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
     CASE(GGML_TYPE_Q1_0, 256, 2, 128,  64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
 
+    CASE(GGML_TYPE_Q2_0, 256, 2, 128,   8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 128,  16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 128,  32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 128,  64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 128,   8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 128,  16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 128,  24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 128,  32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 128,  40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 128,  48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 128,  64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+
     CASE(GGML_TYPE_Q4_0, 256, 2, 128,   8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
     CASE(GGML_TYPE_Q4_0, 256, 2, 128,  16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
     CASE(GGML_TYPE_Q4_0, 256, 2, 128,  32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);

ggml/src/ggml-cuda/mmq-config-rdna3-5.cuh

@@ -11,6 +11,18 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf
     CASE(GGML_TYPE_Q1_0, 256, 2, 128,  96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
     CASE(GGML_TYPE_Q1_0, 256, 2, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
     CASE(GGML_TYPE_Q1_0, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+    CASE(GGML_TYPE_Q2_0, 128, 2,  64,  16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
+    CASE(GGML_TYPE_Q2_0, 128, 2,  64,  32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 128,  64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
+    CASE(GGML_TYPE_Q2_0, 128, 2,  64,  16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+    CASE(GGML_TYPE_Q2_0, 128, 2,  64,  32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 128,  48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 128,  64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 128,  80, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 128,  96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
 
     CASE(GGML_TYPE_Q4_0, 128, 2,  64,  16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
     CASE(GGML_TYPE_Q4_0, 128, 2,  64,  32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);

ggml/src/ggml-cuda/mmq-config-rdna3.cuh

@@ -11,6 +11,18 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf
     CASE(GGML_TYPE_Q1_0, 256, 2, 128,  96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
     CASE(GGML_TYPE_Q1_0, 256, 2, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
     CASE(GGML_TYPE_Q1_0, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+    CASE(GGML_TYPE_Q2_0, 128, 2,  64,  16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
+    CASE(GGML_TYPE_Q2_0, 128, 2,  64,  32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 128,  64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
+    CASE(GGML_TYPE_Q2_0, 128, 2,  64,  16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+    CASE(GGML_TYPE_Q2_0, 128, 2,  64,  32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 128,  48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 128,  64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 128,  80, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 128,  96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
 
     CASE(GGML_TYPE_Q4_0, 128, 2,  64,  16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
     CASE(GGML_TYPE_Q4_0, 128, 2,  64,  32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);

ggml/src/ggml-cuda/mmq-config-rdna4.cuh

@@ -11,6 +11,18 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf
     CASE(GGML_TYPE_Q1_0, 256, 2, 128,  96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
     CASE(GGML_TYPE_Q1_0, 256, 2, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
     CASE(GGML_TYPE_Q1_0, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+    CASE(GGML_TYPE_Q2_0, 128, 2,  64,  16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
+    CASE(GGML_TYPE_Q2_0, 128, 2,  64,  32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
+    CASE(GGML_TYPE_Q2_0, 128, 2,  64,  64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
+    CASE(GGML_TYPE_Q2_0, 128, 2,  64,  16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+    CASE(GGML_TYPE_Q2_0, 128, 2,  64,  32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+    CASE(GGML_TYPE_Q2_0, 128, 2,  64,  48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+    CASE(GGML_TYPE_Q2_0, 128, 2,  64,  64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 128,  80, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 128,  96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
+    CASE(GGML_TYPE_Q2_0, 256, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
 
     CASE(GGML_TYPE_Q4_0, 128, 2,  64,  16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
     CASE(GGML_TYPE_Q4_0, 128, 2,  64,  32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);

ggml/src/ggml-cuda/mmq-load-tiles.cuh

@@ -95,6 +95,103 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
     }
 }
 
+template <ggml_type type, int J, bool fallback> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q2_0(
+        const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) {
+    constexpr int warp_size   = ggml_cuda_get_physical_warp_size();
+    constexpr int nwarps      = ggml_cuda_mmq_get_nthreads(type, J, fallback) / warp_size;
+    constexpr int I           = ggml_cuda_mmq_get_I(type, J, fallback);
+    constexpr int sram_stride = ggml_cuda_mmq_get_sram_stride(type, J, fallback);
+
+#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
+    int   * x_qs = (int   *)  x_tile;
+    float * x_df = (float *) (x_qs + 2*MMQ_TILE_NE_K);
+#else
+    constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_Q8_0, I);
+    int   * x_qs = (int   *)  x_tile;
+    float * x_df = (float *) (x_qs + txs.qs);
+#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
+
+    constexpr int blocks_per_iter = MMQ_ITER_K / QK2_0;
+    constexpr int threads_per_row = blocks_per_iter * QI2_0;
+    constexpr int nrows = warp_size / threads_per_row;
+    constexpr int scale_entries_per_block = QK2_0 / QK8_1;
+    constexpr int scale_entries_per_row = blocks_per_iter * scale_entries_per_block;
+
+    const int txi  = threadIdx.x % threads_per_row;
+    const int kbx  = txi / QI2_0;
+    const int kqsx = txi % QI2_0;
+
+#pragma unroll
+    for (int i0 = 0; i0 < I; i0 += nrows*nwarps) {
+        int i = i0 + threadIdx.y*nrows + threadIdx.x/threads_per_row;
+
+        if (fallback) {
+            i = min(i, i_max);
+        }
+
+        const block_q2_0 * bxi = (const block_q2_0 *) x + kbx0 + i*stride + kbx;
+        // Each 32-element chunk occupies 8 bytes of qs (32 elements * 2 bits = 64 bits)
+        const int qs_offset = 8*kqsx;
+        const int qs0 = bxi->qs[qs_offset + 0] | (bxi->qs[qs_offset + 1] << 8) |
+                        (bxi->qs[qs_offset + 2] << 16) | (bxi->qs[qs_offset + 3] << 24);
+        const int qs1 = bxi->qs[qs_offset + 4] | (bxi->qs[qs_offset + 5] << 8) |
+                        (bxi->qs[qs_offset + 6] << 16) | (bxi->qs[qs_offset + 7] << 24);
+
+        // Unpack 32 2-bit codes into 8 int32s, each holding 4 signed int8s in {-1,0,1,2}.
+        int unpacked_bytes[8];
+#pragma unroll
+        for (int j = 0; j < 4; ++j) {
+            const int shift = j * 8;
+            const int codes = (qs0 >> shift) & 0xFF;
+            const int c0 = ((codes >> 0) & 0x3) - 1;
+            const int c1 = ((codes >> 2) & 0x3) - 1;
+            const int c2 = ((codes >> 4) & 0x3) - 1;
+            const int c3 = ((codes >> 6) & 0x3) - 1;
+            unpacked_bytes[j] = (c0 & 0xFF) | ((c1 & 0xFF) << 8) | ((c2 & 0xFF) << 16) | ((c3 & 0xFF) << 24);
+        }
+#pragma unroll
+        for (int j = 0; j < 4; ++j) {
+            const int shift = j * 8;
+            const int codes = (qs1 >> shift) & 0xFF;
+            const int c0 = ((codes >> 0) & 0x3) - 1;
+            const int c1 = ((codes >> 2) & 0x3) - 1;
+            const int c2 = ((codes >> 4) & 0x3) - 1;
+            const int c3 = ((codes >> 6) & 0x3) - 1;
+            unpacked_bytes[4 + j] = (c0 & 0xFF) | ((c1 & 0xFF) << 8) | ((c2 & 0xFF) << 16) | ((c3 & 0xFF) << 24);
+        }
+
+        const int dst_offset = kbx*(scale_entries_per_block*QI8_0) + kqsx*QI8_0;
+#pragma unroll
+        for (int j = 0; j < 8; ++j) {
+#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
+            x_qs[i*sram_stride           + dst_offset + j] = unpacked_bytes[j];
+#else
+            x_qs[i*(2*MMQ_TILE_NE_K + 1) + dst_offset + j] = unpacked_bytes[j];
+#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
+        }
+    }
+
+    const int ksx = threadIdx.x % scale_entries_per_row;
+    const int scale_block = ksx / scale_entries_per_block;
+
+#pragma unroll
+    for (int i0 = 0; i0 < I; i0 += nwarps) {
+        int i = i0 + threadIdx.y;
+
+        if (fallback) {
+            i = min(i, i_max);
+        }
+
+        const block_q2_0 * bxi = (const block_q2_0 *) x + kbx0 + i*stride + scale_block;
+
+#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
+        x_df[i*sram_stride                           + ksx] = bxi->d;
+#else
+        x_df[i*(2*MMQ_TILE_NE_K/QI8_0) + i/(QI8_0/2) + ksx] = bxi->d;
+#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
+    }
+}
+
 template <ggml_type type, int J, bool fallback> static __device__ __forceinline__ void ggml_cuda_mmq_load_tiles_q4_0(
         const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) {
     constexpr int warp_size   = ggml_cuda_get_physical_warp_size();

ggml/src/ggml-cuda/mmq.cu

@@ -10,6 +10,9 @@ static void ggml_cuda_mul_mat_q_switch_type(ggml_backend_cuda_context & ctx, con
         case GGML_TYPE_Q1_0:
             mul_mat_q_case<GGML_TYPE_Q1_0>(ctx, args, stream);
             break;
+        case GGML_TYPE_Q2_0:
+            mul_mat_q_case<GGML_TYPE_Q2_0>(ctx, args, stream);
+            break;
         case GGML_TYPE_Q4_0:
             mul_mat_q_case<GGML_TYPE_Q4_0>(ctx, args, stream);
             break;
@@ -262,6 +265,7 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t
 
     switch (type) {
         case GGML_TYPE_Q1_0:
+        case GGML_TYPE_Q2_0:
         case GGML_TYPE_Q4_0:
         case GGML_TYPE_Q4_1:
         case GGML_TYPE_Q5_0:

ggml/src/ggml-cuda/mmq.cuh

@@ -60,6 +60,7 @@ static_assert(sizeof(block_fp4_mmq)  == sizeof(block_q8_1_mmq),    "Unexpected b
 static mmq_q8_1_ds_layout mmq_get_q8_1_ds_layout(const ggml_type type_x) {
     switch (type_x) {
         case GGML_TYPE_Q1_0:
+        case GGML_TYPE_Q2_0:
             return MMQ_Q8_1_DS_LAYOUT_D4;
         case GGML_TYPE_Q4_0:
         case GGML_TYPE_Q4_1:
@@ -385,6 +386,7 @@ static constexpr __device__ int ggml_cuda_mmq_get_rows_per_warp(ggml_type type,
 static constexpr __host__ __device__ tile_x_sizes mmq_get_dp4a_tile_x_sizes(ggml_type type, int I) {
     switch (type) {
         case GGML_TYPE_Q1_0:    return MMQ_DP4A_TXS_Q8_0;
+        case GGML_TYPE_Q2_0:    return MMQ_DP4A_TXS_Q8_0;
         case GGML_TYPE_Q4_0:    return MMQ_DP4A_TXS_Q4_0;
         case GGML_TYPE_Q4_1:    return MMQ_DP4A_TXS_Q4_1;
         case GGML_TYPE_Q5_0:    return MMQ_DP4A_TXS_Q8_0;
@@ -542,6 +544,12 @@ static constexpr __device__ ggml_cuda_mmq_util_funcs ggml_cuda_mmq_get_util_func
                     ggml_cuda_mmq_load_tiles_q1_0<type, J, fallback>,
                     ggml_cuda_mmq_vec_dot_q8_0_q8_1_dp4a<type, J, fallback>,
                     ggml_cuda_mmq_write_back_dp4a<type, J, fallback>);
+            case GGML_TYPE_Q2_0:
+                return ggml_cuda_mmq_util_funcs(
+                    VDR_Q2_0_Q8_1_MMQ,
+                    ggml_cuda_mmq_load_tiles_q2_0<type, J, fallback>,
+                    ggml_cuda_mmq_vec_dot_q8_0_q8_1_dp4a<type, J, fallback>,
+                    ggml_cuda_mmq_write_back_dp4a<type, J, fallback>);
             case GGML_TYPE_Q4_0:
                 return ggml_cuda_mmq_util_funcs(
                     VDR_Q4_0_Q8_1_MMQ,
@@ -700,6 +708,12 @@ static constexpr __device__ ggml_cuda_mmq_util_funcs ggml_cuda_mmq_get_util_func
                 ggml_cuda_mmq_load_tiles_q1_0<type, J, fallback>,
                 ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma<type, J, fallback, MMQ_Q8_1_DS_LAYOUT_D4>,
                 ggml_cuda_mmq_write_back_mma<type, J, fallback>);
+        case GGML_TYPE_Q2_0:
+            return ggml_cuda_mmq_util_funcs(
+                -1,
+                ggml_cuda_mmq_load_tiles_q2_0<type, J, fallback>,
+                ggml_cuda_mmq_vec_dot_q8_0_q8_1_mma<type, J, fallback, MMQ_Q8_1_DS_LAYOUT_D4>,
+                ggml_cuda_mmq_write_back_mma<type, J, fallback>);
         case GGML_TYPE_Q4_0:
             return ggml_cuda_mmq_util_funcs(
                 -1,
@@ -1550,6 +1564,7 @@ void mul_mat_q_case(ggml_backend_cuda_context & ctx, const mmq_args & args, cuda
     template void mul_mat_q_case<type>(ggml_backend_cuda_context & ctx, const mmq_args & args, cudaStream_t stream) \
 
 extern DECL_MMQ_CASE(GGML_TYPE_Q1_0);
+extern DECL_MMQ_CASE(GGML_TYPE_Q2_0);
 extern DECL_MMQ_CASE(GGML_TYPE_Q4_0);
 extern DECL_MMQ_CASE(GGML_TYPE_Q4_1);
 extern DECL_MMQ_CASE(GGML_TYPE_Q5_0);

ggml/src/ggml-cuda/mmvq.cu

@@ -10,6 +10,7 @@ typedef float (*vec_dot_q_cuda_t)(const void * __restrict__ vbq, const block_q8_
 static constexpr __device__ vec_dot_q_cuda_t get_vec_dot_q_cuda(ggml_type type) {
     switch (type) {
         case GGML_TYPE_Q1_0:    return vec_dot_q1_0_q8_1;
+        case GGML_TYPE_Q2_0:    return vec_dot_q2_0_q8_1;
         case GGML_TYPE_Q4_0:    return vec_dot_q4_0_q8_1;
         case GGML_TYPE_Q4_1:    return vec_dot_q4_1_q8_1;
         case GGML_TYPE_Q5_0:    return vec_dot_q5_0_q8_1;
@@ -38,6 +39,7 @@ static constexpr __device__ vec_dot_q_cuda_t get_vec_dot_q_cuda(ggml_type type)
 static constexpr __host__ __device__ int get_vdr_mmvq(ggml_type type) {
     switch (type) {
         case GGML_TYPE_Q1_0:    return VDR_Q1_0_Q8_1_MMVQ;
+        case GGML_TYPE_Q2_0:    return VDR_Q2_0_Q8_1_MMVQ;
         case GGML_TYPE_Q4_0:    return VDR_Q4_0_Q8_1_MMVQ;
         case GGML_TYPE_Q4_1:    return VDR_Q4_1_Q8_1_MMVQ;
         case GGML_TYPE_Q5_0:    return VDR_Q5_0_Q8_1_MMVQ;
@@ -1010,6 +1012,12 @@ static void mul_mat_vec_q_switch_type(
                  nchannels_x, nchannels_y, nchannels_dst, stride_channel_x, stride_channel_y, stride_channel_dst,
                  nsamples_x, nsamples_dst, stride_sample_x, stride_sample_y, stride_sample_dst, ids_stride, stream);
             break;
+        case GGML_TYPE_Q2_0:
+            mul_mat_vec_q_switch_ncols_dst<GGML_TYPE_Q2_0>
+                (vx, vy, ids, fusion, dst, ncols_x, nrows_x, ncols_dst, stride_row_x, stride_col_y, stride_col_dst,
+                 nchannels_x, nchannels_y, nchannels_dst, stride_channel_x, stride_channel_y, stride_channel_dst,
+                 nsamples_x, nsamples_dst, stride_sample_x, stride_sample_y, stride_sample_dst, ids_stride, stream);
+            break;
         case GGML_TYPE_Q4_0:
             mul_mat_vec_q_switch_ncols_dst<GGML_TYPE_Q4_0>
                 (vx, vy, ids, fusion, dst, ncols_x, nrows_x, ncols_dst, stride_row_x, stride_col_y, stride_col_dst,

ggml/src/ggml-cuda/template-instances/generate_cu_files.py

@@ -36,6 +36,7 @@ SOURCE_FATTN_MMA_CASE = "DECL_FATTN_MMA_F16_CASE({head_size_kq}, {head_size_v},
 
 TYPES_MMQ = [
     "GGML_TYPE_Q1_0",
+    "GGML_TYPE_Q2_0",
     "GGML_TYPE_Q4_0", "GGML_TYPE_Q4_1", "GGML_TYPE_Q5_0", "GGML_TYPE_Q5_1", "GGML_TYPE_Q8_0",
     "GGML_TYPE_Q2_K", "GGML_TYPE_Q3_K", "GGML_TYPE_Q4_K", "GGML_TYPE_Q5_K", "GGML_TYPE_Q6_K",
     "GGML_TYPE_IQ2_XXS", "GGML_TYPE_IQ2_XS", "GGML_TYPE_IQ2_S", "GGML_TYPE_IQ3_XXS", "GGML_TYPE_IQ3_S",

ggml/src/ggml-cuda/template-instances/mmq-instance-q2_0.cu

@@ -0,0 +1,5 @@
+// This file has been autogenerated by generate_cu_files.py, do not edit manually.
+
+#include "../mmq.cuh"
+
+DECL_MMQ_CASE(GGML_TYPE_Q2_0);

ggml/src/ggml-cuda/vecdotq.cuh

@@ -109,6 +109,9 @@ static __device__ __forceinline__ uint32_t unpack_ksigns(const uint8_t v) {
 #define VDR_Q1_0_Q8_1_MMVQ 1  // Process one 32-element chunk at a time for parallelism
 #define VDR_Q1_0_Q8_1_MMQ  4  // Q1_0 has 128 bits (4 ints) per block
 
+#define VDR_Q2_0_Q8_1_MMVQ 1  // Process one 32-element chunk at a time for parallelism
+#define VDR_Q2_0_Q8_1_MMQ  2  // Q2_0 group 64: 128 bits (4 ints) per block, 2 32-element chunks
+
 #define VDR_Q4_0_Q8_1_MMVQ 2
 #define VDR_Q4_0_Q8_1_MMQ  4
 
@@ -722,6 +725,64 @@ static __device__ __forceinline__ float vec_dot_q1_0_q8_1(
     return d1 * d8 * sumi;
 }
 
+static __device__ __forceinline__ float vec_dot_q2_0_q8_1(
+    const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) {
+
+    const block_q2_0 * bq2_0 = (const block_q2_0 *) vbq + kbx;
+
+    // Q2_0 (group 64): 64 elements with ONE scale, 2 bits per element (4 elements per byte)
+    // Q8_1: 32 elements per block with individual scales
+    // iqs selects which of the 2 chunks of 32 elements to process (0-1)
+
+    const float d2 = bq2_0->d;
+
+    // Process only the chunk specified by iqs
+    const block_q8_1 * bq8_1_chunk = bq8_1 + iqs;
+
+    // Load 64 bits (8 bytes) for this chunk from Q2_0: bytes [8*iqs, 8*iqs+8)
+    const int offset = iqs * 8;
+    const int v0 = bq2_0->qs[offset + 0] | (bq2_0->qs[offset + 1] << 8) |
+                   (bq2_0->qs[offset + 2] << 16) | (bq2_0->qs[offset + 3] << 24);
+    const int v1 = bq2_0->qs[offset + 4] | (bq2_0->qs[offset + 5] << 8) |
+                   (bq2_0->qs[offset + 6] << 16) | (bq2_0->qs[offset + 7] << 24);
+
+    // Unpack 32 2-bit codes into 8 int32s, each holding 4 signed int8 symbols in {-1,0,1,2}.
+    // Stored code c in {0,1,2,3} -> symbol s = c - 1.
+    int vi_bytes[8];
+#pragma unroll
+    for (int j = 0; j < 4; ++j) {
+        const int shift = j * 8;
+        const int codes = (v0 >> shift) & 0xFF;
+        const int c0 = ((codes >> 0) & 0x3) - 1;
+        const int c1 = ((codes >> 2) & 0x3) - 1;
+        const int c2 = ((codes >> 4) & 0x3) - 1;
+        const int c3 = ((codes >> 6) & 0x3) - 1;
+        vi_bytes[j] = (c0 & 0xFF) | ((c1 & 0xFF) << 8) | ((c2 & 0xFF) << 16) | ((c3 & 0xFF) << 24);
+    }
+#pragma unroll
+    for (int j = 0; j < 4; ++j) {
+        const int shift = j * 8;
+        const int codes = (v1 >> shift) & 0xFF;
+        const int c0 = ((codes >> 0) & 0x3) - 1;
+        const int c1 = ((codes >> 2) & 0x3) - 1;
+        const int c2 = ((codes >> 4) & 0x3) - 1;
+        const int c3 = ((codes >> 6) & 0x3) - 1;
+        vi_bytes[4 + j] = (c0 & 0xFF) | ((c1 & 0xFF) << 8) | ((c2 & 0xFF) << 16) | ((c3 & 0xFF) << 24);
+    }
+
+    // Compute dot product for this 32-element chunk
+    int sumi = 0;
+#pragma unroll
+    for (int j = 0; j < 8; ++j) {
+        const int u = get_int_b4(bq8_1_chunk->qs, j);
+        sumi = ggml_cuda_dp4a(vi_bytes[j], u, sumi);
+    }
+
+    // Apply Q2_0's single scale and this chunk's Q8_1 scale
+    const float d8 = __low2float(bq8_1_chunk->ds);
+    return d2 * d8 * sumi;
+}
+
 static __device__ __forceinline__ float vec_dot_q4_0_q8_1(
     const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) {
 

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

@@ -213,7 +213,7 @@ typedef void * ggml_metal_rset_t;
 // a collection of residency sets (non-owning)
 typedef struct ggml_metal_rsets * ggml_metal_rsets_t;
 
-ggml_metal_rsets_t ggml_metal_rsets_init(void);
+ggml_metal_rsets_t ggml_metal_rsets_init(ggml_metal_device_t dev);
 void ggml_metal_rsets_free(ggml_metal_rsets_t rsets);
 
 //

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

@@ -557,7 +557,32 @@ struct ggml_metal_rsets {
     dispatch_group_t d_group;
 };
 
-ggml_metal_rsets_t ggml_metal_rsets_init(void) {
+#if defined(GGML_METAL_HAS_RESIDENCY_SETS)
+static void ggml_metal_dummy_work(ggml_metal_device_t dev) {
+    if (dev->mtl_queue == nil) {
+        return;
+    }
+
+    @autoreleasepool {
+        // perform a minimal dummy operation on the GPU
+        id<MTLBuffer> buf = [dev->mtl_device newBufferWithLength:1 options:MTLResourceStorageModePrivate];
+        id<MTLCommandBuffer> cmd_buf = [dev->mtl_queue commandBuffer];
+
+        {
+            id<MTLBlitCommandEncoder> encoder = [cmd_buf blitCommandEncoder];
+
+            [encoder fillBuffer:buf range:NSMakeRange(0, 1) value:0];
+
+            [encoder endEncoding];
+        }
+
+        [cmd_buf commit];
+        [buf release];
+    }
+}
+#endif
+
+ggml_metal_rsets_t ggml_metal_rsets_init(ggml_metal_device_t dev) {
     ggml_metal_rsets_t res = calloc(1, sizeof(struct ggml_metal_rsets));
 
     res->lock = [[NSLock alloc] init];
@@ -610,6 +635,15 @@ ggml_metal_rsets_t ggml_metal_rsets_init(void) {
 #endif
     });
 
+#if defined(GGML_METAL_HAS_RESIDENCY_SETS)
+    if (@available(macOS 15.0, iOS 18.0, tvOS 18.0, visionOS 2.0, *)) {
+        // workaround for residency set memory not being released if no GPU operation occurs
+        // https://developer.apple.com/forums/thread/839089
+        // https://github.com/ggml-org/llama.cpp/issues/25937
+        ggml_metal_dummy_work(dev);
+    }
+#endif
+
     return res;
 }
 
@@ -864,7 +898,7 @@ ggml_metal_device_t ggml_metal_device_init(int device) {
             }
 
             if (dev->props.use_residency_sets) {
-                dev->rsets = ggml_metal_rsets_init();
+                dev->rsets = ggml_metal_rsets_init(dev);
             } else {
                 dev->rsets = nil;
             }
@@ -1484,6 +1518,7 @@ static void ggml_metal_buffer_rset_free(ggml_metal_buffer_t buf) {
         if (buf->rset) {
             [buf->rset endResidency];
             [buf->rset removeAllAllocations];
+            [buf->rset commit];
             [buf->rset release];
         }
     }

scripts/sync-ggml.last

@@ -1 +1 @@
-9be313313c8ecb9488911bd64550190e3ed80f38
+06ca97616793248fadb410ea8d69c7511b2005e4

src/models/minimax-m3.cpp

@@ -2,7 +2,6 @@
 #include "llama-kv-cache.h"
 #include <cmath>
 #include <vector>
-#include <algorithm>
 #include <cstdint>
 
 // MiniMax-M3: MiniMax-M2 style GQA (per-head QK-norm, partial rotary) with
@@ -126,68 +125,6 @@ public:
     int64_t nblk;
 };
 
-// pooled score of a block with no visible token: -inf from the mask, or -FLT_MAX from the
-// max-pool identity when every element of the block is -inf
-static inline bool msa_score_masked(float x) { return x <= -1e30f; }
-
-// MSA block selection (batch regime)
-// CPU custom op, the token-level expansion and the combination with the causal mask happen on the GPU.
-static void msa_block_mask_op(struct ggml_tensor * dst, int ith, int nth, void * userdata) {
-    const struct ggml_tensor * bs   = dst->src[0];
-    const struct ggml_tensor * bias = dst->src[1];
-    const msa_params * p = (const msa_params *) userdata;
-
-    const int nblk = (int) bs->ne[0];
-    const int Hd   = (int) bs->ne[1];
-    const int S    = (int) bs->ne[2];
-
-    GGML_ASSERT(bs->type   == GGML_TYPE_F32 && ggml_is_contiguous(bs));
-    GGML_ASSERT(bias->type == GGML_TYPE_F32 && ggml_is_contiguous(bias));
-    GGML_ASSERT(dst->type  == GGML_TYPE_F16 && ggml_is_contiguous(dst));
-    GGML_ASSERT(dst->ne[0] == nblk && dst->ne[1] == S && dst->ne[2] == Hd);
-    GGML_ASSERT(bias->ne[0] == nblk && bias->ne[1] == S);
-
-    const int topk = p->topk_blocks < nblk ? p->topk_blocks : nblk;
-
-    const ggml_fp16_t f16_zero = ggml_fp32_to_fp16(0.0f);
-    const ggml_fp16_t f16_ninf = ggml_fp32_to_fp16(-INFINITY);
-
-    std::vector<float> rank(nblk);
-    std::vector<char>  valid(nblk);
-    std::vector<int>   ord(nblk);
-
-    ggml_fp16_t * out = (ggml_fp16_t *) dst->data;
-
-    for (int i = ith; i < S; i += nth) {
-        const float * bias_col = (const float *) bias->data + (size_t) i * nblk;
-        for (int h = 0; h < Hd; ++h) {
-            const float * bs_col = (const float *) bs->data + ((size_t) i * Hd + h) * nblk;
-
-            for (int bk = 0; bk < nblk; ++bk) {
-                // a block is selectable if it has a visible token or is locally forced
-                valid[bk] = !msa_score_masked(bs_col[bk]) || bias_col[bk] > 0.0f;
-                rank [bk] = bias_col[bk] > 0.0f ? bias_col[bk] : bs_col[bk];
-                ord  [bk] = bk;
-            }
-
-            std::partial_sort(ord.begin(), ord.begin() + topk, ord.end(),
-                              [&](int a, int b) { return rank[a] > rank[b]; });
-
-            ggml_fp16_t * dst_col = out + ((size_t) h * S + i) * nblk;
-            for (int bk = 0; bk < nblk; ++bk) {
-                dst_col[bk] = f16_ninf;
-            }
-            for (int t = 0; t < topk; ++t) {
-                const int bk = ord[t];
-                if (!valid[bk]) {
-                    break;   // sorted desc: first invalid -> fewer than topk selectable blocks
-                }
-                dst_col[bk] = f16_zero;
-            }
-        }
-    }
-}
-
 // One FA call for all GQA groups (and at multi-stream decode, all streams) by mapping them onto the FA sequence dim (ne[3])
 ggml_tensor * llama_model_minimax_m3::graph::build_attn_msa_fa(
         ggml_tensor * q_cur,   // [D, HQ, T]
@@ -433,8 +370,8 @@ llama_model_minimax_m3::graph::graph(const llama_model & model, const llm_graph_
                                 msa_mf->nb[1], msa_mf->nb[1], st*msa_mf->nb[3]);
                         ggml_tensor * km_s = ggml_view_3d(ctx0, msa_kqm, n_kv, n_tps, 1,
                                 msa_kqm->nb[1], msa_kqm->nb[3], st*msa_kqm->nb[3]);
-                        ggml_tensor * bias_s = ggml_view_2d(ctx0, msa_loc->bias, nblk, n_tps,
-                                msa_loc->bias->nb[1], st*n_tps*msa_loc->bias->nb[1]);
+                        ggml_tensor * bias_s = ggml_view_3d(ctx0, msa_loc->bias, nblk, 1, n_tps,
+                                msa_loc->bias->nb[1], msa_loc->bias->nb[1], st*n_tps*msa_loc->bias->nb[1]);
                         ggml_tensor * q_s = ggml_view_3d(ctx0, Qcur, D, n_head, n_tps,
                                 Qcur->nb[1], Qcur->nb[2], st*n_tps*Qcur->nb[2]);
                         ggml_tensor * k_s = ggml_view_4d(ctx0, k, D, HKV, n_kv, 1,
@@ -453,15 +390,27 @@ llama_model_minimax_m3::graph::graph(const llama_model & model, const llm_graph_
                         ggml_tensor * bs = ggml_pool_2d(ctx0, sc, GGML_OP_POOL_MAX, blk, 1, blk, 1, 0, 0);
                         cb(bs, "msa_bs", il);
 
-                        // block-level 0/-inf keep mask on the CPU, tiny transfer
-                        ggml_tensor * srcs[2] = { bs, bias_s };
-                        ggml_tensor * bm = ggml_custom_4d(ctx0, GGML_TYPE_F16,
-                                nblk, n_tps, Hd, 1,
-                                srcs, 2, msa_block_mask_op, GGML_N_TASKS_MAX,
-                                const_cast<msa_params *>(&mm.msa_p));
+                        // bias the scores so locally-forced blocks always rank first
+                        ggml_tensor * bsf = ggml_add(ctx0, bs, bias_s);   // [nblk, Hd, n_tps]
+                        cb(bsf, "msa_bsf", il);
+
+                        ggml_tensor * idx = ggml_top_k(ctx0, bsf, K);   // [K, Hd, n_tps] i32
+
+                        ggml_tensor * ninf = ggml_cast(ctx0,
+                                ggml_scale_bias(ctx0, bias_s, 0.0f, -1e30f),
+                                GGML_TYPE_F16);                              // [nblk, 1, n_tps]
+                        ninf = ggml_repeat_4d(ctx0, ninf, nblk, Hd, n_tps, 1);
+                        ggml_tensor * zero = ggml_scale(ctx0,
+                                ggml_cast(ctx0, idx, GGML_TYPE_F32), 0.0f);
+                        ggml_tensor * bm = ggml_set_rows(ctx0,
+                                ggml_reshape_3d(ctx0, ninf, 1, nblk, Hd*n_tps),
+                                ggml_reshape_3d(ctx0, zero, 1, K,    Hd*n_tps),
+                                ggml_reshape_2d(ctx0, idx,     K,    Hd*n_tps));
+                        bm = ggml_reshape_3d(ctx0, bm, nblk, Hd, n_tps);
+                        bm = ggml_cont(ctx0, ggml_permute(ctx0, bm, 0, 2, 1, 3)); // [nblk, n_tps, Hd]
                         cb(bm, "msa_block_mask", il);
 
-                        // expand block -> token granularity on the GPU (j = bk*blk + t),
+                        // expand block -> token granularity (j = bk*blk + t),
                         // then combine with the causal mask in place
                         ggml_tensor * bmx = ggml_repeat_4d(ctx0,
                                 ggml_reshape_3d(ctx0, bm, 1, nblk, n_tps*Hd),

tests/CMakeLists.txt

@@ -278,6 +278,9 @@ set_tests_properties(test-state-restore-fragmented PROPERTIES FIXTURES_REQUIRED
 llama_build_and_test(test-save-load-state.cpp LABEL "model" ARGS -m "${MODEL_DEST}")
 set_tests_properties(test-save-load-state PROPERTIES FIXTURES_REQUIRED test-download-model)
 
+if (APPLE)
+    llama_build(test-rset-release.cpp get-model.cpp)
+endif()
 if (NOT GGML_BACKEND_DL)
     # these tests use the backends directly and cannot be built with dynamic loading
     llama_build_and_test(test-barrier.cpp)

tests/test-backend-ops.cpp

@@ -8013,6 +8013,7 @@ static const ggml_type other_types[] = {
     GGML_TYPE_Q5_0, GGML_TYPE_Q5_1,
     GGML_TYPE_Q8_0,
     GGML_TYPE_Q1_0,
+    GGML_TYPE_Q2_0,
     GGML_TYPE_Q2_K, GGML_TYPE_Q3_K,
     GGML_TYPE_Q5_K,
     GGML_TYPE_Q6_K,
@@ -9563,6 +9564,10 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
     test_cases.emplace_back(new test_flash_attn_ext(128, 64, 4, {1, 1}, 128, 2, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q1_0, GGML_TYPE_Q4_0));
     test_cases.emplace_back(new test_flash_attn_ext(64, 128, 4, {1, 1}, 128, 2, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q4_0, GGML_TYPE_Q1_0));
     test_cases.emplace_back(new test_flash_attn_ext(128, 64, 4, {1, 1}, 64, 2, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q1_0, GGML_TYPE_F16));
+    test_cases.emplace_back(new test_flash_attn_ext(128, 128, 4, {1, 1}, 96, 2, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q2_0, GGML_TYPE_Q2_0));
+    test_cases.emplace_back(new test_flash_attn_ext(128, 64, 4, {1, 1}, 128, 2, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q2_0, GGML_TYPE_Q4_0));
+    test_cases.emplace_back(new test_flash_attn_ext(64, 128, 4, {1, 1}, 128, 2, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q4_0, GGML_TYPE_Q2_0));
+    test_cases.emplace_back(new test_flash_attn_ext(128, 64, 4, {1, 1}, 64, 2, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_Q2_0, GGML_TYPE_F16));
 
     // large-KV F16 cases (Qwen3.6-27B geometry and a llama-class control): the upstream matrix
     // stops at kv=1024, blind to long-context FA bugs (e.g. the oneDNN SDPA ordering race on BMG).

tests/test-rset-release.cpp

@@ -0,0 +1,53 @@
+// ref: https://github.com/ggml-org/llama.cpp/issues/25937
+// only works reliably when run with a large model that occupies 3GB+ of wired memory
+// thus, this test is not run by default
+// example model to run with: google/gemma-4-E4B-it-qat-q4_0-gguf
+
+#include <cstdint>
+#include <mach/mach.h>
+#include <mach/mach_host.h>
+#include <unistd.h>
+
+#include "llama.h"
+#include "get-model.h"
+
+static uint64_t wired_memory() {
+    vm_statistics64_data_t vmstat;
+    mach_msg_type_number_t count = HOST_VM_INFO64_COUNT;
+    if (host_statistics64(mach_host_self(), HOST_VM_INFO64, (host_info64_t)&vmstat, &count) != KERN_SUCCESS) {
+        return UINT64_MAX;
+    }
+    return static_cast<uint64_t>(vmstat.wire_count) * vm_kernel_page_size;
+}
+
+int main(int argc, char ** argv) {
+    auto * model_path = get_model_or_exit(argc, argv);
+
+    llama_backend_init();
+
+    const uint64_t wired_initial = wired_memory();
+
+    llama_model_params params = llama_model_default_params();
+    params.load_mode = LLAMA_LOAD_MODE_NONE;
+    struct llama_model* model = llama_model_load_from_file(model_path, params);
+
+    const uint64_t wired_loaded = wired_memory();
+    const uint64_t wired_delta = wired_loaded - wired_initial;
+    // system memory fluctuates, so we need to allocate enough to reliably detect the release
+    GGML_ASSERT(wired_delta > 2'000'000'000); // 2GB
+
+    llama_model_free(model);
+
+    const uint64_t t_start_ms = ggml_time_ms();
+
+    // expect most of the allocated memory to be released within 10 seconds
+    // we allow for some tolerance due to system-wide memory fluctuations
+    while (wired_memory() > wired_loaded - 0.75 * wired_delta) {
+        GGML_ASSERT(ggml_time_ms() - t_start_ms < 10'000);
+        usleep(100'000); // 100ms
+    }
+
+    llama_backend_free();
+
+    return 0;
+}

tools/ui/src/lib/components/app/chat/ChatForm/ChatFormActions/ChatFormActionAdd/ChatFormActionAddDropdown.svelte

@@ -48,6 +48,9 @@
 	}: Props = $props();
 
 	let dropdownOpen = $state(false);
+	// The system message action moves focus to the message editor, so the menu
+	// must not restore focus to the trigger on close
+	let suppressCloseAutoFocus = false;
 
 	function handleMcpSettingsClick() {
 		dropdownOpen = false;
@@ -96,7 +99,16 @@
 			</Tooltip.Content>
 		</Tooltip.Root>
 
-		<DropdownMenu.Content align="start" class="w-52">
+		<DropdownMenu.Content
+			align="start"
+			class="w-52"
+			onCloseAutoFocus={(e) => {
+				if (suppressCloseAutoFocus) {
+					suppressCloseAutoFocus = false;
+					e.preventDefault();
+				}
+			}}
+		>
 			<ChatFormActionAddReasoningSubmenu />
 
 			<DropdownMenu.Separator />
@@ -148,7 +160,10 @@
 
 			<DropdownMenu.Item
 				class="flex cursor-pointer items-center gap-2"
-				onclick={onSystemPromptClick}
+				onclick={() => {
+					suppressCloseAutoFocus = true;
+					onSystemPromptClick?.();
+				}}
 			>
 				<MessageSquare class={ICON_CLASS_DEFAULT} />
 

tools/ui/src/lib/components/app/chat/ChatMessages/ChatMessage/ChatMessage.svelte

@@ -2,6 +2,7 @@
 	import { goto } from '$app/navigation';
 	import { getChatActionsContext, setMessageEditContext } from '$lib/contexts';
 	import { chatStore, pendingEditMessageId } from '$lib/stores/chat.svelte';
+	import { isMobile } from '$lib/stores/viewport.svelte';
 	import { conversationsStore } from '$lib/stores/conversations.svelte';
 	import { DatabaseService } from '$lib/services/database.service';
 	import { SYSTEM_MESSAGE_PLACEHOLDER } from '$lib/constants';
@@ -46,7 +47,14 @@
 		assistantMessages: number;
 		messageTypes: string[];
 	} | null>(null);
-	let editedContent = $derived(message.content);
+	// The system message placeholder must never surface as editable content; keeping
+	// it in the derived (not just in handleEdit) guards against prop invalidation
+	// reverting the override while editing
+	let editedContent = $derived(
+		message.role === MessageRole.SYSTEM && message.content === SYSTEM_MESSAGE_PLACEHOLDER
+			? ''
+			: message.content
+	);
 
 	let rawEditContent = $derived.by(() => {
 		if (message.role !== MessageRole.ASSISTANT) return undefined;
@@ -265,6 +273,12 @@
 		chatActions.navigateToSibling(siblingId);
 	}
 
+	// After the system message flow ends, hand focus to the main chat form
+	function focusMainChatForm() {
+		if (isMobile.current) return;
+		document.querySelector<HTMLTextAreaElement>('.chat-screen-form-wrapper textarea')?.focus();
+	}
+
 	async function handleSaveEdit() {
 		if (message.role === MessageRole.SYSTEM) {
 			// System messages: update in place without branching
@@ -276,6 +290,8 @@
 				isEditing = false;
 				if (conversationDeleted) {
 					goto(ROUTES.START);
+				} else {
+					focusMainChatForm();
 				}
 				return;
 			}
@@ -285,6 +301,7 @@
 			if (index !== -1) {
 				conversationsStore.updateMessageAtIndex(index, { content: newContent });
 			}
+			focusMainChatForm();
 		} else if (message.role === MessageRole.USER) {
 			const finalExtras = await getMergedExtras();
 			chatActions.editWithBranching(message, editedContent.trim(), finalExtras);

tools/ui/src/lib/components/app/chat/ChatScreen/ChatScreenForm.svelte

@@ -106,15 +106,23 @@
 		onFileRemove?.(fileId);
 	}
 
+	// Auto-focus must not steal focus already claimed elsewhere (e.g. the system
+	// message editor opened just before a navigation)
+	function focusFormUnlessCaptured() {
+		const active = document.activeElement;
+		if (active instanceof HTMLTextAreaElement || active instanceof HTMLInputElement) return;
+		chatFormRef?.focus();
+	}
+
 	onMount(() => {
 		if (!isMobile.current) {
-			setTimeout(() => chatFormRef?.focus(), 100);
+			setTimeout(focusFormUnlessCaptured, 100);
 		}
 	});
 
 	afterNavigate((navigation) => {
 		if (navigation?.from != null && !isMobile.current) {
-			setTimeout(() => chatFormRef?.focus(), 100);
+			setTimeout(focusFormUnlessCaptured, 100);
 		}
 	});
 
@@ -127,7 +135,7 @@
 
 	$effect(() => {
 		if (previousIsLoading && !isLoading) {
-			setTimeout(() => chatFormRef?.focus(), 10);
+			setTimeout(focusFormUnlessCaptured, 10);
 		}
 
 		previousIsLoading = isLoading;

tools/ui/src/lib/services/database.service.ts

@@ -31,14 +31,19 @@ export class DatabaseService {
 	 * Creates a new conversation.
 	 *
 	 * @param name - Name of the conversation
+	 * @param fields - Optional extra fields (e.g. reasoningEffort)
 	 * @returns The created conversation
 	 */
-	static async createConversation(name: string): Promise<DatabaseConversation> {
+	static async createConversation(
+		name: string,
+		fields?: Partial<Omit<DatabaseConversation, 'id' | 'name' | 'lastModified'>>
+	): Promise<DatabaseConversation> {
 		const conversation: DatabaseConversation = {
 			id: uuid(),
 			name,
 			lastModified: Date.now(),
-			currNode: ''
+			currNode: '',
+			...fields
 		};
 
 		await db[IDXDB_TABLES.conversations].add(conversation);
@@ -137,7 +142,7 @@ export class DatabaseService {
 	 * @param systemPrompt - The system prompt content (must be non-empty)
 	 * @param parentId - Parent message ID (typically the root message)
 	 * @returns The created system message
-	 * @throws Error if systemPrompt is empty
+	 * @throws Error if systemPrompt is empty or the parent message does not exist
 	 */
 	static async createSystemMessage(
 		convId: string,
@@ -149,27 +154,30 @@ export class DatabaseService {
 			throw new Error('Cannot create system message with empty content');
 		}
 
-		const systemMessage: DatabaseMessage = {
-			id: uuid(),
-			convId,
-			type: MessageRole.SYSTEM,
-			timestamp: Date.now(),
-			role: MessageRole.SYSTEM,
-			content: trimmedPrompt,
-			parent: parentId,
-			children: []
-		};
+		return await db.transaction('rw', db[IDXDB_TABLES.messages], async () => {
+			const parentMessage = await db[IDXDB_TABLES.messages].get(parentId);
+			if (!parentMessage) {
+				throw new Error(`Parent message ${parentId} not found`);
+			}
 
-		await db[IDXDB_TABLES.messages].add(systemMessage);
+			const systemMessage: DatabaseMessage = {
+				id: uuid(),
+				convId,
+				type: MessageRole.SYSTEM,
+				timestamp: Date.now(),
+				role: MessageRole.SYSTEM,
+				content: trimmedPrompt,
+				parent: parentId,
+				children: []
+			};
 
-		const parentMessage = await db[IDXDB_TABLES.messages].get(parentId);
-		if (parentMessage) {
+			await db[IDXDB_TABLES.messages].add(systemMessage);
 			await db[IDXDB_TABLES.messages].update(parentId, {
 				children: [...parentMessage.children, systemMessage.id]
 			});
-		}
 
-		return systemMessage;
+			return systemMessage;
+		});
 	}
 
 	/**
@@ -442,7 +450,8 @@ export class DatabaseService {
 	}
 
 	/**
-	 * Updates a conversation.
+	 * Updates a conversation. `lastModified` is never stamped implicitly;
+	 * pass it in `updates` to bump the conversation in recency ordering.
 	 *
 	 * @param id - Conversation ID
 	 * @param updates - Partial updates to apply
@@ -452,10 +461,7 @@ export class DatabaseService {
 		id: string,
 		updates: Partial<Omit<DatabaseConversation, 'id'>>
 	): Promise<void> {
-		await db[IDXDB_TABLES.conversations].update(id, {
-			...updates,
-			lastModified: Date.now()
-		});
+		await db[IDXDB_TABLES.conversations].update(id, updates);
 	}
 
 	/**
@@ -473,7 +479,7 @@ export class DatabaseService {
 	 * @returns The new pinned status
 	 */
 	static async toggleConversationPin(id: string): Promise<boolean> {
-		const conversation = await db.conversations.get(id);
+		const conversation = await db[IDXDB_TABLES.conversations].get(id);
 		if (!conversation) {
 			throw new Error(`Conversation ${id} not found`);
 		}
@@ -497,7 +503,6 @@ export class DatabaseService {
 		const result = new Map<string, boolean>();
 		if (cleanIds.length === 0) return result;
 
-		const now = Date.now();
 		await db.transaction('rw', db[IDXDB_TABLES.conversations], async () => {
 			const convs = await db[IDXDB_TABLES.conversations].bulkGet(cleanIds);
 			const updates: DatabaseConversation[] = [];
@@ -505,7 +510,7 @@ export class DatabaseService {
 				const conv = convs[i];
 				if (!conv) continue;
 				const newPinned = !conv.pinned;
-				updates.push({ ...conv, pinned: newPinned, lastModified: now });
+				updates.push({ ...conv, pinned: newPinned });
 				result.set(cleanIds[i], newPinned);
 			}
 			if (updates.length === 0) return;

tools/ui/src/lib/stores/chat.svelte.ts

@@ -1658,7 +1658,8 @@ class ChatStore {
 					generateConversationTitle(newContent, Boolean(config().titleGenerationUseFirstLine))
 				);
 			const messagesToRemove = conversationsStore.activeMessages.slice(messageIndex + 1);
-			for (const message of messagesToRemove) await DatabaseService.deleteMessage(message.id);
+			if (messagesToRemove.length > 0)
+				await DatabaseService.deleteMessageCascading(activeConv.id, messagesToRemove[0].id);
 			conversationsStore.sliceActiveMessages(messageIndex + 1);
 			conversationsStore.updateConversationTimestamp();
 			this.setChatLoading(activeConv.id, true);
@@ -1690,7 +1691,7 @@ class ChatStore {
 		const { index: messageIndex } = result;
 		try {
 			const messagesToRemove = conversationsStore.activeMessages.slice(messageIndex);
-			for (const message of messagesToRemove) await DatabaseService.deleteMessage(message.id);
+			await DatabaseService.deleteMessageCascading(activeConv.id, messagesToRemove[0].id);
 			conversationsStore.sliceActiveMessages(messageIndex);
 			conversationsStore.updateConversationTimestamp();
 			this.setChatLoading(activeConv.id, true);
@@ -2037,7 +2038,7 @@ class ChatStore {
 							timings
 						});
 
-						conversationsStore.updateConversationTimestamp();
+						conversationsStore.updateConversationTimestamp(msg.convId);
 
 						this.setChatLoading(msg.convId, false);
 						this.clearChatStreaming(msg.convId);

tools/ui/src/lib/stores/conversations.svelte.ts

@@ -111,6 +111,9 @@ class ConversationsStore {
 		| ((messageId: string, updates: Partial<DatabaseMessage>) => void)
 		| null = null;
 
+	/** In-flight init run; shared by concurrent callers, reset on failure to allow retry */
+	private initPromise: Promise<void> | null = null;
+
 	/**
 	 *
 	 *
@@ -121,19 +124,25 @@ class ConversationsStore {
 
 	/**
 	 * Initialize the store by loading conversations from database.
-	 * Must be called once after app startup.
+	 * Safe to call multiple times: concurrent callers share a single run,
+	 * and a failed run can be retried by calling again.
 	 */
-	async init(): Promise<void> {
-		if (!browser) return;
-		if (this.isInitialized) return;
+	init(): Promise<void> {
+		if (!browser) return Promise.resolve();
+		if (this.initPromise) return this.initPromise;
 
-		try {
-			await MigrationService.runAllMigrations();
-			await this.loadConversations();
-			this.isInitialized = true;
-		} catch (error) {
-			console.error('Failed to initialize conversations:', error);
-		}
+		this.initPromise = (async () => {
+			try {
+				await MigrationService.runAllMigrations();
+				await this.loadConversations();
+				this.isInitialized = true;
+			} catch (error) {
+				console.error('Failed to initialize conversations:', error);
+				this.initPromise = null;
+			}
+		})();
+
+		return this.initPromise;
 	}
 
 	/**
@@ -237,15 +246,11 @@ class ConversationsStore {
 	 */
 	async createConversation(name?: string): Promise<string> {
 		const conversationName = name || `Chat ${new Date().toLocaleString()}`;
-		const conversation = await DatabaseService.createConversation(conversationName);
 
 		// No MCP override list is seeded: getAllMcpServerOverrides resolves
 		// servers without a per-conversation override to `mcpServers[i].enabled`,
 		// and only explicit toggles are stored on the conversation.
-
-		// Inherit the global reasoning default into the new conversation
-		conversation.reasoningEffort = this.pendingReasoningEffort;
-		await DatabaseService.updateConversation(conversation.id, {
+		const conversation = await DatabaseService.createConversation(conversationName, {
 			reasoningEffort: this.pendingReasoningEffort
 		});
 
@@ -358,10 +363,7 @@ class ConversationsStore {
 	async deleteAll(): Promise<void> {
 		try {
 			const allConversations = await DatabaseService.getAllConversations();
-
-			for (const conv of allConversations) {
-				await DatabaseService.deleteConversation(conv.id);
-			}
+			await DatabaseService.bulkDeleteConversations(allConversations.map((c) => c.id));
 
 			this.clearActiveConversation();
 			this.conversations = [];
@@ -412,7 +414,9 @@ class ConversationsStore {
 			}
 
 			toast.success(
-				convIds.length === 1 ? 'Conversation deleted' : `${convIds.length} conversations deleted`
+				idsToRemove.size === 1
+					? 'Conversation deleted'
+					: `${idsToRemove.size} conversations deleted`
 			);
 		} catch (error) {
 			console.error('Failed to bulk delete conversations:', error);
@@ -443,7 +447,6 @@ class ConversationsStore {
 				const newPinned = updates.get(this.conversations[i].id);
 				if (newPinned !== undefined) this.conversations[i].pinned = newPinned;
 			}
-			this.conversations = [...this.conversations];
 
 			toast.success(
 				convIds.length === 1
@@ -552,7 +555,6 @@ class ConversationsStore {
 
 			if (convIndex !== -1) {
 				this.conversations[convIndex].name = name;
-				this.conversations = [...this.conversations];
 			}
 
 			if (this.activeConversation?.id === convId) {
@@ -576,7 +578,6 @@ class ConversationsStore {
 
 			if (convIndex !== -1) {
 				this.conversations[convIndex].pinned = newPinnedState;
-				this.conversations = [...this.conversations];
 			}
 
 			if (this.activeConversation?.id === convId) {
@@ -591,18 +592,33 @@ class ConversationsStore {
 	}
 
 	/**
-	 * Updates conversation lastModified timestamp and moves it to top of list
+	 * Marks a conversation as recently active: stamps lastModified (persisted)
+	 * and moves it to the top of the list. Only message-activity flows call
+	 * this; metadata updates (rename, pin, settings) do not.
+	 *
+	 * @param convId - Conversation that produced the activity, defaults to the active one
 	 */
-	updateConversationTimestamp(): void {
-		if (!this.activeConversation) return;
+	updateConversationTimestamp(convId?: string): void {
+		const targetId = convId ?? this.activeConversation?.id;
+		if (!targetId) return;
 
-		const chatIndex = this.conversations.findIndex((c) => c.id === this.activeConversation!.id);
+		const now = Date.now();
+
+		const chatIndex = this.conversations.findIndex((c) => c.id === targetId);
 
 		if (chatIndex !== -1) {
-			this.conversations[chatIndex].lastModified = Date.now();
+			this.conversations[chatIndex].lastModified = now;
 			const updatedConv = this.conversations.splice(chatIndex, 1)[0];
 			this.conversations = [updatedConv, ...this.conversations];
 		}
+
+		if (this.activeConversation?.id === targetId) {
+			this.activeConversation = { ...this.activeConversation, lastModified: now };
+		}
+
+		DatabaseService.updateConversation(targetId, { lastModified: now }).catch((error) =>
+			console.error('Failed to update conversation timestamp:', error)
+		);
 	}
 
 	/**
@@ -773,7 +789,6 @@ class ConversationsStore {
 		if (convIndex !== -1) {
 			this.conversations[convIndex].mcpServerOverrides =
 				newOverrides.length > 0 ? newOverrides : undefined;
-			this.conversations = [...this.conversations];
 		}
 	}
 
@@ -837,7 +852,6 @@ class ConversationsStore {
 		const convIndex = this.conversations.findIndex((c) => c.id === this.activeConversation!.id);
 		if (convIndex !== -1) {
 			this.conversations[convIndex].reasoningEffort = effort;
-			this.conversations = [...this.conversations];
 		}
 	}
 

tools/ui/src/lib/utils/legacy-migration.ts

@@ -1,364 +0,0 @@
-/**
- * @deprecated Legacy migration utility — remove at some point in the future once all users have migrated to the new structured agentic message format.
- *
- * Converts old marker-based agentic messages to the new structured format
- * with separate messages per turn.
- *
- * Old format: Single assistant message with markers in content:
- *   <<<reasoning_content_start>>>...<<<reasoning_content_end>>>
- *   <<<AGENTIC_TOOL_CALL_START>>>...<<<AGENTIC_TOOL_CALL_END>>>
- *
- * New format: Separate messages per turn:
- *   - assistant (content + reasoningContent + toolCalls)
- *   - tool (toolCallId + content)
- *   - assistant (next turn)
- *   - ...
- */
-
-import { LEGACY_AGENTIC_REGEX, LEGACY_REASONING_TAGS } from '$lib/constants';
-import { DatabaseService } from '$lib/services/database.service';
-import { MessageRole, MessageType } from '$lib/enums';
-import type { DatabaseMessage } from '$lib/types/database';
-
-const MIGRATION_DONE_KEY = 'llama-ui-migration-v2-done';
-/** @deprecated Use {@link MIGRATION_DONE_KEY} instead */
-const DEPRECATED_MIGRATION_DONE_KEY = 'llama-webui-migration-v2-done';
-
-/**
- * @deprecated Part of legacy migration — remove with the migration module.
- * Check if migration has been performed.
- */
-export function isMigrationNeeded(): boolean {
-	try {
-		// Check new key first, fall back to deprecated old key
-		if (localStorage.getItem(MIGRATION_DONE_KEY)) return false;
-		if (localStorage.getItem(DEPRECATED_MIGRATION_DONE_KEY)) {
-			// Migrate to new key
-			try {
-				localStorage.setItem(MIGRATION_DONE_KEY, String(Date.now()));
-				localStorage.removeItem(DEPRECATED_MIGRATION_DONE_KEY);
-			} catch {
-				// Ignore storage errors
-			}
-			return false;
-		}
-		return true;
-	} catch {
-		return false;
-	}
-}
-
-/**
- * Mark migration as done.
- */
-function markMigrationDone(): void {
-	try {
-		localStorage.setItem(MIGRATION_DONE_KEY, String(Date.now()));
-	} catch {
-		// Ignore localStorage errors
-	}
-}
-
-/**
- * Check if a message has legacy markers in its content.
- */
-function hasLegacyMarkers(message: DatabaseMessage): boolean {
-	if (!message.content) return false;
-	return LEGACY_AGENTIC_REGEX.HAS_LEGACY_MARKERS.test(message.content);
-}
-
-/**
- * Extract reasoning content from legacy marker format.
- */
-function extractLegacyReasoning(content: string): { reasoning: string; cleanContent: string } {
-	let reasoning = '';
-	let cleanContent = content;
-
-	// Extract all reasoning blocks
-	const re = new RegExp(LEGACY_AGENTIC_REGEX.REASONING_EXTRACT.source, 'g');
-	let match;
-	while ((match = re.exec(content)) !== null) {
-		reasoning += match[1];
-	}
-
-	// Remove reasoning tags from content
-	cleanContent = cleanContent
-		.replace(new RegExp(LEGACY_AGENTIC_REGEX.REASONING_BLOCK.source, 'g'), '')
-		.replace(LEGACY_AGENTIC_REGEX.REASONING_OPEN, '');
-
-	return { reasoning, cleanContent };
-}
-
-/**
- * Parse legacy content with tool call markers into structured turns.
- */
-interface ParsedTurn {
-	textBefore: string;
-	toolCalls: Array<{
-		name: string;
-		args: string;
-		result: string;
-	}>;
-}
-
-function parseLegacyToolCalls(content: string): ParsedTurn[] {
-	const turns: ParsedTurn[] = [];
-	const regex = new RegExp(LEGACY_AGENTIC_REGEX.COMPLETED_TOOL_CALL.source, 'g');
-
-	let lastIndex = 0;
-	let currentTurn: ParsedTurn = { textBefore: '', toolCalls: [] };
-	let match;
-
-	while ((match = regex.exec(content)) !== null) {
-		const textBefore = content.slice(lastIndex, match.index).trim();
-
-		// If there's text between tool calls and we already have tool calls,
-		// that means a new turn started (text after tool results = new LLM turn)
-		if (textBefore && currentTurn.toolCalls.length > 0) {
-			turns.push(currentTurn);
-			currentTurn = { textBefore, toolCalls: [] };
-		} else if (textBefore && currentTurn.toolCalls.length === 0) {
-			currentTurn.textBefore = textBefore;
-		}
-
-		currentTurn.toolCalls.push({
-			name: match[1],
-			args: match[2],
-			result: match[3].replace(/^\n+|\n+$/g, '')
-		});
-
-		lastIndex = match.index + match[0].length;
-	}
-
-	// Any remaining text after the last tool call
-	const remainingText = content.slice(lastIndex).trim();
-
-	if (currentTurn.toolCalls.length > 0) {
-		turns.push(currentTurn);
-	}
-
-	// If there's text after all tool calls, it's the final assistant response
-	if (remainingText) {
-		// Remove any partial/open markers
-		const cleanRemaining = remainingText
-			.replace(LEGACY_AGENTIC_REGEX.AGENTIC_TOOL_CALL_OPEN, '')
-			.trim();
-		if (cleanRemaining) {
-			turns.push({ textBefore: cleanRemaining, toolCalls: [] });
-		}
-	}
-
-	// If no tool calls found at all, return the original content as a single turn
-	if (turns.length === 0) {
-		turns.push({ textBefore: content.trim(), toolCalls: [] });
-	}
-
-	return turns;
-}
-
-/**
- * Migrate a single conversation's messages from legacy format to new format.
- */
-async function migrateConversation(convId: string): Promise<number> {
-	const allMessages = await DatabaseService.getConversationMessages(convId);
-	let migratedCount = 0;
-
-	for (const message of allMessages) {
-		if (message.role !== MessageRole.ASSISTANT) continue;
-		if (!hasLegacyMarkers(message)) {
-			// Still check for reasoning-only markers (no tool calls)
-			if (message.content?.includes(LEGACY_REASONING_TAGS.START)) {
-				const { reasoning, cleanContent } = extractLegacyReasoning(message.content);
-				await DatabaseService.updateMessage(message.id, {
-					content: cleanContent.trim(),
-					reasoningContent: reasoning || undefined
-				});
-				migratedCount++;
-			}
-			continue;
-		}
-
-		// Has agentic markers - full migration needed
-		const { reasoning, cleanContent } = extractLegacyReasoning(message.content);
-		const turns = parseLegacyToolCalls(cleanContent);
-
-		// Parse existing toolCalls JSON to try to match IDs
-		let existingToolCalls: Array<{
-			id: string;
-			function?: { name: string; arguments: string };
-		}> = [];
-		if (message.toolCalls) {
-			try {
-				existingToolCalls = JSON.parse(message.toolCalls);
-			} catch {
-				// Ignore
-			}
-		}
-
-		// First turn uses the existing message
-		const firstTurn = turns[0];
-		if (!firstTurn) continue;
-
-		// Match tool calls from the first turn to existing IDs
-		const firstTurnToolCalls = firstTurn.toolCalls.map((tc, i) => {
-			const existing =
-				existingToolCalls.find((e) => e.function?.name === tc.name) || existingToolCalls[i];
-			return {
-				id: existing?.id || `legacy_tool_${i}`,
-				type: 'function' as const,
-				function: { name: tc.name, arguments: tc.args }
-			};
-		});
-
-		// Update the existing message for the first turn
-		await DatabaseService.updateMessage(message.id, {
-			content: firstTurn.textBefore,
-			reasoningContent: reasoning || undefined,
-			toolCalls: firstTurnToolCalls.length > 0 ? JSON.stringify(firstTurnToolCalls) : ''
-		});
-
-		let currentParentId = message.id;
-		let toolCallIdCounter = existingToolCalls.length;
-
-		// Create tool result messages for the first turn
-		for (let i = 0; i < firstTurn.toolCalls.length; i++) {
-			const tc = firstTurn.toolCalls[i];
-			const toolCallId = firstTurnToolCalls[i]?.id || `legacy_tool_${i}`;
-
-			const toolMsg = await DatabaseService.createMessageBranch(
-				{
-					convId,
-					type: MessageType.TEXT,
-					role: MessageRole.TOOL,
-					content: tc.result,
-					toolCallId,
-					timestamp: message.timestamp + i + 1,
-					toolCalls: '',
-					children: []
-				},
-				currentParentId
-			);
-			currentParentId = toolMsg.id;
-		}
-
-		// Create messages for subsequent turns
-		for (let turnIdx = 1; turnIdx < turns.length; turnIdx++) {
-			const turn = turns[turnIdx];
-
-			const turnToolCalls = turn.toolCalls.map((tc, i) => {
-				const idx = toolCallIdCounter + i;
-				const existing = existingToolCalls[idx];
-				return {
-					id: existing?.id || `legacy_tool_${idx}`,
-					type: 'function' as const,
-					function: { name: tc.name, arguments: tc.args }
-				};
-			});
-			toolCallIdCounter += turn.toolCalls.length;
-
-			// Create assistant message for this turn
-			const assistantMsg = await DatabaseService.createMessageBranch(
-				{
-					convId,
-					type: MessageType.TEXT,
-					role: MessageRole.ASSISTANT,
-					content: turn.textBefore,
-					timestamp: message.timestamp + turnIdx * 100,
-					toolCalls: turnToolCalls.length > 0 ? JSON.stringify(turnToolCalls) : '',
-					children: [],
-					model: message.model
-				},
-				currentParentId
-			);
-			currentParentId = assistantMsg.id;
-
-			// Create tool result messages for this turn
-			for (let i = 0; i < turn.toolCalls.length; i++) {
-				const tc = turn.toolCalls[i];
-				const toolCallId = turnToolCalls[i]?.id || `legacy_tool_${toolCallIdCounter + i}`;
-
-				const toolMsg = await DatabaseService.createMessageBranch(
-					{
-						convId,
-						type: MessageType.TEXT,
-						role: MessageRole.TOOL,
-						content: tc.result,
-						toolCallId,
-						timestamp: message.timestamp + turnIdx * 100 + i + 1,
-						toolCalls: '',
-						children: []
-					},
-					currentParentId
-				);
-				currentParentId = toolMsg.id;
-			}
-		}
-
-		// Re-parent any children of the original message to the last created message
-		// (the original message's children list was the next user message or similar)
-		if (message.children.length > 0 && currentParentId !== message.id) {
-			for (const childId of message.children) {
-				// Skip children we just created (they were already properly parented)
-				const child = allMessages.find((m) => m.id === childId);
-				if (!child) continue;
-				// Only re-parent non-tool messages that were original children
-				if (child.role !== MessageRole.TOOL) {
-					await DatabaseService.updateMessage(childId, { parent: currentParentId });
-					// Add to new parent's children
-					const newParent = await DatabaseService.getConversationMessages(convId).then((msgs) =>
-						msgs.find((m) => m.id === currentParentId)
-					);
-					if (newParent && !newParent.children.includes(childId)) {
-						await DatabaseService.updateMessage(currentParentId, {
-							children: [...newParent.children, childId]
-						});
-					}
-				}
-			}
-			// Clear re-parented children from the original message
-			await DatabaseService.updateMessage(message.id, { children: [] });
-		}
-
-		migratedCount++;
-	}
-
-	return migratedCount;
-}
-
-/**
- * @deprecated Part of legacy migration — remove with the migration module.
- * Run the full migration across all conversations.
- * This should be called once at app startup if migration is needed.
- */
-export async function runLegacyMigration(): Promise<void> {
-	if (!isMigrationNeeded()) return;
-
-	if (import.meta.env.DEV && import.meta.env.VITE_DEBUG)
-		console.log('[Migration] Starting legacy message format migration...');
-
-	try {
-		const conversations = await DatabaseService.getAllConversations();
-		let totalMigrated = 0;
-
-		for (const conv of conversations) {
-			const count = await migrateConversation(conv.id);
-			totalMigrated += count;
-		}
-
-		if (import.meta.env.DEV && import.meta.env.VITE_DEBUG) {
-			if (totalMigrated > 0) {
-				console.log(
-					`[Migration] Migrated ${totalMigrated} messages across ${conversations.length} conversations`
-				);
-			} else {
-				console.log('[Migration] No legacy messages found, marking as done');
-			}
-		}
-
-		markMigrationDone();
-	} catch (error) {
-		console.error('[Migration] Failed to migrate legacy messages:', error);
-		// Still mark as done to avoid infinite retry loops
-		markMigrationDone();
-	}
-}