mirror of
https://github.com/ggml-org/llama.cpp.git
synced 2026-09-04 02:37:27 +02:00
opencl: quant lm_head / decode GEMV and medium-batch GEMM optimizations (speculative decoding/MTP) (#26477)
* opencl: quant lm_head / decode GEMV and medium-batch GEMM optimizations * opencl: guard q4_K/q6_K tiled_ns convert-kernel registration for non-Adreno build * opencl: gate q4_K MUL_MAT+GLU fusion dispatch to Adreno * opencl: require the noshuffle weight layout in the q4_K GLU fusion gate * opencl: do not take the vectorized f16 mrow GEMV path on an unaligned row stride * opencl: pass the new get_scale_min_k4 stride argument at the row-major call sites * opencl: enable the q4_K split-K decode GEMV only where it is measured to win * opencl: record the X1-85 split-K datapoint (neutral, exclusion confirmed) * opencl: restrict the tiled lm_head/embed GEMV default to X2E/A8X * opencl: fix q4_K variant kernels to read the transposed scales layout * opencl: keep the flat-GEMV large-m escape opt-in * opencl: guard the o4 GEMV store against the rounded-up dispatch tail * opencl: restore the tiled q4_K/q6_K layout on tensor read-back * opencl: split-K for the q8_0 decode GEMV at small M * opencl: keep the q6_K noshuffle correctness escape ahead of the opt-in gate
This commit is contained in:
Hongqiang Wang
@@ -85,6 +85,7 @@ set(GGML_OPENCL_KERNELS
|
||||
mul_mv_f16_f32_1row
|
||||
mul_mv_f16_f32_l4
|
||||
mul_mv_f16_f32
|
||||
mul_mv_f16_f32_mrow
|
||||
mul_mv_f32_f32
|
||||
mul_mv_q1_0_f32
|
||||
mul_mv_q1_0_f32_flat
|
||||
@@ -180,9 +181,14 @@ set(GGML_OPENCL_KERNELS
|
||||
gemv_noshuffle_q8_0_f32
|
||||
gemm_noshuffle_q8_0_f32
|
||||
gemv_noshuffle_q4_k_f32
|
||||
gemv_noshuffle_q4_k_f32_o4
|
||||
gemv_noshuffle_q4_k_f32_tiled
|
||||
gemm_noshuffle_q4_k_f32
|
||||
gemv_noshuffle_q6_k_f32
|
||||
gemv_noshuffle_q6_k_f32_o4
|
||||
gemv_noshuffle_q6_k_f32_tiled
|
||||
gemm_noshuffle_q6_k_f32
|
||||
gemm_noshuffle_q6_k_f32_tiled
|
||||
gemv_noshuffle_q5_k_f32
|
||||
gemm_noshuffle_q5_k_f32
|
||||
mul
|
||||
|
||||
+1510
-101
File diff suppressed because it is too large
Load Diff
@@ -1110,6 +1110,78 @@ kernel void kernel_restore_block_q4_k_trans4_ns(
|
||||
}
|
||||
}
|
||||
|
||||
//------------------------------------------------------------------------------
|
||||
// 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,
|
||||
@@ -1494,6 +1566,105 @@ kernel void kernel_restore_block_mxfp4_trans(
|
||||
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,
|
||||
|
||||
@@ -4,6 +4,7 @@
|
||||
#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
|
||||
#define ADRENO_GPU 1
|
||||
#define REQD_SUBGROUP_SIZE_128 __attribute__((qcom_reqd_sub_group_size("full")))
|
||||
#define REQD_SUBGROUP_SIZE_64 __attribute__((qcom_reqd_sub_group_size("half")))
|
||||
#endif
|
||||
#define QK_K 256
|
||||
#define K_SCALE_SIZE 12
|
||||
@@ -171,3 +172,319 @@ kernel void kernel_gemm_noshuffle_q4_k_f32(
|
||||
vstore4((float4)(c0.s7, c1.s7, c2.s7, c3.s7), 0, dst + idx);
|
||||
}
|
||||
}
|
||||
|
||||
// 1x8 per-WI tile (1 output row x 8 output cols). For the small-batch
|
||||
// (medium n_q, e.g. MTP/spec verify) path where the 2x8 kernel is starved:
|
||||
// at ne1<=8 the grid is (1, ceil(M/2)) -> only ~M/256 workgroups, leaving
|
||||
// the SP under-occupied. 1 row per WI doubles the M-axis workgroup count
|
||||
// (ceil(M/1)/128 vs ceil(M/2)/128) AND collapses the accumulators to a
|
||||
// single half8 (16 regs, no spill), so more waves co-reside. Same weight
|
||||
// traffic as 2x8 (rows never share weights); the win is pure occupancy.
|
||||
#ifdef ADRENO_GPU
|
||||
REQD_SUBGROUP_SIZE_128
|
||||
#endif
|
||||
kernel void kernel_gemm_noshuffle_q4_k_f32_r1(
|
||||
global const ushort * src0_q,
|
||||
global const uchar * src0_s,
|
||||
global const half * src0_d,
|
||||
global const half * src0_dm,
|
||||
read_only image1d_buffer_t src1,
|
||||
global float * dst,
|
||||
ulong offsetd,
|
||||
int m,
|
||||
int n,
|
||||
int k,
|
||||
int n_no_padding,
|
||||
uchar mask_d6,
|
||||
uchar mask_d4,
|
||||
uchar mask_hi2
|
||||
) {
|
||||
dst = (global float *)((global char *)dst + offsetd);
|
||||
int n_4 = n >> 2;
|
||||
int gy = get_global_id(0);
|
||||
int gx = get_global_id(1); // 1 row per WI
|
||||
|
||||
half8 c0 = 0;
|
||||
half8 B;
|
||||
half dq;
|
||||
|
||||
int num_blocks_K = k / QK_K;
|
||||
|
||||
global const ushort * weight_ptr = src0_q + gx;
|
||||
global const half * d_ptr = src0_d + gx;
|
||||
global const half * dm_ptr = src0_dm + gx;
|
||||
|
||||
for (int i = 0; i < k; i += 32) {
|
||||
int sb_idx = i / QK_K;
|
||||
int sub_idx = (i / 32) % 8;
|
||||
|
||||
half dd = d_ptr [sb_idx * m];
|
||||
half dmm = dm_ptr[sb_idx * m];
|
||||
|
||||
global const uchar * sc0 = src0_s + sb_idx * K_SCALE_SIZE * m + gx;
|
||||
|
||||
uchar sv0, mn0;
|
||||
get_scale_min_k4(sub_idx, sc0, m, &sv0, &mn0, mask_d6, mask_d4, mask_hi2);
|
||||
|
||||
half scale = convert_half(convert_float(dd) * (float)sv0);
|
||||
half mval = convert_half(convert_float(dmm) * (float)mn0);
|
||||
|
||||
for (int l = 0; l < 32; l += 4) {
|
||||
int ki = i + l;
|
||||
ushort bits = weight_ptr[(ki/4) * m];
|
||||
|
||||
B.s0123 = read_imageh(src1, gy*2 + (ki+0) * n_4);
|
||||
B.s4567 = read_imageh(src1, gy*2+1 + (ki+0) * n_4);
|
||||
dq = (bits & 0x000F) * scale - mval;
|
||||
c0 += B * dq;
|
||||
|
||||
B.s0123 = read_imageh(src1, gy*2 + (ki+1) * n_4);
|
||||
B.s4567 = read_imageh(src1, gy*2+1 + (ki+1) * n_4);
|
||||
dq = ((bits & 0x00F0) >> 4) * scale - mval;
|
||||
c0 += B * dq;
|
||||
|
||||
B.s0123 = read_imageh(src1, gy*2 + (ki+2) * n_4);
|
||||
B.s4567 = read_imageh(src1, gy*2+1 + (ki+2) * n_4);
|
||||
dq = ((bits & 0x0F00) >> 8) * scale - mval;
|
||||
c0 += B * dq;
|
||||
|
||||
B.s0123 = read_imageh(src1, gy*2 + (ki+3) * n_4);
|
||||
B.s4567 = read_imageh(src1, gy*2+1 + (ki+3) * n_4);
|
||||
dq = ((bits & 0xF000) >> 12) * scale - mval;
|
||||
c0 += B * dq;
|
||||
}
|
||||
}
|
||||
|
||||
// Output: 8 cols, 1 row per col-step. Scalar store, coalesced across
|
||||
// neighbouring WIs (consecutive gx -> consecutive dst addresses).
|
||||
int idx = (gy<<3)*m + gx;
|
||||
if (idx < m*n_no_padding) { dst[idx] = c0.s0; idx += m; }
|
||||
if (idx < m*n_no_padding) { dst[idx] = c0.s1; idx += m; }
|
||||
if (idx < m*n_no_padding) { dst[idx] = c0.s2; idx += m; }
|
||||
if (idx < m*n_no_padding) { dst[idx] = c0.s3; idx += m; }
|
||||
if (idx < m*n_no_padding) { dst[idx] = c0.s4; idx += m; }
|
||||
if (idx < m*n_no_padding) { dst[idx] = c0.s5; idx += m; }
|
||||
if (idx < m*n_no_padding) { dst[idx] = c0.s6; idx += m; }
|
||||
if (idx < m*n_no_padding) { dst[idx] = c0.s7; }
|
||||
}
|
||||
|
||||
// 2x8 tile, but weights read through an image1d_buffer (CL_R/UINT32 over the
|
||||
// same packed-q buffer) instead of a plain global buffer. The ne1==1 GEMV
|
||||
// already does this and is much faster per weight byte than this GEMM at
|
||||
// small n_q; the structural difference is the image path hits the dedicated
|
||||
// TPL1 weight cache (L1) while the global path only reaches L2. At small n_q
|
||||
// the forward is weight-read-bound, so L1-cached weights is the lever.
|
||||
// The 2 adjacent rows the 2x8 tile reads as a ushort2 are exactly one uint32,
|
||||
// so the vload2 becomes a single read_imageui at index gx + (ki/4)*(m/2).
|
||||
#ifdef ADRENO_GPU
|
||||
REQD_SUBGROUP_SIZE_128
|
||||
#endif
|
||||
kernel void kernel_gemm_noshuffle_q4_k_f32_kimg(
|
||||
read_only image1d_buffer_t src0_q_img,
|
||||
global const uchar * src0_s,
|
||||
global const half * src0_d,
|
||||
global const half * src0_dm,
|
||||
read_only image1d_buffer_t src1,
|
||||
global float * dst,
|
||||
ulong offsetd,
|
||||
int m,
|
||||
int n,
|
||||
int k,
|
||||
int n_no_padding,
|
||||
uchar mask_d6,
|
||||
uchar mask_d4,
|
||||
uchar mask_hi2
|
||||
) {
|
||||
dst = (global float *)((global char *)dst + offsetd);
|
||||
int n_4 = n >> 2;
|
||||
int m_2 = m >> 1;
|
||||
int gy = get_global_id(0);
|
||||
int gx = get_global_id(1);
|
||||
int gx_2 = gx << 1;
|
||||
|
||||
half8 c0 = 0, c1 = 0;
|
||||
half8 B;
|
||||
half2 dequantized_weights;
|
||||
|
||||
int num_blocks_K = k / QK_K;
|
||||
|
||||
global const half * d_ptr = src0_d + gx_2;
|
||||
global const half * dm_ptr = src0_dm + gx_2;
|
||||
|
||||
for (int i = 0; i < k; i += 32) {
|
||||
int sb_idx = i / QK_K;
|
||||
int sub_idx = (i / 32) % 8;
|
||||
|
||||
half2 d = vload2(0, d_ptr + sb_idx * m);
|
||||
half2 dm = vload2(0, dm_ptr + sb_idx * m);
|
||||
|
||||
global const uchar * sc0 = src0_s + sb_idx * K_SCALE_SIZE * m + (gx_2+0);
|
||||
global const uchar * sc1 = sc0 + 1;
|
||||
|
||||
uchar sv0, mn0, sv1, mn1;
|
||||
get_scale_min_k4(sub_idx, sc0, m, &sv0, &mn0, mask_d6, mask_d4, mask_hi2);
|
||||
get_scale_min_k4(sub_idx, sc1, m, &sv1, &mn1, mask_d6, mask_d4, mask_hi2);
|
||||
|
||||
half2 scale = convert_half2(convert_float2(d) * convert_float2((uchar2)(sv0, sv1)));
|
||||
half2 mval = convert_half2(convert_float2(dm) * convert_float2((uchar2)(mn0, mn1)));
|
||||
|
||||
for (int l = 0; l < 32; l += 4) {
|
||||
int ki = i + l;
|
||||
uint wpacked = read_imageui(src0_q_img, gx + (ki/4) * m_2).x;
|
||||
ushort2 bits2 = (ushort2)((ushort)(wpacked & 0xFFFFu), (ushort)(wpacked >> 16));
|
||||
|
||||
// j=0
|
||||
B.s0123 = read_imageh(src1, gy*2 + (ki+0) * n_4);
|
||||
B.s4567 = read_imageh(src1, gy*2+1 + (ki+0) * n_4);
|
||||
dequantized_weights.s0 = (bits2.s0 & 0x000F) * scale.s0 - mval.s0;
|
||||
dequantized_weights.s1 = (bits2.s1 & 0x000F) * scale.s1 - mval.s1;
|
||||
c0 += B * dequantized_weights.s0;
|
||||
c1 += B * dequantized_weights.s1;
|
||||
|
||||
// j=1
|
||||
B.s0123 = read_imageh(src1, gy*2 + (ki+1) * n_4);
|
||||
B.s4567 = read_imageh(src1, gy*2+1 + (ki+1) * n_4);
|
||||
dequantized_weights.s0 = ((bits2.s0 & 0x00F0) >> 4) * scale.s0 - mval.s0;
|
||||
dequantized_weights.s1 = ((bits2.s1 & 0x00F0) >> 4) * scale.s1 - mval.s1;
|
||||
c0 += B * dequantized_weights.s0;
|
||||
c1 += B * dequantized_weights.s1;
|
||||
|
||||
// j=2
|
||||
B.s0123 = read_imageh(src1, gy*2 + (ki+2) * n_4);
|
||||
B.s4567 = read_imageh(src1, gy*2+1 + (ki+2) * n_4);
|
||||
dequantized_weights.s0 = ((bits2.s0 & 0x0F00) >> 8) * scale.s0 - mval.s0;
|
||||
dequantized_weights.s1 = ((bits2.s1 & 0x0F00) >> 8) * scale.s1 - mval.s1;
|
||||
c0 += B * dequantized_weights.s0;
|
||||
c1 += B * dequantized_weights.s1;
|
||||
|
||||
// j=3
|
||||
B.s0123 = read_imageh(src1, gy*2 + (ki+3) * n_4);
|
||||
B.s4567 = read_imageh(src1, gy*2+1 + (ki+3) * n_4);
|
||||
dequantized_weights.s0 = ((bits2.s0 & 0xF000) >> 12) * scale.s0 - mval.s0;
|
||||
dequantized_weights.s1 = ((bits2.s1 & 0xF000) >> 12) * scale.s1 - mval.s1;
|
||||
c0 += B * dequantized_weights.s0;
|
||||
c1 += B * dequantized_weights.s1;
|
||||
}
|
||||
}
|
||||
|
||||
int idx = (gy<<3)*m + (gx<<1);
|
||||
if (idx+1 < m*n_no_padding) { vstore2((float2)(c0.s0, c1.s0), 0, dst + idx); idx += m; }
|
||||
if (idx+1 < m*n_no_padding) { vstore2((float2)(c0.s1, c1.s1), 0, dst + idx); idx += m; }
|
||||
if (idx+1 < m*n_no_padding) { vstore2((float2)(c0.s2, c1.s2), 0, dst + idx); idx += m; }
|
||||
if (idx+1 < m*n_no_padding) { vstore2((float2)(c0.s3, c1.s3), 0, dst + idx); idx += m; }
|
||||
if (idx+1 < m*n_no_padding) { vstore2((float2)(c0.s4, c1.s4), 0, dst + idx); idx += m; }
|
||||
if (idx+1 < m*n_no_padding) { vstore2((float2)(c0.s5, c1.s5), 0, dst + idx); idx += m; }
|
||||
if (idx+1 < m*n_no_padding) { vstore2((float2)(c0.s6, c1.s6), 0, dst + idx); idx += m; }
|
||||
if (idx+1 < m*n_no_padding) { vstore2((float2)(c0.s7, c1.s7), 0, dst + idx); }
|
||||
}
|
||||
|
||||
// Cooperative-K GEMM for the small-batch (n_q in [2..8]) path. Mirrors the
|
||||
// ne1==1 GEMV's structure: a WG is (COK_SG lanes x COK_NSG subgroups); each
|
||||
// lane owns ONE output row and computes its 8 (padded) columns, and the
|
||||
// COK_NSG subgroups SPLIT the K reduction round-robin, combining via a
|
||||
// __local reduction. This is the thing the per-WI GEMM lacked — at small n_q
|
||||
// the old kernel had ~M/256 workgroups each walking all of K serially; this
|
||||
// has M/64 workgroups AND COK_NSG-way K parallelism. Uses REQD_SUBGROUP_SIZE_64
|
||||
// + barrier (same safe reduction pattern as the GEMV; never sub_group_reduce
|
||||
// at full width on X2 per the GDN miscompile note).
|
||||
#define COK_NSG 8
|
||||
#define COK_SG 64
|
||||
#ifdef ADRENO_GPU
|
||||
REQD_SUBGROUP_SIZE_64
|
||||
#endif
|
||||
kernel void kernel_gemm_noshuffle_q4_k_f32_cok(
|
||||
global const ushort * src0_q,
|
||||
global const uchar * src0_s,
|
||||
global const half * src0_d,
|
||||
global const half * src0_dm,
|
||||
read_only image1d_buffer_t src1,
|
||||
global float * dst,
|
||||
ulong offsetd,
|
||||
int m,
|
||||
int n,
|
||||
int k,
|
||||
int n_no_padding,
|
||||
uchar mask_d6,
|
||||
uchar mask_d4,
|
||||
uchar mask_hi2
|
||||
) {
|
||||
dst = (global float *)((global char *)dst + offsetd);
|
||||
int n_4 = n >> 2;
|
||||
int gx = get_global_id(0); // output row
|
||||
int sg = get_local_id(1); // subgroup index (K-split lane)
|
||||
int lane = get_local_id(0); // lane within subgroup (0..COK_SG-1)
|
||||
|
||||
int num_blocks_K = k / QK_K;
|
||||
int num_32blk = k / 32;
|
||||
|
||||
global const ushort * weight_ptr = src0_q + gx;
|
||||
global const half * d_ptr = src0_d + gx;
|
||||
global const half * dm_ptr = src0_dm + gx;
|
||||
|
||||
half8 acc = 0;
|
||||
half8 B;
|
||||
half dq;
|
||||
|
||||
for (int blk = sg; blk < num_32blk; blk += COK_NSG) {
|
||||
int i = blk << 5; // blk * 32
|
||||
int sb_idx = blk >> 3; // (blk*32) / QK_K (QK_K = 256 = 32*8)
|
||||
int sub_idx = blk & 7; // (i/32) % 8
|
||||
|
||||
half dd = d_ptr [sb_idx * m];
|
||||
half dmm = dm_ptr[sb_idx * m];
|
||||
|
||||
global const uchar * sc0 = src0_s + sb_idx * K_SCALE_SIZE * m + gx;
|
||||
uchar sv0, mn0;
|
||||
get_scale_min_k4(sub_idx, sc0, m, &sv0, &mn0, mask_d6, mask_d4, mask_hi2);
|
||||
half scale = convert_half(convert_float(dd) * (float)sv0);
|
||||
half mval = convert_half(convert_float(dmm) * (float)mn0);
|
||||
|
||||
for (int l = 0; l < 32; l += 4) {
|
||||
int ki = i + l;
|
||||
ushort bits = weight_ptr[(ki>>2) * m];
|
||||
|
||||
B.s0123 = read_imageh(src1, (ki+0) * n_4);
|
||||
B.s4567 = read_imageh(src1, 1 + (ki+0) * n_4);
|
||||
dq = (bits & 0x000F) * scale - mval;
|
||||
acc += B * dq;
|
||||
|
||||
B.s0123 = read_imageh(src1, (ki+1) * n_4);
|
||||
B.s4567 = read_imageh(src1, 1 + (ki+1) * n_4);
|
||||
dq = ((bits & 0x00F0) >> 4) * scale - mval;
|
||||
acc += B * dq;
|
||||
|
||||
B.s0123 = read_imageh(src1, (ki+2) * n_4);
|
||||
B.s4567 = read_imageh(src1, 1 + (ki+2) * n_4);
|
||||
dq = ((bits & 0x0F00) >> 8) * scale - mval;
|
||||
acc += B * dq;
|
||||
|
||||
B.s0123 = read_imageh(src1, (ki+3) * n_4);
|
||||
B.s4567 = read_imageh(src1, 1 + (ki+3) * n_4);
|
||||
dq = ((bits & 0xF000) >> 12) * scale - mval;
|
||||
acc += B * dq;
|
||||
}
|
||||
}
|
||||
|
||||
// cross-subgroup reduction over the K-split (float for accuracy)
|
||||
local float8 reduceLM[COK_SG * (COK_NSG - 1)];
|
||||
if (sg > 0) {
|
||||
reduceLM[(sg - 1) * COK_SG + lane] = convert_float8(acc);
|
||||
}
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
if (sg == 0) {
|
||||
float8 sum = convert_float8(acc);
|
||||
for (int s = 0; s < COK_NSG - 1; s++) {
|
||||
sum += reduceLM[s * COK_SG + lane];
|
||||
}
|
||||
int idx = gx;
|
||||
if (idx < m*n_no_padding) { dst[idx] = sum.s0; idx += m; }
|
||||
if (idx < m*n_no_padding) { dst[idx] = sum.s1; idx += m; }
|
||||
if (idx < m*n_no_padding) { dst[idx] = sum.s2; idx += m; }
|
||||
if (idx < m*n_no_padding) { dst[idx] = sum.s3; idx += m; }
|
||||
if (idx < m*n_no_padding) { dst[idx] = sum.s4; idx += m; }
|
||||
if (idx < m*n_no_padding) { dst[idx] = sum.s5; idx += m; }
|
||||
if (idx < m*n_no_padding) { dst[idx] = sum.s6; idx += m; }
|
||||
if (idx < m*n_no_padding) { dst[idx] = sum.s7; }
|
||||
}
|
||||
}
|
||||
|
||||
@@ -5,6 +5,7 @@
|
||||
#pragma OPENCL EXTENSION cl_qcom_reqd_sub_group_size : enable
|
||||
#define ADRENO_GPU 1
|
||||
#define REQD_SUBGROUP_SIZE_128 __attribute__((qcom_reqd_sub_group_size("full")))
|
||||
#define REQD_SUBGROUP_SIZE_64 __attribute__((qcom_reqd_sub_group_size("half")))
|
||||
#endif
|
||||
|
||||
#ifdef ADRENO_GPU
|
||||
@@ -138,3 +139,107 @@ kernel void kernel_gemm_noshuffle_q6_K_f32(
|
||||
vstore4((float4)(c0.s7, c1.s7, c2.s7, c3.s7), 0, dst + idx);
|
||||
}
|
||||
}
|
||||
|
||||
// Cooperative-K q6_K GEMM for the small-batch (n_q in [2..8]) path. Same idea
|
||||
// as the q4_K _cok kernel: WG = (COK_SG lanes x COK_NSG subgroups), each lane
|
||||
// owns ONE output row (half8 over the 8 padded cols), and the COK_NSG
|
||||
// subgroups split the K iterations round-robin and combine via a __local
|
||||
// reduction. Replaces the default 4-row-per-WI tile that walked all of K alone
|
||||
// (~M/512 WGs + serial reduction) at small n_q. REQD_SUBGROUP_SIZE_64 +
|
||||
// barrier (never sub_group_reduce at full width on X2).
|
||||
#define COK_NSG 8
|
||||
#define COK_SG 64
|
||||
#ifdef ADRENO_GPU
|
||||
REQD_SUBGROUP_SIZE_64
|
||||
#endif
|
||||
kernel void kernel_gemm_noshuffle_q6_K_f32_cok(
|
||||
global const ushort * src0_ql,
|
||||
global const uchar * src0_qh,
|
||||
global const ushort * src0_s,
|
||||
global const half * src0_d,
|
||||
read_only image1d_buffer_t src1,
|
||||
global float * dst,
|
||||
ulong offsetd,
|
||||
int m,
|
||||
int n,
|
||||
int k,
|
||||
int n_no_padding,
|
||||
ushort mask_f000,
|
||||
uchar mask_c0
|
||||
) {
|
||||
dst = (global float *)( (global char *)dst + offsetd );
|
||||
|
||||
int n_4 = n >> 2;
|
||||
int gx = get_global_id(0); // output row
|
||||
int sg = get_local_id(1); // subgroup index (K-split)
|
||||
int lane = get_local_id(0); // lane within subgroup
|
||||
|
||||
global const ushort * ptr_ql = src0_ql + gx;
|
||||
global const uchar * ptr_qh = src0_qh + gx;
|
||||
global const ushort * ptr_s = src0_s + gx;
|
||||
global const half * ptr_d = src0_d + gx;
|
||||
|
||||
half8 acc = 0;
|
||||
half8 B;
|
||||
half dq;
|
||||
|
||||
int num_iter = k >> 2; // k/4 iterations, 4 k-values each
|
||||
|
||||
for (int ib = sg; ib < num_iter; ib += COK_NSG) {
|
||||
int i = ib << 2; // ib * 4
|
||||
|
||||
ushort bits4 = ptr_ql[ib * m]; // ql for row gx at this 4-block
|
||||
uchar bits2 = ptr_qh[ib * m]; // qh
|
||||
|
||||
ushort s_packed = ptr_s[(i >> 5) * m]; // (i/16/2) = i/32
|
||||
char2 sc2 = as_char2(s_packed);
|
||||
char scale_s = (((i >> 4) & 1) == 0) ? sc2.s0 : sc2.s1; // (i/16)%2
|
||||
half scale_d = ptr_d[(i >> 8) * m]; // i/256
|
||||
|
||||
// j=0
|
||||
B.s0123 = read_imageh(src1, (i + 0)*n_4 + 0);
|
||||
B.s4567 = read_imageh(src1, (i + 0)*n_4 + 1);
|
||||
dq = (convert_half((bits4 & 0x000F) | ((bits2 & 0x03) << 4)) - 32.f) * scale_s * scale_d;
|
||||
acc += B * dq;
|
||||
|
||||
// j=1
|
||||
B.s0123 = read_imageh(src1, (i + 1)*n_4 + 0);
|
||||
B.s4567 = read_imageh(src1, (i + 1)*n_4 + 1);
|
||||
dq = (convert_half(((bits4 & 0x00F0) >> 4) | ((bits2 & 0x0C) << 2)) - 32.f) * scale_s * scale_d;
|
||||
acc += B * dq;
|
||||
|
||||
// j=2
|
||||
B.s0123 = read_imageh(src1, (i + 2)*n_4 + 0);
|
||||
B.s4567 = read_imageh(src1, (i + 2)*n_4 + 1);
|
||||
dq = (convert_half(((bits4 & 0x0F00) >> 8) | (bits2 & 0x30)) - 32.f) * scale_s * scale_d;
|
||||
acc += B * dq;
|
||||
|
||||
// j=3
|
||||
B.s0123 = read_imageh(src1, (i + 3)*n_4 + 0);
|
||||
B.s4567 = read_imageh(src1, (i + 3)*n_4 + 1);
|
||||
dq = (convert_half(((bits4 & mask_f000) >> 12) | ((bits2 & mask_c0) >> 2)) - 32.f) * scale_s * scale_d;
|
||||
acc += B * dq;
|
||||
}
|
||||
|
||||
local float8 reduceLM[COK_SG * (COK_NSG - 1)];
|
||||
if (sg > 0) {
|
||||
reduceLM[(sg - 1) * COK_SG + lane] = convert_float8(acc);
|
||||
}
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
if (sg == 0) {
|
||||
float8 sum = convert_float8(acc);
|
||||
for (int s = 0; s < COK_NSG - 1; s++) {
|
||||
sum += reduceLM[s * COK_SG + lane];
|
||||
}
|
||||
int idx = gx;
|
||||
if (idx < m*n_no_padding) { dst[idx] = sum.s0; idx += m; }
|
||||
if (idx < m*n_no_padding) { dst[idx] = sum.s1; idx += m; }
|
||||
if (idx < m*n_no_padding) { dst[idx] = sum.s2; idx += m; }
|
||||
if (idx < m*n_no_padding) { dst[idx] = sum.s3; idx += m; }
|
||||
if (idx < m*n_no_padding) { dst[idx] = sum.s4; idx += m; }
|
||||
if (idx < m*n_no_padding) { dst[idx] = sum.s5; idx += m; }
|
||||
if (idx < m*n_no_padding) { dst[idx] = sum.s6; idx += m; }
|
||||
if (idx < m*n_no_padding) { dst[idx] = sum.s7; }
|
||||
}
|
||||
}
|
||||
|
||||
@@ -0,0 +1,136 @@
|
||||
// Batched (N>1) q6_K GEMM over the 64-row-TILED canonical layout produced by
|
||||
// kernel_convert_block_q6_k_tiled_ns (cvt.cl). Companion to the decode kernel
|
||||
// kernel_gemv_noshuffle_q6_K_f32_tiled: SAME pack, SAME canonical e-order
|
||||
// dequant (correct by construction vs reference ggml q6_K), extended to N output
|
||||
// columns. Makes the batched lm_head/embed (perplexity, spec-decode verify,
|
||||
// batched serving) correct on GPU while keeping the tiled convert the fast decode
|
||||
// GEMV depends on.
|
||||
//
|
||||
// One work-item owns one output ROW for a block of BN columns. A work-group is
|
||||
// {64 lanes, NTILES subgroups} = NTILES*64 rows; the global z dimension tiles the
|
||||
// N columns by BN. Each work-item computes its row's FULL K (no K-split, so no
|
||||
// cross-subgroup reduction), which lets the whole work-group share one staged
|
||||
// activation block:
|
||||
//
|
||||
// __local activation staging — the BN columns of the current superblock (BN*256
|
||||
// floats) are loaded into __local once per superblock, cooperatively by all
|
||||
// NTILES*64 work-items, then every row reads its activation from __local. This
|
||||
// removes the ~Nrows-fold redundant image reads of the first version (each lane
|
||||
// re-read the activation), which made the batched GEMM ~2x slower than the plain
|
||||
// noshuffle GEMM.
|
||||
//
|
||||
// Weights are read from __global (coalesced) — matching the decode kernel; the
|
||||
// lm_head weight is streamed with little reuse where coalesced global beats the
|
||||
// Adreno texture cache.
|
||||
|
||||
#pragma OPENCL EXTENSION cl_khr_fp16 : 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 NTILES 4 // 64-row tiles per work-group (NTILES*64 = 256 rows)
|
||||
#define TILE_ROWS 64
|
||||
#define BN 16 // output columns handled per work-group (global z step)
|
||||
#define WG_THREADS (NTILES * TILE_ROWS)
|
||||
|
||||
#if defined(ADRENO_GPU)
|
||||
REQD_SUBGROUP_SIZE_64
|
||||
#endif
|
||||
kernel void kernel_gemm_noshuffle_q6_K_f32_tiled(
|
||||
__global uint4 * src0_ql, // tiled: 8 uint4 granules / superblock
|
||||
__global uint4 * src0_qh, // tiled: 4 uint4 granules / superblock
|
||||
__global char * src0_s, // tiled: 16 chars / superblock
|
||||
__global half * src0_d, // tiled: 1 half / superblock
|
||||
read_only image1d_buffer_t src1, // activation [ne00, ne11] f32 (RGBA), column-major
|
||||
global float * dst,
|
||||
ulong offsetd,
|
||||
int ne00,
|
||||
int ne01,
|
||||
int ne11
|
||||
) {
|
||||
int rit = get_local_id(0); // 0..63 (lane within a tile; coalesces weight loads)
|
||||
int sg = get_local_id(1); // 0..NTILES-1
|
||||
int lid = sg * TILE_ROWS + rit; // 0..WG_THREADS-1 (flat local id)
|
||||
int row = get_group_id(0) * WG_THREADS + lid;
|
||||
int rt = row / TILE_ROWS; // global 64-row tile index
|
||||
int col0 = get_global_id(2) * BN; // first output column of this block
|
||||
|
||||
int nb = ne00 / 256; // superblocks per row
|
||||
int act_col_stride = ne00 / 4; // activation float4 pixels per column
|
||||
|
||||
const bool row_ok = row < ne01;
|
||||
|
||||
// staged activation: BN columns x 256 elements for the current superblock
|
||||
__local float lact[BN * 256];
|
||||
|
||||
float acc[BN];
|
||||
#pragma unroll
|
||||
for (int j = 0; j < BN; ++j) acc[j] = 0.0f;
|
||||
|
||||
for (int sb = 0; sb < nb; ++sb) {
|
||||
// cooperatively stage BN columns' 256 activation elements (= BN*64 float4)
|
||||
for (int p = lid; p < BN * 64; p += WG_THREADS) {
|
||||
int j = p >> 6; // column within the BN block (p / 64)
|
||||
int e4 = p & 63; // element-quad within the column (p % 64)
|
||||
int c = col0 + j;
|
||||
float4 v = (c < ne11)
|
||||
? read_imagef(src1, c * act_col_stride + sb * 64 + e4)
|
||||
: (float4)(0.0f);
|
||||
lact[p * 4 + 0] = v.x;
|
||||
lact[p * 4 + 1] = v.y;
|
||||
lact[p * 4 + 2] = v.z;
|
||||
lact[p * 4 + 3] = v.w; // lact[j*256 + e], e = e4*4 + t
|
||||
}
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
if (row_ok) {
|
||||
int tile_blk = rt * nb + sb; // ne02 == 1 for lm_head/embed
|
||||
|
||||
float dval = (float)src0_d[tile_blk * TILE_ROWS + rit];
|
||||
__global char * sc = src0_s + (tile_blk * TILE_ROWS + rit) * 16;
|
||||
|
||||
uint ql[32];
|
||||
uint qh[16];
|
||||
#pragma unroll
|
||||
for (int g = 0; g < 8; ++g) {
|
||||
uint4 v = src0_ql[(tile_blk * 8 + g) * TILE_ROWS + rit];
|
||||
ql[g*4+0] = v.x; ql[g*4+1] = v.y; ql[g*4+2] = v.z; ql[g*4+3] = v.w;
|
||||
}
|
||||
#pragma unroll
|
||||
for (int g = 0; g < 4; ++g) {
|
||||
uint4 v = src0_qh[(tile_blk * 4 + g) * TILE_ROWS + rit];
|
||||
qh[g*4+0] = v.x; qh[g*4+1] = v.y; qh[g*4+2] = v.z; qh[g*4+3] = v.w;
|
||||
}
|
||||
|
||||
// NOTE: the e loop (256) is deliberately NOT unrolled. Fully unrolling
|
||||
// 256*BN MACs overflows the in-process Adreno compiler (host stack
|
||||
// overflow at clBuildProgram, same class as the FA DK=512 OOM).
|
||||
for (int e = 0; e < 256; ++e) {
|
||||
uint low4 = (ql[e >> 3] >> ((e & 7) * 4)) & 0xF;
|
||||
uint hi2 = (qh[e >> 4] >> ((e & 15) * 2)) & 0x3;
|
||||
int code = (int)(low4 | (hi2 << 4)) - 32;
|
||||
int sidx = ((e >> 7) << 3) + (((e >> 5) & 3) << 1) + ((e >> 4) & 1);
|
||||
float cs = (float)code * (float)sc[sidx] * dval;
|
||||
#pragma unroll
|
||||
for (int j = 0; j < BN; ++j) {
|
||||
acc[j] += cs * lact[j * 256 + e];
|
||||
}
|
||||
}
|
||||
}
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
if (row_ok) {
|
||||
dst = (global float*)((global char*)dst + offsetd);
|
||||
#pragma unroll
|
||||
for (int j = 0; j < BN; ++j) {
|
||||
int c = col0 + j;
|
||||
if (c < ne11) {
|
||||
dst[(ulong)c * ne01 + row] = acc[j];
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -277,3 +277,107 @@ __kernel void kernel_gemv_noshuffle_q4_0_f32(
|
||||
}
|
||||
|
||||
}
|
||||
|
||||
// Multi-column (N in [2..4]) variant of the q4_0 decode GEMV, for the speculative
|
||||
// / MTP verify batch (n_cols = 2..4 = drafted + bonus positions). Routes the small-
|
||||
// batch verify OFF the transposed-GEMM dead-zone (gemm_noshuffle_q4_0) onto the
|
||||
// efficient GEMV path. Each K-block's weights (regA hi+lo) are loaded ONCE and
|
||||
// reused across the n_cols activation columns. Per-column accumulation is
|
||||
// independent and identical to n_cols standalone GEMVs. n_cols==3 is byte-identical
|
||||
// to the original mc3 (col3 disabled, slots 6/7 stay zero). Kept the _mc3 name.
|
||||
#ifdef VECTOR_SUB_GROUP_BROADCAST
|
||||
#define MC_DQ_HI dequantizeBlockAccum_ns_sgbroadcast_8_hi
|
||||
#define MC_DQ_LO dequantizeBlockAccum_ns_sgbroadcast_8_lo
|
||||
#else
|
||||
#define MC_DQ_HI dequantizeBlockAccum_ns_sgbroadcast_1_hi
|
||||
#define MC_DQ_LO dequantizeBlockAccum_ns_sgbroadcast_1_lo
|
||||
#endif
|
||||
// One column c: load this column's activation (own brace scope so the macros'
|
||||
// `shared_y` decl is re-scoped), then dequant (hi+lo) against the shared weights.
|
||||
#define MC_COL_Q40(ts, c) \
|
||||
{ if (slid < 4) { regB.s0123 = read_imagef(src1, (c)*COL_STRIDE + slid*2 + k*8); \
|
||||
regB.s4567 = read_imagef(src1, (c)*COL_STRIDE + 1 + slid*2 + k*8); } \
|
||||
MC_DQ_HI(ts, as_ushort8(regA_hi), regS, regB); \
|
||||
MC_DQ_LO(ts, as_ushort8(regA_lo), regS, regB); }
|
||||
|
||||
#ifdef ADRENO_GPU
|
||||
REQD_SUBGROUP_SIZE_64
|
||||
#endif
|
||||
__kernel void kernel_gemv_noshuffle_q4_0_f32_mc3(
|
||||
__read_only image1d_buffer_t src0_q, // quantized A
|
||||
global half2 * src0_d, // A scales
|
||||
__read_only image1d_buffer_t src1, // B (n_cols columns, col-major image)
|
||||
global float * dst, // C (column-major [M x n_cols])
|
||||
ulong offsetd,
|
||||
int ne00, // K
|
||||
int ne01, // M
|
||||
int n_cols) // N (2..4)
|
||||
{
|
||||
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;
|
||||
// BLOCK_STRIDE_A is the LAYOUT stride between consecutive K-blocks = 4 uints
|
||||
// per q4_0 block * M (set by the trans4_ns convert). The "4" is uints/block, NOT
|
||||
// the subgroup count — keep it fixed so the K-split count (nsg) can vary.
|
||||
uint BLOCK_STRIDE_A = N_SIMDGROUP * M; // = 4 * M (N_SIMDGROUP is the #define 4)
|
||||
uint COL_STRIDE = K / 4; // float4 pixels per activation column
|
||||
uint nsg = get_local_size(1); // runtime K-split (4 default, 8 small-M)
|
||||
|
||||
__private uint4 regA_hi, regA_lo;
|
||||
__private half2 regS;
|
||||
__private float8 regB;
|
||||
|
||||
__private float2 ts0 = (float2)(0.0f);
|
||||
__private float2 ts1 = (float2)(0.0f);
|
||||
__private float2 ts2 = (float2)(0.0f);
|
||||
__private float2 ts3 = (float2)(0.0f);
|
||||
|
||||
for (uint k = groupId; k < (K / QK4_0); k += nsg) {
|
||||
regS = src0_d[gid + k * LINE_STRIDE_A];
|
||||
|
||||
// weights loaded ONCE, reused across the 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;
|
||||
|
||||
MC_COL_Q40(ts0, 0);
|
||||
MC_COL_Q40(ts1, 1);
|
||||
if (n_cols > 2) MC_COL_Q40(ts2, 2);
|
||||
if (n_cols > 3) MC_COL_Q40(ts3, 3);
|
||||
}
|
||||
|
||||
// cross-subgroup reduce over nsg subgroups: pack the (up to 4) columns' float2
|
||||
// into a float8. Generalized to runtime nsg (4 default, 8 for small-M). Each
|
||||
// subgroup writes its partial; subgroup 0 sums the rest into its own acc. At
|
||||
// nsg==4 this is byte-identical to the original (sums subgroups 1,2,3 in order).
|
||||
__local float8 reduceLM[SIMDGROUP_WIDTH * 8];
|
||||
float8 acc = (float8)(ts0.s0, ts0.s1, ts1.s0, ts1.s1, ts2.s0, ts2.s1, ts3.s0, ts3.s1);
|
||||
reduceLM[groupId * SIMDGROUP_WIDTH + slid] = acc;
|
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
if (groupId == 0) {
|
||||
for (uint g = 1; g < nsg; g++) {
|
||||
acc += reduceLM[g * SIMDGROUP_WIDTH + slid];
|
||||
}
|
||||
dst = (global float*)((global char*)dst + offsetd);
|
||||
// dst is column-major [M rows x n_cols 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]));
|
||||
if (n_cols > 2) vstore2((float2)(acc.s4, acc.s5), 0, &(dst[2 * M + gid * 2]));
|
||||
if (n_cols > 3) vstore2((float2)(acc.s6, acc.s7), 0, &(dst[3 * M + gid * 2]));
|
||||
}
|
||||
}
|
||||
#undef MC_COL_Q40
|
||||
#undef MC_DQ_HI
|
||||
#undef MC_DQ_LO
|
||||
|
||||
@@ -286,3 +286,99 @@ kernel void kernel_gemv_noshuffle_q4_1_f32(
|
||||
}
|
||||
|
||||
}
|
||||
|
||||
// Multi-column (N in [2..4]) variant of the q4_1 decode GEMV (spec/MTP verify) =
|
||||
// q4_0 mc3 + the q4_1 per-block min (regM; dequant = q*scale + minv). n_cols=2..4;
|
||||
// routes the small-batch verify OFF the gemm_noshuffle_q4_1 dead-zone. n_cols==3 is
|
||||
// byte-identical to the original mc3. NB: this file spells the vec-broadcast define
|
||||
// BROADCAT (no S) — match it so the fast _8 path compiles.
|
||||
#ifdef VECTOR_SUB_GROUP_BROADCAT
|
||||
#define MC_DQ1_HI dequantizeBlockAccum_ns_sgbroadcast_8_hi
|
||||
#define MC_DQ1_LO dequantizeBlockAccum_ns_sgbroadcast_8_lo
|
||||
#else
|
||||
#define MC_DQ1_HI dequantizeBlockAccum_ns_sgbroadcast_1_hi
|
||||
#define MC_DQ1_LO dequantizeBlockAccum_ns_sgbroadcast_1_lo
|
||||
#endif
|
||||
#define MC_COL_Q41(ts, c) \
|
||||
{ if (slid < 4) { regB.s0123 = read_imagef(src1, (c)*COL_STRIDE + slid*2 + k*8); \
|
||||
regB.s4567 = read_imagef(src1, (c)*COL_STRIDE + 1 + slid*2 + k*8); } \
|
||||
MC_DQ1_HI(ts, as_ushort8(regA_hi), regS, regM, regB); \
|
||||
MC_DQ1_LO(ts, as_ushort8(regA_lo), regS, regM, regB); }
|
||||
#ifdef ADRENO_GPU
|
||||
REQD_SUBGROUP_SIZE_64
|
||||
#endif
|
||||
kernel void kernel_gemv_noshuffle_q4_1_f32_mc3(
|
||||
read_only image1d_buffer_t src0_q,
|
||||
global half2 * src0_d,
|
||||
global half2 * src0_m,
|
||||
read_only image1d_buffer_t src1,
|
||||
global float * dst,
|
||||
ulong offsetd,
|
||||
int ne00,
|
||||
int ne01,
|
||||
int n_cols)
|
||||
{
|
||||
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; // float4 pixels per activation column
|
||||
|
||||
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);
|
||||
private float2 ts3 = (float2)(0.0f);
|
||||
|
||||
for (uint k = groupId; k < (K / QK4_0); k += NSUBGROUPS) {
|
||||
regS = src0_d[gid + k * LINE_STRIDE_A];
|
||||
regM = src0_m[gid + k * LINE_STRIDE_A];
|
||||
|
||||
// weights loaded ONCE, reused across the 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;
|
||||
|
||||
MC_COL_Q41(ts0, 0);
|
||||
MC_COL_Q41(ts1, 1);
|
||||
if (n_cols > 2) MC_COL_Q41(ts2, 2);
|
||||
if (n_cols > 3) MC_COL_Q41(ts3, 3);
|
||||
}
|
||||
|
||||
// cross-subgroup reduce: pack the (up to 4) columns' float2 into a float8.
|
||||
local float8 reduceLM[SUBGROUP_SIZE * 3];
|
||||
float8 acc = (float8)(ts0.s0, ts0.s1, ts1.s0, ts1.s1, ts2.s0, ts2.s1, ts3.s0, ts3.s1);
|
||||
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 n_cols 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]));
|
||||
if (n_cols > 2) vstore2((float2)(acc.s4, acc.s5), 0, &(dst[2 * M + gid * 2]));
|
||||
if (n_cols > 3) vstore2((float2)(acc.s6, acc.s7), 0, &(dst[3 * M + gid * 2]));
|
||||
}
|
||||
}
|
||||
#undef MC_COL_Q41
|
||||
#undef MC_DQ1_HI
|
||||
#undef MC_DQ1_LO
|
||||
|
||||
@@ -228,12 +228,20 @@ kernel void kernel_gemv_noshuffle_q4_k_f32(
|
||||
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;
|
||||
uint BLOCK_STRIDE_A = NSUBGROUPS * M;
|
||||
// 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
|
||||
@@ -259,7 +267,7 @@ kernel void kernel_gemv_noshuffle_q4_k_f32(
|
||||
|
||||
private float2 totalSum = (float2)(0.0f);
|
||||
|
||||
for (uint k = groupId; k < (K / 32); k += NSUBGROUPS) {
|
||||
for (uint k = groupId; k < (K / 32); k += nsg) {
|
||||
uint sb = k / 8;
|
||||
uint j = k % 8;
|
||||
|
||||
@@ -303,28 +311,21 @@ kernel void kernel_gemv_noshuffle_q4_k_f32(
|
||||
#endif // VECTOR_SUB_GROUP_BROADCAST
|
||||
}
|
||||
|
||||
// reduction in local memory, assumes #wave=4
|
||||
local float2 reduceLM[SUBGROUP_SIZE * 3];
|
||||
if (groupId == 1) {
|
||||
reduceLM[SUBGROUP_SIZE * 0 + slid] = totalSum;
|
||||
}
|
||||
if (groupId == 2) {
|
||||
reduceLM[SUBGROUP_SIZE * 1 + slid] = totalSum;
|
||||
}
|
||||
if (groupId == 3) {
|
||||
reduceLM[SUBGROUP_SIZE * 2 + slid] = totalSum;
|
||||
// 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) {
|
||||
totalSum += reduceLM[SUBGROUP_SIZE * 0 + slid];
|
||||
}
|
||||
if (groupId == 0) {
|
||||
totalSum += reduceLM[SUBGROUP_SIZE * 1 + slid];
|
||||
}
|
||||
if (groupId == 0) {
|
||||
totalSum += reduceLM[SUBGROUP_SIZE * 2 + slid];
|
||||
for (uint i = 0; i < nsg - 1; ++i) {
|
||||
totalSum += reduceLM[SUBGROUP_SIZE * i + slid];
|
||||
}
|
||||
}
|
||||
|
||||
// 2 outputs per fiber in wave 0
|
||||
@@ -339,3 +340,484 @@ kernel void kernel_gemv_noshuffle_q4_k_f32(
|
||||
}
|
||||
|
||||
}
|
||||
|
||||
// --- 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]); src0u_* = up (= GLU src[1]).
|
||||
#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; k < (K / 32); k += nsg) { \
|
||||
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; summation order matches the base GEMV's
|
||||
// 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; i < nsg - 1; ++i) {
|
||||
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; a WG runs on one Adreno
|
||||
// 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; kernel_gemv_splitk_reduce_f32
|
||||
// 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; // physical, independent of the K-split
|
||||
|
||||
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; k < (K / 32); k += ksplit * nsg) {
|
||||
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; i < nsg - 1; ++i) {
|
||||
totalSum += reduceLM[SUBGROUP_SIZE * i + slid];
|
||||
}
|
||||
vstore2(totalSum, 0, &(partial[kslice * M + gid * 2]));
|
||||
}
|
||||
}
|
||||
|
||||
// Sum the per-slice partials [ksplit * M] into dst[M]; applies the dst byte offset.
|
||||
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; s < (uint)ksplit; ++s) {
|
||||
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); enabled opt-in on Adreno.
|
||||
#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[1].s0*sy.s1; ts.s0 += wq[2].s0*sy.s2; ts.s0 += wq[3].s0*sy.s3; \
|
||||
ts.s0 += wq[4].s0*sy.s4; ts.s0 += wq[5].s0*sy.s5; ts.s0 += wq[6].s0*sy.s6; ts.s0 += wq[7].s0*sy.s7; \
|
||||
ts.s1 += wq[0].s1*sy.s0; ts.s1 += wq[1].s1*sy.s1; ts.s1 += wq[2].s1*sy.s2; ts.s1 += wq[3].s1*sy.s3; \
|
||||
ts.s1 += wq[4].s1*sy.s4; ts.s1 += wq[5].s1*sy.s5; ts.s1 += wq[6].s1*sy.s6; ts.s1 += wq[7].s1*sy.s7; \
|
||||
sy = sub_group_broadcast(y, ln1); \
|
||||
ts.s0 += wq[8].s0*sy.s0; ts.s0 += wq[9].s0*sy.s1; ts.s0 += wq[10].s0*sy.s2; ts.s0 += wq[11].s0*sy.s3; \
|
||||
ts.s0 += wq[12].s0*sy.s4; ts.s0 += wq[13].s0*sy.s5; ts.s0 += wq[14].s0*sy.s6; ts.s0 += wq[15].s0*sy.s7; \
|
||||
ts.s1 += wq[8].s1*sy.s0; ts.s1 += wq[9].s1*sy.s1; ts.s1 += wq[10].s1*sy.s2; ts.s1 += wq[11].s1*sy.s3; \
|
||||
ts.s1 += wq[12].s1*sy.s4; ts.s1 += wq[13].s1*sy.s5; ts.s1 += wq[14].s1*sy.s6; ts.s1 += wq[15].s1*sy.s7; \
|
||||
}
|
||||
|
||||
// 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; // float4 pixels per activation column
|
||||
|
||||
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; each WI owns its own slot
|
||||
// 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; k < (K / 32); k += NSUBGROUPS) {
|
||||
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); MAC_Q4K_BLOCK(ts0, wq_lo, 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, wq_hi, regB, 0, 1); MAC_Q4K_BLOCK(ts1, wq_lo, 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, wq_hi, regB, 0, 1); MAC_Q4K_BLOCK(ts2, wq_lo, regB, 2, 3); }
|
||||
#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); only one regB + 0 weight
|
||||
// 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]));
|
||||
}
|
||||
}
|
||||
|
||||
@@ -0,0 +1,349 @@
|
||||
#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_o4(
|
||||
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); // 4-output quad index
|
||||
ushort slid = get_sub_group_local_id();
|
||||
|
||||
// Two consecutive pair-indices (each the same access pattern the 2-output
|
||||
// kernel uses); together they cover 4 consecutive output rows.
|
||||
uint gid_a = gid * 2;
|
||||
uint gid_b = gid * 2 + 1;
|
||||
|
||||
uint K = ne00;
|
||||
uint M = ne01;
|
||||
|
||||
uint LINE_STRIDE_A = M / 2;
|
||||
uint BLOCK_STRIDE_A = NSUBGROUPS * M;
|
||||
|
||||
private uint4 regA;
|
||||
private half2 regS_a, regS_b;
|
||||
private half2 regM_a, regM_b;
|
||||
private float8 regB;
|
||||
|
||||
private float2 totalSum_a = (float2)(0.0f);
|
||||
private float2 totalSum_b = (float2)(0.0f);
|
||||
|
||||
for (uint k = groupId; k < (K / 32); k += NSUBGROUPS) {
|
||||
uint sb = k / 8;
|
||||
uint j = k % 8;
|
||||
|
||||
// pair a scales/mins
|
||||
half2 d_a = src0_d[gid_a + sb * LINE_STRIDE_A];
|
||||
half2 dm_a = src0_m[gid_a + sb * LINE_STRIDE_A];
|
||||
global const uchar * sc0a = src0_s + sb * 12 * M + 2 * gid_a;
|
||||
global const uchar * sc1a = sc0a + 1;
|
||||
uchar sv0a, mn0a, sv1a, mn1a;
|
||||
get_scale_min_k4(j, sc0a, M, &sv0a, &mn0a, mask_d6, mask_d4, mask_hi2);
|
||||
get_scale_min_k4(j, sc1a, M, &sv1a, &mn1a, mask_d6, mask_d4, mask_hi2);
|
||||
regS_a = convert_half2(convert_float2(d_a) * convert_float2((uchar2)(sv0a, sv1a)));
|
||||
regM_a = convert_half2(convert_float2(dm_a) * convert_float2((uchar2)(mn0a, mn1a)));
|
||||
|
||||
// pair b scales/mins
|
||||
half2 d_b = src0_d[gid_b + sb * LINE_STRIDE_A];
|
||||
half2 dm_b = src0_m[gid_b + sb * LINE_STRIDE_A];
|
||||
global const uchar * sc0b = src0_s + sb * 12 * M + 2 * gid_b;
|
||||
global const uchar * sc1b = sc0b + 1;
|
||||
uchar sv0b, mn0b, sv1b, mn1b;
|
||||
get_scale_min_k4(j, sc0b, M, &sv0b, &mn0b, mask_d6, mask_d4, mask_hi2);
|
||||
get_scale_min_k4(j, sc1b, M, &sv1b, &mn1b, mask_d6, mask_d4, mask_hi2);
|
||||
regS_b = convert_half2(convert_float2(d_b) * convert_float2((uchar2)(sv0b, sv1b)));
|
||||
regM_b = convert_half2(convert_float2(dm_b) * convert_float2((uchar2)(mn0b, mn1b)));
|
||||
|
||||
// activation: load once, reuse for both pairs
|
||||
if (slid < 4) {
|
||||
regB.s0123 = read_imagef(src1, (slid * 2 + k * 8));
|
||||
regB.s4567 = read_imagef(src1, (1 + slid * 2 + k * 8));
|
||||
}
|
||||
|
||||
// pair a (own block so _lo sees the shared_y declared by _hi)
|
||||
{
|
||||
regA.s0 = read_imageui(src0_q, (gid_a + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 0)).x;
|
||||
regA.s1 = read_imageui(src0_q, (gid_a + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 1)).x;
|
||||
regA.s2 = read_imageui(src0_q, (gid_a + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 2)).x;
|
||||
regA.s3 = read_imageui(src0_q, (gid_a + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 3)).x;
|
||||
#ifdef VECTOR_SUB_GROUP_BROADCAST
|
||||
dequantizeBlockAccum_ns_sgbroadcast_8_hi(totalSum_a, as_ushort8(regA), regS_a, regM_a, regB);
|
||||
#else
|
||||
dequantizeBlockAccum_ns_sgbroadcast_1_hi(totalSum_a, as_ushort8(regA), regS_a, regM_a, regB);
|
||||
#endif
|
||||
regA.s0 = read_imageui(src0_q, (gid_a + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 4)).x;
|
||||
regA.s1 = read_imageui(src0_q, (gid_a + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 5)).x;
|
||||
regA.s2 = read_imageui(src0_q, (gid_a + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 6)).x;
|
||||
regA.s3 = read_imageui(src0_q, (gid_a + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 7)).x;
|
||||
#ifdef VECTOR_SUB_GROUP_BROADCAST
|
||||
dequantizeBlockAccum_ns_sgbroadcast_8_lo(totalSum_a, as_ushort8(regA), regS_a, regM_a, regB);
|
||||
#else
|
||||
dequantizeBlockAccum_ns_sgbroadcast_1_lo(totalSum_a, as_ushort8(regA), regS_a, regM_a, regB);
|
||||
#endif
|
||||
}
|
||||
|
||||
// pair b
|
||||
{
|
||||
regA.s0 = read_imageui(src0_q, (gid_b + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 0)).x;
|
||||
regA.s1 = read_imageui(src0_q, (gid_b + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 1)).x;
|
||||
regA.s2 = read_imageui(src0_q, (gid_b + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 2)).x;
|
||||
regA.s3 = read_imageui(src0_q, (gid_b + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 3)).x;
|
||||
#ifdef VECTOR_SUB_GROUP_BROADCAST
|
||||
dequantizeBlockAccum_ns_sgbroadcast_8_hi(totalSum_b, as_ushort8(regA), regS_b, regM_b, regB);
|
||||
#else
|
||||
dequantizeBlockAccum_ns_sgbroadcast_1_hi(totalSum_b, as_ushort8(regA), regS_b, regM_b, regB);
|
||||
#endif
|
||||
regA.s0 = read_imageui(src0_q, (gid_b + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 4)).x;
|
||||
regA.s1 = read_imageui(src0_q, (gid_b + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 5)).x;
|
||||
regA.s2 = read_imageui(src0_q, (gid_b + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 6)).x;
|
||||
regA.s3 = read_imageui(src0_q, (gid_b + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 7)).x;
|
||||
#ifdef VECTOR_SUB_GROUP_BROADCAST
|
||||
dequantizeBlockAccum_ns_sgbroadcast_8_lo(totalSum_b, as_ushort8(regA), regS_b, regM_b, regB);
|
||||
#else
|
||||
dequantizeBlockAccum_ns_sgbroadcast_1_lo(totalSum_b, as_ushort8(regA), regS_b, regM_b, regB);
|
||||
#endif
|
||||
}
|
||||
}
|
||||
|
||||
// reduce 4 outputs (a.s0, a.s1, b.s0, b.s1) across the 4 subgroups
|
||||
local float4 reduceLM[SUBGROUP_SIZE * 3];
|
||||
float4 acc = (float4)(totalSum_a.s0, totalSum_a.s1, totalSum_b.s0, totalSum_b.s1);
|
||||
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);
|
||||
// The dispatch rounds ne01/4 up to the subgroup width, so the tail
|
||||
// quads past the last row must not store (they wrote 128 rows past
|
||||
// dst on every ne01 % 256 == 128 vocab, e.g. 151936).
|
||||
if (gid * 4 + 3 < (uint)ne01) {
|
||||
vstore4(acc, 0, &(dst[gid * 4]));
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,118 @@
|
||||
// Tiled-wide q4_K GEMV for the long-vocab lm_head/embed (decode path).
|
||||
//
|
||||
// Pairs with kernel_convert_block_q4_k_tiled_ns (cvt.cl): the weights are laid
|
||||
// out CANONICALLY (4-bit code in element order e in [0,256)) and TILED by 64
|
||||
// output rows so the 64-thread lane group coalesces every weight load. Both the
|
||||
// pack (convert) and the unpack (here) are owned by us -> correct by
|
||||
// construction vs the reference ggml q4_K dequant. Same structure as the q6_K
|
||||
// tiled GEMV; the only differences are the 4-bit dequant and the q4_K
|
||||
// scale/min decode (get_scale_min_k4 from the packed 12-byte block).
|
||||
//
|
||||
// One work-item produces one output row. WG = {64 lanes, 4 subgroups}: the 64
|
||||
// lanes cover the 64 rows of one tile (coalesced uint4 reads), the 4 subgroups
|
||||
// split the K-blocks and reduce through __local at the end. Weights read from
|
||||
// __global (lm_head is streamed once per token; texture cache caps it below the
|
||||
// coalesced-global rate).
|
||||
|
||||
#pragma OPENCL EXTENSION cl_khr_fp16 : 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 TILE_ROWS 64
|
||||
|
||||
// Decode one q4_K sub-block scale + min from the packed 12-byte block.
|
||||
// Identical to the o4 kernel's helper (masks hard-coded: d6=0x3F, d4=0x0F, hi2=0xC0).
|
||||
inline void q4k_scale_min(int j, __global const uchar * q, uchar * d, uchar * m) {
|
||||
if (j < 4) {
|
||||
*d = q[j] & 0x3F;
|
||||
*m = q[j+4] & 0x3F;
|
||||
} else {
|
||||
*d = (q[j+4] & 0x0F) | ((q[j-4] & 0xC0) >> 2);
|
||||
*m = ((q[j+4] >> 4) & 0x0F) | ((q[j] & 0xC0) >> 2);
|
||||
}
|
||||
}
|
||||
|
||||
#if defined(ADRENO_GPU)
|
||||
REQD_SUBGROUP_SIZE_64
|
||||
#endif
|
||||
kernel void kernel_gemv_noshuffle_q4_k_f32_tiled(
|
||||
__global uint4 * src0_q, // tiled: 8 uint4 granules / superblock (4-bit codes)
|
||||
__global half * src0_d, // tiled: 1 half / superblock
|
||||
__global half * src0_dm, // tiled: 1 half / superblock
|
||||
__global uchar * src0_s, // tiled: 12 bytes / superblock (packed scales)
|
||||
read_only image1d_buffer_t src1, // activation (RGBA f32)
|
||||
global float * dst,
|
||||
ulong offsetd,
|
||||
int ne00,
|
||||
int ne01
|
||||
) {
|
||||
int grp = get_local_id(1); // subgroup index 0..3 (splits K)
|
||||
int row = get_global_id(0); // output row along ne01
|
||||
int rt = row / TILE_ROWS;
|
||||
int rit = row % TILE_ROWS;
|
||||
|
||||
int nb = ne00 / QK_K; // superblocks per row
|
||||
|
||||
float acc = 0.0f;
|
||||
|
||||
for (int sb = grp; sb < nb; sb += NSUBGROUPS) {
|
||||
int tile_blk = rt * nb + sb; // ne02 == 1 for lm_head/embed
|
||||
|
||||
float dval = (float)src0_d [tile_blk * TILE_ROWS + rit];
|
||||
float dmval = (float)src0_dm[tile_blk * TILE_ROWS + rit];
|
||||
|
||||
// decode the 8 sub-block (scale, min) pairs
|
||||
__global uchar * sc = src0_s + (tile_blk * TILE_ROWS + rit) * 12;
|
||||
float scale[8], minv[8];
|
||||
#pragma unroll
|
||||
for (int is = 0; is < 8; ++is) {
|
||||
uchar sd, sm;
|
||||
q4k_scale_min(is, sc, &sd, &sm);
|
||||
scale[is] = dval * (float)sd;
|
||||
minv[is] = dmval * (float)sm;
|
||||
}
|
||||
|
||||
// 32 uints of 4-bit codes (8 codes/uint), e-order
|
||||
uint q[32];
|
||||
#pragma unroll
|
||||
for (int g = 0; g < 8; ++g) {
|
||||
uint4 v = src0_q[(tile_blk * 8 + g) * TILE_ROWS + rit];
|
||||
q[g*4+0] = v.x; q[g*4+1] = v.y; q[g*4+2] = v.z; q[g*4+3] = v.w;
|
||||
}
|
||||
|
||||
// dequant 256 codes in canonical e-order, MAC with activation.
|
||||
int act_base = sb * 64; // activation float4 pixel base (256/4)
|
||||
#pragma unroll
|
||||
for (int e4 = 0; e4 < 64; ++e4) {
|
||||
float4 a = read_imagef(src1, act_base + e4);
|
||||
#pragma unroll
|
||||
for (int t = 0; t < 4; ++t) {
|
||||
int e = e4 * 4 + t;
|
||||
uint code = (q[e >> 3] >> ((e & 7) * 4)) & 0xF;
|
||||
int is = e >> 5; // sub-block index = e/32
|
||||
float av = (t == 0) ? a.x : (t == 1) ? a.y : (t == 2) ? a.z : a.w;
|
||||
acc += ((float)code * scale[is] - minv[is]) * av;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
// reduce across the NSUBGROUPS subgroups (same rit, different K-subset)
|
||||
local float reduce_lm[NSUBGROUPS * TILE_ROWS];
|
||||
reduce_lm[grp * TILE_ROWS + rit] = acc;
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
if (grp == 0) {
|
||||
float total = reduce_lm[0 * TILE_ROWS + rit]
|
||||
+ reduce_lm[1 * TILE_ROWS + rit]
|
||||
+ reduce_lm[2 * TILE_ROWS + rit]
|
||||
+ reduce_lm[3 * TILE_ROWS + rit];
|
||||
dst = (global float*)((global char*)dst + offsetd);
|
||||
dst[row] = total;
|
||||
}
|
||||
}
|
||||
@@ -329,3 +329,125 @@ kernel void kernel_gemv_noshuffle_q5_k_f32(
|
||||
if (gid * 2 + 1 < M) dst[gid * 2 + 1] = totalSum.s1;
|
||||
}
|
||||
}
|
||||
|
||||
// Multi-column (N in [2..4]) variant of the q5_K decode GEMV (spec/MTP verify) =
|
||||
// q4_K mc3 + the high-bit qh plane (regH). n_cols = 2..4 (drafted + bonus); routes
|
||||
// the small-batch verify OFF the gemm_noshuffle_q5_k dead-zone. n_cols==3 is byte-
|
||||
// identical to the original mc3 (col3 disabled, float8 slots 6/7 stay zero).
|
||||
#ifdef VECTOR_SUB_GROUP_BROADCAST
|
||||
#define MC_DQ5_HI dequantizeBlockAccum_ns_sgbroadcast_8_hi
|
||||
#define MC_DQ5_LO dequantizeBlockAccum_ns_sgbroadcast_8_lo
|
||||
#else
|
||||
#define MC_DQ5_HI dequantizeBlockAccum_ns_sgbroadcast_1_hi
|
||||
#define MC_DQ5_LO dequantizeBlockAccum_ns_sgbroadcast_1_lo
|
||||
#endif
|
||||
#define MC_COL_Q5K(ts, c) \
|
||||
{ if (slid < 4) { regB.s0123 = read_imagef(src1, (c)*COL_STRIDE + slid*2 + k*8); \
|
||||
regB.s4567 = read_imagef(src1, (c)*COL_STRIDE + 1 + slid*2 + k*8); } \
|
||||
MC_DQ5_HI(ts, as_ushort8(regA_hi), as_uchar8(regH), regS, regM, regB); \
|
||||
MC_DQ5_LO(ts, as_ushort8(regA_lo), as_uchar8(regH), regS, regM, regB); }
|
||||
#ifdef ADRENO_GPU
|
||||
REQD_SUBGROUP_SIZE_64
|
||||
#endif
|
||||
kernel void kernel_gemv_noshuffle_q5_k_f32_mc3(
|
||||
read_only image1d_buffer_t src0_q,
|
||||
read_only image1d_buffer_t src0_qh,
|
||||
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,
|
||||
int n_cols)
|
||||
{
|
||||
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 LINE_STRIDE_A_QH = M / 2;
|
||||
uint BLOCK_STRIDE_A_QH = NSUBGROUPS * M / 2;
|
||||
uint scales_per_row = (K / QK_K) * 12;
|
||||
uint COL_STRIDE = K / 4; // float4 pixels per activation column
|
||||
|
||||
private uint4 regA_hi, regA_lo;
|
||||
private ushort4 regH;
|
||||
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);
|
||||
private float2 ts3 = (float2)(0.0f);
|
||||
|
||||
for (uint k = groupId; k < (K / 32); k += NSUBGROUPS) {
|
||||
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 + 2 * gid * scales_per_row + sb * 12;
|
||||
global const uchar * sc1 = src0_s + (2 * gid + 1) * scales_per_row + sb * 12;
|
||||
|
||||
uchar sv0, mn0, sv1, mn1;
|
||||
get_scale_min_k4(j, sc0, &sv0, &mn0, mask_d6, mask_d4, mask_hi2);
|
||||
get_scale_min_k4(j, sc1, &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)));
|
||||
|
||||
// high-bit plane + weights loaded ONCE, reused across the columns
|
||||
regH.s0 = as_ushort(read_imageh(src0_qh, (gid + k * BLOCK_STRIDE_A_QH + LINE_STRIDE_A_QH * 0)).x);
|
||||
regH.s1 = as_ushort(read_imageh(src0_qh, (gid + k * BLOCK_STRIDE_A_QH + LINE_STRIDE_A_QH * 1)).x);
|
||||
regH.s2 = as_ushort(read_imageh(src0_qh, (gid + k * BLOCK_STRIDE_A_QH + LINE_STRIDE_A_QH * 2)).x);
|
||||
regH.s3 = as_ushort(read_imageh(src0_qh, (gid + k * BLOCK_STRIDE_A_QH + LINE_STRIDE_A_QH * 3)).x);
|
||||
|
||||
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;
|
||||
|
||||
MC_COL_Q5K(ts0, 0);
|
||||
MC_COL_Q5K(ts1, 1);
|
||||
if (n_cols > 2) MC_COL_Q5K(ts2, 2);
|
||||
if (n_cols > 3) MC_COL_Q5K(ts3, 3);
|
||||
}
|
||||
|
||||
// cross-subgroup reduce: pack the (up to 4) columns' float2 into a float8.
|
||||
local float8 reduceLM[SUBGROUP_SIZE * 3];
|
||||
float8 acc = (float8)(ts0.s0, ts0.s1, ts1.s0, ts1.s1, ts2.s0, ts2.s1, ts3.s0, ts3.s1);
|
||||
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 n_cols 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]));
|
||||
if (n_cols > 2) vstore2((float2)(acc.s4, acc.s5), 0, &(dst[2 * M + gid * 2]));
|
||||
if (n_cols > 3) vstore2((float2)(acc.s6, acc.s7), 0, &(dst[3 * M + gid * 2]));
|
||||
}
|
||||
}
|
||||
#undef MC_COL_Q5K
|
||||
#undef MC_DQ5_HI
|
||||
#undef MC_DQ5_LO
|
||||
|
||||
@@ -296,3 +296,114 @@ kernel void kernel_gemv_noshuffle_q6_K_f32(
|
||||
if (gid * 2 + 1 < ne01) dst[gid * 2 + 1] = total_sum.s1;
|
||||
}
|
||||
}
|
||||
|
||||
// Multi-column (N=3) q6_K decode GEMV for the spec/MTP verify batch. Same idea
|
||||
// as the q4_K mc3: stay on the efficient GEMV path (subgroup broadcast, no
|
||||
// transpose) instead of the transposed-GEMM dead-zone. Each K-block's weights
|
||||
// (ql/qh, hi+lo) are loaded ONCE and reused across all 3 activation columns.
|
||||
// Per-column accumulation is independent and identical to 3 standalone GEMVs
|
||||
// => byte-identical; does NOT perturb the lm_head logits / spec accept rate.
|
||||
#if defined(ADRENO_GPU)
|
||||
REQD_SUBGROUP_SIZE_64
|
||||
#endif
|
||||
kernel void kernel_gemv_noshuffle_q6_K_f32_mc3(
|
||||
read_only image1d_buffer_t src0_ql,
|
||||
read_only image1d_buffer_t src0_qh,
|
||||
global half2 * src0_s,
|
||||
global half2 * src0_d,
|
||||
read_only image1d_buffer_t src1,
|
||||
global float * dst,
|
||||
ulong offsetd,
|
||||
int ne00,
|
||||
int ne01
|
||||
) {
|
||||
int grp = get_local_id(1);
|
||||
int gid = get_global_id(0);
|
||||
ushort slid = get_sub_group_local_id();
|
||||
|
||||
int nb = ne00 / 32;
|
||||
int line_stride_a = ne01 / 2;
|
||||
int block_stride_a = NSUBGROUPS * ne01;
|
||||
int COL_STRIDE = ne00 / 4; // float4 pixels per activation column
|
||||
|
||||
uint4 ql_hi, ql_lo;
|
||||
ushort4 qh_hi, qh_lo;
|
||||
half2 reg_d;
|
||||
char4 reg_s;
|
||||
float8 reg_b;
|
||||
|
||||
float2 ts0 = 0.0f, ts1 = 0.0f, ts2 = 0.0f;
|
||||
|
||||
for (int k = grp; k < nb; k += NSUBGROUPS) {
|
||||
reg_d = src0_d[gid + k/8 * line_stride_a];
|
||||
reg_s = as_char4(src0_s[gid + k * line_stride_a]);
|
||||
|
||||
// weights loaded ONCE (hi: blocks 0-3, lo: blocks 4-7), reused x3 cols
|
||||
ql_hi.s0 = read_imageui(src0_ql, gid + k*block_stride_a + line_stride_a*0).x;
|
||||
ql_hi.s1 = read_imageui(src0_ql, gid + k*block_stride_a + line_stride_a*1).x;
|
||||
ql_hi.s2 = read_imageui(src0_ql, gid + k*block_stride_a + line_stride_a*2).x;
|
||||
ql_hi.s3 = read_imageui(src0_ql, gid + k*block_stride_a + line_stride_a*3).x;
|
||||
qh_hi.s0 = as_ushort(read_imageh(src0_qh, gid + k*block_stride_a + line_stride_a*0).x);
|
||||
qh_hi.s1 = as_ushort(read_imageh(src0_qh, gid + k*block_stride_a + line_stride_a*1).x);
|
||||
qh_hi.s2 = as_ushort(read_imageh(src0_qh, gid + k*block_stride_a + line_stride_a*2).x);
|
||||
qh_hi.s3 = as_ushort(read_imageh(src0_qh, gid + k*block_stride_a + line_stride_a*3).x);
|
||||
|
||||
ql_lo.s0 = read_imageui(src0_ql, gid + k*block_stride_a + line_stride_a*4).x;
|
||||
ql_lo.s1 = read_imageui(src0_ql, gid + k*block_stride_a + line_stride_a*5).x;
|
||||
ql_lo.s2 = read_imageui(src0_ql, gid + k*block_stride_a + line_stride_a*6).x;
|
||||
ql_lo.s3 = read_imageui(src0_ql, gid + k*block_stride_a + line_stride_a*7).x;
|
||||
qh_lo.s0 = as_ushort(read_imageh(src0_qh, gid + k*block_stride_a + line_stride_a*4).x);
|
||||
qh_lo.s1 = as_ushort(read_imageh(src0_qh, gid + k*block_stride_a + line_stride_a*5).x);
|
||||
qh_lo.s2 = as_ushort(read_imageh(src0_qh, gid + k*block_stride_a + line_stride_a*6).x);
|
||||
qh_lo.s3 = as_ushort(read_imageh(src0_qh, gid + k*block_stride_a + line_stride_a*7).x);
|
||||
|
||||
// Per-column: load only this column's activation (single reg_b live) ->
|
||||
// 1/3 the activation register pressure, cutting the private-mem spill.
|
||||
#ifdef VECTOR_SUB_GROUP_BROADCAT
|
||||
{ if (slid < 4) { reg_b.s0123 = read_imagef(src1, 0*COL_STRIDE + 0 + slid*2 + k*8);
|
||||
reg_b.s4567 = read_imagef(src1, 0*COL_STRIDE + 1 + slid*2 + k*8); }
|
||||
dequantize_block_acc_bcast_8_hi(ts0, as_ushort8(ql_hi), as_uchar8(qh_hi), reg_d, reg_s, reg_b);
|
||||
dequantize_block_acc_bcast_8_lo(ts0, as_ushort8(ql_lo), as_uchar8(qh_lo), reg_d, reg_s, reg_b); }
|
||||
{ if (slid < 4) { reg_b.s0123 = read_imagef(src1, 1*COL_STRIDE + 0 + slid*2 + k*8);
|
||||
reg_b.s4567 = read_imagef(src1, 1*COL_STRIDE + 1 + slid*2 + k*8); }
|
||||
dequantize_block_acc_bcast_8_hi(ts1, as_ushort8(ql_hi), as_uchar8(qh_hi), reg_d, reg_s, reg_b);
|
||||
dequantize_block_acc_bcast_8_lo(ts1, as_ushort8(ql_lo), as_uchar8(qh_lo), reg_d, reg_s, reg_b); }
|
||||
{ if (slid < 4) { reg_b.s0123 = read_imagef(src1, 2*COL_STRIDE + 0 + slid*2 + k*8);
|
||||
reg_b.s4567 = read_imagef(src1, 2*COL_STRIDE + 1 + slid*2 + k*8); }
|
||||
dequantize_block_acc_bcast_8_hi(ts2, as_ushort8(ql_hi), as_uchar8(qh_hi), reg_d, reg_s, reg_b);
|
||||
dequantize_block_acc_bcast_8_lo(ts2, as_ushort8(ql_lo), as_uchar8(qh_lo), reg_d, reg_s, reg_b); }
|
||||
#else
|
||||
{ if (slid < 4) { reg_b.s0123 = read_imagef(src1, 0*COL_STRIDE + 0 + slid*2 + k*8);
|
||||
reg_b.s4567 = read_imagef(src1, 0*COL_STRIDE + 1 + slid*2 + k*8); }
|
||||
dequantize_block_acc_bcast_1_hi(ts0, as_ushort8(ql_hi), as_uchar8(qh_hi), reg_d, reg_s, reg_b);
|
||||
dequantize_block_acc_bcast_1_lo(ts0, as_ushort8(ql_lo), as_uchar8(qh_lo), reg_d, reg_s, reg_b); }
|
||||
{ if (slid < 4) { reg_b.s0123 = read_imagef(src1, 1*COL_STRIDE + 0 + slid*2 + k*8);
|
||||
reg_b.s4567 = read_imagef(src1, 1*COL_STRIDE + 1 + slid*2 + k*8); }
|
||||
dequantize_block_acc_bcast_1_hi(ts1, as_ushort8(ql_hi), as_uchar8(qh_hi), reg_d, reg_s, reg_b);
|
||||
dequantize_block_acc_bcast_1_lo(ts1, as_ushort8(ql_lo), as_uchar8(qh_lo), reg_d, reg_s, reg_b); }
|
||||
{ if (slid < 4) { reg_b.s0123 = read_imagef(src1, 2*COL_STRIDE + 0 + slid*2 + k*8);
|
||||
reg_b.s4567 = read_imagef(src1, 2*COL_STRIDE + 1 + slid*2 + k*8); }
|
||||
dequantize_block_acc_bcast_1_hi(ts2, as_ushort8(ql_hi), as_uchar8(qh_hi), reg_d, reg_s, reg_b);
|
||||
dequantize_block_acc_bcast_1_lo(ts2, as_ushort8(ql_lo), as_uchar8(qh_lo), reg_d, reg_s, reg_b); }
|
||||
#endif
|
||||
}
|
||||
|
||||
local float8 reduce_lm[SUBGROUP_SIZE * 3];
|
||||
float8 acc = (float8)(ts0.s0, ts0.s1, ts1.s0, ts1.s1, ts2.s0, ts2.s1, 0.0f, 0.0f);
|
||||
if (grp == 1) { reduce_lm[SUBGROUP_SIZE*0 + slid] = acc; }
|
||||
if (grp == 2) { reduce_lm[SUBGROUP_SIZE*1 + slid] = acc; }
|
||||
if (grp == 3) { reduce_lm[SUBGROUP_SIZE*2 + slid] = acc; }
|
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
if (grp == 0) {
|
||||
acc += reduce_lm[SUBGROUP_SIZE*0 + slid];
|
||||
acc += reduce_lm[SUBGROUP_SIZE*1 + slid];
|
||||
acc += reduce_lm[SUBGROUP_SIZE*2 + slid];
|
||||
dst = (global float*)((global char*)dst + offsetd);
|
||||
// dst column-major [ne01 rows x 3 cols]: (row, col) at col*ne01 + row
|
||||
vstore2((float2)(acc.s0, acc.s1), 0, &(dst[0*ne01 + gid*2]));
|
||||
vstore2((float2)(acc.s2, acc.s3), 0, &(dst[1*ne01 + gid*2]));
|
||||
vstore2((float2)(acc.s4, acc.s5), 0, &(dst[2*ne01 + gid*2]));
|
||||
}
|
||||
}
|
||||
|
||||
@@ -0,0 +1,372 @@
|
||||
// 4-output-per-WI variant of kernel_gemv_noshuffle_q6_K_f32.
|
||||
// Each WI now produces 4 consecutive outputs (output quad). The activation
|
||||
// fetch (reg_b) is shared across all 4 outputs, doubling per-WI ALU per
|
||||
// activation broadcast and halving the WG count vs the 2-output kernel.
|
||||
//
|
||||
// Implementation: each K-block we fetch TWO sets of (scales + ql + qh)
|
||||
// — one for the low pair (rows 0,1 of the quad) and one for the high pair
|
||||
// (rows 2,3) — and invoke the existing 2-output dequant macros twice
|
||||
// against the *same* reg_b. Identical data layout to the 2-output kernel,
|
||||
// so the host only needs to halve the grid and double the gid-to-output
|
||||
// mapping.
|
||||
//
|
||||
// Opt-in via the host dispatch when GGML_OPENCL_Q6K_GEMV_O4=1.
|
||||
|
||||
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
|
||||
#pragma OPENCL EXTENSION cl_khr_subgroups : 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 NSUBGROUPS 4
|
||||
#define SUBGROUP_SIZE 64
|
||||
|
||||
// Macros are identical to the 2-output kernel — they accept `total_sum` as
|
||||
// a parameter so we can call them twice (once per pair) against different
|
||||
// accumulators against the same reg_b.
|
||||
#define dequantize_block_acc_bcast_8_hi(total_sum, bits4, bits2, cs, y) \
|
||||
float8 shared_y; \
|
||||
shared_y = sub_group_broadcast(y, 0); \
|
||||
total_sum.s0 += ((float)(((bits4.s0 & 0x000F) ) | ((bits2.s0 & 0x03) << 4)) - 32.f) * cs.s0 * shared_y.s0; \
|
||||
total_sum.s0 += ((float)(((bits4.s0 & 0x00F0) >> 4) | ((bits2.s0 & 0x0C) << 2)) - 32.f) * cs.s0 * shared_y.s1; \
|
||||
total_sum.s0 += ((float)(((bits4.s0 & 0x0F00) >> 8) | ((bits2.s0 & 0x30) )) - 32.f) * cs.s0 * shared_y.s2; \
|
||||
total_sum.s0 += ((float)(((bits4.s0 & 0xF000) >> 12) | ((bits2.s0 & 0xC0) >> 2)) - 32.f) * cs.s0 * shared_y.s3; \
|
||||
total_sum.s0 += ((float)(((bits4.s2 & 0x000F) ) | ((bits2.s2 & 0x03) << 4)) - 32.f) * cs.s0 * shared_y.s4; \
|
||||
total_sum.s0 += ((float)(((bits4.s2 & 0x00F0) >> 4) | ((bits2.s2 & 0x0C) << 2)) - 32.f) * cs.s0 * shared_y.s5; \
|
||||
total_sum.s0 += ((float)(((bits4.s2 & 0x0F00) >> 8) | ((bits2.s2 & 0x30) )) - 32.f) * cs.s0 * shared_y.s6; \
|
||||
total_sum.s0 += ((float)(((bits4.s2 & 0xF000) >> 12) | ((bits2.s2 & 0xC0) >> 2)) - 32.f) * cs.s0 * shared_y.s7; \
|
||||
total_sum.s1 += ((float)(((bits4.s1 & 0x000F) ) | ((bits2.s1 & 0x03) << 4)) - 32.f) * cs.s2 * shared_y.s0; \
|
||||
total_sum.s1 += ((float)(((bits4.s1 & 0x00F0) >> 4) | ((bits2.s1 & 0x0C) << 2)) - 32.f) * cs.s2 * shared_y.s1; \
|
||||
total_sum.s1 += ((float)(((bits4.s1 & 0x0F00) >> 8) | ((bits2.s1 & 0x30) )) - 32.f) * cs.s2 * shared_y.s2; \
|
||||
total_sum.s1 += ((float)(((bits4.s1 & 0xF000) >> 12) | ((bits2.s1 & 0xC0) >> 2)) - 32.f) * cs.s2 * shared_y.s3; \
|
||||
total_sum.s1 += ((float)(((bits4.s3 & 0x000F) ) | ((bits2.s3 & 0x03) << 4)) - 32.f) * cs.s2 * shared_y.s4; \
|
||||
total_sum.s1 += ((float)(((bits4.s3 & 0x00F0) >> 4) | ((bits2.s3 & 0x0C) << 2)) - 32.f) * cs.s2 * shared_y.s5; \
|
||||
total_sum.s1 += ((float)(((bits4.s3 & 0x0F00) >> 8) | ((bits2.s3 & 0x30) )) - 32.f) * cs.s2 * shared_y.s6; \
|
||||
total_sum.s1 += ((float)(((bits4.s3 & 0xF000) >> 12) | ((bits2.s3 & 0xC0) >> 2)) - 32.f) * cs.s2 * shared_y.s7; \
|
||||
shared_y = sub_group_broadcast(y, 1); \
|
||||
total_sum.s0 += ((float)(((bits4.s4 & 0x000F) ) | ((bits2.s4 & 0x03) << 4)) - 32.f) * cs.s0 * shared_y.s0; \
|
||||
total_sum.s0 += ((float)(((bits4.s4 & 0x00F0) >> 4) | ((bits2.s4 & 0x0C) << 2)) - 32.f) * cs.s0 * shared_y.s1; \
|
||||
total_sum.s0 += ((float)(((bits4.s4 & 0x0F00) >> 8) | ((bits2.s4 & 0x30) )) - 32.f) * cs.s0 * shared_y.s2; \
|
||||
total_sum.s0 += ((float)(((bits4.s4 & 0xF000) >> 12) | ((bits2.s4 & 0xC0) >> 2)) - 32.f) * cs.s0 * shared_y.s3; \
|
||||
total_sum.s0 += ((float)(((bits4.s6 & 0x000F) ) | ((bits2.s6 & 0x03) << 4)) - 32.f) * cs.s0 * shared_y.s4; \
|
||||
total_sum.s0 += ((float)(((bits4.s6 & 0x00F0) >> 4) | ((bits2.s6 & 0x0C) << 2)) - 32.f) * cs.s0 * shared_y.s5; \
|
||||
total_sum.s0 += ((float)(((bits4.s6 & 0x0F00) >> 8) | ((bits2.s6 & 0x30) )) - 32.f) * cs.s0 * shared_y.s6; \
|
||||
total_sum.s0 += ((float)(((bits4.s6 & 0xF000) >> 12) | ((bits2.s6 & 0xC0) >> 2)) - 32.f) * cs.s0 * shared_y.s7; \
|
||||
total_sum.s1 += ((float)(((bits4.s5 & 0x000F) ) | ((bits2.s5 & 0x03) << 4)) - 32.f) * cs.s2 * shared_y.s0; \
|
||||
total_sum.s1 += ((float)(((bits4.s5 & 0x00F0) >> 4) | ((bits2.s5 & 0x0C) << 2)) - 32.f) * cs.s2 * shared_y.s1; \
|
||||
total_sum.s1 += ((float)(((bits4.s5 & 0x0F00) >> 8) | ((bits2.s5 & 0x30) )) - 32.f) * cs.s2 * shared_y.s2; \
|
||||
total_sum.s1 += ((float)(((bits4.s5 & 0xF000) >> 12) | ((bits2.s5 & 0xC0) >> 2)) - 32.f) * cs.s2 * shared_y.s3; \
|
||||
total_sum.s1 += ((float)(((bits4.s7 & 0x000F) ) | ((bits2.s7 & 0x03) << 4)) - 32.f) * cs.s2 * shared_y.s4; \
|
||||
total_sum.s1 += ((float)(((bits4.s7 & 0x00F0) >> 4) | ((bits2.s7 & 0x0C) << 2)) - 32.f) * cs.s2 * shared_y.s5; \
|
||||
total_sum.s1 += ((float)(((bits4.s7 & 0x0F00) >> 8) | ((bits2.s7 & 0x30) )) - 32.f) * cs.s2 * shared_y.s6; \
|
||||
total_sum.s1 += ((float)(((bits4.s7 & 0xF000) >> 12) | ((bits2.s7 & 0xC0) >> 2)) - 32.f) * cs.s2 * shared_y.s7; \
|
||||
|
||||
#define dequantize_block_acc_bcast_8_lo(total_sum, bits4, bits2, cs, y) \
|
||||
shared_y = sub_group_broadcast(y, 2); \
|
||||
total_sum.s0 += ((float)(((bits4.s0 & 0x000F) ) | ((bits2.s0 & 0x03) << 4)) - 32.f) * cs.s1 * shared_y.s0; \
|
||||
total_sum.s0 += ((float)(((bits4.s0 & 0x00F0) >> 4) | ((bits2.s0 & 0x0C) << 2)) - 32.f) * cs.s1 * shared_y.s1; \
|
||||
total_sum.s0 += ((float)(((bits4.s0 & 0x0F00) >> 8) | ((bits2.s0 & 0x30) )) - 32.f) * cs.s1 * shared_y.s2; \
|
||||
total_sum.s0 += ((float)(((bits4.s0 & 0xF000) >> 12) | ((bits2.s0 & 0xC0) >> 2)) - 32.f) * cs.s1 * shared_y.s3; \
|
||||
total_sum.s0 += ((float)(((bits4.s2 & 0x000F) ) | ((bits2.s2 & 0x03) << 4)) - 32.f) * cs.s1 * shared_y.s4; \
|
||||
total_sum.s0 += ((float)(((bits4.s2 & 0x00F0) >> 4) | ((bits2.s2 & 0x0C) << 2)) - 32.f) * cs.s1 * shared_y.s5; \
|
||||
total_sum.s0 += ((float)(((bits4.s2 & 0x0F00) >> 8) | ((bits2.s2 & 0x30) )) - 32.f) * cs.s1 * shared_y.s6; \
|
||||
total_sum.s0 += ((float)(((bits4.s2 & 0xF000) >> 12) | ((bits2.s2 & 0xC0) >> 2)) - 32.f) * cs.s1 * shared_y.s7; \
|
||||
total_sum.s1 += ((float)(((bits4.s1 & 0x000F) ) | ((bits2.s1 & 0x03) << 4)) - 32.f) * cs.s3 * shared_y.s0; \
|
||||
total_sum.s1 += ((float)(((bits4.s1 & 0x00F0) >> 4) | ((bits2.s1 & 0x0C) << 2)) - 32.f) * cs.s3 * shared_y.s1; \
|
||||
total_sum.s1 += ((float)(((bits4.s1 & 0x0F00) >> 8) | ((bits2.s1 & 0x30) )) - 32.f) * cs.s3 * shared_y.s2; \
|
||||
total_sum.s1 += ((float)(((bits4.s1 & 0xF000) >> 12) | ((bits2.s1 & 0xC0) >> 2)) - 32.f) * cs.s3 * shared_y.s3; \
|
||||
total_sum.s1 += ((float)(((bits4.s3 & 0x000F) ) | ((bits2.s3 & 0x03) << 4)) - 32.f) * cs.s3 * shared_y.s4; \
|
||||
total_sum.s1 += ((float)(((bits4.s3 & 0x00F0) >> 4) | ((bits2.s3 & 0x0C) << 2)) - 32.f) * cs.s3 * shared_y.s5; \
|
||||
total_sum.s1 += ((float)(((bits4.s3 & 0x0F00) >> 8) | ((bits2.s3 & 0x30) )) - 32.f) * cs.s3 * shared_y.s6; \
|
||||
total_sum.s1 += ((float)(((bits4.s3 & 0xF000) >> 12) | ((bits2.s3 & 0xC0) >> 2)) - 32.f) * cs.s3 * shared_y.s7; \
|
||||
shared_y = sub_group_broadcast(y, 3); \
|
||||
total_sum.s0 += ((float)(((bits4.s4 & 0x000F) ) | ((bits2.s4 & 0x03) << 4)) - 32.f) * cs.s1 * shared_y.s0; \
|
||||
total_sum.s0 += ((float)(((bits4.s4 & 0x00F0) >> 4) | ((bits2.s4 & 0x0C) << 2)) - 32.f) * cs.s1 * shared_y.s1; \
|
||||
total_sum.s0 += ((float)(((bits4.s4 & 0x0F00) >> 8) | ((bits2.s4 & 0x30) )) - 32.f) * cs.s1 * shared_y.s2; \
|
||||
total_sum.s0 += ((float)(((bits4.s4 & 0xF000) >> 12) | ((bits2.s4 & 0xC0) >> 2)) - 32.f) * cs.s1 * shared_y.s3; \
|
||||
total_sum.s0 += ((float)(((bits4.s6 & 0x000F) ) | ((bits2.s6 & 0x03) << 4)) - 32.f) * cs.s1 * shared_y.s4; \
|
||||
total_sum.s0 += ((float)(((bits4.s6 & 0x00F0) >> 4) | ((bits2.s6 & 0x0C) << 2)) - 32.f) * cs.s1 * shared_y.s5; \
|
||||
total_sum.s0 += ((float)(((bits4.s6 & 0x0F00) >> 8) | ((bits2.s6 & 0x30) )) - 32.f) * cs.s1 * shared_y.s6; \
|
||||
total_sum.s0 += ((float)(((bits4.s6 & 0xF000) >> 12) | ((bits2.s6 & 0xC0) >> 2)) - 32.f) * cs.s1 * shared_y.s7; \
|
||||
total_sum.s1 += ((float)(((bits4.s5 & 0x000F) ) | ((bits2.s5 & 0x03) << 4)) - 32.f) * cs.s3 * shared_y.s0; \
|
||||
total_sum.s1 += ((float)(((bits4.s5 & 0x00F0) >> 4) | ((bits2.s5 & 0x0C) << 2)) - 32.f) * cs.s3 * shared_y.s1; \
|
||||
total_sum.s1 += ((float)(((bits4.s5 & 0x0F00) >> 8) | ((bits2.s5 & 0x30) )) - 32.f) * cs.s3 * shared_y.s2; \
|
||||
total_sum.s1 += ((float)(((bits4.s5 & 0xF000) >> 12) | ((bits2.s5 & 0xC0) >> 2)) - 32.f) * cs.s3 * shared_y.s3; \
|
||||
total_sum.s1 += ((float)(((bits4.s7 & 0x000F) ) | ((bits2.s7 & 0x03) << 4)) - 32.f) * cs.s3 * shared_y.s4; \
|
||||
total_sum.s1 += ((float)(((bits4.s7 & 0x00F0) >> 4) | ((bits2.s7 & 0x0C) << 2)) - 32.f) * cs.s3 * shared_y.s5; \
|
||||
total_sum.s1 += ((float)(((bits4.s7 & 0x0F00) >> 8) | ((bits2.s7 & 0x30) )) - 32.f) * cs.s3 * shared_y.s6; \
|
||||
total_sum.s1 += ((float)(((bits4.s7 & 0xF000) >> 12) | ((bits2.s7 & 0xC0) >> 2)) - 32.f) * cs.s3 * shared_y.s7; \
|
||||
|
||||
#define dequantize_block_acc_bcast_1_hi(total_sum, bits4, bits2, cs, y) \
|
||||
float shared_y; \
|
||||
shared_y = sub_group_broadcast(y.s0, 0); \
|
||||
total_sum.s0 += ((float)(((bits4.s0 & 0x000F) ) | ((bits2.s0 & 0x03) << 4)) - 32.f) * cs.s0 * shared_y; \
|
||||
total_sum.s1 += ((float)(((bits4.s1 & 0x000F) ) | ((bits2.s1 & 0x03) << 4)) - 32.f) * cs.s2 * shared_y; \
|
||||
shared_y = sub_group_broadcast(y.s1, 0); \
|
||||
total_sum.s0 += ((float)(((bits4.s0 & 0x00F0) >> 4) | ((bits2.s0 & 0x0C) << 2)) - 32.f) * cs.s0 * shared_y; \
|
||||
total_sum.s1 += ((float)(((bits4.s1 & 0x00F0) >> 4) | ((bits2.s1 & 0x0C) << 2)) - 32.f) * cs.s2 * shared_y; \
|
||||
shared_y = sub_group_broadcast(y.s2, 0); \
|
||||
total_sum.s0 += ((float)(((bits4.s0 & 0x0F00) >> 8) | ((bits2.s0 & 0x30) )) - 32.f) * cs.s0 * shared_y; \
|
||||
total_sum.s1 += ((float)(((bits4.s1 & 0x0F00) >> 8) | ((bits2.s1 & 0x30) )) - 32.f) * cs.s2 * shared_y; \
|
||||
shared_y = sub_group_broadcast(y.s3, 0); \
|
||||
total_sum.s0 += ((float)(((bits4.s0 & 0xF000) >> 12) | ((bits2.s0 & 0xC0) >> 2)) - 32.f) * cs.s0 * shared_y; \
|
||||
total_sum.s1 += ((float)(((bits4.s1 & 0xF000) >> 12) | ((bits2.s1 & 0xC0) >> 2)) - 32.f) * cs.s2 * shared_y; \
|
||||
shared_y = sub_group_broadcast(y.s4, 0); \
|
||||
total_sum.s0 += ((float)(((bits4.s2 & 0x000F) ) | ((bits2.s2 & 0x03) << 4)) - 32.f) * cs.s0 * shared_y; \
|
||||
total_sum.s1 += ((float)(((bits4.s3 & 0x000F) ) | ((bits2.s3 & 0x03) << 4)) - 32.f) * cs.s2 * shared_y; \
|
||||
shared_y = sub_group_broadcast(y.s5, 0); \
|
||||
total_sum.s0 += ((float)(((bits4.s2 & 0x00F0) >> 4) | ((bits2.s2 & 0x0C) << 2)) - 32.f) * cs.s0 * shared_y; \
|
||||
total_sum.s1 += ((float)(((bits4.s3 & 0x00F0) >> 4) | ((bits2.s3 & 0x0C) << 2)) - 32.f) * cs.s2 * shared_y; \
|
||||
shared_y = sub_group_broadcast(y.s6, 0); \
|
||||
total_sum.s0 += ((float)(((bits4.s2 & 0x0F00) >> 8) | ((bits2.s2 & 0x30) )) - 32.f) * cs.s0 * shared_y; \
|
||||
total_sum.s1 += ((float)(((bits4.s3 & 0x0F00) >> 8) | ((bits2.s3 & 0x30) )) - 32.f) * cs.s2 * shared_y; \
|
||||
shared_y = sub_group_broadcast(y.s7, 0); \
|
||||
total_sum.s0 += ((float)(((bits4.s2 & 0xF000) >> 12) | ((bits2.s2 & 0xC0) >> 2)) - 32.f) * cs.s0 * shared_y; \
|
||||
total_sum.s1 += ((float)(((bits4.s3 & 0xF000) >> 12) | ((bits2.s3 & 0xC0) >> 2)) - 32.f) * cs.s2 * shared_y; \
|
||||
shared_y = sub_group_broadcast(y.s0, 1); \
|
||||
total_sum.s0 += ((float)(((bits4.s4 & 0x000F) ) | ((bits2.s4 & 0x03) << 4)) - 32.f) * cs.s0 * shared_y; \
|
||||
total_sum.s1 += ((float)(((bits4.s5 & 0x000F) ) | ((bits2.s5 & 0x03) << 4)) - 32.f) * cs.s2 * shared_y; \
|
||||
shared_y = sub_group_broadcast(y.s1, 1); \
|
||||
total_sum.s0 += ((float)(((bits4.s4 & 0x00F0) >> 4) | ((bits2.s4 & 0x0C) << 2)) - 32.f) * cs.s0 * shared_y; \
|
||||
total_sum.s1 += ((float)(((bits4.s5 & 0x00F0) >> 4) | ((bits2.s5 & 0x0C) << 2)) - 32.f) * cs.s2 * shared_y; \
|
||||
shared_y = sub_group_broadcast(y.s2, 1); \
|
||||
total_sum.s0 += ((float)(((bits4.s4 & 0x0F00) >> 8) | ((bits2.s4 & 0x30) )) - 32.f) * cs.s0 * shared_y; \
|
||||
total_sum.s1 += ((float)(((bits4.s5 & 0x0F00) >> 8) | ((bits2.s5 & 0x30) )) - 32.f) * cs.s2 * shared_y; \
|
||||
shared_y = sub_group_broadcast(y.s3, 1); \
|
||||
total_sum.s0 += ((float)(((bits4.s4 & 0xF000) >> 12) | ((bits2.s4 & 0xC0) >> 2)) - 32.f) * cs.s0 * shared_y; \
|
||||
total_sum.s1 += ((float)(((bits4.s5 & 0xF000) >> 12) | ((bits2.s5 & 0xC0) >> 2)) - 32.f) * cs.s2 * shared_y; \
|
||||
shared_y = sub_group_broadcast(y.s4, 1); \
|
||||
total_sum.s0 += ((float)(((bits4.s6 & 0x000F) ) | ((bits2.s6 & 0x03) << 4)) - 32.f) * cs.s0 * shared_y; \
|
||||
total_sum.s1 += ((float)(((bits4.s7 & 0x000F) ) | ((bits2.s7 & 0x03) << 4)) - 32.f) * cs.s2 * shared_y; \
|
||||
shared_y = sub_group_broadcast(y.s5, 1); \
|
||||
total_sum.s0 += ((float)(((bits4.s6 & 0x00F0) >> 4) | ((bits2.s6 & 0x0C) << 2)) - 32.f) * cs.s0 * shared_y; \
|
||||
total_sum.s1 += ((float)(((bits4.s7 & 0x00F0) >> 4) | ((bits2.s7 & 0x0C) << 2)) - 32.f) * cs.s2 * shared_y; \
|
||||
shared_y = sub_group_broadcast(y.s6, 1); \
|
||||
total_sum.s0 += ((float)(((bits4.s6 & 0x0F00) >> 8) | ((bits2.s6 & 0x30) )) - 32.f) * cs.s0 * shared_y; \
|
||||
total_sum.s1 += ((float)(((bits4.s7 & 0x0F00) >> 8) | ((bits2.s7 & 0x30) )) - 32.f) * cs.s2 * shared_y; \
|
||||
shared_y = sub_group_broadcast(y.s7, 1); \
|
||||
total_sum.s0 += ((float)(((bits4.s6 & 0xF000) >> 12) | ((bits2.s6 & 0xC0) >> 2)) - 32.f) * cs.s0 * shared_y; \
|
||||
total_sum.s1 += ((float)(((bits4.s7 & 0xF000) >> 12) | ((bits2.s7 & 0xC0) >> 2)) - 32.f) * cs.s2 * shared_y; \
|
||||
|
||||
#define dequantize_block_acc_bcast_1_lo(total_sum, bits4, bits2, cs, y) \
|
||||
shared_y = sub_group_broadcast(y.s0, 2); \
|
||||
total_sum.s0 += ((float)(((bits4.s0 & 0x000F) ) | ((bits2.s0 & 0x03) << 4)) - 32.f) * cs.s1 * shared_y; \
|
||||
total_sum.s1 += ((float)(((bits4.s1 & 0x000F) ) | ((bits2.s1 & 0x03) << 4)) - 32.f) * cs.s3 * shared_y; \
|
||||
shared_y = sub_group_broadcast(y.s1, 2); \
|
||||
total_sum.s0 += ((float)(((bits4.s0 & 0x00F0) >> 4) | ((bits2.s0 & 0x0C) << 2)) - 32.f) * cs.s1 * shared_y; \
|
||||
total_sum.s1 += ((float)(((bits4.s1 & 0x00F0) >> 4) | ((bits2.s1 & 0x0C) << 2)) - 32.f) * cs.s3 * shared_y; \
|
||||
shared_y = sub_group_broadcast(y.s2, 2); \
|
||||
total_sum.s0 += ((float)(((bits4.s0 & 0x0F00) >> 8) | ((bits2.s0 & 0x30) )) - 32.f) * cs.s1 * shared_y; \
|
||||
total_sum.s1 += ((float)(((bits4.s1 & 0x0F00) >> 8) | ((bits2.s1 & 0x30) )) - 32.f) * cs.s3 * shared_y; \
|
||||
shared_y = sub_group_broadcast(y.s3, 2); \
|
||||
total_sum.s0 += ((float)(((bits4.s0 & 0xF000) >> 12) | ((bits2.s0 & 0xC0) >> 2)) - 32.f) * cs.s1 * shared_y; \
|
||||
total_sum.s1 += ((float)(((bits4.s1 & 0xF000) >> 12) | ((bits2.s1 & 0xC0) >> 2)) - 32.f) * cs.s3 * shared_y; \
|
||||
shared_y = sub_group_broadcast(y.s4, 2); \
|
||||
total_sum.s0 += ((float)(((bits4.s2 & 0x000F) ) | ((bits2.s2 & 0x03) << 4)) - 32.f) * cs.s1 * shared_y; \
|
||||
total_sum.s1 += ((float)(((bits4.s3 & 0x000F) ) | ((bits2.s3 & 0x03) << 4)) - 32.f) * cs.s3 * shared_y; \
|
||||
shared_y = sub_group_broadcast(y.s5, 2); \
|
||||
total_sum.s0 += ((float)(((bits4.s2 & 0x00F0) >> 4) | ((bits2.s2 & 0x0C) << 2)) - 32.f) * cs.s1 * shared_y; \
|
||||
total_sum.s1 += ((float)(((bits4.s3 & 0x00F0) >> 4) | ((bits2.s3 & 0x0C) << 2)) - 32.f) * cs.s3 * shared_y; \
|
||||
shared_y = sub_group_broadcast(y.s6, 2); \
|
||||
total_sum.s0 += ((float)(((bits4.s2 & 0x0F00) >> 8) | ((bits2.s2 & 0x30) )) - 32.f) * cs.s1 * shared_y; \
|
||||
total_sum.s1 += ((float)(((bits4.s3 & 0x0F00) >> 8) | ((bits2.s3 & 0x30) )) - 32.f) * cs.s3 * shared_y; \
|
||||
shared_y = sub_group_broadcast(y.s7, 2); \
|
||||
total_sum.s0 += ((float)(((bits4.s2 & 0xF000) >> 12) | ((bits2.s2 & 0xC0) >> 2)) - 32.f) * cs.s1 * shared_y; \
|
||||
total_sum.s1 += ((float)(((bits4.s3 & 0xF000) >> 12) | ((bits2.s3 & 0xC0) >> 2)) - 32.f) * cs.s3 * shared_y; \
|
||||
shared_y = sub_group_broadcast(y.s0, 3); \
|
||||
total_sum.s0 += ((float)(((bits4.s4 & 0x000F) ) | ((bits2.s4 & 0x03) << 4)) - 32.f) * cs.s1 * shared_y; \
|
||||
total_sum.s1 += ((float)(((bits4.s5 & 0x000F) ) | ((bits2.s5 & 0x03) << 4)) - 32.f) * cs.s3 * shared_y; \
|
||||
shared_y = sub_group_broadcast(y.s1, 3); \
|
||||
total_sum.s0 += ((float)(((bits4.s4 & 0x00F0) >> 4) | ((bits2.s4 & 0x0C) << 2)) - 32.f) * cs.s1 * shared_y; \
|
||||
total_sum.s1 += ((float)(((bits4.s5 & 0x00F0) >> 4) | ((bits2.s5 & 0x0C) << 2)) - 32.f) * cs.s3 * shared_y; \
|
||||
shared_y = sub_group_broadcast(y.s2, 3); \
|
||||
total_sum.s0 += ((float)(((bits4.s4 & 0x0F00) >> 8) | ((bits2.s4 & 0x30) )) - 32.f) * cs.s1 * shared_y; \
|
||||
total_sum.s1 += ((float)(((bits4.s5 & 0x0F00) >> 8) | ((bits2.s5 & 0x30) )) - 32.f) * cs.s3 * shared_y; \
|
||||
shared_y = sub_group_broadcast(y.s3, 3); \
|
||||
total_sum.s0 += ((float)(((bits4.s4 & 0xF000) >> 12) | ((bits2.s4 & 0xC0) >> 2)) - 32.f) * cs.s1 * shared_y; \
|
||||
total_sum.s1 += ((float)(((bits4.s5 & 0xF000) >> 12) | ((bits2.s5 & 0xC0) >> 2)) - 32.f) * cs.s3 * shared_y; \
|
||||
shared_y = sub_group_broadcast(y.s4, 3); \
|
||||
total_sum.s0 += ((float)(((bits4.s6 & 0x000F) ) | ((bits2.s6 & 0x03) << 4)) - 32.f) * cs.s1 * shared_y; \
|
||||
total_sum.s1 += ((float)(((bits4.s7 & 0x000F) ) | ((bits2.s7 & 0x03) << 4)) - 32.f) * cs.s3 * shared_y; \
|
||||
shared_y = sub_group_broadcast(y.s5, 3); \
|
||||
total_sum.s0 += ((float)(((bits4.s6 & 0x00F0) >> 4) | ((bits2.s6 & 0x0C) << 2)) - 32.f) * cs.s1 * shared_y; \
|
||||
total_sum.s1 += ((float)(((bits4.s7 & 0x00F0) >> 4) | ((bits2.s7 & 0x0C) << 2)) - 32.f) * cs.s3 * shared_y; \
|
||||
shared_y = sub_group_broadcast(y.s6, 3); \
|
||||
total_sum.s0 += ((float)(((bits4.s6 & 0x0F00) >> 8) | ((bits2.s6 & 0x30) )) - 32.f) * cs.s1 * shared_y; \
|
||||
total_sum.s1 += ((float)(((bits4.s7 & 0x0F00) >> 8) | ((bits2.s7 & 0x30) )) - 32.f) * cs.s3 * shared_y; \
|
||||
shared_y = sub_group_broadcast(y.s7, 3); \
|
||||
total_sum.s0 += ((float)(((bits4.s6 & 0xF000) >> 12) | ((bits2.s6 & 0xC0) >> 2)) - 32.f) * cs.s1 * shared_y; \
|
||||
total_sum.s1 += ((float)(((bits4.s7 & 0xF000) >> 12) | ((bits2.s7 & 0xC0) >> 2)) - 32.f) * cs.s3 * shared_y; \
|
||||
|
||||
#if defined(ADRENO_GPU)
|
||||
REQD_SUBGROUP_SIZE_64
|
||||
#endif
|
||||
// Q6K_O4_GLOBAL: read the (read-once-per-token, no-reuse) lm_head/embed weights
|
||||
// from __global coalesced instead of image1d_buffer. The texture cache caps the
|
||||
// streaming (no-reuse) lm_head read bandwidth; global coalesced reaches the
|
||||
// higher rate the rest of the model gets. src1 (activation) stays an image (it IS reused via
|
||||
// the cross-subgroup broadcast).
|
||||
#ifdef Q6K_O4_GLOBAL
|
||||
#define Q6K_O4_NAME kernel_gemv_noshuffle_q6_K_f32_o4_global
|
||||
#define QL_ARG __global uint * src0_ql
|
||||
#define QH_ARG __global half * src0_qh
|
||||
#define RD_QL(b,i) (b[i])
|
||||
#define RD_QH(b,i) as_ushort(b[i])
|
||||
#else
|
||||
#define Q6K_O4_NAME kernel_gemv_noshuffle_q6_K_f32_o4
|
||||
#define QL_ARG read_only image1d_buffer_t src0_ql
|
||||
#define QH_ARG read_only image1d_buffer_t src0_qh
|
||||
#define RD_QL(b,i) (read_imageui(b,i).x)
|
||||
#define RD_QH(b,i) as_ushort(read_imageh(b,i).x)
|
||||
#endif
|
||||
kernel void Q6K_O4_NAME(
|
||||
QL_ARG,
|
||||
QH_ARG,
|
||||
global half2 * src0_s,
|
||||
global half2 * src0_d,
|
||||
read_only image1d_buffer_t src1,
|
||||
global float * dst,
|
||||
ulong offsetd,
|
||||
int ne00,
|
||||
int ne01
|
||||
) {
|
||||
int grp = get_local_id(1);
|
||||
int gid = get_global_id(0); // 4-output-quad index
|
||||
ushort slid = get_sub_group_local_id();
|
||||
|
||||
// Map quad index to the two pair-indices the existing 2-output access
|
||||
// pattern uses (consecutive output pairs along ne01). NB: the two pairs are
|
||||
// kept ADJACENT (gid*2, gid*2+1) on purpose -- a "stride-1" split (pairs
|
||||
// ne01/4 apart) is slower because two distant cache-line streams have worse
|
||||
// locality than the adjacent pair whose reads interleave into the same lines
|
||||
// each iteration.
|
||||
int gid_a = gid * 2;
|
||||
int gid_b = gid * 2 + 1;
|
||||
|
||||
int nb = ne00 / 32;
|
||||
|
||||
uint4 reg_a_l_a, reg_a_l_b;
|
||||
ushort4 reg_a_h_a, reg_a_h_b;
|
||||
half2 reg_d_a, reg_d_b;
|
||||
char4 reg_s_a, reg_s_b;
|
||||
float8 reg_b;
|
||||
|
||||
float2 total_sum_a = 0.0f;
|
||||
float2 total_sum_b = 0.0f;
|
||||
|
||||
int line_stride_a = ne01 / 2;
|
||||
int block_stride_a = NSUBGROUPS * ne01;
|
||||
|
||||
for (int k = grp; k < nb; k += NSUBGROUPS) {
|
||||
reg_d_a = src0_d[gid_a + k/8 * line_stride_a];
|
||||
reg_d_b = src0_d[gid_b + k/8 * line_stride_a];
|
||||
reg_s_a = as_char4(src0_s[gid_a + k * line_stride_a]);
|
||||
reg_s_b = as_char4(src0_s[gid_b + k * line_stride_a]);
|
||||
// Precompute the loop-invariant combined scale (sub-block scale * super-block d)
|
||||
// once per pair instead of re-multiplying it for every one of the 256 elements.
|
||||
float4 cs_a = (float4)((float)reg_s_a.s0*(float)reg_d_a.s0, (float)reg_s_a.s1*(float)reg_d_a.s0,
|
||||
(float)reg_s_a.s2*(float)reg_d_a.s1, (float)reg_s_a.s3*(float)reg_d_a.s1);
|
||||
float4 cs_b = (float4)((float)reg_s_b.s0*(float)reg_d_b.s0, (float)reg_s_b.s1*(float)reg_d_b.s0,
|
||||
(float)reg_s_b.s2*(float)reg_d_b.s1, (float)reg_s_b.s3*(float)reg_d_b.s1);
|
||||
|
||||
if (slid < 4) {
|
||||
reg_b.s0123 = read_imagef(src1, 0 + slid*2 + k*8);
|
||||
reg_b.s4567 = read_imagef(src1, 1 + slid*2 + k*8);
|
||||
}
|
||||
|
||||
// Pair a (output rows gid_a*2, gid_a*2+1): read hi+lo then dequant
|
||||
// both in one block so the `_lo` macro can see the `shared_y` that
|
||||
// `_hi` declared. Pair b follows in its own block — fresh shared_y.
|
||||
{
|
||||
reg_a_l_a.s0 = RD_QL(src0_ql, gid_a + k*block_stride_a + line_stride_a*0);
|
||||
reg_a_l_a.s1 = RD_QL(src0_ql, gid_a + k*block_stride_a + line_stride_a*1);
|
||||
reg_a_l_a.s2 = RD_QL(src0_ql, gid_a + k*block_stride_a + line_stride_a*2);
|
||||
reg_a_l_a.s3 = RD_QL(src0_ql, gid_a + k*block_stride_a + line_stride_a*3);
|
||||
reg_a_h_a.s0 = RD_QH(src0_qh, gid_a + k*block_stride_a + line_stride_a*0);
|
||||
reg_a_h_a.s1 = RD_QH(src0_qh, gid_a + k*block_stride_a + line_stride_a*1);
|
||||
reg_a_h_a.s2 = RD_QH(src0_qh, gid_a + k*block_stride_a + line_stride_a*2);
|
||||
reg_a_h_a.s3 = RD_QH(src0_qh, gid_a + k*block_stride_a + line_stride_a*3);
|
||||
#ifdef VECTOR_SUB_GROUP_BROADCAT
|
||||
dequantize_block_acc_bcast_8_hi(total_sum_a, as_ushort8(reg_a_l_a), as_uchar8(reg_a_h_a), cs_a, reg_b);
|
||||
#else
|
||||
dequantize_block_acc_bcast_1_hi(total_sum_a, as_ushort8(reg_a_l_a), as_uchar8(reg_a_h_a), cs_a, reg_b);
|
||||
#endif
|
||||
|
||||
reg_a_l_a.s0 = RD_QL(src0_ql, gid_a + k*block_stride_a + line_stride_a*4);
|
||||
reg_a_l_a.s1 = RD_QL(src0_ql, gid_a + k*block_stride_a + line_stride_a*5);
|
||||
reg_a_l_a.s2 = RD_QL(src0_ql, gid_a + k*block_stride_a + line_stride_a*6);
|
||||
reg_a_l_a.s3 = RD_QL(src0_ql, gid_a + k*block_stride_a + line_stride_a*7);
|
||||
reg_a_h_a.s0 = RD_QH(src0_qh, gid_a + k*block_stride_a + line_stride_a*4);
|
||||
reg_a_h_a.s1 = RD_QH(src0_qh, gid_a + k*block_stride_a + line_stride_a*5);
|
||||
reg_a_h_a.s2 = RD_QH(src0_qh, gid_a + k*block_stride_a + line_stride_a*6);
|
||||
reg_a_h_a.s3 = RD_QH(src0_qh, gid_a + k*block_stride_a + line_stride_a*7);
|
||||
#ifdef VECTOR_SUB_GROUP_BROADCAT
|
||||
dequantize_block_acc_bcast_8_lo(total_sum_a, as_ushort8(reg_a_l_a), as_uchar8(reg_a_h_a), cs_a, reg_b);
|
||||
#else
|
||||
dequantize_block_acc_bcast_1_lo(total_sum_a, as_ushort8(reg_a_l_a), as_uchar8(reg_a_h_a), cs_a, reg_b);
|
||||
#endif
|
||||
}
|
||||
|
||||
{
|
||||
reg_a_l_b.s0 = RD_QL(src0_ql, gid_b + k*block_stride_a + line_stride_a*0);
|
||||
reg_a_l_b.s1 = RD_QL(src0_ql, gid_b + k*block_stride_a + line_stride_a*1);
|
||||
reg_a_l_b.s2 = RD_QL(src0_ql, gid_b + k*block_stride_a + line_stride_a*2);
|
||||
reg_a_l_b.s3 = RD_QL(src0_ql, gid_b + k*block_stride_a + line_stride_a*3);
|
||||
reg_a_h_b.s0 = RD_QH(src0_qh, gid_b + k*block_stride_a + line_stride_a*0);
|
||||
reg_a_h_b.s1 = RD_QH(src0_qh, gid_b + k*block_stride_a + line_stride_a*1);
|
||||
reg_a_h_b.s2 = RD_QH(src0_qh, gid_b + k*block_stride_a + line_stride_a*2);
|
||||
reg_a_h_b.s3 = RD_QH(src0_qh, gid_b + k*block_stride_a + line_stride_a*3);
|
||||
#ifdef VECTOR_SUB_GROUP_BROADCAT
|
||||
dequantize_block_acc_bcast_8_hi(total_sum_b, as_ushort8(reg_a_l_b), as_uchar8(reg_a_h_b), cs_b, reg_b);
|
||||
#else
|
||||
dequantize_block_acc_bcast_1_hi(total_sum_b, as_ushort8(reg_a_l_b), as_uchar8(reg_a_h_b), cs_b, reg_b);
|
||||
#endif
|
||||
|
||||
reg_a_l_b.s0 = RD_QL(src0_ql, gid_b + k*block_stride_a + line_stride_a*4);
|
||||
reg_a_l_b.s1 = RD_QL(src0_ql, gid_b + k*block_stride_a + line_stride_a*5);
|
||||
reg_a_l_b.s2 = RD_QL(src0_ql, gid_b + k*block_stride_a + line_stride_a*6);
|
||||
reg_a_l_b.s3 = RD_QL(src0_ql, gid_b + k*block_stride_a + line_stride_a*7);
|
||||
reg_a_h_b.s0 = RD_QH(src0_qh, gid_b + k*block_stride_a + line_stride_a*4);
|
||||
reg_a_h_b.s1 = RD_QH(src0_qh, gid_b + k*block_stride_a + line_stride_a*5);
|
||||
reg_a_h_b.s2 = RD_QH(src0_qh, gid_b + k*block_stride_a + line_stride_a*6);
|
||||
reg_a_h_b.s3 = RD_QH(src0_qh, gid_b + k*block_stride_a + line_stride_a*7);
|
||||
#ifdef VECTOR_SUB_GROUP_BROADCAT
|
||||
dequantize_block_acc_bcast_8_lo(total_sum_b, as_ushort8(reg_a_l_b), as_uchar8(reg_a_h_b), cs_b, reg_b);
|
||||
#else
|
||||
dequantize_block_acc_bcast_1_lo(total_sum_b, as_ushort8(reg_a_l_b), as_uchar8(reg_a_h_b), cs_b, reg_b);
|
||||
#endif
|
||||
}
|
||||
}
|
||||
|
||||
// Cross-subgroup reduce. Same shape as the 2-output kernel but with the
|
||||
// pair-a and pair-b accumulators concatenated into a single float4.
|
||||
local float4 reduce_lm[SUBGROUP_SIZE * 3];
|
||||
float4 acc = (float4)(total_sum_a.s0, total_sum_a.s1, total_sum_b.s0, total_sum_b.s1);
|
||||
if (grp == 1) { reduce_lm[SUBGROUP_SIZE*0 + slid] = acc; }
|
||||
if (grp == 2) { reduce_lm[SUBGROUP_SIZE*1 + slid] = acc; }
|
||||
if (grp == 3) { reduce_lm[SUBGROUP_SIZE*2 + slid] = acc; }
|
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
if (grp == 0) {
|
||||
acc += reduce_lm[SUBGROUP_SIZE*0 + slid];
|
||||
acc += reduce_lm[SUBGROUP_SIZE*1 + slid];
|
||||
acc += reduce_lm[SUBGROUP_SIZE*2 + slid];
|
||||
dst = (global float*)((global char*)dst + offsetd);
|
||||
// The dispatch rounds ne01/4 up to the subgroup width, so the tail
|
||||
// quads past the last row must not store (they wrote 128 rows past
|
||||
// dst on every ne01 % 256 == 128 vocab, e.g. 151936).
|
||||
if (gid * 4 + 3 < (uint)ne01) {
|
||||
vstore4(acc, 0, &(dst[gid * 4]));
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,196 @@
|
||||
// Tiled-wide q6_K GEMV for the long-vocab lm_head/embed (decode path).
|
||||
//
|
||||
// Pairs with kernel_convert_block_q6_k_tiled_ns (cvt.cl): the weights are laid
|
||||
// out CANONICALLY (6-bit code in element order e in [0,256)) and TILED by 64
|
||||
// output rows so the 64-thread lane group coalesces every weight load. Both the
|
||||
// pack (convert) and the unpack (here) are owned by us — correct by construction
|
||||
// against the reference ggml q6_K dequant, no bit-interleave reverse-engineering.
|
||||
//
|
||||
// One work-item produces one output row. A work-group is {64 lanes, 4 subgroups}:
|
||||
// the 64 lanes cover the 64 rows of one tile (coalesced reads), the 4 subgroups
|
||||
// split the K-blocks and reduce through __local at the end.
|
||||
//
|
||||
// Weights are read from __global (coalesced) rather than image1d_buffer: the
|
||||
// lm_head is read once per token with no reuse, and the Adreno texture cache
|
||||
// caps such a streaming read well below the coalesced-global rate
|
||||
// (see opencl_q6k_gemv_o4_shipped / x2-90 roofline notes).
|
||||
|
||||
#pragma OPENCL EXTENSION cl_khr_fp16 : 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 NSUBGROUPS 4
|
||||
#define TILE_ROWS 64
|
||||
|
||||
#if defined(ADRENO_GPU)
|
||||
REQD_SUBGROUP_SIZE_64
|
||||
#endif
|
||||
kernel void kernel_gemv_noshuffle_q6_K_f32_tiled(
|
||||
__global uint4 * src0_ql, // tiled: 8 uint4 granules / superblock
|
||||
__global uint4 * src0_qh, // tiled: 4 uint4 granules / superblock
|
||||
__global char * src0_s, // tiled: 16 chars / superblock
|
||||
__global half * src0_d, // tiled: 1 half / superblock
|
||||
read_only image1d_buffer_t src1, // activation (RGBA f32)
|
||||
global float * dst,
|
||||
ulong offsetd,
|
||||
int ne00,
|
||||
int ne01
|
||||
) {
|
||||
int grp = get_local_id(1); // subgroup index 0..3 (splits K)
|
||||
int row = get_global_id(0); // output row along ne01
|
||||
int rt = row / TILE_ROWS;
|
||||
int rit = row % TILE_ROWS;
|
||||
|
||||
int nb = ne00 / 256; // superblocks per row
|
||||
|
||||
float acc = 0.0f;
|
||||
|
||||
for (int sb = grp; sb < nb; sb += NSUBGROUPS) {
|
||||
int tile_blk = rt * nb + sb; // ne02 == 1 for lm_head/embed
|
||||
|
||||
// d + 16 scales for this (row, superblock)
|
||||
float dval = (float)src0_d[tile_blk * TILE_ROWS + rit];
|
||||
__global char * sc = src0_s + (tile_blk * TILE_ROWS + rit) * 16;
|
||||
|
||||
// 32 ql-uints (8 codes/uint) + 16 qh-uints (16 codes/uint)
|
||||
uint ql[32];
|
||||
uint qh[16];
|
||||
#pragma unroll
|
||||
for (int g = 0; g < 8; ++g) {
|
||||
uint4 v = src0_ql[(tile_blk * 8 + g) * TILE_ROWS + rit];
|
||||
ql[g*4+0] = v.x; ql[g*4+1] = v.y; ql[g*4+2] = v.z; ql[g*4+3] = v.w;
|
||||
}
|
||||
#pragma unroll
|
||||
for (int g = 0; g < 4; ++g) {
|
||||
uint4 v = src0_qh[(tile_blk * 4 + g) * TILE_ROWS + rit];
|
||||
qh[g*4+0] = v.x; qh[g*4+1] = v.y; qh[g*4+2] = v.z; qh[g*4+3] = v.w;
|
||||
}
|
||||
|
||||
// dequant 256 codes in canonical e-order, MAC with activation.
|
||||
int act_base = sb * 64; // activation float4 pixel base (256/4)
|
||||
#pragma unroll
|
||||
for (int e4 = 0; e4 < 64; ++e4) {
|
||||
float4 a = read_imagef(src1, act_base + e4);
|
||||
#pragma unroll
|
||||
for (int t = 0; t < 4; ++t) {
|
||||
int e = e4 * 4 + t;
|
||||
uint low4 = (ql[e >> 3] >> ((e & 7) * 4)) & 0xF;
|
||||
uint hi2 = (qh[e >> 4] >> ((e & 15) * 2)) & 0x3;
|
||||
int code = (int)(low4 | (hi2 << 4)) - 32;
|
||||
int sidx = ((e >> 7) << 3) + (((e >> 5) & 3) << 1) + ((e >> 4) & 1);
|
||||
float scale = (float)sc[sidx] * dval;
|
||||
float av = (t == 0) ? a.x : (t == 1) ? a.y : (t == 2) ? a.z : a.w;
|
||||
acc += (float)code * scale * av;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
// reduce across the NSUBGROUPS subgroups (same rit, different K-subset)
|
||||
local float reduce_lm[NSUBGROUPS * TILE_ROWS];
|
||||
reduce_lm[grp * TILE_ROWS + rit] = acc;
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
if (grp == 0) {
|
||||
float total = reduce_lm[0 * TILE_ROWS + rit]
|
||||
+ reduce_lm[1 * TILE_ROWS + rit]
|
||||
+ reduce_lm[2 * TILE_ROWS + rit]
|
||||
+ reduce_lm[3 * TILE_ROWS + rit];
|
||||
dst = (global float*)((global char*)dst + offsetd);
|
||||
dst[row] = total;
|
||||
}
|
||||
}
|
||||
|
||||
// Multi-column (N=3) variant of the tiled q6_K decode GEMV, for the speculative/
|
||||
// MTP VERIFY lm_head/embed (ne1=3 = 2 drafts + 1 bonus). Identical tiled weight
|
||||
// layout + unpack as the ne1=1 kernel above; each WI computes 3 output columns,
|
||||
// streaming the (large) lm_head weight ONCE per superblock and reusing it across
|
||||
// the 3 verify activation columns (dequant once per code, MAC into 3 accs). This
|
||||
// is the lm_head analogue of the per-layer mc3 GEMV; the multiply order matches
|
||||
// the ne1=1 kernel, so each column is byte-identical to a standalone tiled GEMV.
|
||||
#if defined(ADRENO_GPU)
|
||||
REQD_SUBGROUP_SIZE_64
|
||||
#endif
|
||||
kernel void kernel_gemv_noshuffle_q6_K_f32_tiled_mc3(
|
||||
__global uint4 * src0_ql,
|
||||
__global uint4 * src0_qh,
|
||||
__global char * src0_s,
|
||||
__global half * src0_d,
|
||||
read_only image1d_buffer_t src1,
|
||||
global float * dst,
|
||||
ulong offsetd,
|
||||
int ne00,
|
||||
int ne01
|
||||
) {
|
||||
int grp = get_local_id(1);
|
||||
int row = get_global_id(0);
|
||||
int rt = row / TILE_ROWS;
|
||||
int rit = row % TILE_ROWS;
|
||||
|
||||
int nb = ne00 / 256;
|
||||
int col_stride = ne00 / 4; // activation float4 pixels per column
|
||||
|
||||
float acc0 = 0.0f, acc1 = 0.0f, acc2 = 0.0f;
|
||||
|
||||
for (int sb = grp; sb < nb; sb += NSUBGROUPS) {
|
||||
int tile_blk = rt * nb + sb;
|
||||
|
||||
float dval = (float)src0_d[tile_blk * TILE_ROWS + rit];
|
||||
__global char * sc = src0_s + (tile_blk * TILE_ROWS + rit) * 16;
|
||||
|
||||
uint ql[32];
|
||||
uint qh[16];
|
||||
#pragma unroll
|
||||
for (int g = 0; g < 8; ++g) {
|
||||
uint4 v = src0_ql[(tile_blk * 8 + g) * TILE_ROWS + rit];
|
||||
ql[g*4+0] = v.x; ql[g*4+1] = v.y; ql[g*4+2] = v.z; ql[g*4+3] = v.w;
|
||||
}
|
||||
#pragma unroll
|
||||
for (int g = 0; g < 4; ++g) {
|
||||
uint4 v = src0_qh[(tile_blk * 4 + g) * TILE_ROWS + rit];
|
||||
qh[g*4+0] = v.x; qh[g*4+1] = v.y; qh[g*4+2] = v.z; qh[g*4+3] = v.w;
|
||||
}
|
||||
|
||||
int act_base = sb * 64;
|
||||
#pragma unroll
|
||||
for (int e4 = 0; e4 < 64; ++e4) {
|
||||
float4 a0 = read_imagef(src1, 0*col_stride + act_base + e4);
|
||||
float4 a1 = read_imagef(src1, 1*col_stride + act_base + e4);
|
||||
float4 a2 = read_imagef(src1, 2*col_stride + act_base + e4);
|
||||
#pragma unroll
|
||||
for (int t = 0; t < 4; ++t) {
|
||||
int e = e4 * 4 + t;
|
||||
uint low4 = (ql[e >> 3] >> ((e & 7) * 4)) & 0xF;
|
||||
uint hi2 = (qh[e >> 4] >> ((e & 15) * 2)) & 0x3;
|
||||
int code = (int)(low4 | (hi2 << 4)) - 32;
|
||||
int sidx = ((e >> 7) << 3) + (((e >> 5) & 3) << 1) + ((e >> 4) & 1);
|
||||
float w = (float)code * ((float)sc[sidx] * dval); // dequant+scale once
|
||||
float av0 = (t == 0) ? a0.x : (t == 1) ? a0.y : (t == 2) ? a0.z : a0.w;
|
||||
float av1 = (t == 0) ? a1.x : (t == 1) ? a1.y : (t == 2) ? a1.z : a1.w;
|
||||
float av2 = (t == 0) ? a2.x : (t == 1) ? a2.y : (t == 2) ? a2.z : a2.w;
|
||||
acc0 += w * av0;
|
||||
acc1 += w * av1;
|
||||
acc2 += w * av2;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
local float4 reduce_lm[NSUBGROUPS * TILE_ROWS];
|
||||
reduce_lm[grp * TILE_ROWS + rit] = (float4)(acc0, acc1, acc2, 0.0f);
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
if (grp == 0) {
|
||||
float4 total = reduce_lm[0 * TILE_ROWS + rit]
|
||||
+ reduce_lm[1 * TILE_ROWS + rit]
|
||||
+ reduce_lm[2 * TILE_ROWS + rit]
|
||||
+ reduce_lm[3 * TILE_ROWS + rit];
|
||||
dst = (global float*)((global char*)dst + offsetd);
|
||||
// dst column-major [ne01 rows x 3 cols]: (row, col) at col*ne01 + row
|
||||
dst[0*ne01 + row] = total.x;
|
||||
dst[1*ne01 + row] = total.y;
|
||||
dst[2*ne01 + row] = total.z;
|
||||
}
|
||||
}
|
||||
@@ -118,6 +118,87 @@
|
||||
elem = (char)((bits8.s7 & 0xFF000000) >> 24); \
|
||||
total_sums += convert_int(elem) * scale * shared_y; \
|
||||
|
||||
// ============================================================================
|
||||
// Split-K variant for small-M decode GEMVs.
|
||||
// ----------------------------------------------------------------------------
|
||||
// The base kernel below puts one output row per lane and splits K only across
|
||||
// the N_SIMDGROUP subgroups of a single workgroup, so M=512 yields M/64 = 8
|
||||
// workgroups -- half the compute units on a 16-CU X2 sit idle, and the kernel
|
||||
// measures ~48 GB/s against the ~122 GB/s the larger projections reach in the
|
||||
// same graph. Here each (kslice, subgroup) pair reduces a disjoint set of
|
||||
// K-blocks into partial[kslice * M + row]; kernel_gemv_splitk_reduce_f32 (in
|
||||
// gemv_noshuffle_q4_k_f32.cl) sums the slices. Same operand order within a
|
||||
// slice as the base kernel; only the cross-slice grouping differs.
|
||||
//
|
||||
// Placed BEFORE the base kernel deliberately: on A6X no kernel may be defined
|
||||
// after one that uses a subgroup builtin, or it silently miscompiles.
|
||||
// ============================================================================
|
||||
#ifdef ADRENO_GPU
|
||||
REQD_SUBGROUP_SIZE_64
|
||||
#endif
|
||||
__kernel void kernel_gemv_noshuffle_q8_0_f32_splitk(
|
||||
__read_only image1d_buffer_t src0_q, // quantized A (weights)
|
||||
global half * src0_d, // A scales
|
||||
__read_only image1d_buffer_t src1, // B (activations)
|
||||
global float * partial, // [ksplit * M], slice-major
|
||||
int ne00, // K
|
||||
int ne01) // M
|
||||
{
|
||||
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;
|
||||
uint BLOCK_STRIDE_A = 8 * M; // physical, independent of the K-split
|
||||
|
||||
__private uint8 regA;
|
||||
__private half regS;
|
||||
__private float8 regB;
|
||||
__private float totalSum = (float)(0.0f);
|
||||
|
||||
#pragma unroll 1
|
||||
for (uint k = kslice * nsg + groupId; k < (K / QK8_0); k += ksplit * nsg) {
|
||||
regS = src0_d[gid + k * LINE_STRIDE_A];
|
||||
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;
|
||||
regA.s4 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 4)).x;
|
||||
regA.s5 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 5)).x;
|
||||
regA.s6 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 6)).x;
|
||||
regA.s7 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 7)).x;
|
||||
|
||||
dequantizeBlockAccum_ns_sgbroadcast_1(totalSum, regA, convert_float(regS), regB);
|
||||
}
|
||||
|
||||
// Intra-workgroup reduce across this K-slice's subgroups. Sized for
|
||||
// nsg <= 8; the host never dispatches more.
|
||||
__local float reduceLM[SIMDGROUP_WIDTH * 7];
|
||||
if (groupId > 0) {
|
||||
reduceLM[SIMDGROUP_WIDTH * (groupId - 1) + slid] = totalSum;
|
||||
}
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
if (groupId == 0) {
|
||||
for (uint i = 0; i < nsg - 1; ++i) {
|
||||
totalSum += reduceLM[SIMDGROUP_WIDTH * i + slid];
|
||||
}
|
||||
// x-grid is padded to CEIL_DIV(M,wave)*wave; guard the tail rows.
|
||||
if (gid < M) {
|
||||
partial[kslice * M + gid] = totalSum;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
#ifdef ADRENO_GPU
|
||||
REQD_SUBGROUP_SIZE_64
|
||||
#endif
|
||||
|
||||
@@ -145,3 +145,52 @@ kernel void kernel_mul_mm_f32_f32_l4_lm(
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
// Multi-column f32 GEMV for the small-N (spec/MTP verify) batch. The tiled GEMM
|
||||
// above always computes a full BM x BN = 64 x 64 output tile, so at ne11=3 with a
|
||||
// skinny weight (e.g. GDN ssm_alpha/ssm_beta, M=32) it launches ONE under-occupied
|
||||
// workgroup at ~2.3% tile utilization. This kernel assigns one 64-thread workgroup
|
||||
// per output element (m,n): the 64 threads split the K reduction (float4) and
|
||||
// tree-reduce in __local (no subgroup ops -> portable). ne01*ne11 workgroups.
|
||||
// Weight row is re-read per column (N small -> negligible). Summation order differs
|
||||
// from the tiled GEMM (lane-strided + tree) -> f32-exact-ish, not bit-identical.
|
||||
kernel void kernel_gemv_f32_f32_mc(
|
||||
global float * src0, ulong offset0, // weight: row m at m*stride_a (elements)
|
||||
global float * src1, ulong offset1, // activations: col n at n*stride_b
|
||||
global float * dst, ulong offsetd, // dst [M x N] col-major: (m,n) at n*stride_d+m
|
||||
int ne00, // K
|
||||
int ne01, // M
|
||||
int ne11, // N
|
||||
int stride_a, // weight row stride (elements) = K
|
||||
int stride_b, // activation col stride (elements) = K
|
||||
int stride_d) // dst column stride (elements) = M
|
||||
{
|
||||
src0 = (global float*)((global char*)src0 + offset0);
|
||||
src1 = (global float*)((global char*)src1 + offset1);
|
||||
dst = (global float*)((global char*)dst + offsetd);
|
||||
|
||||
uint lane = get_local_id(0); // 0..63
|
||||
uint out = get_global_id(1); // 0 .. ne01*ne11 - 1
|
||||
uint m = out % (uint)ne01;
|
||||
uint n = out / (uint)ne01;
|
||||
|
||||
global float4 * wrow = (global float4*)(src0 + (ulong)m * (uint)stride_a);
|
||||
global float4 * xcol = (global float4*)(src1 + (ulong)n * (uint)stride_b);
|
||||
uint k4 = (uint)ne00 >> 2;
|
||||
|
||||
float acc = 0.0f;
|
||||
for (uint k = lane; k < k4; k += 64) {
|
||||
float4 w = wrow[k];
|
||||
float4 x = xcol[k];
|
||||
acc += w.s0*x.s0 + w.s1*x.s1 + w.s2*x.s2 + w.s3*x.s3;
|
||||
}
|
||||
|
||||
local float red[64];
|
||||
red[lane] = acc;
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
for (uint s = 32; s > 0; s >>= 1) {
|
||||
if (lane < s) red[lane] += red[lane + s];
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
}
|
||||
if (lane == 0) dst[(ulong)n * (uint)stride_d + m] = red[0];
|
||||
}
|
||||
|
||||
@@ -0,0 +1,306 @@
|
||||
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
|
||||
|
||||
#ifdef cl_intel_subgroups
|
||||
#pragma OPENCL EXTENSION cl_intel_subgroups : enable
|
||||
#else
|
||||
#pragma OPENCL EXTENSION cl_khr_subgroups : enable
|
||||
#endif
|
||||
|
||||
#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
|
||||
|
||||
// Multi-row f16xf32 GEMV for the DECODE path (single token, ne11*ne12 small).
|
||||
// The legacy kernel_mul_mat_f16_f32_1row runs ONE 64-lane subgroup per workgroup =
|
||||
// one output row per WG, which caps memory-level parallelism at roughly half of
|
||||
// LPDDR5x peak. This variant packs MROW subgroups per workgroup, each
|
||||
// computing a distinct output row, so a WG keeps 64*MROW loads in flight. The
|
||||
// activation column y (shared by every output row) is staged into __local ONCE per
|
||||
// WG and reused across the MROW rows, cutting redundant activation reads. Used for
|
||||
// the f16 attention projections (Q/K/V/O) and lm_head, which dominate decode.
|
||||
// Numerically equivalent to _1row (same f16->f32 widening, same float4 partial sums,
|
||||
// same subgroup-reduce order), so byte-identical to the per-op path.
|
||||
|
||||
#define MROW 16
|
||||
|
||||
#ifdef ADRENO_GPU
|
||||
REQD_SUBGROUP_SIZE_64
|
||||
#endif
|
||||
kernel void kernel_mul_mat_f16_f32_mrow(
|
||||
global char * src0,
|
||||
ulong offset0,
|
||||
global char * src1,
|
||||
ulong offset1,
|
||||
global float * dst,
|
||||
ulong offsetd,
|
||||
int ne00,
|
||||
int ne01,
|
||||
int ne02,
|
||||
ulong nb00,
|
||||
ulong nb01,
|
||||
ulong nb02,
|
||||
ulong nb03,
|
||||
int ne10,
|
||||
int ne11,
|
||||
int ne12,
|
||||
ulong nb10,
|
||||
ulong nb11,
|
||||
ulong nb12,
|
||||
ulong nb13,
|
||||
int ne0,
|
||||
int ne1,
|
||||
int r2,
|
||||
int r3,
|
||||
__local float * ysh
|
||||
) {
|
||||
src0 = (global char*)((global char*)src0 + offset0);
|
||||
src1 = (global char*)((global char*)src1 + offset1);
|
||||
dst = (global float*)((global char*)dst + offsetd);
|
||||
|
||||
int r0 = get_group_id(0) * MROW + get_local_id(1); // output row
|
||||
int r1 = get_group_id(1); // token (ne11)
|
||||
int im = get_group_id(2);
|
||||
int lid = get_sub_group_local_id(); // 0..63
|
||||
int nsg = get_local_size(1); // == MROW
|
||||
|
||||
int i12 = im % ne12;
|
||||
int i13 = im / ne12;
|
||||
|
||||
ulong offset_src1 = r1*nb11 + (i12)*nb12 + (i13)*nb13;
|
||||
global float * y = (global float *) (src1 + offset_src1);
|
||||
|
||||
// Cooperatively stage the activation column (ne00 floats) into __local once per
|
||||
// WG and reuse across the MROW rows. Staging is the actual win here: dropping it
|
||||
// (each subgroup re-reading y from global) regresses below the 1-row kernel.
|
||||
for (int i = get_local_id(1)*get_sub_group_size() + lid; i < ne00; i += nsg*get_sub_group_size()) {
|
||||
ysh[i] = y[i];
|
||||
}
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
if (r0 >= ne01) {
|
||||
return;
|
||||
}
|
||||
|
||||
ulong offset_src0 = r0*nb01 + (i12/r2)*nb02 + (i13/r3)*nb03;
|
||||
global half * x = (global half *) (src0 + offset_src0);
|
||||
|
||||
// The vector path below casts the row pointer to half4, which must be 8-byte aligned.
|
||||
// A row address is r0*nb01 + ..., and a permuted or strided src0 leaves nb01/nb02/nb03
|
||||
// unconstrained -- ne00 % 4 == 0 bounds the element count per row, not the byte stride
|
||||
// between rows. Take the vector path only when this work-item's row is actually
|
||||
// aligned; the scalar loop below has no such requirement.
|
||||
const bool row_aligned = (((ulong) x) & 7) == 0;
|
||||
|
||||
float sumf = 0.0f;
|
||||
if (ne00 < 128 || !row_aligned) {
|
||||
for (int i = lid; i < ne00; i += get_sub_group_size()) {
|
||||
sumf += (float) x[i] * ysh[i];
|
||||
}
|
||||
float all_sum = sub_group_reduce_add(sumf);
|
||||
if (lid == 0) {
|
||||
dst[im*ne1*ne0 + r1*ne0 + r0] = all_sum;
|
||||
}
|
||||
} else {
|
||||
global half4 * x4 = (global half4 *) x;
|
||||
__local float4 * ysh4 = (__local float4 *) ysh;
|
||||
for (int i = lid; i < ne00/4; i += get_sub_group_size()) {
|
||||
float4 yv = ysh4[i];
|
||||
sumf += (float) x4[i].s0 * yv.s0;
|
||||
sumf += (float) x4[i].s1 * yv.s1;
|
||||
sumf += (float) x4[i].s2 * yv.s2;
|
||||
sumf += (float) x4[i].s3 * yv.s3;
|
||||
}
|
||||
float all_sum = sub_group_reduce_add(sumf);
|
||||
if (lid == 0) {
|
||||
for (int i = 4*(ne00/4); i < ne00; ++i) {
|
||||
all_sum += (float) x[i] * ysh[i];
|
||||
}
|
||||
dst[im*ne1*ne0 + r1*ne0 + r0] = all_sum;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
// Register-blocked variant: each 64-lane subgroup accumulates RPT consecutive
|
||||
// output rows instead of one. The staged activation is reused across all RPT rows,
|
||||
// and each lane keeps RPT independent weight loads in flight per column step ->
|
||||
// more memory-level parallelism on the streaming f16 weight read (the BW limiter),
|
||||
// plus RPT fewer staging barriers per output row. Per-row reduction order is
|
||||
// identical to _mrow, so byte-identical to the per-op path. Dispatch guarantees
|
||||
// ne00 >= 128 and ne00 % 4 == 0, so only the half4 path is needed (no tail).
|
||||
#define MROW_RB_BODY(RPT) \
|
||||
src0 = (global char*)((global char*)src0 + offset0); \
|
||||
src1 = (global char*)((global char*)src1 + offset1); \
|
||||
dst = (global float*)((global char*)dst + offsetd); \
|
||||
int r0b = (get_group_id(0) * get_local_size(1) + get_local_id(1)) * (RPT); \
|
||||
int r1 = get_group_id(1); \
|
||||
int im = get_group_id(2); \
|
||||
int lid = get_sub_group_local_id(); \
|
||||
int nsg = get_local_size(1); \
|
||||
int i12 = im % ne12; \
|
||||
int i13 = im / ne12; \
|
||||
ulong off_y = r1*nb11 + i12*nb12 + i13*nb13; \
|
||||
global float * y = (global float *) (src1 + off_y); \
|
||||
for (int i = get_local_id(1)*get_sub_group_size() + lid; i < ne00; \
|
||||
i += nsg*get_sub_group_size()) { \
|
||||
ysh[i] = y[i]; \
|
||||
} \
|
||||
barrier(CLK_LOCAL_MEM_FENCE); \
|
||||
__local float4 * ysh4 = (__local float4 *) ysh; \
|
||||
global half4 * xr[RPT]; \
|
||||
_Pragma("unroll") \
|
||||
for (int rr = 0; rr < (RPT); ++rr) { \
|
||||
int row = r0b + rr; \
|
||||
if (row > ne01 - 1) row = ne01 - 1; \
|
||||
xr[rr] = (global half4 *) (src0 + (ulong)row*nb01 + (i12/r2)*nb02 + (i13/r3)*nb03); \
|
||||
} \
|
||||
float sumf[RPT]; \
|
||||
_Pragma("unroll") \
|
||||
for (int rr = 0; rr < (RPT); ++rr) sumf[rr] = 0.0f; \
|
||||
for (int i = lid; i < ne00/4; i += get_sub_group_size()) { \
|
||||
float4 yv = ysh4[i]; \
|
||||
_Pragma("unroll") \
|
||||
for (int rr = 0; rr < (RPT); ++rr) { \
|
||||
half4 xv = xr[rr][i]; \
|
||||
sumf[rr] += (float) xv.s0 * yv.s0 + (float) xv.s1 * yv.s1 \
|
||||
+ (float) xv.s2 * yv.s2 + (float) xv.s3 * yv.s3; \
|
||||
} \
|
||||
} \
|
||||
_Pragma("unroll") \
|
||||
for (int rr = 0; rr < (RPT); ++rr) { \
|
||||
float s = sub_group_reduce_add(sumf[rr]); \
|
||||
int row = r0b + rr; \
|
||||
if (lid == 0 && row < ne01) { \
|
||||
dst[im*ne1*ne0 + r1*ne0 + row] = s; \
|
||||
} \
|
||||
}
|
||||
|
||||
// half8 (128-bit) load variant: Adreno's load/store unit issues 128-bit
|
||||
// transactions, so half4 (64-bit) loads may leave the load path half-idle. This
|
||||
// processes 8 weight elements per lane per step via half8. Accumulation groups
|
||||
// elements in 8s rather than 4s, so it is NOT bit-identical to _1row (float add is
|
||||
// non-associative) -- experimental BW probe, gate on ne00 % 8 == 0.
|
||||
#define MROW_H8_BODY(RPT) \
|
||||
src0 = (global char*)((global char*)src0 + offset0); \
|
||||
src1 = (global char*)((global char*)src1 + offset1); \
|
||||
dst = (global float*)((global char*)dst + offsetd); \
|
||||
int r0b = (get_group_id(0) * get_local_size(1) + get_local_id(1)) * (RPT); \
|
||||
int r1 = get_group_id(1); \
|
||||
int im = get_group_id(2); \
|
||||
int lid = get_sub_group_local_id(); \
|
||||
int nsg = get_local_size(1); \
|
||||
int i12 = im % ne12; \
|
||||
int i13 = im / ne12; \
|
||||
ulong off_y = r1*nb11 + i12*nb12 + i13*nb13; \
|
||||
global float * y = (global float *) (src1 + off_y); \
|
||||
for (int i = get_local_id(1)*get_sub_group_size() + lid; i < ne00; \
|
||||
i += nsg*get_sub_group_size()) { \
|
||||
ysh[i] = y[i]; \
|
||||
} \
|
||||
barrier(CLK_LOCAL_MEM_FENCE); \
|
||||
__local float4 * ysh4 = (__local float4 *) ysh; \
|
||||
global half8 * xr[RPT]; \
|
||||
_Pragma("unroll") \
|
||||
for (int rr = 0; rr < (RPT); ++rr) { \
|
||||
int row = r0b + rr; \
|
||||
if (row > ne01 - 1) row = ne01 - 1; \
|
||||
xr[rr] = (global half8 *) (src0 + (ulong)row*nb01 + (i12/r2)*nb02 + (i13/r3)*nb03); \
|
||||
} \
|
||||
float sumf[RPT]; \
|
||||
_Pragma("unroll") \
|
||||
for (int rr = 0; rr < (RPT); ++rr) sumf[rr] = 0.0f; \
|
||||
for (int i = lid; i < ne00/8; i += get_sub_group_size()) { \
|
||||
float4 y0 = ysh4[2*i]; \
|
||||
float4 y1 = ysh4[2*i + 1]; \
|
||||
_Pragma("unroll") \
|
||||
for (int rr = 0; rr < (RPT); ++rr) { \
|
||||
half8 xv = xr[rr][i]; \
|
||||
sumf[rr] += (float) xv.s0 * y0.s0 + (float) xv.s1 * y0.s1 \
|
||||
+ (float) xv.s2 * y0.s2 + (float) xv.s3 * y0.s3 \
|
||||
+ (float) xv.s4 * y1.s0 + (float) xv.s5 * y1.s1 \
|
||||
+ (float) xv.s6 * y1.s2 + (float) xv.s7 * y1.s3; \
|
||||
} \
|
||||
} \
|
||||
_Pragma("unroll") \
|
||||
for (int rr = 0; rr < (RPT); ++rr) { \
|
||||
float s = sub_group_reduce_add(sumf[rr]); \
|
||||
int row = r0b + rr; \
|
||||
if (lid == 0 && row < ne01) { \
|
||||
dst[im*ne1*ne0 + r1*ne0 + row] = s; \
|
||||
} \
|
||||
}
|
||||
|
||||
#ifdef ADRENO_GPU
|
||||
REQD_SUBGROUP_SIZE_64
|
||||
#endif
|
||||
kernel void kernel_mul_mat_f16_f32_mrow_h8(
|
||||
global char * src0, ulong offset0,
|
||||
global char * src1, ulong offset1,
|
||||
global float * dst, ulong offsetd,
|
||||
int ne00, int ne01, int ne02,
|
||||
ulong nb00, ulong nb01, ulong nb02, ulong nb03,
|
||||
int ne10, int ne11, int ne12,
|
||||
ulong nb10, ulong nb11, ulong nb12, ulong nb13,
|
||||
int ne0, int ne1, int r2, int r3,
|
||||
__local float * ysh
|
||||
) {
|
||||
MROW_H8_BODY(1)
|
||||
}
|
||||
|
||||
#ifdef ADRENO_GPU
|
||||
REQD_SUBGROUP_SIZE_64
|
||||
#endif
|
||||
kernel void kernel_mul_mat_f16_f32_mrow_h8r2(
|
||||
global char * src0, ulong offset0,
|
||||
global char * src1, ulong offset1,
|
||||
global float * dst, ulong offsetd,
|
||||
int ne00, int ne01, int ne02,
|
||||
ulong nb00, ulong nb01, ulong nb02, ulong nb03,
|
||||
int ne10, int ne11, int ne12,
|
||||
ulong nb10, ulong nb11, ulong nb12, ulong nb13,
|
||||
int ne0, int ne1, int r2, int r3,
|
||||
__local float * ysh
|
||||
) {
|
||||
MROW_H8_BODY(2)
|
||||
}
|
||||
|
||||
#ifdef ADRENO_GPU
|
||||
REQD_SUBGROUP_SIZE_64
|
||||
#endif
|
||||
kernel void kernel_mul_mat_f16_f32_mrow_r2(
|
||||
global char * src0, ulong offset0,
|
||||
global char * src1, ulong offset1,
|
||||
global float * dst, ulong offsetd,
|
||||
int ne00, int ne01, int ne02,
|
||||
ulong nb00, ulong nb01, ulong nb02, ulong nb03,
|
||||
int ne10, int ne11, int ne12,
|
||||
ulong nb10, ulong nb11, ulong nb12, ulong nb13,
|
||||
int ne0, int ne1, int r2, int r3,
|
||||
__local float * ysh
|
||||
) {
|
||||
MROW_RB_BODY(2)
|
||||
}
|
||||
|
||||
#ifdef ADRENO_GPU
|
||||
REQD_SUBGROUP_SIZE_64
|
||||
#endif
|
||||
kernel void kernel_mul_mat_f16_f32_mrow_r4(
|
||||
global char * src0, ulong offset0,
|
||||
global char * src1, ulong offset1,
|
||||
global float * dst, ulong offsetd,
|
||||
int ne00, int ne01, int ne02,
|
||||
ulong nb00, ulong nb01, ulong nb02, ulong nb03,
|
||||
int ne10, int ne11, int ne12,
|
||||
ulong nb10, ulong nb11, ulong nb12, ulong nb13,
|
||||
int ne0, int ne1, int r2, int r3,
|
||||
__local float * ysh
|
||||
) {
|
||||
MROW_RB_BODY(4)
|
||||
}
|
||||
@@ -188,3 +188,182 @@ kernel void kernel_rms_norm_mul(
|
||||
y[i00] = (x[i00] * scale) * f[i00%(ne10/4)];
|
||||
}
|
||||
}
|
||||
|
||||
//------------------------------------------------------------------------------
|
||||
// rms_norm + mul (norm weight) + add (residual), fused. Mirrors
|
||||
// kernel_rms_norm_mul with an extra residual operand src2: computes
|
||||
// y = (rmsnorm(x) * w) + g
|
||||
// in one dispatch, removing one kernel launch + one global round-trip per
|
||||
// residual block (the dominant per-layer adjacency on Gemma matformers).
|
||||
//------------------------------------------------------------------------------
|
||||
kernel void kernel_rms_norm_mul_add(
|
||||
global char * src0,
|
||||
ulong offset0,
|
||||
global char * src1,
|
||||
ulong offset1,
|
||||
global char * src2,
|
||||
ulong offset2,
|
||||
global char * dst,
|
||||
ulong offsetd,
|
||||
int ne00,
|
||||
int ne01,
|
||||
int ne02,
|
||||
int ne03,
|
||||
ulong nb01,
|
||||
ulong nb02,
|
||||
ulong nb03,
|
||||
int ne10,
|
||||
int ne11,
|
||||
int ne12,
|
||||
int ne13,
|
||||
ulong nb11,
|
||||
ulong nb12,
|
||||
ulong nb13,
|
||||
int ne20,
|
||||
int ne21,
|
||||
int ne22,
|
||||
int ne23,
|
||||
ulong nb21,
|
||||
ulong nb22,
|
||||
ulong nb23,
|
||||
ulong nb1,
|
||||
ulong nb2,
|
||||
ulong nb3,
|
||||
float eps,
|
||||
local float * sum
|
||||
) {
|
||||
src0 = src0 + offset0;
|
||||
src1 = src1 + offset1;
|
||||
src2 = src2 + offset2;
|
||||
dst = dst + offsetd;
|
||||
|
||||
if (get_sub_group_id() == 0) {
|
||||
sum[get_sub_group_local_id()] = 0.0f;
|
||||
}
|
||||
|
||||
int i03 = get_group_id(2);
|
||||
int i02 = get_group_id(1);
|
||||
int i01 = get_group_id(0);
|
||||
|
||||
global float4 * x = (global float4 *) (src0 + i03*nb03 + i02*nb02 + i01*nb01);
|
||||
global float4 * f = (global float4 *) (src1 + (i03%ne13)*nb13 + (i02%ne12)*nb12 + (i01%ne11)*nb11);
|
||||
global float4 * g = (global float4 *) (src2 + (i03%ne23)*nb23 + (i02%ne22)*nb22 + (i01%ne21)*nb21);
|
||||
|
||||
float sumf = 0;
|
||||
|
||||
for (int i00 = get_local_id(0); i00 < ne00/4; i00 += get_local_size(0)) {
|
||||
sumf += dot(x[i00], x[i00]);
|
||||
}
|
||||
sumf = sub_group_reduce_add(sumf);
|
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
if (get_sub_group_local_id() == 0) {
|
||||
sum[get_sub_group_id()] = sumf;
|
||||
}
|
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
sumf = sum[get_sub_group_local_id()];
|
||||
sumf = sub_group_reduce_add(sumf);
|
||||
|
||||
float mean = sumf / ne00;
|
||||
float scale = 1.0f/sqrt(mean + eps);
|
||||
|
||||
global float4 * y = (global float4 *) (dst + i03*nb3 + i02*nb2 + i01*nb1);
|
||||
for (int i00 = get_local_id(0); i00 < ne00/4; i00 += get_local_size(0)) {
|
||||
y[i00] = (x[i00] * scale) * f[i00%(ne10/4)] + g[i00%(ne20/4)];
|
||||
}
|
||||
}
|
||||
|
||||
//------------------------------------------------------------------------------
|
||||
// rms_norm + mul(norm weight) + add(residual) + mul(scalar scale), fused.
|
||||
// Computes y = ((rmsnorm(x) * w) + g) * s, where s is a broadcast SCALAR (e.g.
|
||||
// Gemma-4 layer_output_scale). Folds the trailing per-layer l_out scale-mul into
|
||||
// the residual-norm kernel: one extra dispatch + global round-trip saved per
|
||||
// layer. src3 points at the single scale value.
|
||||
//------------------------------------------------------------------------------
|
||||
kernel void kernel_rms_norm_mul_add_scale(
|
||||
global char * src0,
|
||||
ulong offset0,
|
||||
global char * src1,
|
||||
ulong offset1,
|
||||
global char * src2,
|
||||
ulong offset2,
|
||||
global char * src3,
|
||||
ulong offset3,
|
||||
global char * dst,
|
||||
ulong offsetd,
|
||||
int ne00,
|
||||
int ne01,
|
||||
int ne02,
|
||||
int ne03,
|
||||
ulong nb01,
|
||||
ulong nb02,
|
||||
ulong nb03,
|
||||
int ne10,
|
||||
int ne11,
|
||||
int ne12,
|
||||
int ne13,
|
||||
ulong nb11,
|
||||
ulong nb12,
|
||||
ulong nb13,
|
||||
int ne20,
|
||||
int ne21,
|
||||
int ne22,
|
||||
int ne23,
|
||||
ulong nb21,
|
||||
ulong nb22,
|
||||
ulong nb23,
|
||||
ulong nb1,
|
||||
ulong nb2,
|
||||
ulong nb3,
|
||||
float eps,
|
||||
local float * sum
|
||||
) {
|
||||
src0 = src0 + offset0;
|
||||
src1 = src1 + offset1;
|
||||
src2 = src2 + offset2;
|
||||
src3 = src3 + offset3;
|
||||
dst = dst + offsetd;
|
||||
|
||||
const float sc = *((global float *) src3);
|
||||
|
||||
if (get_sub_group_id() == 0) {
|
||||
sum[get_sub_group_local_id()] = 0.0f;
|
||||
}
|
||||
|
||||
int i03 = get_group_id(2);
|
||||
int i02 = get_group_id(1);
|
||||
int i01 = get_group_id(0);
|
||||
|
||||
global float4 * x = (global float4 *) (src0 + i03*nb03 + i02*nb02 + i01*nb01);
|
||||
global float4 * f = (global float4 *) (src1 + (i03%ne13)*nb13 + (i02%ne12)*nb12 + (i01%ne11)*nb11);
|
||||
global float4 * g = (global float4 *) (src2 + (i03%ne23)*nb23 + (i02%ne22)*nb22 + (i01%ne21)*nb21);
|
||||
|
||||
float sumf = 0;
|
||||
|
||||
for (int i00 = get_local_id(0); i00 < ne00/4; i00 += get_local_size(0)) {
|
||||
sumf += dot(x[i00], x[i00]);
|
||||
}
|
||||
sumf = sub_group_reduce_add(sumf);
|
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
if (get_sub_group_local_id() == 0) {
|
||||
sum[get_sub_group_id()] = sumf;
|
||||
}
|
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE);
|
||||
|
||||
sumf = sum[get_sub_group_local_id()];
|
||||
sumf = sub_group_reduce_add(sumf);
|
||||
|
||||
float mean = sumf / ne00;
|
||||
float scale = 1.0f/sqrt(mean + eps);
|
||||
|
||||
global float4 * y = (global float4 *) (dst + i03*nb3 + i02*nb2 + i01*nb1);
|
||||
for (int i00 = get_local_id(0); i00 < ne00/4; i00 += get_local_size(0)) {
|
||||
y[i00] = ((x[i00] * scale) * f[i00%(ne10/4)] + g[i00%(ne20/4)]) * sc;
|
||||
}
|
||||
}
|
||||
|
||||
Reference in New Issue
Block a user