public_git/whisper.cpp
1
0
Fork 0
mirror of https://github.com/ggml-org/whisper.cpp.git synced 2026-10-12 17:25:39 +02:00
Code Issues Packages Projects Releases Wiki Activity

cuda : move neu_padded to where it is used (llama/29940)

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 7914776a72
commit 25453c0adf
1 changed files with 1 additions and 3 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-cuda/mmid.cu
+1 -3
View File
@@ -37,9 +37,6 @@ static __global__ void mm_ids_helper(
const int n_expert_used = n_expert_used_template == 0 ? n_expert_used_var : n_expert_used_template;
const int expert = blockIdx.x;
// token slots per warp lane group, padded to a power of 2 so a warp divides evenly
constexpr int neu_padded = mm_ids_pow2<n_expert_used_template>::value;
extern __shared__ char data_mm_ids_helper[];
mm_ids_helper_store * store = (mm_ids_helper_store *) data_mm_ids_helper;
@@ -69,6 +66,7 @@ static __global__ void mm_ids_helper(
} else {
// Implementation optimized for specific numbers of experts used:
// a warp holds a whole number of token slots, so the slot count is padded to a power of 2
constexpr int neu_padded = mm_ids_pow2<n_expert_used_template>::value;
static_assert(neu_padded <= warp_size && warp_size % neu_padded == 0, "bad n_expert_used");
for (int it0 = 0; it0 < n_tokens; it0 += warp_size/neu_padded) {
const int it = it0 + threadIdx.x / neu_padded;