public_git/whisper.cpp
1
0
Fork 0
mirror of https://github.com/ggml-org/whisper.cpp.git synced 2026-10-11 08:45:40 +02:00
Code Issues Packages Projects Releases Wiki Activity

cuda : move blocks_per_col to where it is used (llama/29939)

Browse Source 
Signed-off-by: Adrien Gallouët <angt@huggingface.co>
This commit is contained in:
Adrien Gallouët
2026-10-06 10:38:19 +03:00
committed by Georgi Gerganov
parent 6421d2c1f2
commit 7914776a72
1 changed files with 2 additions and 3 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-cuda/quantize.cu
+2 -3
View File
@@ -131,8 +131,6 @@ static __global__ void quantize_mmq_nvfp4(
const int64_t ne0, const int64_t ne1, const int64_t ne2, const int n_expert_used) {
#if defined(BLACKWELL_MMA_AVAILABLE)
const int64_t blocks_per_col = (ne0 + QK_FP4_MMQ - 1) / QK_FP4_MMQ;
int64_t base_idx;
if constexpr (scatter) {
base_idx = (int64_t) blockIdx.x * s02; // one physical row per token
@@ -317,7 +315,8 @@ static __global__ void quantize_mmq_nvfp4(
reinterpret_cast<uint8_t *>(yb->d4)[sub] = fp8_code;
}
} else {
block_fp4_mmq * yb = y + (blockIdx.y * ((int64_t) blocks_per_col * ne1) + k_block * ne1 + blockIdx.x);
const int64_t blocks_per_col = (ne0 + QK_FP4_MMQ - 1) / QK_FP4_MMQ;
block_fp4_mmq * yb = y + (blockIdx.y * (blocks_per_col * ne1) + k_block * ne1 + blockIdx.x);
uint32_t * yqs = reinterpret_cast<uint32_t *>(yb->qs);
yqs[2 * sub + 0] = q0;
yqs[2 * sub + 1] = q1;