metal : pad mm results

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

@@ -2443,7 +2443,7 @@ static bool ggml_metal_encode_node(
 #if 0
                 // cpy to tmp buffer in MTLHeap
 
-                id<MTLBuffer> h_src0 = h_src0 = ggml_metal_mem_pool_alloc(mem_pool, ggml_nbytes(src0));
+                id<MTLBuffer> h_src0 = ggml_metal_mem_pool_alloc(mem_pool, ggml_nbytes(src0));
                 if (!h_src0) {
                     GGML_LOG_ERROR("%s: failed to allocate buffer from memory pool, size = %zu\n", __func__, ggml_nbytes(src0));
                     return false;
@@ -2947,6 +2947,12 @@ static bool ggml_metal_encode_node(
                         default: break;
                     }
 
+                    id<MTLBuffer> h_dst = ggml_metal_mem_pool_alloc(mem_pool, sizeof(float)*GGML_PAD(ne01, 64)*GGML_PAD(ne11, 32)*ne12*ne13);
+                    if (!h_dst) {
+                        GGML_LOG_ERROR("%s: failed to allocate buffer from memory pool, size = %zu\n", __func__, ggml_nbytes(src0));
+                        return false;
+                    }
+
                     id<MTLComputePipelineState> pipeline = nil;
 
                     switch (src0->type) {
@@ -2986,8 +2992,8 @@ static bool ggml_metal_encode_node(
                         /*.nb11 =*/ nb11,
                         /*.nb12 =*/ nb12,
                         /*.nb13 =*/ nb13,
-                        /*.ne0  =*/ ne0,
-                        /*.ne1  =*/ ne1,
+                        /*.ne0  =*/ GGML_PAD(ne01, 64),
+                        /*.ne1  =*/ GGML_PAD(ne11, 32),
                         /*.r2   =*/ r2,
                         /*.r3   =*/ r3,
                     };
@@ -2996,10 +3002,40 @@ static bool ggml_metal_encode_node(
                     [encoder setBytes:&args    length:sizeof(args) atIndex:0];
                     [encoder setBuffer:id_src0 offset:offs_src0    atIndex:1];
                     [encoder setBuffer:id_src1 offset:offs_src1    atIndex:2];
-                    [encoder setBuffer:id_dst  offset:offs_dst     atIndex:3];
+                  //[encoder setBuffer:id_dst  offset:offs_dst     atIndex:3];
+                    [encoder setBuffer:h_dst   offset:0            atIndex:3];
 
                     [encoder setThreadgroupMemoryLength:8192 atIndex:0];
                     [encoder dispatchThreadgroups:MTLSizeMake( (ne11 + 31)/32, (ne01 + 63)/64, ne12*ne13) threadsPerThreadgroup:MTLSizeMake(128, 1, 1)];
+
+                    ggml_metal_kargs_cpy args_cpy = {
+                        /*.ne00 =*/ ne0,
+                        /*.ne01 =*/ ne1,
+                        /*.ne02 =*/ ne2,
+                        /*.ne03 =*/ ne3,
+                        /*.nb00 =*/ nb0,
+                        /*.nb01 =*/ nb0*GGML_PAD(ne01, 64),
+                        /*.nb02 =*/ nb0*GGML_PAD(ne01, 64)*GGML_PAD(ne11, 32),
+                        /*.nb03 =*/ nb0*GGML_PAD(ne01, 64)*GGML_PAD(ne11, 32)*ne12,
+                        /*.ne0  =*/ ne0,
+                        /*.ne1  =*/ ne1,
+                        /*.ne2  =*/ ne2,
+                        /*.ne3  =*/ ne3,
+                        /*.nb0  =*/ nb0,
+                        /*.nb1  =*/ nb1,
+                        /*.nb2  =*/ nb2,
+                        /*.nb3  =*/ nb3,
+                    };
+
+                    [encoder setComputePipelineState:ctx->kernels[GGML_METAL_KERNEL_TYPE_CPY_F32_F32].pipeline];
+
+                    [encoder setBytes:&args_cpy length:sizeof(args_cpy) atIndex:0];
+                    [encoder setBuffer:h_dst  offset:0                  atIndex:1];
+                    [encoder setBuffer:id_dst offset:offs_dst           atIndex:2];
+
+                    const int nth_cpy = MIN(1024, ne0);
+
+                    [encoder dispatchThreadgroups:MTLSizeMake(ne1, ne2, ne3) threadsPerThreadgroup:MTLSizeMake(nth_cpy, 1, 1)];
                 } else {
                     id<MTLComputePipelineState> pipeline = nil;
 

ggml/src/ggml-metal/ggml-metal.metal

@@ -6305,34 +6305,34 @@ kernel void kernel_mul_mm(
         }
     } else {
         // block is smaller than 64x32, we should avoid writing data outside of the matrix
-        threadgroup_barrier(mem_flags::mem_threadgroup);
-        threadgroup float * temp_str = ((threadgroup float *) shmem) \
-                                     + 32*(sgitg&1) + (16*(sgitg >> 1))*BLOCK_SIZE_M;
-        for (short i = 0; i < 8; i++) {
-            simdgroup_store(mc[i], temp_str + 8*(i%4) + 8*BLOCK_SIZE_M*(i/4), BLOCK_SIZE_M);
-        }
+        //threadgroup_barrier(mem_flags::mem_threadgroup);
+        //threadgroup float * temp_str = ((threadgroup float *) shmem) \
+        //                             + 32*(sgitg&1) + (16*(sgitg >> 1))*BLOCK_SIZE_M;
+        //for (short i = 0; i < 8; i++) {
+        //    simdgroup_store(mc[i], temp_str + 8*(i%4) + 8*BLOCK_SIZE_M*(i/4), BLOCK_SIZE_M);
+        //}
 
-        threadgroup_barrier(mem_flags::mem_threadgroup);
+        //threadgroup_barrier(mem_flags::mem_threadgroup);
 
-        if (sgitg == 0) {
-            for (int j = tiitg; j < n_cols; j += BLOCK_SIZE_N) {
-                device float  * D  = (device float  *) dst + (r0*BLOCK_SIZE_M) + (r1*BLOCK_SIZE_N + j)*args.ne0 + im*args.ne1*args.ne0;
-                device float4 * D4 = (device float4 *) D;
+        //if (sgitg == 0) {
+        //    for (int j = tiitg; j < n_cols; j += BLOCK_SIZE_N) {
+        //        device float  * D  = (device float  *) dst + (r0*BLOCK_SIZE_M) + (r1*BLOCK_SIZE_N + j)*args.ne0 + im*args.ne1*args.ne0;
+        //        device float4 * D4 = (device float4 *) D;
 
-                threadgroup float  * C  = temp_str + (j*BLOCK_SIZE_M);
-                threadgroup float4 * C4 = (threadgroup float4 *) C;
+        //        threadgroup float  * C  = temp_str + (j*BLOCK_SIZE_M);
+        //        threadgroup float4 * C4 = (threadgroup float4 *) C;
 
-                int i = 0;
-                for (; i < n_rows/4; i++) {
-                    *(D4 + i) = *(C4 + i);
-                }
+        //        int i = 0;
+        //        for (; i < n_rows/4; i++) {
+        //            *(D4 + i) = *(C4 + i);
+        //        }
 
-                i *= 4;
-                for (; i < n_rows; i++) {
-                    *(D + i) = *(C + i);
-                }
-            }
-        }
+        //        i *= 4;
+        //        for (; i < n_rows; i++) {
+        //            *(D + i) = *(C + i);
+        //        }
+        //    }
+        //}
     }
 }
 

tools/imatrix/imatrix.cpp

@@ -46,7 +46,7 @@ private:
     common_params                          m_params;
     std::mutex                             m_mutex;
     int                                    m_last_call = 0;
-    std::vector<float>                     m_src1_data;
+    std::vector<char>                      m_src1_data;
     std::vector<char>                      m_ids; // the expert ids from ggml_mul_mat_id
 };
 
@@ -93,11 +93,13 @@ bool IMatrixCollector::collect_imatrix(struct ggml_tensor * t, bool ask, void *
     const bool is_host = ggml_backend_buffer_is_host(src1->buffer);
 
     if (!is_host) {
-        m_src1_data.resize(ggml_nelements(src1));
-        ggml_backend_tensor_get(src1, m_src1_data.data(), 0, ggml_nbytes(src1));
+        const size_t src1_nbytes = ggml_nbytes(src1);
+        m_src1_data.resize(src1_nbytes);
+        ggml_backend_tensor_get(src1, m_src1_data.data(), 0, src1_nbytes);
     }
 
-    const float * data = is_host ? (const float *) src1->data : m_src1_data.data();
+    const char * data = is_host ? (const char *) src1->data : m_src1_data.data();
+    GGML_ASSERT(src1->nb[0] == ggml_element_size(src1));
 
     // this has been adapted to the new format of storing merged experts in a single 3d tensor
     // ref: https://github.com/ggml-org/llama.cpp/pull/6387
@@ -144,7 +146,7 @@ bool IMatrixCollector::collect_imatrix(struct ggml_tensor * t, bool ask, void *
 
                     const int64_t i11 = idx % src1->ne[1];
                     const int64_t i12 = row;
-                    const float * x = (const float *)((const char *)data + i11*src1->nb[1] + i12*src1->nb[2]);
+                    const float * x = (const float *)(data + i11*src1->nb[1] + i12*src1->nb[2]);
 
                     for (int j = 0; j < (int)src1->ne[0]; ++j) {
                         e.values[e_start + j] += x[j]*x[j];
@@ -180,7 +182,7 @@ bool IMatrixCollector::collect_imatrix(struct ggml_tensor * t, bool ask, void *
         ++e.ncall;
         LOG_DBGV(2, "%s[%d]: %32s, %s, %5d x %5d, %d\n", __func__, m_last_call, wname.c_str(), ggml_op_name(t->op), (int)src1->ne[0], (int)src1->ne[1], (int)src1->type);
         for (int row = 0; row < (int)src1->ne[1]; ++row) {
-            const float * x = data + row * src1->ne[0];
+            const float * x = (const float *) (data + row * src1->nb[1]);
             for (int j = 0; j < (int)src1->ne[0]; ++j) {
                 e.values[j] += x[j]*x[j];
                 e.counts[j]++;