clean up

Author: shawngu-quic

ggml/src/ggml-opencl/ggml-opencl.cpp

@@ -7284,7 +7284,6 @@ static void ggml_cl_mul_mat_id(ggml_backend_t backend, const ggml_tensor * src0,
             CL_CHECK(clSetKernelArg(kernel, 21, sizeof(int),      &ne1));
             CL_CHECK(clSetKernelArg(kernel, 22, sizeof(int),      &r2));
             CL_CHECK(clSetKernelArg(kernel, 23, sizeof(int),      &r3));
-            CL_CHECK(clSetKernelArg(kernel, 24, sizeof(float)*sgs,nullptr));
 #endif // GGML_OPENCL_SOA_Q
             break;
         }

ggml/src/ggml-opencl/kernels/cvt.cl

@@ -146,11 +146,6 @@ kernel void kernel_convert_block_mxfp4(
     for (int i = 0; i < QK_MXFP4 / 2; ++i) {
         q[i] = b->qs[i];
     }
-    // vstore4(*(uint4 *)(&(b->qs[0])), 0, (uint *)(q));
-
-    // if (get_global_id(0) == 0) {
-    //     printf("[kernel_convert_block_mxfp4] e: %d, q[0]: %d, q[15]: %d\n", b->e, b->qs[0], b->qs[15]);
-    // }
 }
 
 kernel void kernel_restore_block_mxfp4(

ggml/src/ggml-opencl/kernels/mul_mv_id_mxfp4_f32.cl

@@ -24,15 +24,6 @@ typedef struct {
     uchar qs[QK_MXFP4/2];
 } block_mxfp4;
 
-
-// static inline half mxfp4_to_fp16(uchar fp4) {
-//     ushort sign = (fp4 >> 3) & 0x1;
-//     ushort d = (fp4 >> 2) & 0x1;
-//     ushort a = fp4 & 0x7;
-    
-//     return (1 - sign * 2) * ((1-d) * a * 0.5f + d * (ushort)(1.2999f * a - 3.0799f));
-// }
-
 // single ushort contains 4 mxfp4 as input
 static inline half4 mxfp4_to_fp16_packed(ushort fp4x4) {
     ushort2 fp16_packed_a, fp16_packed_b, bias_a, bias_b, sign_a, sign_b;
@@ -93,10 +84,8 @@ inline void mul_mv_mxfp4_f32(
     int ne0,
     int ne1,
     int r2,
-    int r3,
-    local  char * shmem
+    int r3
 ) {
-    // local float * shmem_f32 = (local float *) shmem;
     int nb = ne00/QK_MXFP4;
 
     int r0 = get_group_id(0);
@@ -117,35 +106,25 @@ inline void mul_mv_mxfp4_f32(
     const short ix = get_sub_group_local_id()/2;  // 0...15
     const short it = get_sub_group_local_id()%2;  // 0 or 1
 
-    float4 yl[4];
     float sumf[N_R0_MXFP4] = {0.f};
 
     global float * yb = y + ix * QK_MXFP4 + it * 8;
 
     for (int ib = ix; ib < nb; ib += N_SIMDWIDTH/2) {
         global float4 * y4 = (global float4 *)yb;
-        yl[0] = y4[0];
-        yl[1] = y4[4];
-        yl[2] = y4[1];
-        yl[3] = y4[5];
 
         for (short row = 0; row < N_R0_MXFP4; row++) {
             global block_mxfp4 * xb = x + row*nb + ib;
             global ushort       * q2 = (global ushort *)(xb->qs + 8*it);
 
             half4 fp16x4_0 = mxfp4_to_fp16_packed(q2[0]);
             half4 fp16x4_1 = mxfp4_to_fp16_packed(q2[1]);
-            float4 acc1 = yl[0]*(float4)(fp16x4_0.s0, fp16x4_0.s2, fp16x4_1.s0, fp16x4_1.s2);
-            acc1 += yl[1]*(float4)(fp16x4_0.s1, fp16x4_0.s3, fp16x4_1.s1, fp16x4_1.s3);
+            float4 acc1 = y4[0]*(float4)(fp16x4_0.s0, fp16x4_0.s2, fp16x4_1.s0, fp16x4_1.s2);
+            acc1 += y4[4]*(float4)(fp16x4_0.s1, fp16x4_0.s3, fp16x4_1.s1, fp16x4_1.s3);
             fp16x4_0 = mxfp4_to_fp16_packed(q2[2]);
             fp16x4_1 = mxfp4_to_fp16_packed(q2[3]);
-            acc1 += yl[2]*(float4)(fp16x4_0.s0, fp16x4_0.s2, fp16x4_1.s0, fp16x4_1.s2);
-            acc1 += yl[3]*(float4)(fp16x4_0.s1, fp16x4_0.s3, fp16x4_1.s1, fp16x4_1.s3);
-
-            // float4 acc1 = yl[0]*(float4)(mxfp4_to_fp16(q2[0] &  0x0F), mxfp4_to_fp16(q2[1] &  0x0F), mxfp4_to_fp16(q2[2] &  0x0F), mxfp4_to_fp16(q2[3] &  0x0F));
-            // acc1 += yl[1]*(float4)(mxfp4_to_fp16(q2[0] >> 4   ), mxfp4_to_fp16(q2[1] >> 4   ), mxfp4_to_fp16(q2[2] >> 4   ), mxfp4_to_fp16(q2[3] >> 4   ));
-            // acc1 += yl[2]*(float4)(mxfp4_to_fp16(q2[4] &  0x0F), mxfp4_to_fp16(q2[5] &  0x0F), mxfp4_to_fp16(q2[6] &  0x0F), mxfp4_to_fp16(q2[7] &  0x0F));
-            // acc1 += yl[3]*(float4)(mxfp4_to_fp16(q2[4] >> 4   ), mxfp4_to_fp16(q2[5] >> 4   ), mxfp4_to_fp16(q2[6] >> 4   ), mxfp4_to_fp16(q2[7] >> 4   ));
+            acc1 += y4[1]*(float4)(fp16x4_0.s0, fp16x4_0.s2, fp16x4_1.s0, fp16x4_1.s2);
+            acc1 += y4[5]*(float4)(fp16x4_0.s1, fp16x4_0.s3, fp16x4_1.s1, fp16x4_1.s3);
 
             sumf[row] += e8m0_to_fp32(xb->e) * ((acc1.s0 + acc1.s1) + (acc1.s2 + acc1.s3));
         }
@@ -192,8 +171,7 @@ kernel void kernel_mul_mv_id_mxfp4_f32(
     int           ne0,
     int           ne1,
     int           r2,
-    int           r3,
-    local  char * shmem
+    int           r3
 ) {
     src0 = (global char *)((global char *)src0 + offset0);
     src1 = (global char *)((global char *)src1 + offset1);
@@ -214,14 +192,8 @@ kernel void kernel_mul_mv_id_mxfp4_f32(
     global char * src0_cur = src0 + i02*nb02;
     global char * src1_cur = src1 + i11*nb11 + i12*nb12;
 
-    // if (get_global_id(0) == 0 && get_global_id(1) == 0 && get_global_id(2) == 0) {
-    //     printf("[kernel_mul_mv_id_mxfp4_f32_flat] src1(%lu): %f, src2(%lu): %d\n", offset1, ((global float*)src1)[0], offset2, ((global int*)src2)[0]);
-    //     global block_mxfp4 * block = (global block_mxfp4 *)(src0);
-    //     printf("[kernel_mul_mv_id_mxfp4_f32] i02: %d, offset0: %d, e: %d, q[0]: %d, q[16]: %d\n", i02, offset0, block->e, block->qs[0], block->qs[15]);
-    // }
-
     global char * dst_cur = dst + (i1*ne0 + i2*ne1*ne0)*sizeof(float);
 
     mul_mv_mxfp4_f32(src0_cur, src1_cur, dst_cur,
-        ne00, nb01, nb02, nb03, ne12, nb11, nb12, nb13, ne0, ne1, r2, r3, shmem);
+        ne00, nb01, nb02, nb03, ne12, nb11, nb12, nb13, ne0, ne1, r2, r3);
 }