bug fixes for running on metal

ggml/src/ggml-metal/CMakeLists.txt

@@ -100,6 +100,9 @@ if (GGML_METAL_EMBED_LIBRARY)
             COMMAND echo ".globl _ggml_metallib_${kind_sym}_start"               >> "${ASM}"
             COMMAND echo "_ggml_metallib_${kind_sym}_start:"                     >> "${ASM}"
             COMMAND echo .incbin "\"${EMBED}\""                                  >> "${ASM}"
+            # sccache hashes only the .s contents, not the incbin'd file, so
+            # embed the source digest here to bust the cache when kernels change
+            COMMAND echo "/* src-md5: `md5 -q ${EMBED}` */"                      >> "${ASM}"
             COMMAND echo ".globl _ggml_metallib_${kind_sym}_end"                 >> "${ASM}"
             COMMAND echo "_ggml_metallib_${kind_sym}_end:"                       >> "${ASM}"
             DEPENDS ../ggml-common.h ggml-metal-impl.h

ggml/src/ggml-metal/kernels/dequantize.h

@@ -756,13 +756,15 @@ void dequantize_iq4_nl(device const block_iq4_nl * xb, short il, thread type4x4
     const float d = xb->d;
     uint32_t aux32;
     thread const uint8_t * q8 = (thread const uint8_t *)&aux32;
+    float4x4 tmp;
     for (int i = 0; i < 4; ++i) {
         aux32 = ((q4[2*i] | (q4[2*i+1] << 16)) >> 4*il) & 0x0f0f0f0f;
-        reg[i][0] = d * kvalues_iq4nl_f[q8[0]];
-        reg[i][1] = d * kvalues_iq4nl_f[q8[1]];
-        reg[i][2] = d * kvalues_iq4nl_f[q8[2]];
-        reg[i][3] = d * kvalues_iq4nl_f[q8[3]];
+        tmp[i][0] = d * kvalues_iq4nl_f[q8[0]];
+        tmp[i][1] = d * kvalues_iq4nl_f[q8[1]];
+        tmp[i][2] = d * kvalues_iq4nl_f[q8[2]];
+        tmp[i][3] = d * kvalues_iq4nl_f[q8[3]];
     }
+    reg = (type4x4) tmp;
 }
 
 template <typename type4>

ggml/src/ggml-metal/kernels/quantize.h

@@ -19,6 +19,59 @@ void quantize_q1_0(device const float * src, device block_q1_0 & dst) {
     }
 }
 
+// 2-bit: byte j of qs holds elements j, j+8, j+16, j+24, two bits each
+void quantize_q2_0s(device const float * src, device block_q2_0s & dst) {
+    float amax = 0.0f;
+    float vmax = 0.0f;
+    for (int j = 0; j < QK2_0S; ++j) {
+        if (fabs(src[j]) > amax) {
+            amax = fabs(src[j]);
+            vmax = src[j];
+        }
+    }
+    const float d  = vmax / -2.0f;
+    const float id = d != 0.0f ? 1.0f / d : 0.0f;
+
+    dst.d = d;
+
+    for (int j = 0; j < QK2_0S / 4; ++j) {
+        dst.qs[j] = 0;
+    }
+    for (int j = 0; j < QK2_0S; ++j) {
+        const float x0 = src[j] * id;
+        const uint8_t xi0 = (uint8_t) clamp((int) floor(x0 + 2.5f), 0, 3);
+        dst.qs[j % (QK2_0S / 4)] |= (xi0 & 0x03) << (2 * (j / (QK2_0S / 4)));
+    }
+}
+
+void quantize_q2_1(device const float * src, device block_q2_1 & dst) {
+    float vmin = FLT_MAX;
+    float vmax = -FLT_MAX;
+    for (int j = 0; j < QK2_1; ++j) {
+        const float v = src[j];
+        if (v < vmin) {
+            vmin = v;
+        }
+        if (v > vmax) {
+            vmax = v;
+        }
+    }
+    const float d  = (vmax - vmin) / 3.0f;
+    const float id = d != 0.0f ? 1.0f / d : 0.0f;
+
+    dst.d = d;
+    dst.m = vmin;
+
+    for (int j = 0; j < QK2_1 / 4; ++j) {
+        dst.qs[j] = 0;
+    }
+    for (int j = 0; j < QK2_1; ++j) {
+        const float x0 = (src[j] - vmin) * id;
+        const uint8_t xi0 = (uint8_t) clamp((int) floor(x0 + 0.5f), 0, 3);
+        dst.qs[j % (QK2_1 / 4)] |= (xi0 & 0x03) << (2 * (j / (QK2_1 / 4)));
+    }
+}
+
 void quantize_q6_0(device const float * src, device block_q6_0 & dst) {
     float amax = 0.0f;
     float vmax = 0.0f;

ggml/src/ggml-metal/kernels/quantize.metal

@@ -296,7 +296,7 @@ kernel void kernel_get_rows_f(
     auto pdst = (      device T  *) ((      device char *)  dst + i12*args.nb3  + i11*args.nb2  + i10*args.nb1);
 
     for (int ind = iw0*ntg.x + tiitg; ind < args.ne00t;) {
-        pdst[ind] = psrc[ind];
+        pdst[ind] = (T) psrc[ind];
 
         break;
     }