char-fork/llama.cpp

git clone https://git.t4t.associates/char-fork/llama.cpp

Hongqiang Wangopencl: quant lm_head / decode GEMV and medium-batch GEMM optimizations (speculative decoding/MTP) (#26477)95ef7fc162026-09-03

//------------------------------------------------------------------------------
// This file is contains kernels for data conversion.
// These kernels are used when loading the model, so its performance is less
// important.
//------------------------------------------------------------------------------
#pragma OPENCL EXTENSION cl_khr_fp16 : enable

#ifdef cl_intel_required_subgroup_size
#pragma OPENCL EXTENSION cl_intel_required_subgroup_size : enable
#define INTEL_GPU 1
#define REQD_SUBGROUP_SIZE_16 __attribute__((intel_reqd_sub_group_size(16)))
#define REQD_SUBGROUP_SIZE_32 __attribute__((intel_reqd_sub_group_size(32)))
#elif defined(cl_qcom_reqd_sub_group_size)
#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
#define ADRENO_GPU 1
#define REQD_SUBGROUP_SIZE_64  __attribute__((qcom_reqd_sub_group_size("half")))
#define REQD_SUBGROUP_SIZE_128 __attribute__((qcom_reqd_sub_group_size("full")))
#endif

#define QK4_0                   32
#define QR4_0                   2
#define QK4_1                   32
#define QR4_1                   2
#define QK5_0                   32
#define QR5_0                   2
#define QK5_1                   32
#define QR5_1                   2
#define QK8_0                   32
#define QR8_0                   1
#define QK1_0                   128
#define QR1_0                   1
#define QK_K                    256
#define K_SCALE_SIZE            (3 * QK_K / 64)
#define K_QUANTS_PER_ITERATION  2

typedef char int8_t;
typedef uchar uint8_t;
typedef short int16_t;
typedef ushort uint16_t;
typedef int int32_t;
typedef uint uint32_t;

//------------------------------------------------------------------------------
// block_q1_0
//------------------------------------------------------------------------------
typedef struct {
    half d;             // delta
    uchar qs[QK1_0/8];  // 1-bit signs (16 bytes)
} block_q1_0;

//------------------------------------------------------------------------------
// block_q4_0
//------------------------------------------------------------------------------
struct block_q4_0
{
    half d;
    uint8_t qs[QK4_0 / 2];
};

//------------------------------------------------------------------------------
// block_q4_1
//------------------------------------------------------------------------------
struct block_q4_1 {
    half d; // delta
    half m; // min
    uchar qs[QK4_1 / 2]; // nibbles / quants
};

//------------------------------------------------------------------------------
// block_q5_0
//------------------------------------------------------------------------------
struct block_q5_0 {
    half d; // delta
    uchar qh[4]; // 5-th bit of quants
    uchar qs[QK5_0 / 2]; // nibbles / quants
};

//------------------------------------------------------------------------------
// block_q5_1
//------------------------------------------------------------------------------
struct block_q5_1 {
    half d; // delta
    half m; // min
    uchar qh[4]; // 5-th bit of quants
    uchar qs[QK5_1 / 2]; // nibbles / quants
};

//------------------------------------------------------------------------------
// block_q4_k
//------------------------------------------------------------------------------
struct block_q4_K {
    half d; // delta
    half dm; // min
    uchar s[K_SCALE_SIZE];
    uchar q[QK_K / 2]; // nibbles / quants
};

//------------------------------------------------------------------------------
// block_q5_k
//------------------------------------------------------------------------------
struct block_q5_K {
    half d; // delta
    half dm; // min
    uchar s[K_SCALE_SIZE];
    uchar qh[QK_K / 8];
    uchar qs[QK_K / 2]; // nibbles / quants
};

//------------------------------------------------------------------------------
// block_q6_K
//------------------------------------------------------------------------------
struct block_q6_K {
    uint8_t ql[QK_K/2];      // quants, lower 4 bits
    uint8_t qh[QK_K/4];      // quants, upper 2 bits
    int8_t  scales[QK_K/16]; // scales, quantized with 8 bits
    half d;                  // super-block scale
};

//------------------------------------------------------------------------------
// block_iq4_nl
//------------------------------------------------------------------------------
#define QK4_NL 32

struct block_iq4_nl
{
    half d;
    uint8_t qs[QK4_NL / 2];
};

//------------------------------------------------------------------------------
// bf16 to f16
//------------------------------------------------------------------------------
kernel void kernel_convert_bf16_to_f16(
    global const ushort * src,
    global half * dst,
    ulong off_dst,
    ulong n
) {
    uint i = get_global_id(0);
    if (i >= n) {
        return;
    }

    dst[i + off_dst] = (half) as_float((uint) src[i] << 16);
}

//------------------------------------------------------------------------------
// f16 to bf16
//------------------------------------------------------------------------------
kernel void kernel_convert_f16_to_bf16(
    global const half * src,
    ulong off_src,
    global ushort * dst,
    ulong n
) {
    uint i = get_global_id(0);
    if (i >= n) {
        return;
    }

    float f = (float) src[i + off_src];
    uint bits = as_uint(f);
    if ((bits & 0x7fffffffu) > 0x7f800000u) {
        // nan to quiet nan
        dst[i] = (ushort)((bits >> 16) | 0x40u);
    } else {
        uint rounded = bits + 0x7fffu + ((bits >> 16) & 1u);
        dst[i] = (ushort)(rounded >> 16);
    }
}

//------------------------------------------------------------------------------
// kernel_convert_block_q1_0
// Convert block_q1_0 (AOS) to 2 separate arrays (SOA): quant bytes + scales.
// q1_0 bits are stored in natural order (bit j of byte i -> weight 8*i + j)
//------------------------------------------------------------------------------
kernel void kernel_convert_block_q1_0(
    global block_q1_0 * src0,
    global uchar * dst_q,
    global half  * dst_d
) {
    global block_q1_0 * b = (global block_q1_0 *) src0 + get_global_id(0);
    global uchar      * q = (global uchar *) dst_q + (QK1_0/8)*get_global_id(0);
    global half       * d = (global half *) dst_d + get_global_id(0);

    *d = b->d;

    for (int i = 0; i < QK1_0/8; ++i) {
        q[i] = b->qs[i];
    }
}

kernel void kernel_restore_block_q1_0(
    global uchar * src_q,
    global half  * src_d,
    global block_q1_0 * dst
) {
    global block_q1_0 * b = (global block_q1_0 *) dst + get_global_id(0);
    global uchar      * q = (global uchar *) src_q + (QK1_0/8)*get_global_id(0);
    global half       * d = (global half *) src_d + get_global_id(0);

    b->d = *d;
    for (int i = 0; i < QK1_0/8; ++i) {
        b->qs[i] = q[i];
    }
}

//------------------------------------------------------------------------------
// kernel_convert_block_q4_0
// Convert the block_q4_0 format to 2 separate arrays (AOS -> SOA).
// This kernel does not deshuffle the bits.
//------------------------------------------------------------------------------
kernel void kernel_convert_block_q4_0(
    global struct block_q4_0 * src0,
    global uchar * dst_q,
    global half  * dst_d
) {
    global struct block_q4_0 * b = (global struct block_q4_0 *) src0 + get_global_id(0);
    global uchar * q = (global uchar *) dst_q + QK4_0/2*get_global_id(0);
    global half  * d = (global half *) dst_d + get_global_id(0);

    *d = b->d;

    for (int i = 0; i < QK4_0/2; ++i) {
        q[i] = b->qs[i];
    }
}

kernel void kernel_restore_block_q4_0(
    global uchar * src_q,
    global half  * src_d,
    global struct block_q4_0 * dst
) {
    global struct block_q4_0 * b = (global struct block_q4_0 *) dst + get_global_id(0);
    global uchar * q = (global uchar *) src_q + QK4_0/2*get_global_id(0);
    global half  * d = (global half *) src_d + get_global_id(0);

    b->d = *d;
    for (int i = 0; i < QK4_0/2; ++i) {
        b->qs[i] = q[i];
    }
}

//------------------------------------------------------------------------------
// kernel_convert_block_q4_0_noshuffle
// Flatten q4_0 weights and unshuffle the bits
//------------------------------------------------------------------------------

kernel void kernel_convert_block_q4_0_noshuffle(
    global struct block_q4_0 * src0,
    global uchar * dst_q,
    global half  * dst_d
) {
    global struct block_q4_0 * b = (global struct block_q4_0 *) src0 + get_global_id(0);
    global uchar * q = (global uchar *) dst_q + QK4_0/2*get_global_id(0);
    global half  * d = (global half *) dst_d + get_global_id(0);

    *d = b->d;
    for (int i = 0; i < QK4_0/4; ++i) {
        uchar x0 = b->qs[2*i + 0];
        uchar x1 = b->qs[2*i + 1];

        q[i + 0      ] = convert_uchar(x0 & 0x0F) | convert_uchar((x1 & 0x0F) << 4);
        q[i + QK4_0/4] = convert_uchar((x0 & 0xF0) >> 4) | convert_uchar(x1 & 0xF0);

#ifdef ADRENO_GPU
        // Workaround for adreno - must have the following printf statement for
        // the kernel to work properly. Otherwise it produces incorrect result.
        // convert_uchar above also seems necessary.
        // Compare against a large number so that it does not print anything.
        // get_sub_group_local_id() also works.
        if (get_global_id(0) == 65536*4096) {
            printf("%04x - %02x\n", *(global ushort*)d, ((x0 & 0xF0) >> 4) | (x1 & 0xF0));
        }
#endif
    }
}

kernel void kernel_restore_block_q4_0_noshuffle(
    global uchar * src_q,
    global half  * src_d,
    global struct block_q4_0 * dst,
    uchar mask_0F,
    uchar mask_F0
) {
    global struct block_q4_0 * b = (global struct block_q4_0 *) dst + get_global_id(0);
    global uchar * q = (global uchar *) src_q + QK4_0/2*get_global_id(0);
    global half  * d = (global half *) src_d + get_global_id(0);

    b->d = *d;
    for (int i = 0; i < QK4_0/4; ++i) {
        uchar x0 = q[i + 0      ] ;
        uchar x1 = q[i + QK4_0/4];

        b->qs[2*i + 0] = convert_uchar((x0 & mask_0F) | ((x1 & mask_0F) << 4));
        b->qs[2*i + 1] = convert_uchar(((x0 & mask_F0) >> 4) | (x1 & mask_F0));
    }
}

kernel void kernel_convert_block_q4_0_trans4_ns(
    global struct block_q4_0 * src0,
    __global uint * dst_q,
    __global half * dst_d,
    uint ne00,
    uint ne01
) {
    uint i00 = get_global_id(1);
    uint i01 = get_global_id(0);
    uint i02 = get_global_id(2);

    if (i01 >= ne01) {
        return;
    }

    uint ne00_blk = ne00 / QK4_0;
    uint src_blk_offset = i00 + i01 * ne00_blk + i02 * ne00_blk * ne01;
    uint dst_blk_offset = i01 + i00 * ne01 + i02 * ne00_blk * ne01;

    global struct block_q4_0 * b = src0 + src_blk_offset;
    dst_d[dst_blk_offset] = b->d;

    // extract quantization and unshuffle
    ushort8 pre_block = ((global ushort8 *)(&(b->qs[0])))[0];

    ushort8 post_block = (ushort8)(0);

    uchar * pre_block_ptr = (uchar *)(&pre_block);
    uchar * post_block_ptr = (uchar *)(&post_block);

    for (int i = 0; i < QK4_0 / 4; ++i) {
        uchar x0 = pre_block_ptr[2*i + 0];
        uchar x1 = pre_block_ptr[2*i + 1];

        post_block_ptr[i + 0        ] = convert_uchar(x0 & 0x0F) | convert_uchar((x1 & 0x0F) << 4);
        post_block_ptr[i + QK4_0 / 4] = convert_uchar((x0 & 0xF0) >> 4) | convert_uchar(x1 & 0xF0);
    }

    uint4 q_block = as_uint4(post_block);

    uint offset = i02 * ne00_blk * ne01 * 4 + i00 * ne01 * 4 + i01;
    dst_q[offset] = q_block.x;
    dst_q[offset + ne01] = q_block.y;
    dst_q[offset + ne01 * 2] = q_block.z;
    dst_q[offset + ne01 * 3] = q_block.w;
}

kernel void kernel_restore_block_q4_0_trans4_ns(
    __global uint * src_q,
    __global half * src_d,
    __global struct block_q4_0 * dst0,
    uint ne00,
    uint ne01
) {
    uint i00 = get_global_id(1);
    uint i01 = get_global_id(0);
    uint i02 = get_global_id(2);

    if (i01 >= ne01) {
        return;
    }

    uint ne00_blk = ne00 / QK4_0;
    uint dst_blk_offset = i00 + i01 * ne00_blk + i02 * ne00_blk * ne01;
    uint src_d_offset = i01 + i00 * ne01 + i02 * ne00_blk * ne01;

    __global struct block_q4_0 * b = dst0 + dst_blk_offset;
    b->d = src_d[src_d_offset];

    // collect transposed quantization parts for a block
    uint src_q_offset = i02 * ne00_blk * ne01 * 4 + i00 * ne01 * 4 + i01;
    uint4 q_block;
    q_block.x = src_q[src_q_offset];
    q_block.y = src_q[src_q_offset + ne01];
    q_block.z = src_q[src_q_offset + ne01 * 2];
    q_block.w = src_q[src_q_offset + ne01 * 3];

    ushort8 post_block = as_ushort8(q_block);
    ushort8 pre_block = (ushort8)(0);

    uchar * pre_block_ptr = (uchar *)(&pre_block);
    uchar * post_block_ptr = (uchar *)(&post_block);

    for (int i = 0; i < QK4_0 / 4; ++i) {
        uchar x0 = post_block_ptr[i + 0];
        uchar x1 = post_block_ptr[i + QK4_0 / 4];

        pre_block_ptr[2 * i + 0] = convert_uchar(x0 & 0x0F) | convert_uchar((x1 & 0x0F) << 4);
        pre_block_ptr[2 * i + 1] = convert_uchar((x0 & 0xF0) >> 4) | convert_uchar(x1 & 0xF0);
    }

    ((__global ushort8 *)(&(b->qs[0])))[0] = pre_block;
}

//------------------------------------------------------------------------------
// kernel_convert_block_q4_1
// Convert the block_q4_1 format to 2 separate arrays (AOS -> SOA).
// This kernel does not deshuffle the bits.
//------------------------------------------------------------------------------
kernel void kernel_convert_block_q4_1(
    global struct block_q4_1 * src0,
    global uchar * dst_q,
    global half  * dst_d,
    global half  * dst_m
) {
    global struct block_q4_1 * b = (global struct block_q4_1 *) src0 + get_global_id(0);
    global uchar * q = (global uchar *) dst_q + QK4_1/2*get_global_id(0);
    global half  * d = (global half *) dst_d + get_global_id(0);
    global half  * m = (global half *) dst_m + get_global_id(0);

    *d = b->d;
    *m = b->m;

    for (int i = 0; i < QK4_1/2; ++i) {
        q[i] = b->qs[i];
    }
}

kernel void kernel_restore_block_q4_1(
    global uchar * src_q,
    global half  * src_d,
    global half  * src_m,
    global struct block_q4_1 * dst
) {
    global struct block_q4_1 * b = (global struct block_q4_1 *) dst + get_global_id(0);
    global uchar * q = (global uchar *) src_q + QK4_1/2*get_global_id(0);
    global half  * d = (global half *) src_d + get_global_id(0);
    global half  * m = (global half *) src_m + get_global_id(0);

    b->d = *d;
    b->m = *m;
    for (int i = 0; i < QK4_1/2; ++i) {
        b->qs[i] = q[i];
    }
}

kernel void kernel_convert_block_q4_1_noshuffle(
    global struct block_q4_1 * src0,
    global uchar * dst_q,
    global half  * dst_d,
    global half  * dst_m
) {
    global struct block_q4_1 * b = (global struct block_q4_1 *) src0 + get_global_id(0);
    global uchar * q = (global uchar *) dst_q + QK4_1/2*get_global_id(0);
    global half  * d = (global half *) dst_d + get_global_id(0);
    global half  * m = (global half *) dst_m + get_global_id(0);

    *d = b->d;
    *m = b->m;
    for (int i = 0; i < QK4_1/4; ++i) {
        uchar x0 = b->qs[2*i + 0];
        uchar x1 = b->qs[2*i + 1];

        q[i + 0      ] = convert_uchar(x0 & 0x0F) | convert_uchar((x1 & 0x0F) << 4);
        q[i + QK4_1/4] = convert_uchar((x0 & 0xF0) >> 4) | convert_uchar(x1 & 0xF0);

#ifdef ADRENO_GPU
        if (get_global_id(0) == 65536*4096) {
            printf("%04x - %02x\n", *(global ushort*)d, ((x0 & 0xF0) >> 4) | (x1 & 0xF0));
        }
#endif
    }
}

kernel void kernel_restore_block_q4_1_noshuffle(
    global uchar * src_q,
    global half  * src_d,
    global half  * src_m,
    global struct block_q4_1 * dst,
    uchar mask_0F,
    uchar mask_F0
) {
    global struct block_q4_1 * b = (global struct block_q4_1 *) dst + get_global_id(0);
    global uchar * q = (global uchar *) src_q + QK4_1/2*get_global_id(0);
    global half  * d = (global half *) src_d + get_global_id(0);
    global half  * m = (global half *) src_m + get_global_id(0);

    b->d = *d;
    b->m = *m;
    for (int i = 0; i < QK4_1/4; ++i) {
        uchar x0 = q[i + 0      ] ;
        uchar x1 = q[i + QK4_1/4];

        b->qs[2*i + 0] = convert_uchar((x0 & mask_0F) | ((x1 & mask_0F) << 4));
        b->qs[2*i + 1] = convert_uchar(((x0 & mask_F0) >> 4) | (x1 & mask_F0));
    }
}

kernel void kernel_convert_block_q4_1_trans4_ns(
    __global struct block_q4_1 * src0,
    __global uint * dst_q,
    __global half * dst_d,
    __global half * dst_m,
    uint ne00,
    uint ne01
) {
    uint i00 = get_global_id(1);
    uint i01 = get_global_id(0);
    uint i02 = get_global_id(2);

    if (i01 >= ne01) {
        return;
    }

    uint ne00_blk = ne00 / QK4_1;
    uint src_blk_offset = i00 + i01 * ne00_blk + i02 * ne00_blk * ne01;
    uint dst_blk_offset = i01 + i00 * ne01 + i02 * ne00_blk * ne01;

    global struct block_q4_1 * b = src0 + src_blk_offset;
    dst_d[dst_blk_offset] = b->d;
    dst_m[dst_blk_offset] = b->m;

    // extract quantization and unshuffle
    ushort8 pre_block = ((global ushort8 *)(&(b->qs[0])))[0];

    ushort8 post_block = (ushort8)(0);

    uchar * pre_block_ptr = (uchar *)(&pre_block);
    uchar * post_block_ptr = (uchar *)(&post_block);

    for (int i = 0; i < QK4_1 / 4; ++i) {
        uchar x0 = pre_block_ptr[2*i + 0];
        uchar x1 = pre_block_ptr[2*i + 1];

        post_block_ptr[i + 0        ] = convert_uchar(x0 & 0x0F) | convert_uchar((x1 & 0x0F) << 4);
        post_block_ptr[i + QK4_1 / 4] = convert_uchar((x0 & 0xF0) >> 4) | convert_uchar(x1 & 0xF0);
    }

    uint4 q_block = as_uint4(post_block);

    uint offset = i02 * ne00_blk * ne01 * 4 + i00 * ne01 * 4 + i01;
    dst_q[offset] = q_block.x;
    dst_q[offset + ne01] = q_block.y;
    dst_q[offset + ne01 * 2] = q_block.z;
    dst_q[offset + ne01 * 3] = q_block.w;
}

kernel void kernel_restore_block_q4_1_trans4_ns(
    __global uint * src_q,
    __global half * src_d,
    __global half * src_m,
    __global struct block_q4_1 * dst0,
    uint ne00,
    uint ne01
) {
    int i00 = get_global_id(1);
    uint i01 = get_global_id(0);
    uint i02 = get_global_id(2);

    if (i01 >= ne01) {
        return;
    }

    uint ne00_blk = ne00 / QK4_1;
    uint dst_blk_offset = i00 + i01 * ne00_blk + i02 * ne00_blk * ne01;
    uint src_dm_offset = i01 + i00 * ne01 + i02 * ne00_blk * ne01;

    __global struct block_q4_1 * b = dst0 + dst_blk_offset;
    b->d = src_d[src_dm_offset];
    b->m = src_m[src_dm_offset];

    // collect transposed quantization parts for a block
    uint src_q_offset = i02 * ne00_blk * ne01 * 4 + i00 * ne01 * 4 + i01;
    uint4 q_block;
    q_block.x = src_q[src_q_offset];
    q_block.y = src_q[src_q_offset + ne01];
    q_block.z = src_q[src_q_offset + ne01 * 2];
    q_block.w = src_q[src_q_offset + ne01 * 3];

    ushort8 post_block = as_ushort8(q_block);
    ushort8 pre_block = (ushort8)(0);

    uchar * pre_block_ptr = (uchar *)(&pre_block);
    uchar * post_block_ptr = (uchar *)(&post_block);

    for (int i = 0; i < QK4_0 / 4; ++i) {
        uchar x0 = post_block_ptr[i + 0];
        uchar x1 = post_block_ptr[i + QK4_0 / 4];

        pre_block_ptr[2 * i + 0] = convert_uchar(x0 & 0x0F) | convert_uchar((x1 & 0x0F) << 4);
        pre_block_ptr[2 * i + 1] = convert_uchar((x0 & 0xF0) >> 4) | convert_uchar(x1 & 0xF0);
    }

    ((__global ushort8 *)(&(b->qs[0])))[0] = pre_block;
}

//------------------------------------------------------------------------------
// kernel_convert_block_q5_0
// Convert the block_q5_0 format to 3 separate arrays (AOS -> SOA).
// This kernel does not deshuffle the bits.
//------------------------------------------------------------------------------
kernel void kernel_convert_block_q5_0(
    global struct block_q5_0 * src0,
    global uchar * dst_qs,
    global uint  * dst_qh,
    global half  * dst_d,
    ulong n_blk
) {
    if (get_global_id(0) >= n_blk) {
        return;
    }

    global struct block_q5_0 * b = (global struct block_q5_0 *) src0 + get_global_id(0);
    global uchar * qs = (global uchar *) dst_qs + (QK5_0/2)*get_global_id(0);
    global uint  * qh = (global uint  *) dst_qh + get_global_id(0);
    global half  * d  = (global half  *) dst_d  + get_global_id(0);

    *d = b->d;
    *qh = *((global uint *)(b->qh));

    for (int i = 0; i < QK5_0/2; ++i) {
        qs[i] = b->qs[i];
    }
}

kernel void kernel_restore_block_q5_0(
    global uchar * src_qs,
    global uint  * src_qh,
    global half  * src_d,
    global struct block_q5_0 * dst
) {
    global struct block_q5_0 * b = (global struct block_q5_0 *) dst + get_global_id(0);
    global uchar * qs = (global uchar *) src_qs + (QK5_0/2)*get_global_id(0);
    global uint  * qh = (global uint  *) src_qh + get_global_id(0);
    global half  * d  = (global half  *) src_d  + get_global_id(0);

    b->d = *d;
    *((global uint *)(b->qh)) = *qh;
    for (int i = 0; i < QK5_0/2; ++i) {
        b->qs[i] = qs[i];
    }
}

kernel void kernel_convert_block_q5_0_noshuffle(
    global struct block_q5_0 * src0,
    global uchar * dst_q,
    global uint  * dst_qh,
    global half  * dst_d
) {
    global struct block_q5_0 * b = (global struct block_q5_0 *) src0 + get_global_id(0);
    global uchar * q  = (global uchar *) dst_q + QK5_0/2*get_global_id(0);
    global uint  * qh = (global uint  *) dst_qh + get_global_id(0);
    global half  * d  = (global half  *) dst_d  + get_global_id(0);

    *d = b->d;
    *qh = *((global uint *)(b->qh));

    for (int i = 0; i < QK5_0/4; ++i) {
        uchar x0 = b->qs[2*i + 0];
        uchar x1 = b->qs[2*i + 1];

        q[i + 0      ] = convert_uchar(x0 & 0x0F) | convert_uchar((x1 & 0x0F) << 4);
        q[i + QK5_0/4] = convert_uchar((x0 & 0xF0) >> 4) | convert_uchar(x1 & 0xF0);

#ifdef ADRENO_GPU
        if (get_global_id(0) == 65536*4096) {
            printf("%04x - %02x\n", *(global ushort*)d, ((x0 & 0xF0) >> 4) | (x1 & 0xF0));
        }
#endif
    }
}

kernel void kernel_restore_block_q5_0_noshuffle(
    global uchar * src_q,
    global uint  * src_qh,
    global half  * src_d,
    global struct block_q5_0 * dst,
    uchar mask_0F,
    uchar mask_F0
) {
    global struct block_q5_0 * b = (global struct block_q5_0 *) dst + get_global_id(0);
    global uchar * q  = (global uchar *) src_q + QK5_0/2*get_global_id(0);
    global uint  * qh = (global uint  *) src_qh + get_global_id(0);
    global half  * d  = (global half  *) src_d  + get_global_id(0);

    b->d = *d;
    *((global uint *)(b->qh)) = *qh;

    for (int i = 0; i < QK5_0/4; ++i) {
        uchar x0 = q[i + 0      ];
        uchar x1 = q[i + QK5_0/4];

        b->qs[2*i + 0] = convert_uchar((x0 & mask_0F) | ((x1 & mask_0F) << 4));
        b->qs[2*i + 1] = convert_uchar(((x0 & mask_F0) >> 4) | (x1 & mask_F0));
    }
}

kernel void kernel_convert_block_q5_0_trans4_ns(
    __global struct block_q5_0 * src0,
    __global uint * dst_qs,
    __global uint * dst_qh,
    __global half * dst_d,
    uint ne00,
    uint ne01
) {
    uint i00 = get_global_id(1);
    uint i01 = get_global_id(0);
    uint i02 = get_global_id(2);

    if (i01 >= ne01) {
        return;
    }

    uint ne00_blk = ne00 / QK5_0;
    uint src_blk_offset = i00 + i01 * ne00_blk + i02 * ne00_blk * ne01;
    uint dst_blk_offset = i01 + i00 * ne01 + i02 * ne00_blk * ne01;

    global struct block_q5_0 * b = src0 + src_blk_offset;
    dst_d[dst_blk_offset] = b->d;

    dst_qh[dst_blk_offset] = ((global uint *)(&(b->qh[0])))[0];

    // extract quantization and unshuffle
    ushort8 pre_block = ((global ushort8 *)(&(b->qs[0])))[0];
    ushort8 post_block = (ushort8)(0);

    uchar * pre_block_ptr = (uchar *)(&pre_block);
    uchar * post_block_ptr = (uchar *)(&post_block);

    for (int i = 0; i < QK5_0 / 4; ++i) {
        uchar x0 = pre_block_ptr[2*i + 0];
        uchar x1 = pre_block_ptr[2*i + 1];

        post_block_ptr[i + 0        ] = convert_uchar(x0 & 0x0F) | convert_uchar((x1 & 0x0F) << 4);
        post_block_ptr[i + QK5_0 / 4] = convert_uchar((x0 & 0xF0) >> 4) | convert_uchar(x1 & 0xF0);
    }

    uint4 q_block = as_uint4(post_block);

    uint offset = i02 * ne00_blk * ne01 * 4 + i00 * ne01 * 4 + i01;
    dst_qs[offset] = q_block.x;
    dst_qs[offset + ne01] = q_block.y;
    dst_qs[offset + ne01 * 2] = q_block.z;
    dst_qs[offset + ne01 * 3] = q_block.w;
}

kernel void kernel_restore_block_q5_0_trans4_ns(
    __global uint * src_qs,
    __global uint * src_qh,
    __global half * src_d,
    __global struct block_q5_0 * dst0,
    uint ne00,
    uint ne01
) {
    int i00 = get_global_id(1);
    uint i01 = get_global_id(0);
    uint i02 = get_global_id(2);

    if (i01 >= ne01) {
        return;
    }

    uint ne00_blk = ne00 / QK5_0;
    uint dst_blk_offset = i00 + i01 * ne00_blk + i02 * ne00_blk * ne01;
    uint src_blk_offset = i01 + i00 * ne01 + i02 * ne00_blk * ne01;

    __global struct block_q5_0 * b = dst0 + dst_blk_offset;
    b->d = src_d[src_blk_offset];

    ((__global uint *)(&(b->qh[0])))[0] = src_qh[src_blk_offset];

    // collect transposed quantization parts for a block
    uint src_q_offset = i02 * ne00_blk * ne01 * 4 + i00 * ne01 * 4 + i01;
    uint4 q_block;
    q_block.x = src_qs[src_q_offset];
    q_block.y = src_qs[src_q_offset + ne01];
    q_block.z = src_qs[src_q_offset + ne01 * 2];
    q_block.w = src_qs[src_q_offset + ne01 * 3];

    ushort8 post_block = as_ushort8(q_block);
    ushort8 pre_block = (ushort8)(0);

    uchar * pre_block_ptr = (uchar *)(&pre_block);
    uchar * post_block_ptr = (uchar *)(&post_block);

    for (int i = 0; i < QK5_0 / 4; ++i) {
        uchar x0 = post_block_ptr[i + 0];
        uchar x1 = post_block_ptr[i + QK5_0 / 4];

        pre_block_ptr[2 * i + 0] = convert_uchar(x0 & 0x0F) | convert_uchar((x1 & 0x0F) << 4);
        pre_block_ptr[2 * i + 1] = convert_uchar((x0 & 0xF0) >> 4) | convert_uchar(x1 & 0xF0);
    }

    ((__global ushort8 *)(&(b->qs[0])))[0] = pre_block;
}

//------------------------------------------------------------------------------
// kernel_convert_block_q5_1
// Convert the block_q5_1 format to 4 separate arrays (AOS -> SOA).
// This kernel does not deshuffle the bits.
//------------------------------------------------------------------------------
kernel void kernel_convert_block_q5_1(
    global struct block_q5_1 * src0,
    global uchar * dst_qs,
    global uint  * dst_qh,
    global half  * dst_d,
    global half  * dst_m,
    ulong n_blk
) {
    if (get_global_id(0) >= n_blk) {
        return;
    }

    global struct block_q5_1 * b = (global struct block_q5_1 *) src0 + get_global_id(0);
    global uchar * qs = (global uchar *) dst_qs + (QK5_1/2)*get_global_id(0);
    global uint  * qh = (global uint  *) dst_qh + get_global_id(0);
    global half  * d  = (global half  *) dst_d  + get_global_id(0);
    global half  * m  = (global half  *) dst_m  + get_global_id(0);

    *d = b->d;
    *m = b->m;
    *qh = *((global uint *)(b->qh));

    for (int i = 0; i < QK5_1/2; ++i) {
        qs[i] = b->qs[i];
    }
}

kernel void kernel_restore_block_q5_1(
    global uchar * src_qs,
    global uint  * src_qh,
    global half  * src_d,
    global half  * src_m,
    global struct block_q5_1 * dst
) {
    global struct block_q5_1 * b = (global struct block_q5_1 *) dst + get_global_id(0);
    global uchar * qs = (global uchar *) src_qs + (QK5_1/2)*get_global_id(0);
    global uint  * qh = (global uint  *) src_qh + get_global_id(0);
    global half  * d  = (global half  *) src_d  + get_global_id(0);
    global half  * m  = (global half  *) src_m  + get_global_id(0);

    b->d = *d;
    b->m = *m;
    *((global uint *)(b->qh)) = *qh;
    for (int i = 0; i < QK5_1/2; ++i) {
        b->qs[i] = qs[i];
    }
}

kernel void kernel_convert_block_q5_1_noshuffle(
    global struct block_q5_1 * src0,
    global uchar * dst_q,
    global uint  * dst_qh,
    global half  * dst_d,
    global half  * dst_m
) {
    global struct block_q5_1 * b = (global struct block_q5_1 *) src0 + get_global_id(0);
    global uchar * q  = (global uchar *) dst_q + QK5_1/2*get_global_id(0);
    global uint  * qh = (global uint  *) dst_qh + get_global_id(0);
    global half  * d  = (global half  *) dst_d  + get_global_id(0);
    global half  * m  = (global half  *) dst_m  + get_global_id(0);

    *d = b->d;
    *m = b->m;
    *qh = *((global uint *)(b->qh));

    for (int i = 0; i < QK5_1/4; ++i) {
        uchar x0 = b->qs[2*i + 0];
        uchar x1 = b->qs[2*i + 1];

        q[i + 0      ] = convert_uchar(x0 & 0x0F) | convert_uchar((x1 & 0x0F) << 4);
        q[i + QK5_1/4] = convert_uchar((x0 & 0xF0) >> 4) | convert_uchar(x1 & 0xF0);

#ifdef ADRENO_GPU
        if (get_global_id(0) == 65536*4096) {
            printf("%04x - %02x\n", *(global ushort*)d, ((x0 & 0xF0) >> 4) | (x1 & 0xF0));
        }
#endif
    }
}

kernel void kernel_restore_block_q5_1_noshuffle(
    global uchar * src_q,
    global uint  * src_qh,
    global half  * src_d,
    global half  * src_m,
    global struct block_q5_1 * dst,
    uchar mask_0F,
    uchar mask_F0
) {
    global struct block_q5_1 * b = (global struct block_q5_1 *) dst + get_global_id(0);
    global uchar * q  = (global uchar *) src_q + QK5_1/2*get_global_id(0);
    global uint  * qh = (global uint  *) src_qh + get_global_id(0);
    global half  * d  = (global half  *) src_d  + get_global_id(0);
    global half  * m  = (global half  *) src_m  + get_global_id(0);

    b->d = *d;
    b->m = *m;
    *((global uint *)(b->qh)) = *qh;

    for (int i = 0; i < QK5_1/4; ++i) {
        uchar x0 = q[i + 0      ];
        uchar x1 = q[i + QK5_1/4];

        b->qs[2*i + 0] = convert_uchar((x0 & mask_0F) | ((x1 & mask_0F) << 4));
        b->qs[2*i + 1] = convert_uchar(((x0 & mask_F0) >> 4) | (x1 & mask_F0));
    }
}

kernel void kernel_convert_block_q5_1_trans4_ns(
    __global struct block_q5_1 * src0,
    __global uint * dst_qs,
    __global uint * dst_qh,
    __global half * dst_d,
    __global half * dst_m,
    uint ne00,
    uint ne01
) {
    uint i00 = get_global_id(1);
    uint i01 = get_global_id(0);
    uint i02 = get_global_id(2);

    if (i01 >= ne01) {
        return;
    }

    uint ne00_blk = ne00 / QK5_1;
    uint src_blk_offset = i00 + i01 * ne00_blk + i02 * ne00_blk * ne01;
    uint dst_blk_offset = i01 + i00 * ne01 + i02 * ne00_blk * ne01;

    global struct block_q5_1 * b = src0 + src_blk_offset;
    dst_d[dst_blk_offset] = b->d;
    dst_m[dst_blk_offset] = b->m;

    dst_qh[dst_blk_offset] = ((global uint *)(&(b->qh[0])))[0];

    // extract quantization and unshuffle
    ushort8 pre_block = ((global ushort8 *)(&(b->qs[0])))[0];
    ushort8 post_block = (ushort8)(0);

    uchar * pre_block_ptr = (uchar *)(&pre_block);
    uchar * post_block_ptr = (uchar *)(&post_block);

    for (int i = 0; i < QK5_1 / 4; ++i) {
        uchar x0 = pre_block_ptr[2*i + 0];
        uchar x1 = pre_block_ptr[2*i + 1];

        post_block_ptr[i + 0        ] = convert_uchar(x0 & 0x0F) | convert_uchar((x1 & 0x0F) << 4);
        post_block_ptr[i + QK5_1 / 4] = convert_uchar((x0 & 0xF0) >> 4) | convert_uchar(x1 & 0xF0);
    }

    uint4 q_block = as_uint4(post_block);

    uint offset = i02 * ne00_blk * ne01 * 4 + i00 * ne01 * 4 + i01;
    dst_qs[offset] = q_block.x;
    dst_qs[offset + ne01] = q_block.y;
    dst_qs[offset + ne01 * 2] = q_block.z;
    dst_qs[offset + ne01 * 3] = q_block.w;
}

kernel void kernel_restore_block_q5_1_trans4_ns(
    __global uint * src_qs,
    __global uint * src_qh,
    __global half * src_d,
    __global half * src_m,
    __global struct block_q5_1 * dst0,
    uint ne00,
    uint ne01
) {
    int i00 = get_global_id(1);
    uint i01 = get_global_id(0);
    uint i02 = get_global_id(2);

    if (i01 >= ne01) {
        return;
    }

    uint ne00_blk = ne00 / QK5_1;
    uint dst_blk_offset = i00 + i01 * ne00_blk + i02 * ne00_blk * ne01;
    uint src_blk_offset = i01 + i00 * ne01 + i02 * ne00_blk * ne01;

    __global struct block_q5_1 * b = dst0 + dst_blk_offset;
    b->d = src_d[src_blk_offset];
    b->m = src_m[src_blk_offset];

    ((__global uint *)(&(b->qh[0])))[0] = src_qh[src_blk_offset];

    // collect transposed quantization parts for a block
    uint src_q_offset = i02 * ne00_blk * ne01 * 4 + i00 * ne01 * 4 + i01;
    uint4 q_block;
    q_block.x = src_qs[src_q_offset];
    q_block.y = src_qs[src_q_offset + ne01];
    q_block.z = src_qs[src_q_offset + ne01 * 2];
    q_block.w = src_qs[src_q_offset + ne01 * 3];

    ushort8 post_block = as_ushort8(q_block);
    ushort8 pre_block = (ushort8)(0);

    uchar * pre_block_ptr = (uchar *)(&pre_block);
    uchar * post_block_ptr = (uchar *)(&post_block);

    for (int i = 0; i < QK5_1 / 4; ++i) {
        uchar x0 = post_block_ptr[i + 0];
        uchar x1 = post_block_ptr[i + QK5_1 / 4];

        pre_block_ptr[2 * i + 0] = convert_uchar(x0 & 0x0F) | convert_uchar((x1 & 0x0F) << 4);
        pre_block_ptr[2 * i + 1] = convert_uchar((x0 & 0xF0) >> 4) | convert_uchar(x1 & 0xF0);
    }
    ((__global ushort8 *)(&(b->qs[0])))[0] = pre_block;
}

kernel void kernel_convert_block_q4_k_trans4_ns(
    __global struct block_q4_K * src0,
    __global uint  * dst_q,
    __global half  * dst_d,
    __global half  * dst_dm,
    __global uchar * dst_s,
    uint ne00,
    uint ne01,
    uchar mask_0F,
    uchar mask_F0
) {
    uint i00 = get_global_id(1);
    uint i01 = get_global_id(0);
    uint i02 = get_global_id(2);

    if (i01 >= ne01) {
        return;
    }

    uint ne00_blk = ne00 / QK_K;
    uint src_blk_offset = i00 + i01 * ne00_blk + i02 * ne00_blk * ne01;
    uint dst_blk_offset = i01 + i00 * ne01     + i02 * ne00_blk * ne01;

    __global struct block_q4_K * b = src0 + src_blk_offset;

    dst_d [dst_blk_offset] = b->d;
    dst_dm[dst_blk_offset] = b->dm;

    uint4 qv[8];
    uchar * qv_bytes = (uchar *)qv;
    for (int i = 0; i < QK_K / 64; ++i) {
        for (int j = 0; j < 16; ++j) {
            uchar x0 = b->q[i*32 + 2*j];
            uchar x1 = b->q[i*32 + 2*j + 1];

            qv_bytes[i*32 + j     ] = convert_uchar(x0 & mask_0F) | convert_uchar((x1 & mask_0F) << 4);
            qv_bytes[i*32 + j + 16] = convert_uchar((x0 & mask_F0) >> 4) | convert_uchar(x1 & mask_F0);
        }
    }

    uint base = i02 * ne00_blk * ne01 * 32 + i00 * ne01 * 32 + i01;
    #pragma unroll
    for (int p = 0; p < 8; ++p) {
        uint4 v = qv[p];
        dst_q[base + (p * 4 + 0) * ne01] = v.x;
        dst_q[base + (p * 4 + 1) * ne01] = v.y;
        dst_q[base + (p * 4 + 2) * ne01] = v.z;
        dst_q[base + (p * 4 + 3) * ne01] = v.w;
    }

    __global uchar * s_dst = dst_s + (i02 * ne01 + i01) * ne00_blk * K_SCALE_SIZE + i00 * K_SCALE_SIZE;
    #pragma unroll
    for (int i = 0; i < K_SCALE_SIZE; ++i) {
        s_dst[i] = b->s[i];
    }
}

kernel void kernel_restore_block_q4_k_trans4_ns(
    __global uint  * src_q,
    __global half  * src_d,
    __global half  * src_dm,
    __global uchar * src_s,
    __global struct block_q4_K * dst0,
    uint ne00,
    uint ne01,
    uchar mask_0F,
    uchar mask_F0
) {
    uint i00 = get_global_id(1);  // block index along K
    uint i01 = get_global_id(0);  // row index
    uint i02 = get_global_id(2);  // batch index

    if (i01 >= ne01) {
        return;
    }

    uint ne00_blk = ne00 / QK_K;

    uint src_blk_offset = i01 + i00 * ne01 + i02 * ne00_blk * ne01;
    uint dst_blk_offset = i00 + i01 * ne00_blk + i02 * ne00_blk * ne01;

    __global struct block_q4_K * b = dst0 + dst_blk_offset;

    b->d  = src_d[src_blk_offset];
    b->dm = src_dm[src_blk_offset];

    __global uchar * s_src = src_s + (i02 * ne01 + i01) * ne00_blk * K_SCALE_SIZE + i00 * K_SCALE_SIZE;
    for (int i = 0; i < K_SCALE_SIZE; ++i) {
        b->s[i] = s_src[i];
    }

    uint base = i02 * ne00_blk * ne01 * 32 + i00 * ne01 * 32 + i01;

    uint4 qv[8];
    for (int p = 0; p < 8; ++p) {
        qv[p].x = src_q[base + (p * 4 + 0) * ne01];
        qv[p].y = src_q[base + (p * 4 + 1) * ne01];
        qv[p].z = src_q[base + (p * 4 + 2) * ne01];
        qv[p].w = src_q[base + (p * 4 + 3) * ne01];
    }

    uchar * qv_bytes = (uchar *)qv;
    for (int i = 0; i < QK_K / 64; ++i) {
        for (int j = 0; j < 16; ++j) {
            uchar lo = qv_bytes[i*32 + j];
            uchar hi = qv_bytes[i*32 + j + 16];
            b->q[i*32 + 2*j]     = convert_uchar((lo & mask_0F) | ((hi & mask_0F) << 4));
            b->q[i*32 + 2*j + 1] = convert_uchar(((lo & mask_F0) >> 4) | (hi & mask_F0));
        }
    }
}

//------------------------------------------------------------------------------
// kernel_convert_block_q4_k_tiled_ns
//
// Tiled-wide layout for the long-vocab q4_K lm_head/embed GEMV (decode path).
// Mirror of kernel_convert_block_q6_k_tiled_ns: recovers each weight's 4-bit
// code in CANONICAL ggml element order (e in [0,256)) and re-packs into 32 uints
// (8 codes/uint), stored TILED by 64 output rows so the matching GEMV
// (gemv_noshuffle_q4_k_f32_tiled) coalesces every weight load. The 12-byte
// packed scale block `s` and d/dm are stored per (row, K-block) tiled; the GEMV
// re-derives the 8 (scale,min) pairs via get_scale_min_k4, exactly like the o4
// kernel. Both ends owned here -> correct by construction vs the reference q4_K
// dequant. Requires ne01 % 64 == 0 (gated host-side). Buffer sizes identical to
// the trans4_ns layout.
//
//   q  uint4 granule g of (row r, K-block sb): idx = ((rt*ne00_blk+sb)*8 + g)*64 + rit
//   s  (12 bytes) of (r, sb):                  idx = (rt*ne00_blk+sb)*64 + rit, *12
//   d/dm (half)  of (r, sb):                   idx = (rt*ne00_blk+sb)*64 + rit
//   where rt = r/64, rit = r%64.
//------------------------------------------------------------------------------
kernel void kernel_convert_block_q4_k_tiled_ns(
    __global struct block_q4_K * src0,
    __global uint  * dst_q,    // 32 uints / superblock (4-bit codes, 8 codes/uint)
    __global half  * dst_d,    // 1 half  / superblock
    __global half  * dst_dm,   // 1 half  / superblock
    __global uchar * dst_s,    // K_SCALE_SIZE (12) bytes / superblock
    uint ne00,
    uint ne01
) {
    uint i00 = get_global_id(1);   // K-block index (superblock along ne00)
    uint i01 = get_global_id(0);   // output row index (along ne01)
    uint i02 = get_global_id(2);   // batch

    uint ne00_blk = ne00 / QK_K;

    uint src_blk_offset = i00 + i01 * ne00_blk + i02 * ne00_blk * ne01;
    __global struct block_q4_K * b = src0 + src_blk_offset;

    uint rt  = i01 / 64;
    uint rit = i01 % 64;
    uint tile_blk = (i02 * (ne01 / 64) + rt) * ne00_blk + i00;

    // --- recover canonical 4-bit codes in e-order, pack 8 codes/uint ---
    uint qw[32] = {0};
    for (uint e = 0; e < 256; ++e) {
        uint g    = e >> 6;           // group 0..3 (q advances 32 bytes/group)
        uint within = e & 63u;
        uint hlf  = within >> 5;      // 0 = low nibble, 1 = high nibble
        uint l    = within & 31u;     // 0..31
        uchar byte = b->q[g * 32u + l];
        uint code = (hlf == 0u) ? (uint)(byte & 0x0F) : (uint)(byte >> 4);
        qw[e >> 3] |= code << ((e & 7u) * 4u);
    }

    for (uint gr = 0; gr < 8; ++gr) {
        uint base = (tile_blk * 8u + gr) * 64u + rit;   // uint4 index
        dst_q[base * 4u + 0u] = qw[gr * 4u + 0u];
        dst_q[base * 4u + 1u] = qw[gr * 4u + 1u];
        dst_q[base * 4u + 2u] = qw[gr * 4u + 2u];
        dst_q[base * 4u + 3u] = qw[gr * 4u + 3u];
    }

    // packed scales (12 bytes), tiled per (row, block)
    __global uchar * s_dst = dst_s + (tile_blk * 64u + rit) * K_SCALE_SIZE;
    #pragma unroll
    for (int i = 0; i < K_SCALE_SIZE; ++i) {
        s_dst[i] = b->s[i];
    }

    dst_d [tile_blk * 64u + rit] = b->d;
    dst_dm[tile_blk * 64u + rit] = b->dm;
}

kernel void kernel_convert_block_q5_k_trans4_ns(
    __global struct block_q5_K * src0,
    __global uint  * dst_qs,
    __global uint  * dst_qh,
    __global half  * dst_d,
    __global half  * dst_dm,
    __global uchar * dst_s,
    uint ne00,
    uint ne01,
    uchar mask_0F,
    uchar mask_F0
) {
    uint i00 = get_global_id(1);
    uint i01 = get_global_id(0);
    uint i02 = get_global_id(2);

    if (i01 >= ne01) {
        return;
    }

    uint ne00_blk = ne00 / QK_K;
    uint src_blk_offset = i00 + i01 * ne00_blk + i02 * ne00_blk * ne01;
    uint dst_blk_offset = i01 + i00 * ne01     + i02 * ne00_blk * ne01;

    __global struct block_q5_K * b = src0 + src_blk_offset;

    dst_d [dst_blk_offset] = b->d;
    dst_dm[dst_blk_offset] = b->dm;

    for (int k = 0; k < 8; k++) {
        uchar b0 = 0, b1 = 0, b2 = 0, b3 = 0;
        for (int bit = 0; bit < 8; bit++) {
            b0 |= (uchar)(((b->qh[bit]      >> k) & 1) << bit);
            b1 |= (uchar)(((b->qh[8  + bit] >> k) & 1) << bit);
            b2 |= (uchar)(((b->qh[16 + bit] >> k) & 1) << bit);
            b3 |= (uchar)(((b->qh[24 + bit] >> k) & 1) << bit);
        }
        uint packed = (uint)b0 | ((uint)b1 << 8) | ((uint)b2 << 16) | ((uint)b3 << 24);
        dst_qh[i01 + (i00 * 8 + k) * ne01 + i02 * ne00_blk * 8 * ne01] = packed;
    }

    uint4 qv[8];
    uchar * qv_bytes = (uchar *)qv;
    for (int i = 0; i < QK_K / 64; ++i) {
        for (int j = 0; j < 16; ++j) {
            uchar x0 = b->qs[i*32 + 2*j];
            uchar x1 = b->qs[i*32 + 2*j + 1];

            qv_bytes[i*32 + j     ] = convert_uchar(x0 & mask_0F) | convert_uchar((x1 & mask_0F) << 4);
            qv_bytes[i*32 + j + 16] = convert_uchar((x0 & mask_F0) >> 4) | convert_uchar(x1 & mask_F0);
        }
    }

    uint base = i02 * ne00_blk * ne01 * 32 + i00 * ne01 * 32 + i01;
    #pragma unroll
    for (int p = 0; p < 8; ++p) {
        uint4 v = qv[p];
        dst_qs[base + (p * 4 + 0) * ne01] = v.x;
        dst_qs[base + (p * 4 + 1) * ne01] = v.y;
        dst_qs[base + (p * 4 + 2) * ne01] = v.z;
        dst_qs[base + (p * 4 + 3) * ne01] = v.w;
    }

    __global uchar * s_dst = dst_s + (i02 * ne01 + i01) * ne00_blk * K_SCALE_SIZE + i00 * K_SCALE_SIZE;
    #pragma unroll
    for (int i = 0; i < K_SCALE_SIZE; ++i) {
        s_dst[i] = b->s[i];
    }
}

kernel void kernel_restore_block_q5_k_trans4_ns(
    __global uint  * src_qs,
    __global uint  * src_qh,
    __global half  * src_d,
    __global half  * src_dm,
    __global uchar * src_s,
    __global struct block_q5_K * dst0,
    uint ne00,
    uint ne01,
    uchar mask_0F,
    uchar mask_F0
) {
    uint i00 = get_global_id(1);  // block index along K
    uint i01 = get_global_id(0);  // row index
    uint i02 = get_global_id(2);  // batch index

    if (i01 >= ne01) {
        return;
    }

    uint ne00_blk = ne00 / QK_K;

    uint src_blk_offset = i01 + i00 * ne01 + i02 * ne00_blk * ne01;
    uint dst_blk_offset = i00 + i01 * ne00_blk + i02 * ne00_blk * ne01;

    __global struct block_q5_K * b = dst0 + dst_blk_offset;

    b->d  = src_d[src_blk_offset];
    b->dm = src_dm[src_blk_offset];

    for (int j = 0; j < 32; j++) b->qh[j] = 0;
    for (int k = 0; k < 8; k++) {
        uint packed = src_qh[i01 + (i00 * 8 + k) * ne01 + i02 * ne00_blk * 8 * ne01];
        uchar b0 = (uchar)(packed & 0xFF);
        uchar b1 = (uchar)((packed >> 8) & 0xFF);
        uchar b2 = (uchar)((packed >> 16) & 0xFF);
        uchar b3 = (uchar)((packed >> 24) & 0xFF);
        for (int bit = 0; bit < 8; bit++) {
            b->qh[bit]      |= (uchar)(((b0 >> bit) & 1) << k);
            b->qh[8  + bit] |= (uchar)(((b1 >> bit) & 1) << k);
            b->qh[16 + bit] |= (uchar)(((b2 >> bit) & 1) << k);
            b->qh[24 + bit] |= (uchar)(((b3 >> bit) & 1) << k);
        }
    }

    __global uchar * s_src = src_s + (i02 * ne01 + i01) * ne00_blk * K_SCALE_SIZE + i00 * K_SCALE_SIZE;
    for (int i = 0; i < K_SCALE_SIZE; ++i) {
        b->s[i] = s_src[i];
    }

    uint base = i02 * ne00_blk * ne01 * 32 + i00 * ne01 * 32 + i01;

    uint4 qv[8];
    for (int p = 0; p < 8; ++p) {
        qv[p].x = src_qs[base + (p * 4 + 0) * ne01];
        qv[p].y = src_qs[base + (p * 4 + 1) * ne01];
        qv[p].z = src_qs[base + (p * 4 + 2) * ne01];
        qv[p].w = src_qs[base + (p * 4 + 3) * ne01];
    }

    uchar * qv_bytes = (uchar *)qv;
    for (int i = 0; i < QK_K / 64; ++i) {
        for (int j = 0; j < 16; ++j) {
            uchar lo = qv_bytes[i*32 + j];
            uchar hi = qv_bytes[i*32 + j + 16];
            b->qs[i*32 + 2*j]     = convert_uchar((lo & mask_0F) | ((hi & mask_0F) << 4));
            b->qs[i*32 + 2*j + 1] = convert_uchar(((lo & mask_F0) >> 4) | (hi & mask_F0));
        }
    }
}

kernel void kernel_convert_block_q6_k_trans4_ns(
    __global struct block_q6_K * src0,
    __global uint  * dst_ql,
    __global uint  * dst_qh,
    __global half  * dst_d,
    __global char  * dst_s,
    uint ne00,
    uint ne01,
    uchar mask_0F,
    uchar mask_F0
) {
    uint i00 = get_global_id(1);
    uint i01 = get_global_id(0);
    uint i02 = get_global_id(2);

    if (i01 >= ne01) {
        return;
    }

    uint ne00_blk = ne00 / QK_K;

    uint src_blk_offset = i00 + i01 * ne00_blk + i02 * ne00_blk * ne01;
    uint dst_blk_offset = i01 + i00 * ne01     + i02 * ne00_blk * ne01;

    __global struct block_q6_K * b = src0 + src_blk_offset;

    dst_d[dst_blk_offset] = b->d;

    uint4 qlv[8];
    uchar * qlv_bytes = (uchar *)qlv;
    for (int i = 0; i < 2; ++i) {
        for (int j = 0; j < 16; ++j) {
            uchar x0 = b->ql[i*64 + 2*j];
            uchar x1 = b->ql[i*64 + 2*j + 1];
            uchar x2 = b->ql[i*64 + 32 + 2*j];
            uchar x3 = b->ql[i*64 + 32 + 2*j + 1];
            qlv_bytes[i*64 + j     ] = convert_uchar(x0 & mask_0F) | convert_uchar((x1 & mask_0F) << 4);
            qlv_bytes[i*64 + j + 16] = convert_uchar(x2 & mask_0F) | convert_uchar((x3 & mask_0F) << 4);
            qlv_bytes[i*64 + j + 32] = convert_uchar((x0 & mask_F0) >> 4) | convert_uchar(x1 & mask_F0);
            qlv_bytes[i*64 + j + 48] = convert_uchar((x2 & mask_F0) >> 4) | convert_uchar(x3 & mask_F0);
        }
    }

    uint ql_base = i02 * ne00_blk * ne01 * 32 + i00 * ne01 * 32 + i01;

    #pragma unroll
    for (int p = 0; p < 8; ++p) {
        uint4 v = qlv[p];
        dst_ql[ql_base + (p * 4 + 0) * ne01] = v.x;
        dst_ql[ql_base + (p * 4 + 1) * ne01] = v.y;
        dst_ql[ql_base + (p * 4 + 2) * ne01] = v.z;
        dst_ql[ql_base + (p * 4 + 3) * ne01] = v.w;
    }

    uint qhv[16] = {0};

    for (int n = 0; n < 2; ++n) {
        for (int l = 0; l < 32; ++l) {
            uchar h = b->qh[n*32 + l];
            int u = l / 16;
            int bit_pos = (l % 16) * 2;
            qhv[(n*4 + 0)*2 + u] |= ((uint)((h >> 0) & 0x03)) << bit_pos;
            qhv[(n*4 + 1)*2 + u] |= ((uint)((h >> 2) & 0x03)) << bit_pos;
            qhv[(n*4 + 2)*2 + u] |= ((uint)((h >> 4) & 0x03)) << bit_pos;
            qhv[(n*4 + 3)*2 + u] |= ((uint)((h >> 6) & 0x03)) << bit_pos;
        }
    }

    uint qh_base = i02 * ne00_blk * ne01 * 16 + i00 * ne01 * 16 + i01;

    for (int p = 0; p < 16; ++p) {
        dst_qh[qh_base + p * ne01] = qhv[p];
    }

    __global char * s_dst = dst_s + (i02 * ne01 + i01) * ne00_blk * 16 + i00 * 16;
    #pragma unroll
    for (int i = 0; i < 16; ++i) {
        s_dst[i] = b->scales[i];
    }
}

kernel void kernel_restore_block_q6_k_trans4_ns(
    __global uint  * src_ql,
    __global uint  * src_qh,
    __global half  * src_d,
    __global char  * src_s,
    __global struct block_q6_K * dst0,
    uint ne00,
    uint ne01,
    uchar mask_0F,
    uchar mask_F0
) {
    uint i00 = get_global_id(1);  // block index along K
    uint i01 = get_global_id(0);  // row index
    uint i02 = get_global_id(2);  // batch index

    if (i01 >= ne01) {
        return;
    }

    uint ne00_blk = ne00 / QK_K;

    uint src_blk_offset = i01 + i00 * ne01 + i02 * ne00_blk * ne01;
    uint dst_blk_offset = i00 + i01 * ne00_blk + i02 * ne00_blk * ne01;

    __global struct block_q6_K * b = dst0 + dst_blk_offset;

    b->d = src_d[src_blk_offset];

    uint ql_base = i02 * ne00_blk * ne01 * 32 + i00 * ne01 * 32 + i01;
    uint4 qlv[8];
    for (int p = 0; p < 8; ++p) {
        qlv[p].x = src_ql[ql_base + (p * 4 + 0) * ne01];
        qlv[p].y = src_ql[ql_base + (p * 4 + 1) * ne01];
        qlv[p].z = src_ql[ql_base + (p * 4 + 2) * ne01];
        qlv[p].w = src_ql[ql_base + (p * 4 + 3) * ne01];
    }

    uchar * qlv_bytes = (uchar *)qlv;
    for (int i = 0; i < 2; ++i) {
        for (int j = 0; j < 16; ++j) {
            uchar lo_02 = qlv_bytes[i*64 + j];
            uchar lo_13 = qlv_bytes[i*64 + j + 16];
            uchar hi_02 = qlv_bytes[i*64 + j + 32];
            uchar hi_13 = qlv_bytes[i*64 + j + 48];
            b->ql[i*64 + 2*j]          = convert_uchar((lo_02 & mask_0F) | ((hi_02 & mask_0F) << 4));
            b->ql[i*64 + 2*j + 1]      = convert_uchar(((lo_02 & mask_F0) >> 4) | (hi_02 & mask_F0));
            b->ql[i*64 + 32 + 2*j]     = convert_uchar((lo_13 & mask_0F) | ((hi_13 & mask_0F) << 4));
            b->ql[i*64 + 32 + 2*j + 1] = convert_uchar(((lo_13 & mask_F0) >> 4) | (hi_13 & mask_F0));
        }
    }

    uint qh_base = i02 * ne00_blk * ne01 * 16 + i00 * ne01 * 16 + i01;
    uint qhv[16];
    for (int p = 0; p < 16; ++p) {
        qhv[p] = src_qh[qh_base + p * ne01];
    }

    for (int n = 0; n < 2; ++n) {
        for (int l = 0; l < 32; ++l) {
            int u = l / 16;
            int bit_pos = (l % 16) * 2;
            uchar v0 = (uchar)((qhv[(n*4 + 0)*2 + u] >> bit_pos) & 0x03);
            uchar v1 = (uchar)((qhv[(n*4 + 1)*2 + u] >> bit_pos) & 0x03);
            uchar v2 = (uchar)((qhv[(n*4 + 2)*2 + u] >> bit_pos) & 0x03);
            uchar v3 = (uchar)((qhv[(n*4 + 3)*2 + u] >> bit_pos) & 0x03);
            b->qh[n*32 + l] = v0 | (v1 << 2) | (v2 << 4) | (v3 << 6);
        }
    }

    __global char * s_src = src_s + (i02 * ne01 + i01) * ne00_blk * 16 + i00 * 16;
    for (int i = 0; i < 16; ++i) {
        b->scales[i] = s_src[i];
    }
}

//------------------------------------------------------------------------------
// block_mxfp4
//------------------------------------------------------------------------------
#define QK_MXFP4 32
struct block_mxfp4 {
    uchar e; // E8M0
    uchar qs[QK_MXFP4 / 2];
};

//------------------------------------------------------------------------------
// kernel_convert_block_mxfp4
// Convert the block_mxfp4 format to 2 separate arrays (AOS -> SOA).
// This kernel does not deshuffle the bits.
//------------------------------------------------------------------------------
kernel void kernel_convert_block_mxfp4(
    global struct block_mxfp4 * src0,
    global uchar * dst_q,
    global uchar * dst_e
) {
    global struct block_mxfp4 * b = (global struct block_mxfp4 *) src0 + get_global_id(0);
    global uchar * q = (global uchar *) dst_q + QK_MXFP4 / 2 * get_global_id(0);
    global uchar * e = (global uchar *) dst_e + get_global_id(0);

    *e = b->e;

    for (int i = 0; i < QK_MXFP4 / 2; ++i) {
        q[i] = b->qs[i];
    }
}

kernel void kernel_convert_block_mxfp4_trans(
    global struct block_mxfp4 * src0,
    __global uint4 * dst_q,
    __global uchar * dst_e,
    uint ne00,
    uint ne01
) {
    int i00 = get_global_id(1);
    uint i01 = get_global_id(0);
    uint i02 = get_global_id(2);

    uint ne00_blk = ne00 / QK_MXFP4;
    uint src_blk_offset = i00 + i01 * ne00_blk + i02 * ne00_blk * ne01;
    uint dst_blk_offset = i01 + i00 * ne01 + i02 * ne00_blk * ne01;

    global struct block_mxfp4 * b = src0 + src_blk_offset;

    dst_q[dst_blk_offset] = ((global uint4 *)(&(b->qs[0])))[0];
    dst_e[dst_blk_offset] = b->e;
}

kernel void kernel_restore_block_mxfp4(
    global uchar * src_q,
    global half  * src_e,
    global struct block_mxfp4 * dst
) {
    global struct block_mxfp4 * b = (global struct block_mxfp4 *) dst + get_global_id(0);
    global uchar * q = (global uchar *) src_q + QK_MXFP4 / 2 * get_global_id(0);
    global uchar * e = (global uchar *) src_e + get_global_id(0);

    b->e = *e;
    for (int i = 0; i < QK_MXFP4 / 2; ++i) {
        b->qs[i] = q[i];
    }
}

kernel void kernel_restore_block_mxfp4_trans(
    __global uint4 * src_q,
    __global uchar * src_e,
    global struct block_mxfp4 * dst,
    uint ne00,
    uint ne01
) {
    int i00 = get_global_id(1);
    uint i01 = get_global_id(0);
    uint i02 = get_global_id(2);

    uint ne00_blk = ne00 / QK_MXFP4;
    uint src_blk_offset = i01 + i00 * ne01 + i02 * ne00_blk * ne01;
    uint dst_blk_offset = i00 + i01 * ne00_blk + i02 * ne00_blk * ne01;

    global struct block_mxfp4 * b = dst + dst_blk_offset;

    ((global uint4 *)(&(b->qs[0])))[0] = src_q[src_blk_offset];
    b->e = src_e[src_blk_offset];
}

//------------------------------------------------------------------------------
// kernel_convert_block_q6_k_tiled_ns
//
// Tiled-wide layout for the long-vocab q6_K lm_head/embed GEMV (decode path).
// Unlike *_trans4_ns (which mirrors the bit-interleave the legacy 2-output GEMV
// consumes), this kernel is correct-by-construction against the CANONICAL ggml
// q6_K dequant: it recovers each weight's 6-bit code in element order e in
// [0,256), then re-packs low-4-bits into 32 uints (8 codes/uint) and high-2-bits
// into 16 uints (16 codes/uint). The matching GEMV (gemv_noshuffle_q6_k_f32_tiled)
// unpacks the same order, so both ends are owned here.
//
// Storage is TILED by 64 output rows so the GEMV's 64-thread tile coalesces:
//   ql uint4 granule g of (row r, K-block sb): idx = ((rt*ne00_blk + sb)*8 + g)*64 + rit
//   qh uint4 granule g:                         idx = ((rt*ne00_blk + sb)*4 + g)*64 + rit
//   scales (char16) of (r, sb):                 idx = (rt*ne00_blk + sb)*64 + rit
//   d (half) of (r, sb):                        idx = (rt*ne00_blk + sb)*64 + rit
// where rt = r/64, rit = r%64. Requires ne01 % 64 == 0 (gated host-side).
// Buffer sizes are byte-identical to the trans4_ns layout.
//------------------------------------------------------------------------------
kernel void kernel_convert_block_q6_k_tiled_ns(
    __global struct block_q6_K * src0,
    __global uint  * dst_ql,   // 32 uints / superblock (low 4 bits, 8 codes/uint)
    __global uint  * dst_qh,   // 16 uints / superblock (high 2 bits, 16 codes/uint)
    __global half  * dst_d,    // 1 half  / superblock
    __global char  * dst_s,    // 16 chars/ superblock
    uint ne00,
    uint ne01
) {
    uint i00 = get_global_id(1);   // K-block index (superblock along ne00)
    uint i01 = get_global_id(0);   // output row index (along ne01)
    uint i02 = get_global_id(2);   // batch

    uint ne00_blk = ne00 / QK_K;

    // Source block: row-major over (i02, i01, i00).
    uint src_blk_offset = i00 + i01 * ne00_blk + i02 * ne00_blk * ne01;
    __global struct block_q6_K * b = src0 + src_blk_offset;

    uint rt  = i01 / 64;
    uint rit = i01 % 64;
    uint tile_blk = (i02 * (ne01 / 64) + rt) * ne00_blk + i00;  // tile-major (row-tile, K-block)

    // --- recover canonical 6-bit codes, pack into ql (4b) + qh (2b) in e-order ---
    // 32 ql-uints (8 low-nibbles each) + 16 qh-uints (16 2-bit slots each).
    uint qlw[32] = {0};
    uint qhw[16] = {0};

    for (uint e = 0; e < 256; ++e) {
        uint n   = (e >= 128) ? 1u : 0u;       // which 128-half
        uint within = e - n * 128u;
        uint q   = within / 32u;               // quadrant 0..3
        uint l   = within % 32u;               // 0..31

        uint off_ql = n * 64u;                 // raw ql byte base for this half
        uint off_qh = n * 32u;                 // raw qh byte base for this half

        uchar low4;
        uchar qlb0 = b->ql[off_ql + l];
        uchar qlb1 = b->ql[off_ql + l + 32];
        if (q == 0)      low4 = qlb0 & 0x0F;
        else if (q == 1) low4 = qlb1 & 0x0F;
        else if (q == 2) low4 = (qlb0 >> 4) & 0x0F;
        else             low4 = (qlb1 >> 4) & 0x0F;

        uchar hi2 = (b->qh[off_qh + l] >> (q * 2u)) & 0x03;

        // pack low4 (e-order): uint e/8, nibble (e%8)
        qlw[e >> 3] |= ((uint)low4) << ((e & 7u) * 4u);
        // pack hi2 (e-order): uint e/16, 2-bit slot (e%16)
        qhw[e >> 4] |= ((uint)hi2)  << ((e & 15u) * 2u);
    }

    // --- write tiled ---
    for (uint g = 0; g < 8; ++g) {
        uint base = (tile_blk * 8u + g) * 64u + rit;   // uint4 index
        dst_ql[base * 4u + 0u] = qlw[g * 4u + 0u];
        dst_ql[base * 4u + 1u] = qlw[g * 4u + 1u];
        dst_ql[base * 4u + 2u] = qlw[g * 4u + 2u];
        dst_ql[base * 4u + 3u] = qlw[g * 4u + 3u];
    }
    for (uint g = 0; g < 4; ++g) {
        uint base = (tile_blk * 4u + g) * 64u + rit;   // uint4 index
        dst_qh[base * 4u + 0u] = qhw[g * 4u + 0u];
        dst_qh[base * 4u + 1u] = qhw[g * 4u + 1u];
        dst_qh[base * 4u + 2u] = qhw[g * 4u + 2u];
        dst_qh[base * 4u + 3u] = qhw[g * 4u + 3u];
    }

    // scales: 16 chars contiguous per (row, block), tiled
    __global char * s_dst = dst_s + (tile_blk * 64u + rit) * 16u;
    #pragma unroll
    for (int i = 0; i < 16; ++i) {
        s_dst[i] = b->scales[i];
    }

    // super-block scale
    dst_d[tile_blk * 64u + rit] = b->d;
}

kernel void kernel_convert_block_mxfp4_trans4_ns(
    global struct block_mxfp4 * src0,
    __global uint * dst_q,
    __global uchar * dst_e,
    uint ne00,
    uint ne01
) {
    uint i00 = get_global_id(1);
    uint i01 = get_global_id(0);
    uint i02 = get_global_id(2);

    if (i01 >= ne01) {
        return;
    }

    uint ne00_blk = ne00 / QK_MXFP4;
    uint src_blk_offset = i00 + i01 * ne00_blk + i02 * ne00_blk * ne01;
    uint dst_blk_offset = i01 + i00 * ne01 + i02 * ne00_blk * ne01;

    global struct block_mxfp4 * b = src0 + src_blk_offset;
    dst_e[dst_blk_offset] = b->e;

    // extract quantization and unshuffle
    ushort8 pre_block = ((global ushort8 *)(&(b->qs[0])))[0];

    ushort8 post_block = (ushort8)(0);

    uchar * pre_block_ptr = (uchar *)(&pre_block);
    uchar * post_block_ptr = (uchar *)(&post_block);

    for (int i = 0; i < QK_MXFP4 / 4; ++i) {
        uchar x0 = pre_block_ptr[2*i + 0];
        uchar x1 = pre_block_ptr[2*i + 1];

        post_block_ptr[i + 0        ] = convert_uchar(x0 & 0x0F) | convert_uchar((x1 & 0x0F) << 4);
        post_block_ptr[i + QK_MXFP4 / 4] = convert_uchar((x0 & 0xF0) >> 4) | convert_uchar(x1 & 0xF0);
    }

    uint4 q_block = as_uint4(post_block);

    uint offset = i02 * ne00_blk * ne01 * 4 + i00 * ne01 * 4 + i01;
    dst_q[offset] = q_block.x;
    dst_q[offset + ne01] = q_block.y;
    dst_q[offset + ne01 * 2] = q_block.z;
    dst_q[offset + ne01 * 3] = q_block.w;
}

kernel void kernel_restore_block_mxfp4_trans4_ns(
    __global uint * src_q,
    __global uchar * src_e,
    __global struct block_mxfp4 * dst0,
    uint ne00,
    uint ne01
) {
    uint i00 = get_global_id(1);
    uint i01 = get_global_id(0);
    uint i02 = get_global_id(2);

    if (i01 >= ne01) {
        return;
    }

    uint ne00_blk = ne00 / QK_MXFP4;
    uint dst_blk_offset = i00 + i01 * ne00_blk + i02 * ne00_blk * ne01;
    uint src_d_offset = i01 + i00 * ne01 + i02 * ne00_blk * ne01;

    __global struct block_mxfp4 * b = dst0 + dst_blk_offset;
    b->e = src_e[src_d_offset];

    // collect transposed quantization parts for a block
    uint src_q_offset = i02 * ne00_blk * ne01 * 4 + i00 * ne01 * 4 + i01;
    uint4 q_block;
    q_block.x = src_q[src_q_offset];
    q_block.y = src_q[src_q_offset + ne01];
    q_block.z = src_q[src_q_offset + ne01 * 2];
    q_block.w = src_q[src_q_offset + ne01 * 3];

    ushort8 post_block = as_ushort8(q_block);
    ushort8 pre_block = (ushort8)(0);

    uchar * pre_block_ptr = (uchar *)(&pre_block);
    uchar * post_block_ptr = (uchar *)(&post_block);

    for (int i = 0; i < QK_MXFP4 / 4; ++i) {
        uchar x0 = post_block_ptr[i + 0];
        uchar x1 = post_block_ptr[i + QK_MXFP4 / 4];

        pre_block_ptr[2 * i + 0] = convert_uchar(x0 & 0x0F) | convert_uchar((x1 & 0x0F) << 4);
        pre_block_ptr[2 * i + 1] = convert_uchar((x0 & 0xF0) >> 4) | convert_uchar(x1 & 0xF0);
    }

    ((__global ushort8 *)(&(b->qs[0])))[0] = pre_block;
}


//------------------------------------------------------------------------------
// block_q8_0
//------------------------------------------------------------------------------
typedef struct {
    half d;       // delta
    char qs[QK8_0]; // quants
} block_q8_0;

kernel void kernel_convert_block_q8_0(
    global block_q8_0 * src0,
    global uchar * dst_q,
    global half  * dst_d
) {
    global block_q8_0 * b = (global block_q8_0 *) src0 + get_global_id(0);
    global uchar      * q = (global uchar *) dst_q + QK8_0*get_global_id(0);
    global half       * d = (global half *) dst_d + get_global_id(0);

    *d = b->d;

    for (int i = 0; i < QK8_0; ++i) {
        q[i] = b->qs[i];
    }
}

kernel void kernel_restore_block_q8_0(
    global uchar * src_q,
    global half  * src_d,
    global block_q8_0 * dst
) {
    global block_q8_0 * b = (global block_q8_0 *) dst + get_global_id(0);
    global uchar      * q = (global uchar *) src_q + QK8_0*get_global_id(0);
    global half       * d = (global half *) src_d + get_global_id(0);

    b->d = *d;
    for (int i = 0; i < QK8_0; ++i) {
        b->qs[i] = q[i];
    }
}

// View-aware AoS q8_0 -> f32 dequant (f32/f32 FA path).
kernel void kernel_dequant_q8_0_f32_view_aos(
    global char * src,
    ulong         src_offset,
    ulong         src_nb1,
    ulong         src_nb2,
    ulong         src_nb3,
    int           nblk0,
    int           ne1,
    int           ne2,
    int           ne3,
    global float * dst
) {
    int blk_i0 = get_global_id(0);
    int i1     = get_global_id(1);
    int batch  = get_global_id(2);

    if (blk_i0 >= nblk0) return;
    if (i1     >= ne1)   return;

    int i2 = batch % ne2;
    int i3 = batch / ne2;
    if (i3 >= ne3) return;

    global char * block = src + src_offset + (ulong)i3*src_nb3 + (ulong)i2*src_nb2 + (ulong)i1*src_nb1 + (ulong)blk_i0 * (2 + QK8_0);
    float d = vload_half(0, (global half *)block);
    global char * qs = block + 2;

    ulong dst_row_base = ((ulong)i3 * ne2 * ne1 + (ulong)i2 * ne1 + (ulong)i1) * nblk0;
    global float * out = dst + (dst_row_base + blk_i0) * QK8_0;

    for (int i = 0; i < QK8_0; ++i) {
        out[i] = d * (float)qs[i];
    }
}

// View-aware AoS q8_0 -> f16 dequant. Rows tight, batch strides may be gapped.
kernel void kernel_dequant_q8_0_f16_view_aos(
    global char * src,
    ulong         src_offset,
    ulong         src_nb1,
    ulong         src_nb2,
    ulong         src_nb3,
    int           nblk0,
    int           ne1,
    int           ne2,
    int           ne3,
    global half * dst
) {
    int blk_i0 = get_global_id(0);
    int i1     = get_global_id(1);
    int batch  = get_global_id(2);

    if (blk_i0 >= nblk0) return;
    if (i1     >= ne1)   return;

    int i2 = batch % ne2;
    int i3 = batch / ne2;
    if (i3 >= ne3) return;

    global char * block = src + src_offset + (ulong)i3*src_nb3 + (ulong)i2*src_nb2 + (ulong)i1*src_nb1 + (ulong)blk_i0 * (2 + QK8_0);
    float d = vload_half(0, (global half *)block);
    global char * qs = block + 2;

    ulong dst_row_base = ((ulong)i3 * ne2 * ne1 + (ulong)i2 * ne1 + (ulong)i1) * nblk0;
    global half * out = dst + (dst_row_base + blk_i0) * QK8_0;

    for (int i = 0; i < QK8_0; ++i) {
        out[i] = (half)(d * (float)qs[i]);
    }
}

// View-aware AoS q4_0 -> f32 dequant (mirrors the q8_0 view variant).
kernel void kernel_dequant_q4_0_f32_view_aos(
    global char * src,
    ulong         src_offset,
    ulong         src_nb1,
    ulong         src_nb2,
    ulong         src_nb3,
    int           nblk0,
    int           ne1,
    int           ne2,
    int           ne3,
    global float * dst
) {
    int blk_i0 = get_global_id(0);
    int i1     = get_global_id(1);
    int batch  = get_global_id(2);

    if (blk_i0 >= nblk0) return;
    if (i1     >= ne1)   return;

    int i2 = batch % ne2;
    int i3 = batch / ne2;
    if (i3 >= ne3) return;

    global char * block = src + src_offset + (ulong)i3*src_nb3 + (ulong)i2*src_nb2 + (ulong)i1*src_nb1 + (ulong)blk_i0 * (2 + QK4_0/2);
    float d = vload_half(0, (global half *)block);
    global uchar * qs = (global uchar *)(block + 2);

    ulong dst_row_base = ((ulong)i3 * ne2 * ne1 + (ulong)i2 * ne1 + (ulong)i1) * nblk0;
    global float * out = dst + (dst_row_base + blk_i0) * QK4_0;

    for (int i = 0; i < QK4_0/2; ++i) {
        uchar byte = qs[i];
        int q0 = (int)(byte & 0x0F) - 8;
        int q1 = (int)(byte >> 4)   - 8;
        out[i]            = d * (float)q0;
        out[i + QK4_0/2]  = d * (float)q1;
    }
}

// View-aware AoS q4_0 -> f16 dequant (mirrors the q8_0 view variant).
kernel void kernel_dequant_q4_0_f16_view_aos(
    global char * src,
    ulong         src_offset,
    ulong         src_nb1,
    ulong         src_nb2,
    ulong         src_nb3,
    int           nblk0,
    int           ne1,
    int           ne2,
    int           ne3,
    global half * dst
) {
    int blk_i0 = get_global_id(0);
    int i1     = get_global_id(1);
    int batch  = get_global_id(2);

    if (blk_i0 >= nblk0) return;
    if (i1     >= ne1)   return;

    int i2 = batch % ne2;
    int i3 = batch / ne2;
    if (i3 >= ne3) return;

    global char * block = src + src_offset + (ulong)i3*src_nb3 + (ulong)i2*src_nb2 + (ulong)i1*src_nb1 + (ulong)blk_i0 * (2 + QK4_0/2);
    float d = vload_half(0, (global half *)block);
    global uchar * qs = (global uchar *)(block + 2);

    ulong dst_row_base = ((ulong)i3 * ne2 * ne1 + (ulong)i2 * ne1 + (ulong)i1) * nblk0;
    global half * out = dst + (dst_row_base + blk_i0) * QK4_0;

    for (int i = 0; i < QK4_0/2; ++i) {
        uchar byte = qs[i];
        int q0 = (int)(byte & 0x0F) - 8;
        int q1 = (int)(byte >> 4)   - 8;
        out[i]          = (half)(d * (float)q0);
        out[i + QK4_0/2] = (half)(d * (float)q1);
    }
}

kernel void kernel_restore_block_q8_0_trans(
    global uchar * src_q,
    global half  * src_d,
    global block_q8_0 * dst,
    uint ne00,
    uint ne01
){
    uint num_blk_per_row = ne00 / QK8_0;

    global block_q8_0 * b = (global block_q8_0 *) dst + get_global_id(0) * num_blk_per_row;
    global uchar      * q = (global uchar *) src_q + get_global_id(0) * 4; // 4 8-bit packed
    global half       * d = (global half *) src_d + get_global_id(0);

    for (uint blk = 0; blk < num_blk_per_row; blk++) {
        b->d = *d;

        for (uint i = 0; i < QK8_0; i+=4) {
            b->qs[i]   = q[0];
            b->qs[i+1] = q[1];
            b->qs[i+2] = q[2];
            b->qs[i+3] = q[3];

            q += 4 * ne01; // M stride
        }

        d += ne01;

        b++;
    }
}

//------------------------------------------------------------------------------
// kernel_convert_block_q4_K
// Convert the block_q4_K format to 4 separate arrays (AOS -> SOA).
// This kernel does not deshuffle the bits.
// Each thread processes a super block.
// Mask args are just to keep the signature consistent with the no-shuffle
// version and they are not used in this kernel.
//------------------------------------------------------------------------------
kernel void kernel_convert_block_q4_K(
    global struct block_q4_K * src0,
    global uchar * dst_q,
    global uchar * dst_s,
    global half  * dst_d,
    global half  * dst_dm,
    uchar mask_0F,
    uchar mask_F0
) {
    global struct block_q4_K * b = (global struct block_q4_K *) src0 + get_global_id(0);
    global uchar * q  = (global uchar *) dst_q  + QK_K/2*get_global_id(0);
    global uchar * s  = (global uchar *) dst_s  + K_SCALE_SIZE*get_global_id(0);
    global half  * d  = (global half  *) dst_d  + get_global_id(0);
    global half  * dm = (global half  *) dst_dm + get_global_id(0);

    *d  = b->d;
    *dm = b->dm;

    for (int i = 0; i < QK_K/2; ++i) {
        q[i] = b->q[i];
    }
    for (int i = 0; i < K_SCALE_SIZE; ++i) {
        s[i] = b->s[i];
    }
}

// Restore block_q4_K from flattened arrays.
// Each thread processes a super block.
// Mask args are just to keep the signature consistent with the no-shuffle ones.
kernel void kernel_restore_block_q4_K(
    global uchar * src_q,
    global uchar * src_s,
    global half  * src_d,
    global half  * src_dm,
    global struct block_q4_K * dst,
    uchar mask_0F,
    uchar mask_F0
) {
    global struct block_q4_K * b = (global struct block_q4_K *) dst + get_global_id(0);
    global uchar * q  = (global uchar *) src_q  + QK_K/2*get_global_id(0);
    global uchar * s  = (global uchar *) src_s + K_SCALE_SIZE*get_global_id(0);
    global half  * d  = (global half  *) src_d  + get_global_id(0);
    global half  * dm = (global half  *) src_dm  + get_global_id(0);

    b->d  = *d;
    b->dm = *dm;

    for (int i = 0; i < QK_K/2; ++i) {
        b->q[i] = q[i];
    }
    for (int i = 0; i < K_SCALE_SIZE; ++i) {
        b->s[i] = s[i];
    }
}

kernel void kernel_convert_block_q4_K_noshuffle(
    global struct block_q4_K * src0,
    global uchar * dst_q,
    global uchar * dst_s,
    global half  * dst_d,
    global half  * dst_dm,
    uchar mask_0F,
    uchar mask_F0
) {
    global struct block_q4_K * b = (global struct block_q4_K *) src0 + get_global_id(0);
    global uchar * q  = (global uchar *) dst_q  + QK_K/2 * get_global_id(0);
    global uchar * s  = (global uchar *) dst_s  + K_SCALE_SIZE * get_global_id(0);
    global half  * d  = (global half  *) dst_d  + get_global_id(0);
    global half  * dm = (global half  *) dst_dm + get_global_id(0);

    *d  = b->d;
    *dm = b->dm;

    for (int i = 0; i < QK_K / 64; ++i) {
        for (int j = 0; j < 16; ++j) {
            uchar x0 = b->q[i*32 + 2*j];
            uchar x1 = b->q[i*32 + 2*j + 1];
            q[i*32 + j]      = convert_uchar(x0 & mask_0F) | convert_uchar((x1 & mask_0F) << 4);
            q[i*32 + j + 16] = convert_uchar((x0 & mask_F0) >> 4)   | convert_uchar(x1 & mask_F0);
        }
    }

    for (int i = 0; i < K_SCALE_SIZE; ++i) {
        s[i] = b->s[i];
    }
}

kernel void kernel_restore_block_q4_K_noshuffle(
    global uchar * src_q,
    global uchar * src_s,
    global half  * src_d,
    global half  * src_dm,
    global struct block_q4_K * dst,
    uchar mask_0F,
    uchar mask_F0
) {
    global struct block_q4_K * b = (global struct block_q4_K *) dst + get_global_id(0);
    global uchar * q  = (global uchar *) src_q  + QK_K/2 * get_global_id(0);
    global uchar * s  = (global uchar *) src_s  + K_SCALE_SIZE * get_global_id(0);
    global half  * d  = (global half  *) src_d  + get_global_id(0);
    global half  * dm = (global half  *) src_dm + get_global_id(0);

    b->d  = *d;
    b->dm = *dm;

    for (int i = 0; i < QK_K / 64; ++i) {
        for (int j = 0; j < 16; ++j) {
            uchar lo = q[i*32 + j];
            uchar hi = q[i*32 + j + 16];
            b->q[i*32 + 2*j]     = convert_uchar((lo & mask_0F) | ((hi & mask_0F) << 4));
            b->q[i*32 + 2*j + 1] = convert_uchar(((lo & mask_F0) >> 4) | (hi & mask_F0));
        }
    }

    for (int i = 0; i < K_SCALE_SIZE; ++i) {
        b->s[i] = s[i];
    }
}

//------------------------------------------------------------------------------
// kernel_convert_block_q5_K
// Convert the block_q5_K format to 5 separate arrays (AOS -> SOA).
// Each thread processes a super block.
//------------------------------------------------------------------------------
kernel void kernel_convert_block_q5_K(
    global struct block_q5_K * src0,
    global uchar * dst_q,
    global uchar * dst_qh,
    global uchar * dst_s,
    global half  * dst_d,
    global half  * dst_dm,
    uchar mask_0F,
    uchar mask_F0
) {
    global struct block_q5_K * b  = (global struct block_q5_K *) src0 + get_global_id(0);
    global uchar * q  = (global uchar *) dst_q  + QK_K/2*get_global_id(0);
    global uchar * qh = (global uchar *) dst_qh + QK_K/8*get_global_id(0);
    global uchar * s  = (global uchar *) dst_s  + K_SCALE_SIZE*get_global_id(0);
    global half  * d  = (global half  *) dst_d  + get_global_id(0);
    global half  * dm = (global half  *) dst_dm + get_global_id(0);

    *d  = b->d;
    *dm = b->dm;

    for (int i = 0; i < QK_K/2; ++i) {
        q[i] = b->qs[i];
    }
    for (int i = 0; i < QK_K/8; ++i) {
        qh[i] = b->qh[i];
    }
    for (int i = 0; i < K_SCALE_SIZE; ++i) {
        s[i] = b->s[i];
    }
}

// Restore block_q5_K from flattened arrays.
// Each thread processes a super block.
kernel void kernel_restore_block_q5_K(
    global uchar * src_q,
    global uchar * src_qh,
    global uchar * src_s,
    global half  * src_d,
    global half  * src_dm,
    global struct block_q5_K * dst,
    uchar mask_0F,
    uchar mask_F0
) {
    global struct block_q5_K * b  = (global struct block_q5_K *) dst + get_global_id(0);
    global uchar * q  = (global uchar *) src_q  + QK_K/2*get_global_id(0);
    global uchar * qh = (global uchar *) src_qh + QK_K/8*get_global_id(0);
    global uchar * s  = (global uchar *) src_s  + K_SCALE_SIZE*get_global_id(0);
    global half  * d  = (global half  *) src_d  + get_global_id(0);
    global half  * dm = (global half  *) src_dm + get_global_id(0);

    b->d    = *d;
    b->dm = *dm;

    for (int i = 0; i < QK_K/2; ++i) {
        b->qs[i] = q[i];
    }
    for (int i = 0; i < QK_K/8; ++i) {
        b->qh[i] = qh[i];
    }
    for (int i = 0; i < K_SCALE_SIZE; ++i) {
        b->s[i] = s[i];
    }
}

kernel void kernel_convert_block_q5_K_noshuffle(
    global struct block_q5_K * src0,
    global uchar * dst_q,
    global uchar * dst_qh,
    global uchar * dst_s,
    global half  * dst_d,
    global half  * dst_dm,
    uchar mask_0F,
    uchar mask_F0
) {
    global struct block_q5_K * b  = (global struct block_q5_K *) src0 + get_global_id(0);
    global uchar * q  = (global uchar *) dst_q  + QK_K/2       * get_global_id(0);
    global uchar * qh = (global uchar *) dst_qh + QK_K/8       * get_global_id(0);
    global uchar * s  = (global uchar *) dst_s  + K_SCALE_SIZE * get_global_id(0);
    global half  * d  = (global half  *) dst_d  + get_global_id(0);
    global half  * dm = (global half  *) dst_dm + get_global_id(0);

    *d  = b->d;
    *dm = b->dm;

    for (int i = 0; i < QK_K / 64; ++i) {
        for (int j = 0; j < 16; ++j) {
            uchar x0 = b->qs[i*32 + 2*j];
            uchar x1 = b->qs[i*32 + 2*j + 1];
            q[i*32 + j]      = convert_uchar(x0 & mask_0F) | convert_uchar((x1 & mask_0F) << 4);
            q[i*32 + j + 16] = convert_uchar((x0 & mask_F0) >> 4) | convert_uchar(x1 & mask_F0);
        }
    }

    for (int l = 0; l < QK_K/8; ++l) {
        uchar x0 = 0;
        for (int i = 0; i < 8; ++i) {
            x0 |= ((b->qh[(l%4)*8+i] >> (l/4)) & 0x01) << i;
        }
        qh[l] = x0;
    }

    for (int i = 0; i < K_SCALE_SIZE; ++i) {
        s[i] = b->s[i];
    }
}

kernel void kernel_restore_block_q5_K_noshuffle(
    global uchar * src_q,
    global uchar * src_qh,
    global uchar * src_s,
    global half  * src_d,
    global half  * src_dm,
    global struct block_q5_K * dst,
    uchar mask_0F,
    uchar mask_F0
) {
    global struct block_q5_K * b  = (global struct block_q5_K *) dst + get_global_id(0);
    global uchar * q  = (global uchar *) src_q  + QK_K/2       * get_global_id(0);
    global uchar * qh = (global uchar *) src_qh + QK_K/8       * get_global_id(0);
    global uchar * s  = (global uchar *) src_s  + K_SCALE_SIZE * get_global_id(0);
    global half  * d  = (global half  *) src_d  + get_global_id(0);
    global half  * dm = (global half  *) src_dm + get_global_id(0);

    b->d  = *d;
    b->dm = *dm;

    for (int i = 0; i < QK_K / 64; ++i) {
        for (int j = 0; j < 16; ++j) {
            uchar lo = q[i*32 + j];
            uchar hi = q[i*32 + j + 16];
            b->qs[i*32 + 2*j]     = convert_uchar((lo & mask_0F) | ((hi & mask_0F) << 4));
            b->qs[i*32 + 2*j + 1] = convert_uchar(((lo & mask_F0) >> 4) | (hi & mask_F0));
        }
    }

    for (int g = 0; g < 4; ++g) {
        for (int i = 0; i < 8; ++i) {
            uchar x0 = 0;
            for (int k = 0; k < 8; ++k) {
                x0 |= ((qh[4*k+g] >> i) & 0x01) << k;
            }
            b->qh[g*8+i] = x0;
        }
    }

    for (int i = 0; i < K_SCALE_SIZE; ++i) {
        b->s[i] = s[i];
    }
}

//------------------------------------------------------------------------------
// kernel_convert_block_q6_K
// Convert the block_q6_K format to 3 separate arrays (AOS -> SOA).
// This kernel does not deshuffle the bits.
// Each thread processes a super block.
//------------------------------------------------------------------------------
kernel void kernel_convert_block_q6_K(
    global struct block_q6_K * src0,
    global uchar * dst_ql,
    global uchar * dst_qh,
    global char  * dst_s,
    global half  * dst_d,
    uchar          mask_lsb_8,
    ulong          n_blk
) {
    if (get_global_id(0) >= n_blk) {
        return;
    }
    global struct block_q6_K * b = (global struct block_q6_K *) src0 + get_global_id(0);
    global uchar * ql = (global uchar *) dst_ql + QK_K/2*get_global_id(0);
    global uchar * qh = (global uchar *) dst_qh + QK_K/4*get_global_id(0);
    global char  * s  = (global char  *) dst_s  + QK_K/16*get_global_id(0);
    global half  * d  = (global half  *) dst_d  + get_global_id(0);

    *d = b->d;

    for (int i = 0; i < QK_K/2; ++i) {
        ql[i] = b->ql[i];
    }
    for (int i = 0; i < QK_K/4; ++i) {
        qh[i] = b->qh[i];
    }
    for (int i = 0; i < QK_K/16; ++i) {
        s[i] = b->scales[i];
    }
}

// Restore block_q6_K from flattened arrays.
// Each thread processes a super block.
kernel void kernel_restore_block_q6_K(
    global uchar * dst_ql,
    global uchar * dst_qh,
    global char  * dst_s,
    global half  * dst_d,
    global struct block_q6_K * dst,
    uchar mask_lsb_8,
    ulong n_blk
) {
    if (get_global_id(0) >= n_blk) {
        return;
    }
    global struct block_q6_K * b = (global struct block_q6_K *) dst + get_global_id(0);
    global uchar * ql = (global uchar *) dst_ql + QK_K/2*get_global_id(0);
    global uchar * qh = (global uchar *) dst_qh + QK_K/4*get_global_id(0);
    global char  * s  = (global char  *) dst_s  + QK_K/16*get_global_id(0);
    global half  * d  = (global half  *) dst_d  + get_global_id(0);

    b->d = *d;

    for (int i = 0; i < QK_K/2; ++i) {
        b->ql[i] = ql[i];
    }
    for (int i = 0; i < QK_K/4; ++i) {
        b->qh[i] = qh[i];
    }
    for (int i = 0; i < QK_K/16; ++i) {
        b->scales[i] = s[i];
    }
}

kernel void kernel_convert_block_q6_K_noshuffle(
    global struct block_q6_K * src0,
    global uchar * dst_ql,
    global uchar * dst_qh,
    global char  * dst_s,
    global half  * dst_d,
    uchar          mask_lsb_8,
    ulong          n_blk
) {
    if (get_global_id(0) >= n_blk) {
        return;
    }
    global struct block_q6_K * b = (global struct block_q6_K *) src0 + get_global_id(0);
    global uchar * ql = (global uchar *) dst_ql + QK_K/2*get_global_id(0);
    global uchar * qh = (global uchar *) dst_qh + QK_K/4*get_global_id(0);
    global char  * s  = (global char  *) dst_s  + QK_K/16*get_global_id(0);
    global half  * d  = (global half  *) dst_d  + get_global_id(0);

    *d = b->d;

    for (int i = 0; i < QK_K/2/4; ++i) {
        uchar x0 = b->ql[i*2 + 0] & mask_lsb_8;
        uchar x1 = b->ql[i*2 + 1] & mask_lsb_8;
        ql[i +  0] = (x0 & 0x0F)        | ((x1 & 0x0F) << 4);
        ql[i + 32] = ((x0 & 0xF0) >> 4) | (x1 & 0xF0);

        uchar x2 = b->ql[i*2 + 0 + 64] & mask_lsb_8;
        uchar x3 = b->ql[i*2 + 1 + 64] & mask_lsb_8;
        ql[i + 64] = (x2 & 0x0F)        | ((x3 & 0x0F) << 4);
        ql[i + 96] = ((x2 & 0xF0) >> 4) | (x3 & 0xF0);
    }

    for (int i = 0; i < QK_K/4/8; ++i) {
        uchar x0 = b->qh[i*4 + 0] & mask_lsb_8;
        uchar x1 = b->qh[i*4 + 1] & mask_lsb_8;
        uchar x2 = b->qh[i*4 + 2] & mask_lsb_8;
        uchar x3 = b->qh[i*4 + 3] & mask_lsb_8;
        qh[i +  0] = (x0 & 0x03)        | ((x1 & 0x03) << 2) | ((x2 & 0x03) << 4) | ((x3 & 0x03) << 6);
        qh[i +  8] = ((x0 & 0x0C) >> 2) | (x1 & 0x0C)        | ((x2 & 0x0C) << 2) | ((x3 & 0x0C) << 4);
        qh[i + 16] = ((x0 & 0x30) >> 4) | ((x1 & 0x30) >> 2) | (x2 & 0x30)        | ((x3 & 0x30) << 2);
        qh[i + 24] = ((x0 & 0xC0) >> 6) | ((x1 & 0xC0) >> 4) | ((x2 & 0xC0) >> 2) | (x3 & 0xC0);

        uchar x4 = b->qh[i*4 + 0 + 32] & mask_lsb_8;
        uchar x5 = b->qh[i*4 + 1 + 32] & mask_lsb_8;
        uchar x6 = b->qh[i*4 + 2 + 32] & mask_lsb_8;
        uchar x7 = b->qh[i*4 + 3 + 32] & mask_lsb_8;
        qh[i + 32] = (x4 & 0x03)        | ((x5 & 0x03) << 2) | ((x6 & 0x03) << 4) | ((x7 & 0x03) << 6);
        qh[i + 40] = ((x4 & 0x0C) >> 2) | (x5 & 0x0C)        | ((x6 & 0x0C) << 2) | ((x7 & 0x0C) << 4);
        qh[i + 48] = ((x4 & 0x30) >> 4) | ((x5 & 0x30) >> 2) | (x6 & 0x30)        | ((x7 & 0x30) << 2);
        qh[i + 56] = ((x4 & 0xC0) >> 6) | ((x5 & 0xC0) >> 4) | ((x6 & 0xC0) >> 2) | (x7 & 0xC0);
    }

    for (int i = 0; i < QK_K/16; ++i) {
        s[i] = b->scales[i];
    }
}

kernel void kernel_restore_block_q6_K_noshuffle(
    global uchar * src_ql,
    global uchar * src_qh,
    global char  * src_s,
    global half  * src_d,
    global struct block_q6_K * dst,
    uchar          mask_lsb_8,
    ulong          n_blk
) {
    if (get_global_id(0) >= n_blk) {
        return;
    }
    global struct block_q6_K * b = (global struct block_q6_K *) dst + get_global_id(0);
    global uchar * ql = (global uchar *) src_ql + QK_K/2*get_global_id(0);
    global uchar * qh = (global uchar *) src_qh + QK_K/4*get_global_id(0);
    global char  * s  = (global char  *) src_s  + QK_K/16*get_global_id(0);
    global half  * d  = (global half  *) src_d  + get_global_id(0);

    b->d = *d;

    for (int i = 0; i < QK_K/2/4; ++i) {
        uchar x0   = ql[i +  0] & mask_lsb_8;
        uchar x1   = ql[i + 32] & mask_lsb_8;
        b->ql[i*2 + 0] = (x0 & 0x0F)        | ((x1 & 0x0F) << 4);
        b->ql[i*2 + 1] = ((x0 & 0xF0) >> 4) | (x1 & 0xF0);

        uchar x2   = ql[i + 64] & mask_lsb_8;
        uchar x3   = ql[i + 96] & mask_lsb_8;
        b->ql[i*2 + 0 + 64] = (x2 & 0x0F)        | ((x3 & 0x0F) << 4);
        b->ql[i*2 + 1 + 64] = ((x2 & 0xF0) >> 4) | (x3 & 0xF0);
    }

    for (int i = 0; i < QK_K/4/8; ++i) {
        uchar x0 = qh[i +  0] & mask_lsb_8;
        uchar x1 = qh[i +  8] & mask_lsb_8;
        uchar x2 = qh[i + 16] & mask_lsb_8;
        uchar x3 = qh[i + 24] & mask_lsb_8;
        b->qh[i*4 + 0] = (x0 & 0x03)        | ((x1 & 0x03) << 2) | ((x2 & 0x03) << 4) | ((x3 & 0x03) << 6);
        b->qh[i*4 + 1] = ((x0 & 0x0C) >> 2) | (x1 & 0x0C)        | ((x2 & 0x0C) << 2) | ((x3 & 0x0C) << 4);
        b->qh[i*4 + 2] = ((x0 & 0x30) >> 4) | ((x1 & 0x30) >> 2) | (x2 & 0x30)        | ((x3 & 0x30) << 2);
        b->qh[i*4 + 3] = ((x0 & 0xC0) >> 6) | ((x1 & 0xC0) >> 4) | ((x2 & 0xC0) >> 2) | (x3 & 0xC0);

        uchar x4 = qh[i +  0 + 32] & mask_lsb_8;
        uchar x5 = qh[i +  8 + 32] & mask_lsb_8;
        uchar x6 = qh[i + 16 + 32] & mask_lsb_8;
        uchar x7 = qh[i + 24 + 32] & mask_lsb_8;
        b->qh[i*4 + 0 + 32] = (x4 & 0x03)        | ((x5 & 0x03) << 2) | ((x6 & 0x03) << 4) | ((x7 & 0x03) << 6);
        b->qh[i*4 + 1 + 32] = ((x4 & 0x0C) >> 2) | (x5 & 0x0C)        | ((x6 & 0x0C) << 2) | ((x7 & 0x0C) << 4);
        b->qh[i*4 + 2 + 32] = ((x4 & 0x30) >> 4) | ((x5 & 0x30) >> 2) | (x6 & 0x30)        | ((x7 & 0x30) << 2);
        b->qh[i*4 + 3 + 32] = ((x4 & 0xC0) >> 6) | ((x5 & 0xC0) >> 4) | ((x6 & 0xC0) >> 2) | (x7 & 0xC0);
    }

    for (int i = 0; i < QK_K/16; ++i) {
        b->scales[i] = s[i];
    }
}

//------------------------------------------------------------------------------
// kernel_convert_block_iq4_nl
// Convert the block_iq4_nl format to 2 separate arrays (AOS -> SOA).
//------------------------------------------------------------------------------
kernel void kernel_convert_block_iq4_nl(
    global struct block_iq4_nl * src0,
    global uchar * dst_q,
    global half  * dst_d,
    uchar          mask_0F,
    uchar          mask_F0,
    ulong          n_blk
) {
    if (get_global_id(0) >= n_blk) {
        return;
    }
    global struct block_iq4_nl * b = (global struct block_iq4_nl *) src0 + get_global_id(0);
    global uchar * q = (global uchar *) dst_q + QK4_NL/2*get_global_id(0);
    global half  * d = (global half *) dst_d + get_global_id(0);

    *d = b->d;

    for (int i = 0; i < QK4_NL/2; ++i) {
        q[i] = b->qs[i];
    }
}

kernel void kernel_restore_block_iq4_nl(
    global uchar * src_q,
    global half  * src_d,
    global struct block_iq4_nl * dst,
    ulong          n_blk
) {
    if (get_global_id(0) >= n_blk) {
        return;
    }
    global struct block_iq4_nl * b = (global struct block_iq4_nl *) dst + get_global_id(0);
    global uchar * q = (global uchar *) src_q + QK4_NL/2*get_global_id(0);
    global half  * d = (global half *) src_d + get_global_id(0);

    b->d = *d;

    for (int i = 0; i < QK4_NL/2; ++i) {
        b->qs[i] = q[i];
    }
}

kernel void kernel_convert_block_iq4_nl_noshuffle(
    global struct block_iq4_nl * src0,
    global uchar * dst_q,
    global half  * dst_d,
    uchar          mask_0F,
    uchar          mask_F0,
    ulong          n_blk
) {
    if (get_global_id(0) >= n_blk) {
        return;
    }
    global struct block_iq4_nl * b = (global struct block_iq4_nl *) src0 + get_global_id(0);
    global uchar * q = (global uchar *) dst_q + QK4_NL/2*get_global_id(0);
    global half  * d = (global half *) dst_d + get_global_id(0);

    *d = b->d;
    for (int i = 0; i < QK4_NL/4; ++i) {
        uchar x0 = b->qs[2*i + 0];
        uchar x1 = b->qs[2*i + 1];

        q[i + 0       ] = convert_uchar(x0 & mask_0F) | convert_uchar((x1 & mask_0F) << 4);
        q[i + QK4_NL/4] = convert_uchar((x0 & mask_F0) >> 4) | convert_uchar(x1 & mask_F0);
    }
}

kernel void kernel_restore_block_iq4_nl_noshuffle(
    global uchar * src_q,
    global half  * src_d,
    global struct block_iq4_nl * dst,
    uchar mask_0F,
    uchar mask_F0,
    ulong n_blk
) {
    if (get_global_id(0) >= n_blk) {
        return;
    }
    global struct block_iq4_nl * b = (global struct block_iq4_nl *) dst + get_global_id(0);
    global uchar * q = (global uchar *) src_q + QK4_NL/2*get_global_id(0);
    global half  * d = (global half *) src_d + get_global_id(0);

    b->d = *d;
    for (int i = 0; i < QK4_NL/4; ++i) {
        uchar x0 = q[i + 0       ];
        uchar x1 = q[i + QK4_NL/4];

        b->qs[2*i + 0] = convert_uchar((x0 & mask_0F) | ((x1 & mask_0F) << 4));
        b->qs[2*i + 1] = convert_uchar(((x0 & mask_F0) >> 4) | (x1 & mask_F0));
    }
}

// ---------------------------------------------------------------------------
// kernel_moe_expand_scale_q8_0
//
// Expand the q8_0 per-32-block scale d (one half/block, [expert][row][block]) into
// the UNIFORM scale[16] format the generic dp4a MoE GEMM (kernel_gemm_moe_q8_1_dp4a,
// MOE_QT=80) consumes: 16 f16 per 256-superblock (per-16-element segment), where the
// two segments of each 32-block share the block's d. q8_0 is symmetric -> no min
// buffer (the GEMM runs with has_min=0). The int8 weight codes are reused verbatim
// from the existing flat q8_0 weight buffer (extra0_q8_0->q), so only the scale is
// rebuilt here. One work-item per (row, superblock, expert).
// ---------------------------------------------------------------------------
kernel void kernel_moe_expand_scale_q8_0(
    global const half * src_d,      // [expert][row][block], one scale per 32-block
    global       half * dst_scale,  // [expert][row][block][2] (FLAT per-32-block)
    int ne00,
    int ne01
) {
    int row = get_global_id(0);
    int blk = get_global_id(1);   // 32-block index along K
    int e   = get_global_id(2);
    if (row >= ne01) { return; }

    long nb = ne00 / 32;          // 32-blocks per row (K only needs % 32 == 0)
    half d  = src_d[((long)e*ne01 + row)*nb + blk];
    long b  = (((long)e*ne01 + row)*nb + blk) * 2;
    dst_scale[b + 0] = d;
    dst_scale[b + 1] = d;
}

// ---------------------------------------------------------------------------
// kernel_moe_expand_scale_q5_0
//
// q5_0 = symmetric, value = d*(code-16), code = nibble | (hi<<4) in 0..31. The
// generic dp4a MoE GEMM keeps the unsigned code and centers via the min term:
//   scale*dp4a(code,a) - min*sum(a),  scale = d,  min = d*16.
// Reads the existing q5_0 d ([expert][block][row], one half/32-block, from the
// trans4 convert) and writes the FLAT per-32-block uniform scale[2]/min[1] in
// [expert][row][block] order (a transpose). One work-item per (row, block, expert).
// ---------------------------------------------------------------------------
kernel void kernel_moe_expand_scale_q5_0(
    global const half * src_d,      // [expert][block][row]
    global       half * dst_scale,  // [expert][row][block][2]
    global       half * dst_min,    // [expert][row][block]
    int ne00,
    int ne01
) {
    int row = get_global_id(0);
    int blk = get_global_id(1);
    int e   = get_global_id(2);
    if (row >= ne01) { return; }

    long nb = ne00 / 32;
    half d  = src_d[(long)e*nb*ne01 + (long)blk*ne01 + row];   // [expert][block][row]
    long sb = (((long)e*ne01 + row)*nb + blk) * 2;
    long mb = ((long)e*ne01 + row)*nb + blk;
    dst_scale[sb + 0] = d;
    dst_scale[sb + 1] = d;
    dst_min[mb] = (half)((float)d * 16.0f);
}

// ---------------------------------------------------------------------------
// kernel_moe_expand_scale_q5_K
//
// q5_K value = d*sv*code + (-dm*mn), with the 6-bit packed per-sub-block scale sv
// and min mn (8 sub-blocks of 32 per 256-superblock, decoded by get_scale_min_k4
// from the 12-byte s[]). The generic dp4a MoE GEMM (kernel_gemm_moe_q8_1_dp4a,
// MOE_QT=5) keeps the unsigned 5-bit code and applies scale/min via the uniform
// per-32-block buffers:
//   acc += sc0*a_d*raw1 + sc1*a_d*raw2 - mn_u*a_s,
//   sc0 = sc1 = d*sv (both per-16 segments of a 32-block share the sub-block scale),
//   mn_u = dm*mn (positive; the GEMM subtracts it -> the -dm*mn min term).
// q5_K's q_img (low nibbles) + qh (hi-bit plane) are already in the layout the GEMM
// reads (same trans4_ns convert that feeds gemm_moe_q5_k_f32_ns), so only the scale
// is rebuilt here.
//
// One work-item per (row, superblock, expert); each emits 8 sub-blocks.
// ---------------------------------------------------------------------------
kernel void kernel_moe_expand_scale_q5_K(
    global const uchar * src_s,     // [expert][row][superblock][12]
    global const half  * src_d,     // [expert][superblock][row]
    global const half  * src_dm,    // [expert][superblock][row]
    global       half  * dst_scale, // [expert][row][32block][2]
    global       half  * dst_min,   // [expert][row][32block]
    int ne00,
    int ne01
) {
    int row = get_global_id(0);
    int sb  = get_global_id(1);   // superblock index along K
    int e   = get_global_id(2);
    if (row >= ne01) { return; }

    long nsb    = ne00 / 256;     // superblocks per row
    long nblk32 = ne00 / 32;      // 32-blocks per row

    float d  = (float)src_d [((long)e*nsb + sb)*ne01 + row];
    float dm = (float)src_dm[((long)e*nsb + sb)*ne01 + row];

    __global const uchar * sc = src_s + ((long)e*ne01 + row)*nsb*12 + (long)sb*12;

    for (int j = 0; j < 8; ++j) {
        uchar sv, mn;
        // get_scale_min_k4 (6-bit packed scale/min for sub-block j of 8)
        if (j < 4) {
            sv = sc[j]   & 63;
            mn = sc[j+4] & 63;
        } else {
            sv = (sc[j+4] & 0x0F) | ((sc[j-4] & 0xC0) >> 2);
            mn = ((sc[j+4] >> 4) & 0x0F) | ((sc[j]   & 0xC0) >> 2);
        }
        long sub   = (long)sb*8 + j;
        long sbase = (((long)e*ne01 + row)*nblk32 + sub) * 2;
        half s_val = (half)(d  * (float)sv);
        dst_scale[sbase + 0] = s_val;
        dst_scale[sbase + 1] = s_val;
        dst_min[((long)e*ne01 + row)*nblk32 + sub] = (half)(dm * (float)mn);
    }
}