sleepy/llama.cpp
1
0
Fork 0
Code Issues 16 Pull Requests Actions 31 Packages Projects Releases Wiki Activity

opencl: add initial mxfp4 support via mv (#15270)

Browse Source 
* opencl: add reference `mul_mv_mxfp4_f32`

* opencl: add reference `mul_mv_id` for mxfp4

* Q4_0 tranpose fix for Adreno

---------

Co-authored-by: shawngu-quic <shawngu@qti.qualcomm.com>
This commit is contained in:
lhez
2025-08-16 00:52:14 +08:00
committed by GitHub
parent 5edf1592fd
commit e2c1bfff53
5 changed files with 496 additions and 6 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-opencl/kernels/transpose.cl
+20
View File
@@ -24,6 +24,26 @@ kernel void kernel_transpose_16(
write_imageh(output, (i_2+3)*rows+j, (half4)(temp0.s3, temp1.s3, temp2.s3, temp3.s3));
}
// Padded kernel for irregular shape
kernel void kernel_transpose_16_4x1(
__read_only image1d_buffer_t input,
__write_only image1d_buffer_t output,
const uint rows,
const uint cols
) {
const int i = get_global_id(0);
const int j = get_global_id(1);
const int j_2 = j << 2;
half temp0 = read_imageh(input, (j_2 + 0) * cols + i).x;
half temp1 = read_imageh(input, (j_2 + 1) * cols + i).x;
half temp2 = read_imageh(input, (j_2 + 2) * cols + i).x;
half temp3 = read_imageh(input, (j_2 + 3) * cols + i).x;
write_imageh(output, i * rows + j, (half4)(temp0, temp1, temp2, temp3));
}
// 32-bit transpose, loading/storing a 4x4 tile of elements
kernel void kernel_transpose_32(
__read_only image1d_buffer_t input,