From 9166ceaf16599406366cd2d01dc6a2126bbada4d Mon Sep 17 00:00:00 2001 From: lhez Date: Thu, 15 Oct 2026 00:21:18 -0700 Subject: [PATCH] opencl: fix OOB tail access for adreno noshuffle gemv kernels (#30421) Co-authored-by: Hongqiang Wang --- .../kernels/gemv_noshuffle_iq4_nl_f32.cl | 20 ++-- .../kernels/gemv_noshuffle_q1_0_f32.cl | 12 ++- .../kernels/gemv_noshuffle_q4_0_f32.cl | 50 +++++----- .../kernels/gemv_noshuffle_q4_0_f32_spec.cl | 20 ++-- .../kernels/gemv_noshuffle_q4_1_f32.cl | 54 ++++++----- .../kernels/gemv_noshuffle_q4_k_f32.cl | 91 +++++++++++-------- .../kernels/gemv_noshuffle_q4_k_f32_o4.cl | 5 +- .../kernels/gemv_noshuffle_q5_0_f32.cl | 28 +++--- .../kernels/gemv_noshuffle_q5_1_f32.cl | 30 +++--- .../kernels/gemv_noshuffle_q5_k_f32.cl | 78 ++++++++-------- .../kernels/gemv_noshuffle_q6_k_f32.cl | 84 +++++++++-------- .../kernels/gemv_noshuffle_q6_k_f32_o4.cl | 5 +- .../kernels/gemv_noshuffle_q8_0_f32.cl | 40 ++++---- 13 files changed, 285 insertions(+), 232 deletions(-) diff --git a/ggml/src/ggml-opencl/kernels/gemv_noshuffle_iq4_nl_f32.cl b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_iq4_nl_f32.cl index 1f832cb25..e17ca09ba 100644 --- a/ggml/src/ggml-opencl/kernels/gemv_noshuffle_iq4_nl_f32.cl +++ b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_iq4_nl_f32.cl @@ -230,6 +230,8 @@ kernel void kernel_gemv_noshuffle_iq4_nl_f32( uint M = ne01; uint LINE_STRIDE_A = M / 2; + // Keep tail lanes active for subgroup ops, but read a valid weight row. + uint gid_s = min(gid, LINE_STRIDE_A - 1); uint BLOCK_STRIDE_A = NSUBGROUPS * M; private uint4 regA; @@ -240,7 +242,7 @@ kernel void kernel_gemv_noshuffle_iq4_nl_f32( // loop along K in block granularity, skip 4 blocks every iter for (uint k = groupId; k < (K / QK4_NL); k += NSUBGROUPS) { - regS = src0_d[gid + k * LINE_STRIDE_A]; // each fiber loads scale of two rows + regS = src0_d[gid_s + k * LINE_STRIDE_A]; // each fiber loads scale of two rows // first 4 fibers in each wave load 8 B values to its private scope if (slid < 4) { regB.s0123 = read_imagef(src1, (slid * 2 + k * 8)); @@ -248,20 +250,20 @@ kernel void kernel_gemv_noshuffle_iq4_nl_f32( } // load half weights for two blocks in consecutive rows - 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.s0 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 0)).x; + regA.s1 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 1)).x; + regA.s2 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 2)).x; + regA.s3 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 3)).x; #ifdef VECTOR_SUB_GROUP_BROADCAST dequantizeBlockAccum_ns_sgbroadcast_8_hi(totalSum, as_ushort8(regA), regS, regB); #else dequantizeBlockAccum_ns_sgbroadcast_1_hi(totalSum, as_ushort8(regA), regS, regB); #endif // VECTOR_SUB_GROUP_BROADCAST - 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; + regA.s0 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 4)).x; + regA.s1 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 5)).x; + regA.s2 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 6)).x; + regA.s3 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 7)).x; #ifdef VECTOR_SUB_GROUP_BROADCAST dequantizeBlockAccum_ns_sgbroadcast_8_lo(totalSum, as_ushort8(regA), regS, regB); #else diff --git a/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q1_0_f32.cl b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q1_0_f32.cl index 9efede294..d935a489c 100644 --- a/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q1_0_f32.cl +++ b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q1_0_f32.cl @@ -73,6 +73,8 @@ __kernel void kernel_gemv_noshuffle_q1_0_f32( uint M = ne01; uint LINE_STRIDE_A = M; + // Keep tail lanes active for subgroup ops, but read a valid weight row. + uint gid_s = min(gid, LINE_STRIDE_A - 1); uint BLOCK_STRIDE_A = 4 * M; uint4 regA; @@ -83,7 +85,7 @@ __kernel void kernel_gemv_noshuffle_q1_0_f32( #pragma unroll 1 for (uint kb = groupId; kb < (K / QK1_0); kb += N_SIMDGROUP) { - regS = src0_d[gid + kb * LINE_STRIDE_A]; // each fiber loads its row's scale + regS = src0_d[gid_s + kb * LINE_STRIDE_A]; // each fiber loads its row's scale // first 16 fibers load 8 B values each -> 128 activations for this block if (slid < 16) { @@ -92,10 +94,10 @@ __kernel void kernel_gemv_noshuffle_q1_0_f32( } // load this row's 4 uint32 (128 sign bits) - regA.s0 = read_imageui(src0_q, (gid + kb * BLOCK_STRIDE_A + LINE_STRIDE_A * 0)).x; - regA.s1 = read_imageui(src0_q, (gid + kb * BLOCK_STRIDE_A + LINE_STRIDE_A * 1)).x; - regA.s2 = read_imageui(src0_q, (gid + kb * BLOCK_STRIDE_A + LINE_STRIDE_A * 2)).x; - regA.s3 = read_imageui(src0_q, (gid + kb * BLOCK_STRIDE_A + LINE_STRIDE_A * 3)).x; + regA.s0 = read_imageui(src0_q, (gid_s + kb * BLOCK_STRIDE_A + LINE_STRIDE_A * 0)).x; + regA.s1 = read_imageui(src0_q, (gid_s + kb * BLOCK_STRIDE_A + LINE_STRIDE_A * 1)).x; + regA.s2 = read_imageui(src0_q, (gid_s + kb * BLOCK_STRIDE_A + LINE_STRIDE_A * 2)).x; + regA.s3 = read_imageui(src0_q, (gid_s + kb * BLOCK_STRIDE_A + LINE_STRIDE_A * 3)).x; float scale = (float)regS; dequantizeBlockAccum_q1(totalSum, regA.s0, scale, regB, 0); diff --git a/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q4_0_f32.cl b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q4_0_f32.cl index 0dcc929a6..6066ae3a9 100644 --- a/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q4_0_f32.cl +++ b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q4_0_f32.cl @@ -224,6 +224,8 @@ __kernel void kernel_gemv_noshuffle_q4_0_f32( uint M = ne01; uint LINE_STRIDE_A = M / 2; + // Keep tail lanes active for subgroup ops, but read a valid weight row. + uint gid_s = min(gid, LINE_STRIDE_A - 1); uint BLOCK_STRIDE_A = N_SIMDGROUP * M; __private uint4 regA; @@ -234,7 +236,7 @@ __kernel void kernel_gemv_noshuffle_q4_0_f32( // loop along K in block granularity, skip 4 blocks every iter for (uint k = groupId; k < (K / QK4_0); k += N_SIMDGROUP) { - regS = src0_d[gid + k * LINE_STRIDE_A]; // each fiber loads scale of two rows + regS = src0_d[gid_s + k * LINE_STRIDE_A]; // each fiber loads scale of two rows // first 4 fibers in each wave load 8 B values to its private scope if (slid < 4) { regB.s0123 = read_imagef(src1, (slid * 2 + k * 8)); @@ -242,20 +244,20 @@ __kernel void kernel_gemv_noshuffle_q4_0_f32( } // load half weights for two blocks in consecutive rows - 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.s0 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 0)).x; + regA.s1 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 1)).x; + regA.s2 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 2)).x; + regA.s3 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 3)).x; #ifdef VECTOR_SUB_GROUP_BROADCAST dequantizeBlockAccum_ns_sgbroadcast_8_hi(totalSum, as_ushort8(regA), regS, regB); #else dequantizeBlockAccum_ns_sgbroadcast_1_hi(totalSum, as_ushort8(regA), regS, regB); #endif // VECTOR_SUB_GROUP_BROADCAST - 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; + regA.s0 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 4)).x; + regA.s1 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 5)).x; + regA.s2 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 6)).x; + regA.s3 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 7)).x; #ifdef VECTOR_SUB_GROUP_BROADCAST dequantizeBlockAccum_ns_sgbroadcast_8_lo(totalSum, as_ushort8(regA), regS, regB); #else @@ -329,6 +331,8 @@ __kernel void kernel_gemv_noshuffle_q4_0_f32_mc3( uint M = ne01; uint LINE_STRIDE_A = M / 2; + // Keep tail lanes active for subgroup ops, but read a valid weight row. + uint gid_s = min(gid, LINE_STRIDE_A - 1); // 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. @@ -347,17 +351,17 @@ __kernel void kernel_gemv_noshuffle_q4_0_f32_mc3( __private float2 ts3 = (float2)(0.0f); for (uint k = groupId; k < (K / QK4_0); k += nsg) { - regS = src0_d[gid + k * LINE_STRIDE_A]; + regS = src0_d[gid_s + 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; + regA_hi.s0 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 0)).x; + regA_hi.s1 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 1)).x; + regA_hi.s2 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 2)).x; + regA_hi.s3 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 3)).x; + regA_lo.s0 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 4)).x; + regA_lo.s1 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 5)).x; + regA_lo.s2 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 6)).x; + regA_lo.s3 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 7)).x; MC_COL_Q40(ts0, 0); MC_COL_Q40(ts1, 1); @@ -381,10 +385,12 @@ __kernel void kernel_gemv_noshuffle_q4_0_f32_mc3( } 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])); + if (gid < LINE_STRIDE_A) { + 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 diff --git a/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q4_0_f32_spec.cl b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q4_0_f32_spec.cl index 0dca20f71..63990c414 100644 --- a/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q4_0_f32_spec.cl +++ b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q4_0_f32_spec.cl @@ -210,6 +210,8 @@ __kernel void kernel_gemv_noshuffle_q4_0_f32( { uint groupId = get_local_id(1); uint gid = get_global_id(0); + // Keep tail lanes active for subgroup ops, but read a valid weight row. + uint gid_s = min(gid, (uint)LINE_STRIDE_A - 1u); ushort slid = get_sub_group_local_id(); __private uint4 regA; @@ -220,7 +222,7 @@ __kernel void kernel_gemv_noshuffle_q4_0_f32( // loop along K in block granularity, skip 4 blocks every iter for (uint k = groupId; k < (K / QK4_0); k += N_SIMDGROUP) { - regS = src0_d[gid + k * LINE_STRIDE_A]; // each fiber loads scale of two rows + regS = src0_d[gid_s + k * LINE_STRIDE_A]; // each fiber loads scale of two rows // first 4 fibers in each wave load 8 B values to its private scope if (slid < 4) { regB.s0123 = read_imagef(src1, (slid * 2 + k * 8)); @@ -228,20 +230,20 @@ __kernel void kernel_gemv_noshuffle_q4_0_f32( } // load half weights for two blocks in consecutive rows - 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.s0 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 0)).x; + regA.s1 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 1)).x; + regA.s2 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 2)).x; + regA.s3 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 3)).x; #ifdef VECTOR_SUB_GROUP_BROADCAST dequantizeBlockAccum_ns_sgbroadcast_8_hi(totalSum, as_ushort8(regA), regS, regB); #else dequantizeBlockAccum_ns_sgbroadcast_1_hi(totalSum, as_ushort8(regA), regS, regB); #endif // VECTOR_SUB_GROUP_BROADCAST - 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; + regA.s0 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 4)).x; + regA.s1 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 5)).x; + regA.s2 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 6)).x; + regA.s3 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 7)).x; #ifdef VECTOR_SUB_GROUP_BROADCAST dequantizeBlockAccum_ns_sgbroadcast_8_lo(totalSum, as_ushort8(regA), regS, regB); #else diff --git a/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q4_1_f32.cl b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q4_1_f32.cl index 2ccf4214c..b0b4793d7 100644 --- a/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q4_1_f32.cl +++ b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q4_1_f32.cl @@ -209,6 +209,8 @@ kernel void kernel_gemv_noshuffle_q4_1_f32( uint M = ne01; uint LINE_STRIDE_A = M / 2; + // Keep tail lanes active for subgroup ops, but read a valid weight row. + uint gid_s = min(gid, LINE_STRIDE_A - 1); uint BLOCK_STRIDE_A = NSUBGROUPS * M; private uint4 regA; @@ -220,8 +222,8 @@ kernel void kernel_gemv_noshuffle_q4_1_f32( // loop along K in block granularity, skip 4 blocks every iter for (uint k = groupId; k < (K / QK4_0); k += NSUBGROUPS) { - regS = src0_d[gid + k * LINE_STRIDE_A]; // each fiber loads scale of two rows - regM = src0_m[gid + k * LINE_STRIDE_A]; // each fiber loads min of two rows + regS = src0_d[gid_s + k * LINE_STRIDE_A]; // each fiber loads scale of two rows + regM = src0_m[gid_s + k * LINE_STRIDE_A]; // each fiber loads min of two rows // first 4 fibers in each wave load 8 B values to its private scope if (slid < 4) { regB.s0123 = read_imagef(src1, (slid * 2 + k * 8)); @@ -229,20 +231,20 @@ kernel void kernel_gemv_noshuffle_q4_1_f32( } // load half weights for two blocks in consecutive rows - 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.s0 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 0)).x; + regA.s1 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 1)).x; + regA.s2 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 2)).x; + regA.s3 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 3)).x; #ifdef VECTOR_SUB_GROUP_BROADCAT dequantizeBlockAccum_ns_sgbroadcast_8_hi(totalSum, as_ushort8(regA), regS, regM, regB); #else dequantizeBlockAccum_ns_sgbroadcast_1_hi(totalSum, as_ushort8(regA), regS, regM, regB); #endif // VECTOR_SUB_GROUP_BROADCAT - 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; + regA.s0 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 4)).x; + regA.s1 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 5)).x; + regA.s2 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 6)).x; + regA.s3 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 7)).x; #ifdef VECTOR_SUB_GROUP_BROADCAT dequantizeBlockAccum_ns_sgbroadcast_8_lo(totalSum, as_ushort8(regA), regS, regM, regB); #else @@ -326,6 +328,8 @@ kernel void kernel_gemv_noshuffle_q4_1_f32_mc3( uint M = ne01; uint LINE_STRIDE_A = M / 2; + // Keep tail lanes active for subgroup ops, but read a valid weight row. + uint gid_s = min(gid, LINE_STRIDE_A - 1); uint BLOCK_STRIDE_A = NSUBGROUPS * M; uint COL_STRIDE = K / 4; // float4 pixels per activation column @@ -339,18 +343,18 @@ kernel void kernel_gemv_noshuffle_q4_1_f32_mc3( 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]; + regS = src0_d[gid_s + k * LINE_STRIDE_A]; + regM = src0_m[gid_s + 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; + regA_hi.s0 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 0)).x; + regA_hi.s1 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 1)).x; + regA_hi.s2 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 2)).x; + regA_hi.s3 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 3)).x; + regA_lo.s0 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 4)).x; + regA_lo.s1 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 5)).x; + regA_lo.s2 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 6)).x; + regA_lo.s3 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 7)).x; MC_COL_Q41(ts0, 0); MC_COL_Q41(ts1, 1); @@ -373,10 +377,12 @@ kernel void kernel_gemv_noshuffle_q4_1_f32_mc3( 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])); + if (gid < LINE_STRIDE_A) { + 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 diff --git a/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q4_k_f32.cl b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q4_k_f32.cl index 916e00118..3dc30db58 100644 --- a/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q4_k_f32.cl +++ b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q4_k_f32.cl @@ -415,6 +415,8 @@ kernel void kernel_gemv_noshuffle_q4_k_f32_glu( uint M = ne01; uint LINE_STRIDE_A = M / 2; + // Keep tail lanes active for subgroup ops, but read a valid weight row. + uint gid_s = min(gid, LINE_STRIDE_A - 1); uint BLOCK_STRIDE_A = 4 * M; private uint4 regA; @@ -434,9 +436,9 @@ kernel void kernel_gemv_noshuffle_q4_k_f32_glu( 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; \ + half2 d = DD[gid_s + sb * LINE_STRIDE_A]; \ + half2 dm = MM[gid_s + sb * LINE_STRIDE_A]; \ + global const uchar * sc0 = SS + sb * 12 * M + 2 * gid_s; \ global const uchar * sc1 = sc0 + 1; \ uchar sv0, mn0, sv1, mn1; \ get_scale_min_k4(j, sc0, M, &sv0, &mn0, mask_d6, mask_d4, mask_hi2); \ @@ -447,15 +449,15 @@ kernel void kernel_gemv_noshuffle_q4_k_f32_glu( 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; \ + regA.s0 = read_imageui(Q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 0)).x; \ + regA.s1 = read_imageui(Q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 1)).x; \ + regA.s2 = read_imageui(Q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 2)).x; \ + regA.s3 = read_imageui(Q, (gid_s + 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; \ + regA.s0 = read_imageui(Q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 4)).x; \ + regA.s1 = read_imageui(Q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 5)).x; \ + regA.s2 = read_imageui(Q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 6)).x; \ + regA.s3 = read_imageui(Q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 7)).x; \ DEQ_LO(SUM, as_ushort8(regA), regS, regM, regB); \ } @@ -489,8 +491,10 @@ kernel void kernel_gemv_noshuffle_q4_k_f32_glu( 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); + if (gid < LINE_STRIDE_A) { + dst[gid * 2 + 0] = glu_apply(glu_op, gateSum.s0, upSum.s0); + dst[gid * 2 + 1] = glu_apply(glu_op, gateSum.s1, upSum.s1); + } } } @@ -533,6 +537,9 @@ kernel void kernel_gemv_noshuffle_q4_k_f32_splitk( uint LINE_STRIDE_A = M / 2; uint BLOCK_STRIDE_A = 4 * M; // physical, independent of the K-split + // Keep tail lanes active for subgroup ops, but read a valid weight row. + uint gid_s = min(gid, LINE_STRIDE_A - 1); + private uint4 regA; private half2 regS, regM; private float8 regB; @@ -542,9 +549,9 @@ kernel void kernel_gemv_noshuffle_q4_k_f32_splitk( 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; + half2 d = src0_d[gid_s + sb * LINE_STRIDE_A]; + half2 dm = src0_m[gid_s + sb * LINE_STRIDE_A]; + global const uchar * sc0 = src0_s + sb * 12 * M + 2 * gid_s; global const uchar * sc1 = sc0 + 1; uchar sv0, mn0, sv1, mn1; get_scale_min_k4(j, sc0, M, &sv0, &mn0, mask_d6, mask_d4, mask_hi2); @@ -555,19 +562,19 @@ kernel void kernel_gemv_noshuffle_q4_k_f32_splitk( 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.s0 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 0)).x; + regA.s1 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 1)).x; + regA.s2 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 2)).x; + regA.s3 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 3)).x; #ifdef VECTOR_SUB_GROUP_BROADCAST dequantizeBlockAccum_ns_sgbroadcast_8_hi(totalSum, as_ushort8(regA), regS, regM, regB); #else dequantizeBlockAccum_ns_sgbroadcast_1_hi(totalSum, as_ushort8(regA), regS, regM, regB); #endif - 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; + regA.s0 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 4)).x; + regA.s1 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 5)).x; + regA.s2 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 6)).x; + regA.s3 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 7)).x; #ifdef VECTOR_SUB_GROUP_BROADCAST dequantizeBlockAccum_ns_sgbroadcast_8_lo(totalSum, as_ushort8(regA), regS, regM, regB); #else @@ -584,7 +591,9 @@ kernel void kernel_gemv_noshuffle_q4_k_f32_splitk( for (uint i = 0; i < nsg - 1; ++i) { totalSum += reduceLM[SUBGROUP_SIZE * i + slid]; } - vstore2(totalSum, 0, &(partial[kslice * M + gid * 2])); + // Keep tail stores inside this K-slice. + if (gid * 2 + 0 < M) partial[kslice * M + gid * 2 + 0] = totalSum.s0; + if (gid * 2 + 1 < M) partial[kslice * M + gid * 2 + 1] = totalSum.s1; } } @@ -689,6 +698,8 @@ kernel void kernel_gemv_noshuffle_q4_k_f32_mc3( uint M = ne01; uint LINE_STRIDE_A = M / 2; + // Keep tail lanes active for subgroup ops, but read a valid weight row. + uint gid_s = min(gid, LINE_STRIDE_A - 1); uint BLOCK_STRIDE_A = NSUBGROUPS * M; uint COL_STRIDE = K / 4; // float4 pixels per activation column @@ -713,10 +724,10 @@ kernel void kernel_gemv_noshuffle_q4_k_f32_mc3( 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]; + half2 d = src0_d[gid_s + sb * LINE_STRIDE_A]; + half2 dm = src0_m[gid_s + sb * LINE_STRIDE_A]; - global const uchar * sc0 = src0_s + sb * 12 * M + 2 * gid; + global const uchar * sc0 = src0_s + sb * 12 * M + 2 * gid_s; global const uchar * sc1 = sc0 + 1; uchar sv0, mn0, sv1, mn1; @@ -727,14 +738,14 @@ kernel void kernel_gemv_noshuffle_q4_k_f32_mc3( 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; + regA_hi.s0 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 0)).x; + regA_hi.s1 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 1)).x; + regA_hi.s2 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 2)).x; + regA_hi.s3 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 3)).x; + regA_lo.s0 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 4)).x; + regA_lo.s1 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 5)).x; + regA_lo.s2 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 6)).x; + regA_lo.s3 = read_imageui(src0_q, (gid_s + 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- @@ -827,8 +838,10 @@ kernel void kernel_gemv_noshuffle_q4_k_f32_mc3( 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])); + if (gid < LINE_STRIDE_A) { + 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])); + } } } diff --git a/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q4_k_f32_o4.cl b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q4_k_f32_o4.cl index 02916bb91..2c1c659a3 100644 --- a/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q4_k_f32_o4.cl +++ b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q4_k_f32_o4.cl @@ -231,8 +231,9 @@ kernel void kernel_gemv_noshuffle_q4_k_f32_o4( // 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; + // Keep tail lanes active for subgroup ops, but read valid weight pairs. + uint gid_a = min(gid * 2, (uint)(ne01 / 2) - 1u); + uint gid_b = min(gid * 2 + 1, (uint)(ne01 / 2) - 1u); uint K = ne00; uint M = ne01; diff --git a/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q5_0_f32.cl b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q5_0_f32.cl index 7dbf5a3bb..d8537e158 100644 --- a/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q5_0_f32.cl +++ b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q5_0_f32.cl @@ -209,6 +209,8 @@ __kernel void kernel_gemv_noshuffle_q5_0_f32( uint M = ne01; uint LINE_STRIDE_A = M / 2; + // Keep tail lanes active for subgroup ops, but read a valid weight row. + uint gid_s = min(gid, LINE_STRIDE_A - 1); uint BLOCK_STRIDE_A = NSUBGROUPS * M; private uint4 regA; @@ -218,13 +220,13 @@ __kernel void kernel_gemv_noshuffle_q5_0_f32( private float2 totalSum = (float2)(0.0f); for (uint k = groupId; k < (K / QK5_0); k += NSUBGROUPS) { - regS = src0_d[gid + k * LINE_STRIDE_A]; + regS = src0_d[gid_s + k * LINE_STRIDE_A]; ushort4 qh_raw; - qh_raw.s0 = src0_qh[gid + (4*k + 0) * LINE_STRIDE_A]; - qh_raw.s1 = src0_qh[gid + (4*k + 1) * LINE_STRIDE_A]; - qh_raw.s2 = src0_qh[gid + (4*k + 2) * LINE_STRIDE_A]; - qh_raw.s3 = src0_qh[gid + (4*k + 3) * LINE_STRIDE_A]; + qh_raw.s0 = src0_qh[gid_s + (4*k + 0) * LINE_STRIDE_A]; + qh_raw.s1 = src0_qh[gid_s + (4*k + 1) * LINE_STRIDE_A]; + qh_raw.s2 = src0_qh[gid_s + (4*k + 2) * LINE_STRIDE_A]; + qh_raw.s3 = src0_qh[gid_s + (4*k + 3) * LINE_STRIDE_A]; uchar8 raw = as_uchar8(qh_raw); uchar8 qh_bytes = (uchar8)(raw.s0, raw.s2, raw.s4, raw.s6, @@ -236,10 +238,10 @@ __kernel void kernel_gemv_noshuffle_q5_0_f32( regB.s4567 = read_imagef(src1, (1 + slid * 2 + k * 8)); } - regA.s0 = read_imageui(src0_qs, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 0)).x; - regA.s1 = read_imageui(src0_qs, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 1)).x; - regA.s2 = read_imageui(src0_qs, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 2)).x; - regA.s3 = read_imageui(src0_qs, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 3)).x; + regA.s0 = read_imageui(src0_qs, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 0)).x; + regA.s1 = read_imageui(src0_qs, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 1)).x; + regA.s2 = read_imageui(src0_qs, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 2)).x; + regA.s3 = read_imageui(src0_qs, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 3)).x; #ifdef VECTOR_SUB_GROUP_BROADCAST dequantizeBlockAccum_ns_q5_0_sgbroadcast_8_hi(totalSum, as_ushort8(regA), qh_bytes, regS, regB); @@ -247,10 +249,10 @@ __kernel void kernel_gemv_noshuffle_q5_0_f32( dequantizeBlockAccum_ns_q5_0_sgbroadcast_1_hi(totalSum, as_ushort8(regA), qh_bytes, regS, regB); #endif // VECTOR_SUB_GROUP_BROADCAST - regA.s0 = read_imageui(src0_qs, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 4)).x; - regA.s1 = read_imageui(src0_qs, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 5)).x; - regA.s2 = read_imageui(src0_qs, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 6)).x; - regA.s3 = read_imageui(src0_qs, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 7)).x; + regA.s0 = read_imageui(src0_qs, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 4)).x; + regA.s1 = read_imageui(src0_qs, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 5)).x; + regA.s2 = read_imageui(src0_qs, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 6)).x; + regA.s3 = read_imageui(src0_qs, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 7)).x; #ifdef VECTOR_SUB_GROUP_BROADCAST dequantizeBlockAccum_ns_q5_0_sgbroadcast_8_lo(totalSum, as_ushort8(regA), qh_bytes, regS, regB); #else diff --git a/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q5_1_f32.cl b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q5_1_f32.cl index ba0e2a711..82b4c741e 100644 --- a/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q5_1_f32.cl +++ b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q5_1_f32.cl @@ -210,6 +210,8 @@ __kernel void kernel_gemv_noshuffle_q5_1_f32( uint M = ne01; uint LINE_STRIDE_A = M / 2; + // Keep tail lanes active for subgroup ops, but read a valid weight row. + uint gid_s = min(gid, LINE_STRIDE_A - 1); uint BLOCK_STRIDE_A = NSUBGROUPS * M; __private uint4 regA; @@ -220,14 +222,14 @@ __kernel void kernel_gemv_noshuffle_q5_1_f32( __private float2 totalSum = (float2)(0.0f); for (uint k = groupId; k < (K / QK5_1); k += NSUBGROUPS) { - regS = src0_d[gid + k * LINE_STRIDE_A]; - regM = src0_m[gid + k * LINE_STRIDE_A]; + regS = src0_d[gid_s + k * LINE_STRIDE_A]; + regM = src0_m[gid_s + k * LINE_STRIDE_A]; ushort4 qh_raw; - qh_raw.s0 = src0_qh[gid + (4*k + 0) * LINE_STRIDE_A]; - qh_raw.s1 = src0_qh[gid + (4*k + 1) * LINE_STRIDE_A]; - qh_raw.s2 = src0_qh[gid + (4*k + 2) * LINE_STRIDE_A]; - qh_raw.s3 = src0_qh[gid + (4*k + 3) * LINE_STRIDE_A]; + qh_raw.s0 = src0_qh[gid_s + (4*k + 0) * LINE_STRIDE_A]; + qh_raw.s1 = src0_qh[gid_s + (4*k + 1) * LINE_STRIDE_A]; + qh_raw.s2 = src0_qh[gid_s + (4*k + 2) * LINE_STRIDE_A]; + qh_raw.s3 = src0_qh[gid_s + (4*k + 3) * LINE_STRIDE_A]; uchar8 raw = as_uchar8(qh_raw); uchar8 qh_bytes = (uchar8)(raw.s0, raw.s2, raw.s4, raw.s6, @@ -239,10 +241,10 @@ __kernel void kernel_gemv_noshuffle_q5_1_f32( regB.s4567 = read_imagef(src1, (1 + slid * 2 + k * 8)); } - regA.s0 = read_imageui(src0_qs, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 0)).x; - regA.s1 = read_imageui(src0_qs, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 1)).x; - regA.s2 = read_imageui(src0_qs, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 2)).x; - regA.s3 = read_imageui(src0_qs, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 3)).x; + regA.s0 = read_imageui(src0_qs, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 0)).x; + regA.s1 = read_imageui(src0_qs, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 1)).x; + regA.s2 = read_imageui(src0_qs, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 2)).x; + regA.s3 = read_imageui(src0_qs, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 3)).x; #ifdef VECTOR_SUB_GROUP_BROADCAST dequantizeBlockAccum_ns_q5_1_sgbroadcast_8_hi(totalSum, as_ushort8(regA), qh_bytes, regS, regM, regB); @@ -250,10 +252,10 @@ __kernel void kernel_gemv_noshuffle_q5_1_f32( dequantizeBlockAccum_ns_q5_1_sgbroadcast_1_hi(totalSum, as_ushort8(regA), qh_bytes, regS, regM, regB); #endif // VECTOR_SUB_GROUP_BROADCAST - regA.s0 = read_imageui(src0_qs, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 4)).x; - regA.s1 = read_imageui(src0_qs, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 5)).x; - regA.s2 = read_imageui(src0_qs, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 6)).x; - regA.s3 = read_imageui(src0_qs, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 7)).x; + regA.s0 = read_imageui(src0_qs, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 4)).x; + regA.s1 = read_imageui(src0_qs, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 5)).x; + regA.s2 = read_imageui(src0_qs, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 6)).x; + regA.s3 = read_imageui(src0_qs, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 7)).x; #ifdef VECTOR_SUB_GROUP_BROADCAST dequantizeBlockAccum_ns_q5_1_sgbroadcast_8_lo(totalSum, as_ushort8(regA), qh_bytes, regS, regM, regB); #else diff --git a/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q5_k_f32.cl b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q5_k_f32.cl index ae864b19b..6570167c9 100644 --- a/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q5_k_f32.cl +++ b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q5_k_f32.cl @@ -232,6 +232,8 @@ kernel void kernel_gemv_noshuffle_q5_k_f32( uint M = ne01; uint LINE_STRIDE_A = M / 2; + // Keep tail lanes active for subgroup ops, but read a valid weight row. + uint gid_s = min(gid, LINE_STRIDE_A - 1); uint BLOCK_STRIDE_A = NSUBGROUPS * M; uint LINE_STRIDE_A_QH = M / 2; @@ -250,11 +252,11 @@ kernel void kernel_gemv_noshuffle_q5_k_f32( 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]; + half2 d = src0_d[gid_s + sb * LINE_STRIDE_A]; + half2 dm = src0_m[gid_s + sb * LINE_STRIDE_A]; - global const uchar * sc0 = src0_s + 2 * gid * scales_per_row + sb * 12; - global const uchar * sc1 = src0_s + (2 * gid + 1) * scales_per_row + sb * 12; + global const uchar * sc0 = src0_s + 2 * gid_s * scales_per_row + sb * 12; + global const uchar * sc1 = src0_s + (2 * gid_s + 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); @@ -268,25 +270,25 @@ kernel void kernel_gemv_noshuffle_q5_k_f32( regB.s4567 = read_imagef(src1, (1 + slid * 2 + k * 8)); } - 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); + regH.s0 = as_ushort(read_imageh(src0_qh, (gid_s + k * BLOCK_STRIDE_A_QH + LINE_STRIDE_A_QH * 0)).x); + regH.s1 = as_ushort(read_imageh(src0_qh, (gid_s + k * BLOCK_STRIDE_A_QH + LINE_STRIDE_A_QH * 1)).x); + regH.s2 = as_ushort(read_imageh(src0_qh, (gid_s + k * BLOCK_STRIDE_A_QH + LINE_STRIDE_A_QH * 2)).x); + regH.s3 = as_ushort(read_imageh(src0_qh, (gid_s + k * BLOCK_STRIDE_A_QH + LINE_STRIDE_A_QH * 3)).x); - 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.s0 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 0)).x; + regA.s1 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 1)).x; + regA.s2 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 2)).x; + regA.s3 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 3)).x; #ifdef VECTOR_SUB_GROUP_BROADCAST dequantizeBlockAccum_ns_sgbroadcast_8_hi(totalSum, as_ushort8(regA), as_uchar8(regH), regS, regM, regB); #else dequantizeBlockAccum_ns_sgbroadcast_1_hi(totalSum, as_ushort8(regA), as_uchar8(regH), regS, regM, regB); #endif // VECTOR_SUB_GROUP_BROADCAST - 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; + regA.s0 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 4)).x; + regA.s1 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 5)).x; + regA.s2 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 6)).x; + regA.s3 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 7)).x; #ifdef VECTOR_SUB_GROUP_BROADCAST dequantizeBlockAccum_ns_sgbroadcast_8_lo(totalSum, as_ushort8(regA), as_uchar8(regH), regS, regM, regB); #else @@ -373,6 +375,8 @@ kernel void kernel_gemv_noshuffle_q5_k_f32_mc3( uint M = ne01; uint LINE_STRIDE_A = M / 2; + // Keep tail lanes active for subgroup ops, but read a valid weight row. + uint gid_s = min(gid, LINE_STRIDE_A - 1); uint BLOCK_STRIDE_A = NSUBGROUPS * M; uint LINE_STRIDE_A_QH = M / 2; uint BLOCK_STRIDE_A_QH = NSUBGROUPS * M / 2; @@ -393,11 +397,11 @@ kernel void kernel_gemv_noshuffle_q5_k_f32_mc3( 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]; + half2 d = src0_d[gid_s + sb * LINE_STRIDE_A]; + half2 dm = src0_m[gid_s + sb * LINE_STRIDE_A]; - global const uchar * sc0 = src0_s + 2 * gid * scales_per_row + sb * 12; - global const uchar * sc1 = src0_s + (2 * gid + 1) * scales_per_row + sb * 12; + global const uchar * sc0 = src0_s + 2 * gid_s * scales_per_row + sb * 12; + global const uchar * sc1 = src0_s + (2 * gid_s + 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); @@ -407,19 +411,19 @@ kernel void kernel_gemv_noshuffle_q5_k_f32_mc3( 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); + regH.s0 = as_ushort(read_imageh(src0_qh, (gid_s + k * BLOCK_STRIDE_A_QH + LINE_STRIDE_A_QH * 0)).x); + regH.s1 = as_ushort(read_imageh(src0_qh, (gid_s + k * BLOCK_STRIDE_A_QH + LINE_STRIDE_A_QH * 1)).x); + regH.s2 = as_ushort(read_imageh(src0_qh, (gid_s + k * BLOCK_STRIDE_A_QH + LINE_STRIDE_A_QH * 2)).x); + regH.s3 = as_ushort(read_imageh(src0_qh, (gid_s + 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; + regA_hi.s0 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 0)).x; + regA_hi.s1 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 1)).x; + regA_hi.s2 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 2)).x; + regA_hi.s3 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 3)).x; + regA_lo.s0 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 4)).x; + regA_lo.s1 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 5)).x; + regA_lo.s2 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 6)).x; + regA_lo.s3 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 7)).x; MC_COL_Q5K(ts0, 0); MC_COL_Q5K(ts1, 1); @@ -442,10 +446,12 @@ kernel void kernel_gemv_noshuffle_q5_k_f32_mc3( 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])); + if (gid < LINE_STRIDE_A) { + 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 diff --git a/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q6_k_f32.cl b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q6_k_f32.cl index 32624ac86..a6908a6a7 100644 --- a/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q6_k_f32.cl +++ b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q6_k_f32.cl @@ -219,26 +219,28 @@ kernel void kernel_gemv_noshuffle_q6_K_f32( float2 total_sum = 0.0f; int line_stride_a = ne01 / 2; + // Keep tail lanes active for subgroup ops, but read a valid weight row. + int gid_s = min(gid, line_stride_a - 1); int block_stride_a = NSUBGROUPS * ne01; 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]); + reg_d = src0_d[gid_s + k/8 * line_stride_a]; + reg_s = as_char4(src0_s[gid_s + k * line_stride_a]); 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); } - reg_a_l.s0 = read_imageui(src0_ql, gid + k*block_stride_a + line_stride_a*0).x; - reg_a_l.s1 = read_imageui(src0_ql, gid + k*block_stride_a + line_stride_a*1).x; - reg_a_l.s2 = read_imageui(src0_ql, gid + k*block_stride_a + line_stride_a*2).x; - reg_a_l.s3 = read_imageui(src0_ql, gid + k*block_stride_a + line_stride_a*3).x; + reg_a_l.s0 = read_imageui(src0_ql, gid_s + k*block_stride_a + line_stride_a*0).x; + reg_a_l.s1 = read_imageui(src0_ql, gid_s + k*block_stride_a + line_stride_a*1).x; + reg_a_l.s2 = read_imageui(src0_ql, gid_s + k*block_stride_a + line_stride_a*2).x; + reg_a_l.s3 = read_imageui(src0_ql, gid_s + k*block_stride_a + line_stride_a*3).x; - reg_a_h.s0 = as_ushort(read_imageh(src0_qh, gid + k*block_stride_a + line_stride_a*0).x); - reg_a_h.s1 = as_ushort(read_imageh(src0_qh, gid + k*block_stride_a + line_stride_a*1).x); - reg_a_h.s2 = as_ushort(read_imageh(src0_qh, gid + k*block_stride_a + line_stride_a*2).x); - reg_a_h.s3 = as_ushort(read_imageh(src0_qh, gid + k*block_stride_a + line_stride_a*3).x); + reg_a_h.s0 = as_ushort(read_imageh(src0_qh, gid_s + k*block_stride_a + line_stride_a*0).x); + reg_a_h.s1 = as_ushort(read_imageh(src0_qh, gid_s + k*block_stride_a + line_stride_a*1).x); + reg_a_h.s2 = as_ushort(read_imageh(src0_qh, gid_s + k*block_stride_a + line_stride_a*2).x); + reg_a_h.s3 = as_ushort(read_imageh(src0_qh, gid_s + k*block_stride_a + line_stride_a*3).x); #ifdef VECTOR_SUB_GROUP_BROADCAT dequantize_block_acc_bcast_8_hi(total_sum, as_ushort8(reg_a_l), as_uchar8(reg_a_h), reg_d, reg_s, reg_b); @@ -246,15 +248,15 @@ kernel void kernel_gemv_noshuffle_q6_K_f32( dequantize_block_acc_bcast_1_hi(total_sum, as_ushort8(reg_a_l), as_uchar8(reg_a_h), reg_d, reg_s, reg_b); #endif // VECTOR_SUB_GROUP_BROADCAT - reg_a_l.s0 = read_imageui(src0_ql, gid + k*block_stride_a + line_stride_a*4).x; - reg_a_l.s1 = read_imageui(src0_ql, gid + k*block_stride_a + line_stride_a*5).x; - reg_a_l.s2 = read_imageui(src0_ql, gid + k*block_stride_a + line_stride_a*6).x; - reg_a_l.s3 = read_imageui(src0_ql, gid + k*block_stride_a + line_stride_a*7).x; + reg_a_l.s0 = read_imageui(src0_ql, gid_s + k*block_stride_a + line_stride_a*4).x; + reg_a_l.s1 = read_imageui(src0_ql, gid_s + k*block_stride_a + line_stride_a*5).x; + reg_a_l.s2 = read_imageui(src0_ql, gid_s + k*block_stride_a + line_stride_a*6).x; + reg_a_l.s3 = read_imageui(src0_ql, gid_s + k*block_stride_a + line_stride_a*7).x; - reg_a_h.s0 = as_ushort(read_imageh(src0_qh, gid + k*block_stride_a + line_stride_a*4).x); - reg_a_h.s1 = as_ushort(read_imageh(src0_qh, gid + k*block_stride_a + line_stride_a*5).x); - reg_a_h.s2 = as_ushort(read_imageh(src0_qh, gid + k*block_stride_a + line_stride_a*6).x); - reg_a_h.s3 = as_ushort(read_imageh(src0_qh, gid + k*block_stride_a + line_stride_a*7).x); + reg_a_h.s0 = as_ushort(read_imageh(src0_qh, gid_s + k*block_stride_a + line_stride_a*4).x); + reg_a_h.s1 = as_ushort(read_imageh(src0_qh, gid_s + k*block_stride_a + line_stride_a*5).x); + reg_a_h.s2 = as_ushort(read_imageh(src0_qh, gid_s + k*block_stride_a + line_stride_a*6).x); + reg_a_h.s3 = as_ushort(read_imageh(src0_qh, gid_s + k*block_stride_a + line_stride_a*7).x); #ifdef VECTOR_SUB_GROUP_BROADCAT dequantize_block_acc_bcast_8_lo(total_sum, as_ushort8(reg_a_l), as_uchar8(reg_a_h), reg_d, reg_s, reg_b); @@ -323,6 +325,8 @@ kernel void kernel_gemv_noshuffle_q6_K_f32_mc3( int nb = ne00 / 32; int line_stride_a = ne01 / 2; + // Keep tail lanes active for subgroup ops, but read a valid weight row. + int gid_s = min(gid, line_stride_a - 1); int block_stride_a = NSUBGROUPS * ne01; int COL_STRIDE = ne00 / 4; // float4 pixels per activation column @@ -335,27 +339,27 @@ kernel void kernel_gemv_noshuffle_q6_K_f32_mc3( 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]); + reg_d = src0_d[gid_s + k/8 * line_stride_a]; + reg_s = as_char4(src0_s[gid_s + 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_hi.s0 = read_imageui(src0_ql, gid_s + k*block_stride_a + line_stride_a*0).x; + ql_hi.s1 = read_imageui(src0_ql, gid_s + k*block_stride_a + line_stride_a*1).x; + ql_hi.s2 = read_imageui(src0_ql, gid_s + k*block_stride_a + line_stride_a*2).x; + ql_hi.s3 = read_imageui(src0_ql, gid_s + k*block_stride_a + line_stride_a*3).x; + qh_hi.s0 = as_ushort(read_imageh(src0_qh, gid_s + k*block_stride_a + line_stride_a*0).x); + qh_hi.s1 = as_ushort(read_imageh(src0_qh, gid_s + k*block_stride_a + line_stride_a*1).x); + qh_hi.s2 = as_ushort(read_imageh(src0_qh, gid_s + k*block_stride_a + line_stride_a*2).x); + qh_hi.s3 = as_ushort(read_imageh(src0_qh, gid_s + 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); + ql_lo.s0 = read_imageui(src0_ql, gid_s + k*block_stride_a + line_stride_a*4).x; + ql_lo.s1 = read_imageui(src0_ql, gid_s + k*block_stride_a + line_stride_a*5).x; + ql_lo.s2 = read_imageui(src0_ql, gid_s + k*block_stride_a + line_stride_a*6).x; + ql_lo.s3 = read_imageui(src0_ql, gid_s + k*block_stride_a + line_stride_a*7).x; + qh_lo.s0 = as_ushort(read_imageh(src0_qh, gid_s + k*block_stride_a + line_stride_a*4).x); + qh_lo.s1 = as_ushort(read_imageh(src0_qh, gid_s + k*block_stride_a + line_stride_a*5).x); + qh_lo.s2 = as_ushort(read_imageh(src0_qh, gid_s + k*block_stride_a + line_stride_a*6).x); + qh_lo.s3 = as_ushort(read_imageh(src0_qh, gid_s + 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. @@ -402,8 +406,10 @@ kernel void kernel_gemv_noshuffle_q6_K_f32_mc3( 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])); + if (gid < line_stride_a) { + 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])); + } } } diff --git a/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q6_k_f32_o4.cl b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q6_k_f32_o4.cl index 84447e61b..22efb6c3c 100644 --- a/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q6_k_f32_o4.cl +++ b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q6_k_f32_o4.cl @@ -249,8 +249,9 @@ kernel void Q6K_O4_NAME( // 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; + // Keep tail lanes active for subgroup ops, but read valid weight pairs. + int gid_a = min(gid * 2, ne01 / 2 - 1); + int gid_b = min(gid * 2 + 1, ne01 / 2 - 1); int nb = ne00 / 32; diff --git a/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q8_0_f32.cl b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q8_0_f32.cl index 6f6d7425c..6109c967a 100644 --- a/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q8_0_f32.cl +++ b/ggml/src/ggml-opencl/kernels/gemv_noshuffle_q8_0_f32.cl @@ -155,6 +155,8 @@ __kernel void kernel_gemv_noshuffle_q8_0_f32_splitk( uint M = ne01; uint LINE_STRIDE_A = M; + // Keep tail lanes active for subgroup ops, but read a valid weight row. + uint gid_s = min(gid, LINE_STRIDE_A - 1); uint BLOCK_STRIDE_A = 8 * M; // physical, independent of the K-split __private uint8 regA; @@ -164,19 +166,19 @@ __kernel void kernel_gemv_noshuffle_q8_0_f32_splitk( #pragma unroll 1 for (uint k = kslice * nsg + groupId; k < (K / QK8_0); k += ksplit * nsg) { - regS = src0_d[gid + k * LINE_STRIDE_A]; + regS = src0_d[gid_s + 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; + regA.s0 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 0)).x; + regA.s1 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 1)).x; + regA.s2 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 2)).x; + regA.s3 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 3)).x; + regA.s4 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 4)).x; + regA.s5 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 5)).x; + regA.s6 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 6)).x; + regA.s7 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 7)).x; dequantizeBlockAccum_ns_sgbroadcast_1(totalSum, regA, convert_float(regS), regB); } @@ -227,6 +229,8 @@ __kernel void kernel_gemv_noshuffle_q8_0_f32( uint M = ne01; uint LINE_STRIDE_A = M; + // Keep tail lanes active for subgroup ops, but read a valid weight row. + uint gid_s = min(gid, LINE_STRIDE_A - 1); uint BLOCK_STRIDE_A = 8 * M; // 32 / 4 = 8 __private uint8 regA; @@ -238,7 +242,7 @@ __kernel void kernel_gemv_noshuffle_q8_0_f32( // loop along K in block granularity, skip 4 blocks every iter #pragma unroll 1 /* tell compiler not to unroll */ for (uint k = groupId; k < (K / QK8_0); k += N_SIMDGROUP) { - regS = src0_d[gid + k * LINE_STRIDE_A]; // each fiber loads scale of one rows + regS = src0_d[gid_s + k * LINE_STRIDE_A]; // each fiber loads the scale of one row // first 4 fibers in each wave load 8 B values to its private scope if (slid < 4) { regB.s0123 = read_imagef(src1, (slid * 2 + k * 8)); @@ -246,14 +250,14 @@ __kernel void kernel_gemv_noshuffle_q8_0_f32( } // load weights for one block in consecutive rows - 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; + regA.s0 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 0)).x; + regA.s1 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 1)).x; + regA.s2 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 2)).x; + regA.s3 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 3)).x; + regA.s4 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 4)).x; + regA.s5 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 5)).x; + regA.s6 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 6)).x; + regA.s7 = read_imageui(src0_q, (gid_s + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 7)).x; dequantizeBlockAccum_ns_sgbroadcast_1(totalSum, regA, convert_float(regS), regB); }