#pragma OPENCL EXTENSION cl_khr_fp16 : enable
#pragma OPENCL EXTENSION cl_khr_subgroups : enable
#ifdef 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")))
#endif
#define QK_K 256
#define NSUBGROUPS 4
#define SUBGROUP_SIZE 64
// scales are transposed: consecutive codes of a row are `stride` apart
inline void get_scale_min_k4(
int j,
global const uchar * q,
uint stride,
uchar * d,
uchar * m,
uchar mask_d6,
uchar mask_d4,
uchar mask_hi2
) {
if (j < 4) {
*d = q[j*stride] & mask_d6 *m = q[(j+4)*stride] & mask_d6 } else {
*d = (q[(j+4)*stride] & mask_d4) | ((q[(j-4)*stride] & mask_hi2) >> 2) *m = ((q[(j+4)*stride] >> 4) & mask_d4) | ((q[j*stride] & mask_hi2) >> 2) }
}
#define dequantizeBlockAccum_ns_sgbroadcast_1_hi(total_sums, bits4, scale, minv, y) \
float shared_y shared_y = sub_group_broadcast(y.s0, 0) total_sums.s0 += ((bits4.s0 & 0x000F) * scale.s0 - minv.s0) * shared_y total_sums.s1 += ((bits4.s1 & 0x000F) * scale.s1 - minv.s1) * shared_y shared_y = sub_group_broadcast(y.s1, 0) total_sums.s0 += (((bits4.s0 & 0x00F0) >> 4) * scale.s0 - minv.s0) * shared_y total_sums.s1 += (((bits4.s1 & 0x00F0) >> 4) * scale.s1 - minv.s1) * shared_y shared_y = sub_group_broadcast(y.s2, 0) total_sums.s0 += (((bits4.s0 & 0x0F00) >> 8) * scale.s0 - minv.s0) * shared_y total_sums.s1 += (((bits4.s1 & 0x0F00) >> 8) * scale.s1 - minv.s1) * shared_y shared_y = sub_group_broadcast(y.s3, 0) total_sums.s0 += (((bits4.s0 & 0xF000) >> 12) * scale.s0 - minv.s0) * shared_y total_sums.s1 += (((bits4.s1 & 0xF000) >> 12) * scale.s1 - minv.s1) * shared_y shared_y = sub_group_broadcast(y.s4, 0) total_sums.s0 += ((bits4.s2 & 0x000F) * scale.s0 - minv.s0) * shared_y total_sums.s1 += ((bits4.s3 & 0x000F) * scale.s1 - minv.s1) * shared_y shared_y = sub_group_broadcast(y.s5, 0) total_sums.s0 += (((bits4.s2 & 0x00F0) >> 4) * scale.s0 - minv.s0) * shared_y total_sums.s1 += (((bits4.s3 & 0x00F0) >> 4) * scale.s1 - minv.s1) * shared_y shared_y = sub_group_broadcast(y.s6, 0) total_sums.s0 += (((bits4.s2 & 0x0F00) >> 8) * scale.s0 - minv.s0) * shared_y total_sums.s1 += (((bits4.s3 & 0x0F00) >> 8) * scale.s1 - minv.s1) * shared_y shared_y = sub_group_broadcast(y.s7, 0) total_sums.s0 += (((bits4.s2 & 0xF000) >> 12) * scale.s0 - minv.s0) * shared_y total_sums.s1 += (((bits4.s3 & 0xF000) >> 12) * scale.s1 - minv.s1) * shared_y shared_y = sub_group_broadcast(y.s0, 1) total_sums.s0 += ((bits4.s4 & 0x000F) * scale.s0 - minv.s0) * shared_y total_sums.s1 += ((bits4.s5 & 0x000F) * scale.s1 - minv.s1) * shared_y shared_y = sub_group_broadcast(y.s1, 1) total_sums.s0 += (((bits4.s4 & 0x00F0) >> 4) * scale.s0 - minv.s0) * shared_y total_sums.s1 += (((bits4.s5 & 0x00F0) >> 4) * scale.s1 - minv.s1) * shared_y shared_y = sub_group_broadcast(y.s2, 1) total_sums.s0 += (((bits4.s4 & 0x0F00) >> 8) * scale.s0 - minv.s0) * shared_y total_sums.s1 += (((bits4.s5 & 0x0F00) >> 8) * scale.s1 - minv.s1) * shared_y shared_y = sub_group_broadcast(y.s3, 1) total_sums.s0 += (((bits4.s4 & 0xF000) >> 12) * scale.s0 - minv.s0) * shared_y total_sums.s1 += (((bits4.s5 & 0xF000) >> 12) * scale.s1 - minv.s1) * shared_y shared_y = sub_group_broadcast(y.s4, 1) total_sums.s0 += ((bits4.s6 & 0x000F) * scale.s0 - minv.s0) * shared_y total_sums.s1 += ((bits4.s7 & 0x000F) * scale.s1 - minv.s1) * shared_y shared_y = sub_group_broadcast(y.s5, 1) total_sums.s0 += (((bits4.s6 & 0x00F0) >> 4) * scale.s0 - minv.s0) * shared_y total_sums.s1 += (((bits4.s7 & 0x00F0) >> 4) * scale.s1 - minv.s1) * shared_y shared_y = sub_group_broadcast(y.s6, 1) total_sums.s0 += (((bits4.s6 & 0x0F00) >> 8) * scale.s0 - minv.s0) * shared_y total_sums.s1 += (((bits4.s7 & 0x0F00) >> 8) * scale.s1 - minv.s1) * shared_y shared_y = sub_group_broadcast(y.s7, 1) total_sums.s0 += (((bits4.s6 & 0xF000) >> 12) * scale.s0 - minv.s0) * shared_y total_sums.s1 += (((bits4.s7 & 0xF000) >> 12) * scale.s1 - minv.s1) * shared_y
#define dequantizeBlockAccum_ns_sgbroadcast_1_lo(total_sums, bits4, scale, minv, y) \
shared_y = sub_group_broadcast(y.s0, 2) total_sums.s0 += ((bits4.s0 & 0x000F) * scale.s0 - minv.s0) * shared_y total_sums.s1 += ((bits4.s1 & 0x000F) * scale.s1 - minv.s1) * shared_y shared_y = sub_group_broadcast(y.s1, 2) total_sums.s0 += (((bits4.s0 & 0x00F0) >> 4) * scale.s0 - minv.s0) * shared_y total_sums.s1 += (((bits4.s1 & 0x00F0) >> 4) * scale.s1 - minv.s1) * shared_y shared_y = sub_group_broadcast(y.s2, 2) total_sums.s0 += (((bits4.s0 & 0x0F00) >> 8) * scale.s0 - minv.s0) * shared_y total_sums.s1 += (((bits4.s1 & 0x0F00) >> 8) * scale.s1 - minv.s1) * shared_y shared_y = sub_group_broadcast(y.s3, 2) total_sums.s0 += (((bits4.s0 & 0xF000) >> 12) * scale.s0 - minv.s0) * shared_y total_sums.s1 += (((bits4.s1 & 0xF000) >> 12) * scale.s1 - minv.s1) * shared_y shared_y = sub_group_broadcast(y.s4, 2) total_sums.s0 += ((bits4.s2 & 0x000F) * scale.s0 - minv.s0) * shared_y total_sums.s1 += ((bits4.s3 & 0x000F) * scale.s1 - minv.s1) * shared_y shared_y = sub_group_broadcast(y.s5, 2) total_sums.s0 += (((bits4.s2 & 0x00F0) >> 4) * scale.s0 - minv.s0) * shared_y total_sums.s1 += (((bits4.s3 & 0x00F0) >> 4) * scale.s1 - minv.s1) * shared_y shared_y = sub_group_broadcast(y.s6, 2) total_sums.s0 += (((bits4.s2 & 0x0F00) >> 8) * scale.s0 - minv.s0) * shared_y total_sums.s1 += (((bits4.s3 & 0x0F00) >> 8) * scale.s1 - minv.s1) * shared_y shared_y = sub_group_broadcast(y.s7, 2) total_sums.s0 += (((bits4.s2 & 0xF000) >> 12) * scale.s0 - minv.s0) * shared_y total_sums.s1 += (((bits4.s3 & 0xF000) >> 12) * scale.s1 - minv.s1) * shared_y shared_y = sub_group_broadcast(y.s0, 3) total_sums.s0 += ((bits4.s4 & 0x000F) * scale.s0 - minv.s0) * shared_y total_sums.s1 += ((bits4.s5 & 0x000F) * scale.s1 - minv.s1) * shared_y shared_y = sub_group_broadcast(y.s1, 3) total_sums.s0 += (((bits4.s4 & 0x00F0) >> 4) * scale.s0 - minv.s0) * shared_y total_sums.s1 += (((bits4.s5 & 0x00F0) >> 4) * scale.s1 - minv.s1) * shared_y shared_y = sub_group_broadcast(y.s2, 3) total_sums.s0 += (((bits4.s4 & 0x0F00) >> 8) * scale.s0 - minv.s0) * shared_y total_sums.s1 += (((bits4.s5 & 0x0F00) >> 8) * scale.s1 - minv.s1) * shared_y shared_y = sub_group_broadcast(y.s3, 3) total_sums.s0 += (((bits4.s4 & 0xF000) >> 12) * scale.s0 - minv.s0) * shared_y total_sums.s1 += (((bits4.s5 & 0xF000) >> 12) * scale.s1 - minv.s1) * shared_y shared_y = sub_group_broadcast(y.s4, 3) total_sums.s0 += ((bits4.s6 & 0x000F) * scale.s0 - minv.s0) * shared_y total_sums.s1 += ((bits4.s7 & 0x000F) * scale.s1 - minv.s1) * shared_y shared_y = sub_group_broadcast(y.s5, 3) total_sums.s0 += (((bits4.s6 & 0x00F0) >> 4) * scale.s0 - minv.s0) * shared_y total_sums.s1 += (((bits4.s7 & 0x00F0) >> 4) * scale.s1 - minv.s1) * shared_y shared_y = sub_group_broadcast(y.s6, 3) total_sums.s0 += (((bits4.s6 & 0x0F00) >> 8) * scale.s0 - minv.s0) * shared_y total_sums.s1 += (((bits4.s7 & 0x0F00) >> 8) * scale.s1 - minv.s1) * shared_y shared_y = sub_group_broadcast(y.s7, 3) total_sums.s0 += (((bits4.s6 & 0xF000) >> 12) * scale.s0 - minv.s0) * shared_y total_sums.s1 += (((bits4.s7 & 0xF000) >> 12) * scale.s1 - minv.s1) * shared_y
#define dequantizeBlockAccum_ns_sgbroadcast_8_hi(total_sums, bits4, scale, minv, y) \
float8 shared_y shared_y = sub_group_broadcast(y, 0) total_sums.s0 += ((bits4.s0 & 0x000F) * scale.s0 - minv.s0) * shared_y.s0 total_sums.s0 += (((bits4.s0 & 0x00F0) >> 4) * scale.s0 - minv.s0) * shared_y.s1 total_sums.s0 += (((bits4.s0 & 0x0F00) >> 8) * scale.s0 - minv.s0) * shared_y.s2 total_sums.s0 += (((bits4.s0 & 0xF000) >> 12) * scale.s0 - minv.s0) * shared_y.s3 total_sums.s0 += ((bits4.s2 & 0x000F) * scale.s0 - minv.s0) * shared_y.s4 total_sums.s0 += (((bits4.s2 & 0x00F0) >> 4) * scale.s0 - minv.s0) * shared_y.s5 total_sums.s0 += (((bits4.s2 & 0x0F00) >> 8) * scale.s0 - minv.s0) * shared_y.s6 total_sums.s0 += (((bits4.s2 & 0xF000) >> 12) * scale.s0 - minv.s0) * shared_y.s7 total_sums.s1 += ((bits4.s1 & 0x000F) * scale.s1 - minv.s1) * shared_y.s0 total_sums.s1 += (((bits4.s1 & 0x00F0) >> 4) * scale.s1 - minv.s1) * shared_y.s1 total_sums.s1 += (((bits4.s1 & 0x0F00) >> 8) * scale.s1 - minv.s1) * shared_y.s2 total_sums.s1 += (((bits4.s1 & 0xF000) >> 12) * scale.s1 - minv.s1) * shared_y.s3 total_sums.s1 += ((bits4.s3 & 0x000F) * scale.s1 - minv.s1) * shared_y.s4 total_sums.s1 += (((bits4.s3 & 0x00F0) >> 4) * scale.s1 - minv.s1) * shared_y.s5 total_sums.s1 += (((bits4.s3 & 0x0F00) >> 8) * scale.s1 - minv.s1) * shared_y.s6 total_sums.s1 += (((bits4.s3 & 0xF000) >> 12) * scale.s1 - minv.s1) * shared_y.s7 shared_y = sub_group_broadcast(y, 1) total_sums.s0 += ((bits4.s4 & 0x000F) * scale.s0 - minv.s0) * shared_y.s0 total_sums.s0 += (((bits4.s4 & 0x00F0) >> 4) * scale.s0 - minv.s0) * shared_y.s1 total_sums.s0 += (((bits4.s4 & 0x0F00) >> 8) * scale.s0 - minv.s0) * shared_y.s2 total_sums.s0 += (((bits4.s4 & 0xF000) >> 12) * scale.s0 - minv.s0) * shared_y.s3 total_sums.s0 += ((bits4.s6 & 0x000F) * scale.s0 - minv.s0) * shared_y.s4 total_sums.s0 += (((bits4.s6 & 0x00F0) >> 4) * scale.s0 - minv.s0) * shared_y.s5 total_sums.s0 += (((bits4.s6 & 0x0F00) >> 8) * scale.s0 - minv.s0) * shared_y.s6 total_sums.s0 += (((bits4.s6 & 0xF000) >> 12) * scale.s0 - minv.s0) * shared_y.s7 total_sums.s1 += ((bits4.s5 & 0x000F) * scale.s1 - minv.s1) * shared_y.s0 total_sums.s1 += (((bits4.s5 & 0x00F0) >> 4) * scale.s1 - minv.s1) * shared_y.s1 total_sums.s1 += (((bits4.s5 & 0x0F00) >> 8) * scale.s1 - minv.s1) * shared_y.s2 total_sums.s1 += (((bits4.s5 & 0xF000) >> 12) * scale.s1 - minv.s1) * shared_y.s3 total_sums.s1 += ((bits4.s7 & 0x000F) * scale.s1 - minv.s1) * shared_y.s4 total_sums.s1 += (((bits4.s7 & 0x00F0) >> 4) * scale.s1 - minv.s1) * shared_y.s5 total_sums.s1 += (((bits4.s7 & 0x0F00) >> 8) * scale.s1 - minv.s1) * shared_y.s6 total_sums.s1 += (((bits4.s7 & 0xF000) >> 12) * scale.s1 - minv.s1) * shared_y.s7
#define dequantizeBlockAccum_ns_sgbroadcast_8_lo(total_sums, bits4, scale, minv, y) \
shared_y = sub_group_broadcast(y, 2) total_sums.s0 += ((bits4.s0 & 0x000F) * scale.s0 - minv.s0) * shared_y.s0 total_sums.s0 += (((bits4.s0 & 0x00F0) >> 4) * scale.s0 - minv.s0) * shared_y.s1 total_sums.s0 += (((bits4.s0 & 0x0F00) >> 8) * scale.s0 - minv.s0) * shared_y.s2 total_sums.s0 += (((bits4.s0 & 0xF000) >> 12) * scale.s0 - minv.s0) * shared_y.s3 total_sums.s0 += ((bits4.s2 & 0x000F) * scale.s0 - minv.s0) * shared_y.s4 total_sums.s0 += (((bits4.s2 & 0x00F0) >> 4) * scale.s0 - minv.s0) * shared_y.s5 total_sums.s0 += (((bits4.s2 & 0x0F00) >> 8) * scale.s0 - minv.s0) * shared_y.s6 total_sums.s0 += (((bits4.s2 & 0xF000) >> 12) * scale.s0 - minv.s0) * shared_y.s7 total_sums.s1 += ((bits4.s1 & 0x000F) * scale.s1 - minv.s1) * shared_y.s0 total_sums.s1 += (((bits4.s1 & 0x00F0) >> 4) * scale.s1 - minv.s1) * shared_y.s1 total_sums.s1 += (((bits4.s1 & 0x0F00) >> 8) * scale.s1 - minv.s1) * shared_y.s2 total_sums.s1 += (((bits4.s1 & 0xF000) >> 12) * scale.s1 - minv.s1) * shared_y.s3 total_sums.s1 += ((bits4.s3 & 0x000F) * scale.s1 - minv.s1) * shared_y.s4 total_sums.s1 += (((bits4.s3 & 0x00F0) >> 4) * scale.s1 - minv.s1) * shared_y.s5 total_sums.s1 += (((bits4.s3 & 0x0F00) >> 8) * scale.s1 - minv.s1) * shared_y.s6 total_sums.s1 += (((bits4.s3 & 0xF000) >> 12) * scale.s1 - minv.s1) * shared_y.s7 shared_y = sub_group_broadcast(y, 3) total_sums.s0 += ((bits4.s4 & 0x000F) * scale.s0 - minv.s0) * shared_y.s0 total_sums.s0 += (((bits4.s4 & 0x00F0) >> 4) * scale.s0 - minv.s0) * shared_y.s1 total_sums.s0 += (((bits4.s4 & 0x0F00) >> 8) * scale.s0 - minv.s0) * shared_y.s2 total_sums.s0 += (((bits4.s4 & 0xF000) >> 12) * scale.s0 - minv.s0) * shared_y.s3 total_sums.s0 += ((bits4.s6 & 0x000F) * scale.s0 - minv.s0) * shared_y.s4 total_sums.s0 += (((bits4.s6 & 0x00F0) >> 4) * scale.s0 - minv.s0) * shared_y.s5 total_sums.s0 += (((bits4.s6 & 0x0F00) >> 8) * scale.s0 - minv.s0) * shared_y.s6 total_sums.s0 += (((bits4.s6 & 0xF000) >> 12) * scale.s0 - minv.s0) * shared_y.s7 total_sums.s1 += ((bits4.s5 & 0x000F) * scale.s1 - minv.s1) * shared_y.s0 total_sums.s1 += (((bits4.s5 & 0x00F0) >> 4) * scale.s1 - minv.s1) * shared_y.s1 total_sums.s1 += (((bits4.s5 & 0x0F00) >> 8) * scale.s1 - minv.s1) * shared_y.s2 total_sums.s1 += (((bits4.s5 & 0xF000) >> 12) * scale.s1 - minv.s1) * shared_y.s3 total_sums.s1 += ((bits4.s7 & 0x000F) * scale.s1 - minv.s1) * shared_y.s4 total_sums.s1 += (((bits4.s7 & 0x00F0) >> 4) * scale.s1 - minv.s1) * shared_y.s5 total_sums.s1 += (((bits4.s7 & 0x0F00) >> 8) * scale.s1 - minv.s1) * shared_y.s6 total_sums.s1 += (((bits4.s7 & 0xF000) >> 12) * scale.s1 - minv.s1) * shared_y.s7
#ifdef ADRENO_GPU
REQD_SUBGROUP_SIZE_64
#endif
kernel void kernel_gemv_noshuffle_q4_k_f32(
read_only image1d_buffer_t src0_q,
global half2 * src0_d,
global half2 * src0_m,
global uchar * src0_s,
read_only image1d_buffer_t src1,
global float * dst,
ulong offsetd,
int ne00,
int ne01,
uchar mask_d6,
uchar mask_d4,
uchar mask_hi2)
{
uint groupId = get_local_id(1) uint gid = get_global_id(0) ushort slid = get_sub_group_local_id() // K-split factor = #subgroups in the WG. Read from the launch (NOT a compile
// constant) so small-M projections (Kcur/Vcur/Qcur) can dispatch a wider
// K-split (more waves/SP -> latency hiding) while large-M keeps 4. The
// physical weight layout stride below is INDEPENDENT of this (see BLOCK_STRIDE_A).
uint nsg = get_local_size(1)
uint K = ne00 uint M = ne01
uint LINE_STRIDE_A = M / 2 // Physical per-K-block stride in the packed image: 8 uints/block-row-pair *
// (M/2) row-pairs = 4*M uints. This is a layout constant, not tied to nsg.
uint BLOCK_STRIDE_A = 4 * M uint scales_per_row = (K / QK_K) * 12
// The x-grid is padded to CEIL_DIV(ne01/2,64)*64, so when ne01 % 128 != 0 the
// tail lanes hold gid >= ne01/2. The output stores below are guarded, but the
// input fetches are not: src0_d and src0_m are raw global half2 pointers,
// src0_s is a raw global uchar pointer, and read_imageui on an
// image1d_buffer_t is UNDEFINED out of range -- an image clamps only for
// SAMPLER reads, which these are not. Those lanes therefore read past the end
// of all three allocations. For a [2816, 2112] weight (2112 % 128 == 64) the
// top tail lane is gid = 1087 while only gid < 1056 is backed, and it runs
// 32 half2 past src0_d/src0_m, 31 uints past the quant image, and 63 bytes
// past src0_s.
//
// Clamp the row used for every fetch. The lanes stay ACTIVE, which the
// sub_group_broadcast in the dequant macros requires, and their results are
// still discarded by the existing output guard. No-op and byte-identical
// whenever ne01 % 128 == 0.
uint gid_s = min(gid, LINE_STRIDE_A - 1)
private uint4 regA private half2 regS private half2 regM private float8 regB
private float2 totalSum = (float2)(0.0f)
for (uint k = groupId uint sb = k / 8 uint j = k % 8
half2 d = src0_d[gid_s + sb * LINE_STRIDE_A] half2 dm = src0_m[gid_s + sb * LINE_STRIDE_A]
global const uchar * sc0 = src0_s + sb * 12 * M + 2 * gid_s global const uchar * sc1 = sc0 + 1
uchar sv0, mn0, sv1, mn1 get_scale_min_k4(j, sc0, M, &sv0, &mn0, mask_d6, mask_d4, mask_hi2) get_scale_min_k4(j, sc1, M, &sv1, &mn1, mask_d6, mask_d4, mask_hi2)
regS = convert_half2(convert_float2(d) * convert_float2((uchar2)(sv0, sv1))) regM = convert_half2(convert_float2(dm) * convert_float2((uchar2)(mn0, mn1)))
if (slid < 4) {
regB.s0123 = read_imagef(src1, (slid * 2 + k * 8)) regB.s4567 = read_imagef(src1, (1 + slid * 2 + k * 8)) }
// load half weights for two blocks in consecutive rows
regA.s0 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 0)).x regA.s1 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 1)).x regA.s2 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 2)).x regA.s3 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 3)).x#ifdef VECTOR_SUB_GROUP_BROADCAST
dequantizeBlockAccum_ns_sgbroadcast_8_hi(totalSum, as_ushort8(regA), regS, regM, regB)#else
dequantizeBlockAccum_ns_sgbroadcast_1_hi(totalSum, as_ushort8(regA), regS, regM, regB)#endif // VECTOR_SUB_GROUP_BROADCAST
regA.s0 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 4)).x regA.s1 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 5)).x regA.s2 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 6)).x regA.s3 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 7)).x#ifdef VECTOR_SUB_GROUP_BROADCAST
dequantizeBlockAccum_ns_sgbroadcast_8_lo(totalSum, as_ushort8(regA), regS, regM, regB)#else
dequantizeBlockAccum_ns_sgbroadcast_1_lo(totalSum, as_ushort8(regA), regS, regM, regB)#endif // VECTOR_SUB_GROUP_BROADCAST
}
// Cross-subgroup reduction in local memory. Generalized to nsg subgroups
// (was a hard-coded 4-wave unroll). Sized for up to 16 subgroups (the widest
// K-split we dispatch for small M). At nsg==4 the accumulation order is
// identical to the original unroll -> byte-identical for the large-M path.
local float2 reduceLM[SUBGROUP_SIZE * 15] if (groupId > 0) {
reduceLM[SUBGROUP_SIZE * (groupId - 1) + slid] = totalSum }
barrier(CLK_LOCAL_MEM_FENCE)
if (groupId == 0) {
for (uint i = 0 totalSum += reduceLM[SUBGROUP_SIZE * i + slid] }
}
// 2 outputs per fiber in wave 0
if (groupId == 0) {
dst = (global float*)((global char*)dst + offsetd) // Guard the two output rows. The x-grid is padded to CEIL_DIV(ne01/2,64)*64,
// so when ne01 is not a multiple of 128 the tail row-pairs run past row ne01
// and would overrun dst into the adjacent tensor. No-op / byte-identical when
// ne01 % 128 == 0 (M/2 already a multiple of 64 -> no padding).
if (gid * 2 + 0 < M) dst[gid * 2 + 0] = totalSum.s0 if (gid * 2 + 1 < M) dst[gid * 2 + 1] = totalSum.s1 }
}
// --- Fused gate+up GEMV + GLU epilogue (FFN) ------------------------------------
// Folds the FFN's two decode GEMVs (ffn_gate, ffn_up) and the following GLU into a
// SINGLE dispatch: {MUL_MAT(Wg,x), MUL_MAT(Wu,x), GLU}. Both matmuls share the same
// activation x (ffn_norm), so the activation image read is issued ONCE per K-block
// and reused for the gate and up dot products (the per-op path re-reads it twice and
// also materializes the two full ffn-wide intermediates to global, which the GLU
// then re-reads). The gate/up partial sums are accumulated in the SAME per-fiber
// order and reduced in the SAME cross-subgroup order as the standalone GEMV, and the
// GLU formula is the exact scalar expression from kernels/glu.cl, so the output is
// BYTE-IDENTICAL to the per-op matmul+matmul+glu path -> safe to default on.
// glu_op: REGLU=0, GEGLU=1, SWIGLU=2, GEGLU_ERF=4, GEGLU_QUICK=5 (ggml_glu_op).
// Weights: src0g_* = gate (= GLU src[0])#define GLU_GEGLU_COEF_A 0.044715f
#define GLU_SQRT_2_OVER_PI 0.79788456080286535587989211986876f
#define GLU_SQRT_2_INV 0.70710678118654752440084436210484f
#define GLU_QUICK_COEF -1.702f
inline float glu_apply(int glu_op, float g, float u) {
float act if (glu_op == 1) { // GEGLU (tanh-approx gelu)
act = 0.5f*g*(1.0f + tanh(GLU_SQRT_2_OVER_PI*g*(1.0f + GLU_GEGLU_COEF_A*g*g))) } else if (glu_op == 2) { // SWIGLU (silu)
act = g / (1.0f + exp(-g)) } else if (glu_op == 0) { // REGLU
return g*u*(g > 0.0f) } else if (glu_op == 4) { // GEGLU_ERF
act = 0.5f*g*(1.0f + erf(g*GLU_SQRT_2_INV)) } else { // GEGLU_QUICK (glu_op == 5)
act = g*(1.0f/(1.0f + exp(GLU_QUICK_COEF*g))) }
return act*u}
#ifdef ADRENO_GPU
REQD_SUBGROUP_SIZE_64
#endif
kernel void kernel_gemv_noshuffle_q4_k_f32_glu(
read_only image1d_buffer_t src0g_q,
global half2 * src0g_d,
global half2 * src0g_m,
global uchar * src0g_s,
read_only image1d_buffer_t src0u_q,
global half2 * src0u_d,
global half2 * src0u_m,
global uchar * src0u_s,
read_only image1d_buffer_t src1,
global float * dst,
ulong offsetd,
int ne00,
int ne01,
int glu_op,
uchar mask_d6,
uchar mask_d4,
uchar mask_hi2)
{
uint groupId = get_local_id(1) uint gid = get_global_id(0) ushort slid = get_sub_group_local_id() uint nsg = get_local_size(1)
uint K = ne00 uint M = ne01
uint LINE_STRIDE_A = M / 2 uint BLOCK_STRIDE_A = 4 * M
private uint4 regA private half2 regS, regM private float8 regB
private float2 gateSum = (float2)(0.0f) private float2 upSum = (float2)(0.0f)
// Two SEQUENTIAL K-loops (gate fully, then up). Keeping only one weight's
// working set live at a time holds the kernel's register footprint at ~the
// base single-weight GEMV's, so its max WG stays 1024 (16 subgroups) and the
// per-subgroup K-split matches the standalone wide GEMV exactly -> the gate
// and up partial sums are BYTE-IDENTICAL to the per-op path. The macro body
// is the base kernel's inner loop verbatim, parameterized by weight source.
#define Q4K_GLU_LOOP(SUM, Q, DD, MM, SS) \
for (uint k = groupId uint sb = k / 8 uint j = k % 8 half2 d = DD[gid + sb * LINE_STRIDE_A] half2 dm = MM[gid + sb * LINE_STRIDE_A] global const uchar * sc0 = SS + sb * 12 * M + 2 * gid global const uchar * sc1 = sc0 + 1 uchar sv0, mn0, sv1, mn1 get_scale_min_k4(j, sc0, M, &sv0, &mn0, mask_d6, mask_d4, mask_hi2) get_scale_min_k4(j, sc1, M, &sv1, &mn1, mask_d6, mask_d4, mask_hi2) regS = convert_half2(convert_float2(d) * convert_float2((uchar2)(sv0, sv1))) regM = convert_half2(convert_float2(dm) * convert_float2((uchar2)(mn0, mn1))) if (slid < 4) { \
regB.s0123 = read_imagef(src1, (slid * 2 + k * 8)) regB.s4567 = read_imagef(src1, (1 + slid * 2 + k * 8)) } \
regA.s0 = read_imageui(Q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 0)).x regA.s1 = read_imageui(Q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 1)).x regA.s2 = read_imageui(Q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 2)).x regA.s3 = read_imageui(Q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 3)).x DEQ_HI(SUM, as_ushort8(regA), regS, regM, regB) regA.s0 = read_imageui(Q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 4)).x regA.s1 = read_imageui(Q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 5)).x regA.s2 = read_imageui(Q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 6)).x regA.s3 = read_imageui(Q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 7)).x DEQ_LO(SUM, as_ushort8(regA), regS, regM, regB) }
#ifdef VECTOR_SUB_GROUP_BROADCAST
#define DEQ_HI dequantizeBlockAccum_ns_sgbroadcast_8_hi
#define DEQ_LO dequantizeBlockAccum_ns_sgbroadcast_8_lo
#else
#define DEQ_HI dequantizeBlockAccum_ns_sgbroadcast_1_hi
#define DEQ_LO dequantizeBlockAccum_ns_sgbroadcast_1_lo
#endif
Q4K_GLU_LOOP(gateSum, src0g_q, src0g_d, src0g_m, src0g_s)
Q4K_GLU_LOOP(upSum, src0u_q, src0u_d, src0u_m, src0u_s)
#undef DEQ_HI
#undef DEQ_LO
#undef Q4K_GLU_LOOP
// Cross-subgroup reduction in local memory. Packs gate (xy) + up (zw) into a
// float4 so both reduce in one pass // per-channel loop -> byte-identical partial sums.
local float4 reduceLM[SUBGROUP_SIZE * 15] if (groupId > 0) {
reduceLM[SUBGROUP_SIZE * (groupId - 1) + slid] = (float4)(gateSum, upSum) }
barrier(CLK_LOCAL_MEM_FENCE) if (groupId == 0) {
for (uint i = 0 float4 p = reduceLM[SUBGROUP_SIZE * i + slid] gateSum += p.xy upSum += p.zw }
dst = (global float*)((global char*)dst + offsetd) dst[gid * 2 + 0] = glu_apply(glu_op, gateSum.s0, upSum.s0) dst[gid * 2 + 1] = glu_apply(glu_op, gateSum.s1, upSum.s1) }
}
// --- Split-K-across-workgroups decode GEMV (small-M projections) ----------------
// A single-token GEMV makes only ceil(M/2/64) workgroups// compute unit, so for small M (Kcur/Vcur, M=512 -> 4 WGs) most of the 16 CUs sit
// idle and the matmul is bandwidth-starved even with a wide intra-WG K-split. This
// variant adds a SECOND grid dimension of `ksplit` workgroups that each reduce a
// disjoint slice of K and write a per-slice partial// then sums the partials into dst. Identical math/layout to the base kernel
// (physical block stride 4*M, get_scale_min_k4) -> coherent. Gated host-side to
// M<=1024 (M>=2048
// already fills the CUs and the extra reduce dispatch only hurts).
#ifdef ADRENO_GPU
REQD_SUBGROUP_SIZE_64
#endif
kernel void kernel_gemv_noshuffle_q4_k_f32_splitk(
read_only image1d_buffer_t src0_q,
global half2 * src0_d,
global half2 * src0_m,
global uchar * src0_s,
read_only image1d_buffer_t src1,
global float * partial, // [ksplit * M], slice-major
int ne00,
int ne01,
uchar mask_d6,
uchar mask_d4,
uchar mask_hi2)
{
uint groupId = get_local_id(1) uint gid = get_global_id(0) ushort slid = get_sub_group_local_id() uint nsg = get_local_size(1) uint ksplit = get_num_groups(1) uint kslice = get_group_id(1)
uint K = ne00 uint M = ne01 uint LINE_STRIDE_A = M / 2 uint BLOCK_STRIDE_A = 4 * M
private uint4 regA private half2 regS, regM private float8 regB private float2 totalSum = (float2)(0.0f)
// each (kslice, subgroup) pair owns a disjoint set of K-blocks
for (uint k = kslice * nsg + groupId uint sb = k / 8 uint j = k % 8 half2 d = src0_d[gid + sb * LINE_STRIDE_A] half2 dm = src0_m[gid + sb * LINE_STRIDE_A] global const uchar * sc0 = src0_s + sb * 12 * M + 2 * gid global const uchar * sc1 = sc0 + 1 uchar sv0, mn0, sv1, mn1 get_scale_min_k4(j, sc0, M, &sv0, &mn0, mask_d6, mask_d4, mask_hi2) get_scale_min_k4(j, sc1, M, &sv1, &mn1, mask_d6, mask_d4, mask_hi2) regS = convert_half2(convert_float2(d) * convert_float2((uchar2)(sv0, sv1))) regM = convert_half2(convert_float2(dm) * convert_float2((uchar2)(mn0, mn1))) if (slid < 4) {
regB.s0123 = read_imagef(src1, (slid * 2 + k * 8)) regB.s4567 = read_imagef(src1, (1 + slid * 2 + k * 8)) }
regA.s0 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 0)).x regA.s1 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 1)).x regA.s2 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 2)).x regA.s3 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 3)).x#ifdef VECTOR_SUB_GROUP_BROADCAST
dequantizeBlockAccum_ns_sgbroadcast_8_hi(totalSum, as_ushort8(regA), regS, regM, regB)#else
dequantizeBlockAccum_ns_sgbroadcast_1_hi(totalSum, as_ushort8(regA), regS, regM, regB)#endif
regA.s0 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 4)).x regA.s1 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 5)).x regA.s2 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 6)).x regA.s3 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 7)).x#ifdef VECTOR_SUB_GROUP_BROADCAST
dequantizeBlockAccum_ns_sgbroadcast_8_lo(totalSum, as_ushort8(regA), regS, regM, regB)#else
dequantizeBlockAccum_ns_sgbroadcast_1_lo(totalSum, as_ushort8(regA), regS, regM, regB)#endif
}
local float2 reduceLM[SUBGROUP_SIZE * 15] if (groupId > 0) {
reduceLM[SUBGROUP_SIZE * (groupId - 1) + slid] = totalSum }
barrier(CLK_LOCAL_MEM_FENCE) if (groupId == 0) {
for (uint i = 0 totalSum += reduceLM[SUBGROUP_SIZE * i + slid] }
vstore2(totalSum, 0, &(partial[kslice * M + gid * 2])) }
}
// Sum the per-slice partials [ksplit * M] into dst[M]kernel void kernel_gemv_splitk_reduce_f32(
global float * partial,
global float * dst,
ulong offsetd,
int ne01, // M
int ksplit)
{
uint r = get_global_id(0) if (r >= (uint)ne01) return float acc = 0.0f for (uint s = 0 acc += partial[s * (uint)ne01 + r] }
dst = (global float*)((global char*)dst + offsetd) dst[r] = acc}
// --- Dequant-once macros for the mc3 verify GEMV (Q4K_MC3_DEQUANT_ONCE) ---
// The inline dequantizeBlockAccum_* macros recompute the dequantized weight
// ((code & mask)>>shift)*scale - minv ONCE PER COLUMN (3x), and the flat
// 32-FMA unroll spills ~430 B of temporaries. These macros split the work:
// DEQUANT_Q4K_BLOCK computes the 16 weights/row of one 32-block ONCE into a
// half2[] (row0 in .s0, row1 in .s1) — stored as half, the exact type the
// inline expression yields (int*half-half), so no extra rounding. MAC_Q4K_BLOCK
// then accumulates them against a column's broadcast activation in the SAME
// per-accumulator order as the inline macro. Each weight value and each
// accumulator's add-chain is bit-for-bit identical => byte-identical output,
// while the dequant ALU drops 3x->1x and the live set shrinks. Requires the
// Qualcomm vector sub_group_broadcast (float8)#define DEQ_Q4K_HALF2(b0, b1, msk, sh, scale, minv) \
(half2)( ((b0 & msk) >> sh) * scale.s0 - minv.s0, \
((b1 & msk) >> sh) * scale.s1 - minv.s1 )
#define DEQUANT_Q4K_BLOCK(wq, bits, scale, minv) \
wq[0] = DEQ_Q4K_HALF2(bits.s0, bits.s1, 0x000F, 0, scale, minv) wq[1] = DEQ_Q4K_HALF2(bits.s0, bits.s1, 0x00F0, 4, scale, minv) wq[2] = DEQ_Q4K_HALF2(bits.s0, bits.s1, 0x0F00, 8, scale, minv) wq[3] = DEQ_Q4K_HALF2(bits.s0, bits.s1, 0xF000, 12, scale, minv) wq[4] = DEQ_Q4K_HALF2(bits.s2, bits.s3, 0x000F, 0, scale, minv) wq[5] = DEQ_Q4K_HALF2(bits.s2, bits.s3, 0x00F0, 4, scale, minv) wq[6] = DEQ_Q4K_HALF2(bits.s2, bits.s3, 0x0F00, 8, scale, minv) wq[7] = DEQ_Q4K_HALF2(bits.s2, bits.s3, 0xF000, 12, scale, minv) wq[8] = DEQ_Q4K_HALF2(bits.s4, bits.s5, 0x000F, 0, scale, minv) wq[9] = DEQ_Q4K_HALF2(bits.s4, bits.s5, 0x00F0, 4, scale, minv) wq[10] = DEQ_Q4K_HALF2(bits.s4, bits.s5, 0x0F00, 8, scale, minv) wq[11] = DEQ_Q4K_HALF2(bits.s4, bits.s5, 0xF000, 12, scale, minv) wq[12] = DEQ_Q4K_HALF2(bits.s6, bits.s7, 0x000F, 0, scale, minv) wq[13] = DEQ_Q4K_HALF2(bits.s6, bits.s7, 0x00F0, 4, scale, minv) wq[14] = DEQ_Q4K_HALF2(bits.s6, bits.s7, 0x0F00, 8, scale, minv) wq[15] = DEQ_Q4K_HALF2(bits.s6, bits.s7, 0xF000, 12, scale, minv)
// ln0/ln1 = the two source lanes whose activation float8 this block consumes
// (0,1 for the hi block, 2,3 for the lo block — matching the inline _hi/_lo).
#define MAC_Q4K_BLOCK(ts, wq, y, ln0, ln1) { \
float8 sy = sub_group_broadcast(y, ln0) ts.s0 += wq[0].s0*sy.s0 ts.s0 += wq[4].s0*sy.s4 ts.s1 += wq[0].s1*sy.s0 ts.s1 += wq[4].s1*sy.s4 sy = sub_group_broadcast(y, ln1) ts.s0 += wq[8].s0*sy.s0 ts.s0 += wq[12].s0*sy.s4 ts.s1 += wq[8].s1*sy.s0 ts.s1 += wq[12].s1*sy.s4}
// Multi-column (N=3) variant of the q4_K decode GEMV, for the speculative /
// MTP verify batch (ne1=3 = 2 drafts + 1 bonus). Stays on the efficient GEMV
// path (subgroup-broadcast activation, NSUBGROUPS K-split) instead of the
// transposed-GEMM dead-zone path. Each K-block's weights (regA_hi/regA_lo) are
// loaded ONCE and reused across all 3 activation columns — same weight traffic
// as one decode, ~3x the (cheap) dequant ALU. Per-column accumulation is
// independent and identical to 3 standalone GEMVs => byte-identical, so it does
// NOT perturb the lm_head logits / spec accept rate.
#ifdef ADRENO_GPU
REQD_SUBGROUP_SIZE_64
#endif
kernel void kernel_gemv_noshuffle_q4_k_f32_mc3(
read_only image1d_buffer_t src0_q,
global half2 * src0_d,
global half2 * src0_m,
global uchar * src0_s,
read_only image1d_buffer_t src1,
global float * dst,
ulong offsetd,
int ne00,
int ne01,
uchar mask_d6,
uchar mask_d4,
uchar mask_hi2)
{
uint groupId = get_local_id(1) uint gid = get_global_id(0) ushort slid = get_sub_group_local_id()
uint K = ne00 uint M = ne01
uint LINE_STRIDE_A = M / 2 uint BLOCK_STRIDE_A = NSUBGROUPS * M uint COL_STRIDE = K / 4
private uint4 regA_hi, regA_lo private half2 regS, regM private float8 regB
private float2 ts0 = (float2)(0.0f) private float2 ts1 = (float2)(0.0f) private float2 ts2 = (float2)(0.0f)
#ifdef Q4K_MC3_DEQUANT_LDS
// One 16-half2 block buffer per WI (reused hi->lo): forces the dequantized
// weights into LDS instead of private arrays (which spill to slow global on
// Adreno). 64*NSUBGROUPS WIs * 16 half2 = 16 KB // range (flat*16) -> no cross-lane sharing, no barrier needed.
local half2 wstage[SUBGROUP_SIZE * NSUBGROUPS * 16] local half2 * ws = wstage + (groupId * SUBGROUP_SIZE + slid) * 16#endif
for (uint k = groupId uint sb = k / 8 uint j = k % 8
half2 d = src0_d[gid + sb * LINE_STRIDE_A] half2 dm = src0_m[gid + sb * LINE_STRIDE_A]
global const uchar * sc0 = src0_s + sb * 12 * M + 2 * gid global const uchar * sc1 = sc0 + 1
uchar sv0, mn0, sv1, mn1 get_scale_min_k4(j, sc0, M, &sv0, &mn0, mask_d6, mask_d4, mask_hi2) get_scale_min_k4(j, sc1, M, &sv1, &mn1, mask_d6, mask_d4, mask_hi2)
regS = convert_half2(convert_float2(d) * convert_float2((uchar2)(sv0, sv1))) regM = convert_half2(convert_float2(dm) * convert_float2((uchar2)(mn0, mn1)))
// weights loaded ONCE, reused across the 3 columns
regA_hi.s0 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 0)).x regA_hi.s1 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 1)).x regA_hi.s2 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 2)).x regA_hi.s3 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 3)).x regA_lo.s0 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 4)).x regA_lo.s1 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 5)).x regA_lo.s2 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 6)).x regA_lo.s3 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 7)).x
#ifdef Q4K_MC3_DEQUANT_ONCE
// Dequant the 32 weights/row (16 hi + 16 lo) ONCE into half2[] (byte-
// identical to the inline intermediate), then MAC against each column's
// activation. Drops the dequant ALU 3x->1x and the macro-temp spill.
half2 wq_hi[16], wq_lo[16] DEQUANT_Q4K_BLOCK(wq_hi, as_ushort8(regA_hi), regS, regM) DEQUANT_Q4K_BLOCK(wq_lo, as_ushort8(regA_lo), regS, regM) { if (slid < 4) { regB.s0123 = read_imagef(src1, 0*COL_STRIDE + slid*2 + k*8) regB.s4567 = read_imagef(src1, 0*COL_STRIDE + 1 + slid*2 + k*8) MAC_Q4K_BLOCK(ts0, wq_hi, regB, 0, 1) { if (slid < 4) { regB.s0123 = read_imagef(src1, 1*COL_STRIDE + slid*2 + k*8) regB.s4567 = read_imagef(src1, 1*COL_STRIDE + 1 + slid*2 + k*8) MAC_Q4K_BLOCK(ts1, wq_hi, regB, 0, 1) { if (slid < 4) { regB.s0123 = read_imagef(src1, 2*COL_STRIDE + slid*2 + k*8) regB.s4567 = read_imagef(src1, 2*COL_STRIDE + 1 + slid*2 + k*8) MAC_Q4K_BLOCK(ts2, wq_hi, regB, 0, 1)#elif defined(Q4K_MC3_DEQUANT_LDS)
// LDS-staged dequant: dequant a 32-block ONCE into the per-WI LDS slot
// (hi pass then lo pass, overwriting), MAC each column from LDS. ts*
// receive hi-then-lo in the same order as DEQUANT_ONCE -> byte-identical.
// Activations reloaded per pass (cheap, imaged) // regs live -> the weight working set lives in LDS, not spilled private.
DEQUANT_Q4K_BLOCK(ws, as_ushort8(regA_hi), regS, regM) { if (slid < 4) { regB.s0123 = read_imagef(src1, 0*COL_STRIDE + slid*2 + k*8) regB.s4567 = read_imagef(src1, 0*COL_STRIDE + 1 + slid*2 + k*8) MAC_Q4K_BLOCK(ts0, ws, regB, 0, 1) { if (slid < 4) { regB.s0123 = read_imagef(src1, 1*COL_STRIDE + slid*2 + k*8) regB.s4567 = read_imagef(src1, 1*COL_STRIDE + 1 + slid*2 + k*8) MAC_Q4K_BLOCK(ts1, ws, regB, 0, 1) { if (slid < 4) { regB.s0123 = read_imagef(src1, 2*COL_STRIDE + slid*2 + k*8) regB.s4567 = read_imagef(src1, 2*COL_STRIDE + 1 + slid*2 + k*8) MAC_Q4K_BLOCK(ts2, ws, regB, 0, 1) DEQUANT_Q4K_BLOCK(ws, as_ushort8(regA_lo), regS, regM) { if (slid < 4) { regB.s0123 = read_imagef(src1, 0*COL_STRIDE + slid*2 + k*8) regB.s4567 = read_imagef(src1, 0*COL_STRIDE + 1 + slid*2 + k*8) MAC_Q4K_BLOCK(ts0, ws, regB, 2, 3) { if (slid < 4) { regB.s0123 = read_imagef(src1, 1*COL_STRIDE + slid*2 + k*8) regB.s4567 = read_imagef(src1, 1*COL_STRIDE + 1 + slid*2 + k*8) MAC_Q4K_BLOCK(ts1, ws, regB, 2, 3) { if (slid < 4) { regB.s0123 = read_imagef(src1, 2*COL_STRIDE + slid*2 + k*8) regB.s4567 = read_imagef(src1, 2*COL_STRIDE + 1 + slid*2 + k*8) MAC_Q4K_BLOCK(ts2, ws, regB, 2, 3)#else
// Per-column: load only this column's activation (single regB live at a
// time -> 1/3 the activation register pressure vs holding all 3) then
// dequant against the shared weights. Cuts the private-mem spill.
#ifdef VECTOR_SUB_GROUP_BROADCAST
{ if (slid < 4) { regB.s0123 = read_imagef(src1, 0*COL_STRIDE + slid*2 + k*8) regB.s4567 = read_imagef(src1, 0*COL_STRIDE + 1 + slid*2 + k*8) dequantizeBlockAccum_ns_sgbroadcast_8_hi(ts0, as_ushort8(regA_hi), regS, regM, regB) dequantizeBlockAccum_ns_sgbroadcast_8_lo(ts0, as_ushort8(regA_lo), regS, regM, regB) { if (slid < 4) { regB.s0123 = read_imagef(src1, 1*COL_STRIDE + slid*2 + k*8) regB.s4567 = read_imagef(src1, 1*COL_STRIDE + 1 + slid*2 + k*8) dequantizeBlockAccum_ns_sgbroadcast_8_hi(ts1, as_ushort8(regA_hi), regS, regM, regB) dequantizeBlockAccum_ns_sgbroadcast_8_lo(ts1, as_ushort8(regA_lo), regS, regM, regB) { if (slid < 4) { regB.s0123 = read_imagef(src1, 2*COL_STRIDE + slid*2 + k*8) regB.s4567 = read_imagef(src1, 2*COL_STRIDE + 1 + slid*2 + k*8) dequantizeBlockAccum_ns_sgbroadcast_8_hi(ts2, as_ushort8(regA_hi), regS, regM, regB) dequantizeBlockAccum_ns_sgbroadcast_8_lo(ts2, as_ushort8(regA_lo), regS, regM, regB)#else
{ if (slid < 4) { regB.s0123 = read_imagef(src1, 0*COL_STRIDE + slid*2 + k*8) regB.s4567 = read_imagef(src1, 0*COL_STRIDE + 1 + slid*2 + k*8) dequantizeBlockAccum_ns_sgbroadcast_1_hi(ts0, as_ushort8(regA_hi), regS, regM, regB) dequantizeBlockAccum_ns_sgbroadcast_1_lo(ts0, as_ushort8(regA_lo), regS, regM, regB) { if (slid < 4) { regB.s0123 = read_imagef(src1, 1*COL_STRIDE + slid*2 + k*8) regB.s4567 = read_imagef(src1, 1*COL_STRIDE + 1 + slid*2 + k*8) dequantizeBlockAccum_ns_sgbroadcast_1_hi(ts1, as_ushort8(regA_hi), regS, regM, regB) dequantizeBlockAccum_ns_sgbroadcast_1_lo(ts1, as_ushort8(regA_lo), regS, regM, regB) { if (slid < 4) { regB.s0123 = read_imagef(src1, 2*COL_STRIDE + slid*2 + k*8) regB.s4567 = read_imagef(src1, 2*COL_STRIDE + 1 + slid*2 + k*8) dequantizeBlockAccum_ns_sgbroadcast_1_hi(ts2, as_ushort8(regA_hi), regS, regM, regB) dequantizeBlockAccum_ns_sgbroadcast_1_lo(ts2, as_ushort8(regA_lo), regS, regM, regB)#endif
#endif // Q4K_MC3_DEQUANT_ONCE
}
// cross-subgroup reduce: pack the 3 columns' float2 into a float8 (6 used).
local float8 reduceLM[SUBGROUP_SIZE * 3] float8 acc = (float8)(ts0.s0, ts0.s1, ts1.s0, ts1.s1, ts2.s0, ts2.s1, 0.0f, 0.0f) if (groupId == 1) { reduceLM[SUBGROUP_SIZE * 0 + slid] = acc if (groupId == 2) { reduceLM[SUBGROUP_SIZE * 1 + slid] = acc if (groupId == 3) { reduceLM[SUBGROUP_SIZE * 2 + slid] = acc
barrier(CLK_LOCAL_MEM_FENCE)
if (groupId == 0) {
acc += reduceLM[SUBGROUP_SIZE * 0 + slid] acc += reduceLM[SUBGROUP_SIZE * 1 + slid] acc += reduceLM[SUBGROUP_SIZE * 2 + slid] dst = (global float*)((global char*)dst + offsetd) // dst is column-major [M rows x 3 cols]: (row, col) at col*M + row
vstore2((float2)(acc.s0, acc.s1), 0, &(dst[0 * M + gid * 2])) vstore2((float2)(acc.s2, acc.s3), 0, &(dst[1 * M + gid * 2])) vstore2((float2)(acc.s4, acc.s5), 0, &(dst[2 * M + gid * 2])) }
}