public_git/whisper.cpp
1
0
Fork 0
mirror of https://github.com/ggml-org/whisper.cpp.git synced 2026-10-22 03:48:19 +02:00
Code Issues Packages Projects Releases Wiki Activity

CUDA: Optimize accumulation in mmq for NVFP4 type (llama/29857)

Browse Source 
* ggml_cuda: optimize accumulation in mmq_vec_dot_fp4_fp4_mma for better performance

* remove whitespace

* fix: correct indentation in mma_block_scaled_fp4 loop
This commit is contained in:
Konrad Moren
2026-10-06 10:38:22 +03:00
committed by Georgi Gerganov
parent 82fa9df729
commit 8c6ca4ab76
1 changed files with 2 additions and 5 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-cuda/mmq-vec-dot.cuh
+2 -5
View File
@@ -1225,14 +1225,11 @@ template <ggml_type type, int J, bool fallback> static __device__ __forceinline_
#pragma unroll #pragma unroll
for (int n = 0; n < ntx; ++n) { for (int n = 0; n < ntx; ++n) {
// accumulate in place into the output sum array
tile_C & C = *reinterpret_cast<tile_C *>(sum + (j0 / tile_C::J + n) * tile_C::ne);
#pragma unroll #pragma unroll
for (int frag = 0; frag < nfrags; ++frag) { for (int frag = 0; frag < nfrags; ++frag) {
tile_C C = {};
mma_block_scaled_fp4<type>(C, A[n][frag], B[frag], scaleA[n][frag], scaleB[frag]); mma_block_scaled_fp4<type>(C, A[n][frag], B[frag], scaleA[n][frag], scaleB[frag]);
#pragma unroll
for (int l = 0; l < tile_C::ne; ++l) {
sum[(j0 / tile_C::J + n) * tile_C::ne + l] += C.x[l];
}
} }
} }
} }