nikola/llama_cpp
1
0
Fork 0
Code Issues Pull Requests Actions 18 Packages Projects Releases Wiki Activity

opencl: refactor Adreno q4_0 (#22335)

Browse Source 
* opencl: refactor adreno q4_0 gemm/gemv dispatch

* opencl: refactor q4_0 gemm/gemv loading, use consistent names

* opencl: use consistent name for adreno q8_0 gemm/gemv

* opencl: use consistent names for adreno q4_0 gemm/gemv

* opencl: simplify adreno q4_0 set_tensor

* opencl: refactor q4_0 get_tensor
This commit is contained in:
lhez
2026-05-05 09:38:57 -07:00
committed by GitHub
parent d5003b6e4d
commit ff806a110d
7 changed files with 355 additions and 623 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-opencl/CMakeLists.txt
+5 -5
View File
@@ -66,8 +66,6 @@ set(GGML_OPENCL_KERNELS
diag diag
div div
gelu gelu
gemv_noshuffle_general
gemv_noshuffle
get_rows get_rows
glu glu
group_norm group_norm
@@ -75,7 +73,6 @@ set(GGML_OPENCL_KERNELS
im2col_f32 im2col_f32
im2col_f16 im2col_f16
mean mean
mul_mat_Ab_Bi_8x4
mul_mv_f16_f16 mul_mv_f16_f16
mul_mv_f16_f32_1row mul_mv_f16_f32_1row
mul_mv_f16_f32_l4 mul_mv_f16_f32_l4
@@ -120,12 +117,15 @@ set(GGML_OPENCL_KERNELS
mul_mm_q4_k_f32_l4_lm mul_mm_q4_k_f32_l4_lm
mul_mm_q5_k_f32_l4_lm mul_mm_q5_k_f32_l4_lm
mul_mm_q6_k_f32_l4_lm mul_mm_q6_k_f32_l4_lm
mul_mm_q8_0_f32_8x4 gemv_noshuffle_q4_0_f32
gemv_noshuffle_q4_0_f32_spec
gemm_noshuffle_q4_0_f32
gemv_noshuffle_q4_1_f32 gemv_noshuffle_q4_1_f32
gemm_noshuffle_q4_1_f32 gemm_noshuffle_q4_1_f32
gemv_noshuffle_iq4_nl_f32 gemv_noshuffle_iq4_nl_f32
gemm_noshuffle_iq4_nl_f32 gemm_noshuffle_iq4_nl_f32
gemv_noshuffle_general_q8_0_f32 gemv_noshuffle_q8_0_f32
gemm_noshuffle_q8_0_f32
gemv_noshuffle_q4_k_f32 gemv_noshuffle_q4_k_f32
gemm_noshuffle_q4_k_f32 gemm_noshuffle_q4_k_f32
gemv_noshuffle_q6_k_f32 gemv_noshuffle_q6_k_f32
ggml/src/ggml-opencl/ggml-opencl.cpp
+311 -579
View File
File diff suppressed because it is too large Load Diff
ggml/src/ggml-opencl/kernels/mul_mat_Ab_Bi_8x4.cl → ggml/src/ggml-opencl/kernels/gemm_noshuffle_q4_0_f32.cl
+1 -1
View File
@@ -17,7 +17,7 @@
REQD_SUBGROUP_SIZE_128 REQD_SUBGROUP_SIZE_128
#endif #endif
kernel void kernel_mul_mat_Ab_Bi_8x4( kernel void kernel_gemm_noshuffle_q4_0_f32(
global const ushort * src0_q, // quantized A global const ushort * src0_q, // quantized A
global const half * src0_d, // A scales global const half * src0_d, // A scales
__read_only image1d_buffer_t src1, // B (1d image) __read_only image1d_buffer_t src1, // B (1d image)
ggml/src/ggml-opencl/kernels/mul_mm_q8_0_f32_8x4.cl → ggml/src/ggml-opencl/kernels/gemm_noshuffle_q8_0_f32.cl
+1 -1
View File
@@ -11,7 +11,7 @@
REQD_SUBGROUP_SIZE_128 REQD_SUBGROUP_SIZE_128
#endif #endif
kernel void kernel_mul_mm_q8_0_f32_8x4( kernel void kernel_gemm_noshuffle_q8_0_f32(
global const uint * src0_q, global const uint * src0_q,
global const half * src0_d, global const half * src0_d,
__read_only image1d_buffer_t src1, __read_only image1d_buffer_t src1,
ggml/src/ggml-opencl/kernels/gemv_noshuffle_general.cl → ggml/src/ggml-opencl/kernels/gemv_noshuffle_q4_0_f32.cl
+5 -5
View File
@@ -191,7 +191,7 @@
#ifdef ADRENO_GPU #ifdef ADRENO_GPU
REQD_SUBGROUP_SIZE_64 REQD_SUBGROUP_SIZE_64
#endif #endif
__kernel void kernel_gemv_noshuffle( __kernel void kernel_gemv_noshuffle_q4_0_f32(
__read_only image1d_buffer_t src0_q, // quantized A __read_only image1d_buffer_t src0_q, // quantized A
global half2 * src0_d, // A scales global half2 * src0_d, // A scales
__read_only image1d_buffer_t src1, // B __read_only image1d_buffer_t src1, // B
@@ -238,21 +238,21 @@ __kernel void kernel_gemv_noshuffle(
regA.s1 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 1)).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.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.s3 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 3)).x;
#ifdef VECTOR_SUB_GROUP_BROADCAT #ifdef VECTOR_SUB_GROUP_BROADCAST
dequantizeBlockAccum_ns_sgbroadcast_8_hi(totalSum, as_ushort8(regA), regS, regB); dequantizeBlockAccum_ns_sgbroadcast_8_hi(totalSum, as_ushort8(regA), regS, regB);
#else #else
dequantizeBlockAccum_ns_sgbroadcast_1_hi(totalSum, as_ushort8(regA), regS, regB); dequantizeBlockAccum_ns_sgbroadcast_1_hi(totalSum, as_ushort8(regA), regS, regB);
#endif // VECTOR_SUB_GROUP_BROADCAT #endif // VECTOR_SUB_GROUP_BROADCAST
regA.s0 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 4)).x; 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.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.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.s3 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 7)).x;
#ifdef VECTOR_SUB_GROUP_BROADCAT #ifdef VECTOR_SUB_GROUP_BROADCAST
dequantizeBlockAccum_ns_sgbroadcast_8_lo(totalSum, as_ushort8(regA), regS, regB); dequantizeBlockAccum_ns_sgbroadcast_8_lo(totalSum, as_ushort8(regA), regS, regB);
#else #else
dequantizeBlockAccum_ns_sgbroadcast_1_lo(totalSum, as_ushort8(regA), regS, regB); dequantizeBlockAccum_ns_sgbroadcast_1_lo(totalSum, as_ushort8(regA), regS, regB);
#endif // VECTOR_SUB_GROUP_BROADCAT #endif // VECTOR_SUB_GROUP_BROADCAST
} }
// reduction in local memory, assumes #wave=4 // reduction in local memory, assumes #wave=4
ggml/src/ggml-opencl/kernels/gemv_noshuffle.cl → ggml/src/ggml-opencl/kernels/gemv_noshuffle_q4_0_f32_spec.cl
+5 -5
View File
@@ -191,7 +191,7 @@
#ifdef ADRENO_GPU #ifdef ADRENO_GPU
REQD_SUBGROUP_SIZE_64 REQD_SUBGROUP_SIZE_64
#endif #endif
__kernel void kernel_gemv_noshuffle( __kernel void kernel_gemv_noshuffle_q4_0_f32(
__read_only image1d_buffer_t src0_q, // quantized A __read_only image1d_buffer_t src0_q, // quantized A
global half2 * src0_d, // A scales global half2 * src0_d, // A scales
__read_only image1d_buffer_t src1, // B __read_only image1d_buffer_t src1, // B
@@ -232,21 +232,21 @@ __kernel void kernel_gemv_noshuffle(
regA.s1 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 1)).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.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.s3 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 3)).x;
#ifdef VECTOR_SUB_GROUP_BROADCAT #ifdef VECTOR_SUB_GROUP_BROADCAST
dequantizeBlockAccum_ns_sgbroadcast_8_hi(totalSum, as_ushort8(regA), regS, regB); dequantizeBlockAccum_ns_sgbroadcast_8_hi(totalSum, as_ushort8(regA), regS, regB);
#else #else
dequantizeBlockAccum_ns_sgbroadcast_1_hi(totalSum, as_ushort8(regA), regS, regB); dequantizeBlockAccum_ns_sgbroadcast_1_hi(totalSum, as_ushort8(regA), regS, regB);
#endif // VECTOR_SUB_GROUP_BROADCAT #endif // VECTOR_SUB_GROUP_BROADCAST
regA.s0 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 4)).x; 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.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.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.s3 = read_imageui(src0_q, (gid + k * BLOCK_STRIDE_A + LINE_STRIDE_A * 7)).x;
#ifdef VECTOR_SUB_GROUP_BROADCAT #ifdef VECTOR_SUB_GROUP_BROADCAST
dequantizeBlockAccum_ns_sgbroadcast_8_lo(totalSum, as_ushort8(regA), regS, regB); dequantizeBlockAccum_ns_sgbroadcast_8_lo(totalSum, as_ushort8(regA), regS, regB);
#else #else
dequantizeBlockAccum_ns_sgbroadcast_1_lo(totalSum, as_ushort8(regA), regS, regB); dequantizeBlockAccum_ns_sgbroadcast_1_lo(totalSum, as_ushort8(regA), regS, regB);
#endif // VECTOR_SUB_GROUP_BROADCAT #endif // VECTOR_SUB_GROUP_BROADCAST
} }
// reduction in local memory, assumes #wave=4 // reduction in local memory, assumes #wave=4
ggml/src/ggml-opencl/kernels/gemv_noshuffle_general_q8_0_f32.cl → ggml/src/ggml-opencl/kernels/gemv_noshuffle_q8_0_f32.cl
View File