public_git
/
whisper.cpp
mirror of https://github.com/ggml-org/whisper.cpp.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity
whisper.cpp/ggml/src/ggml-cuda
History
Pascal f74102e4f8
cuda: GET_ROWS quants (llama/25962) 
* cuda: add k-quant support to GET_ROWS

Device-side embedding lookups require GET_ROWS to handle the k-quants
used by common GGUF recipes (Q4_K_M stores token_embd as q6_K). Without
it the backend rejects the op and the scheduler falls back to the host,
copying the full embedding matrix back on every token in single-device
graphs.

Factor the super-block dequantizers out of the dequantize_block kernels
in convert.cu into shared device functions in dequantize.cuh and reuse
them from a new k_get_rows_kq kernel : one thread block dequantizes one
(dst row, super-block) pair with the existing thread layouts, 32 threads
for q4_K and 64 for the other k-quants.

Covers q2_K to q6_K in get_rows_cuda and supports_op. i-quants are left
as a TODO.

* cuda: add i-quant support to GET_ROWS

Extends the shared super-block dequantizers to the nine i-quants and
reuses them from k_get_rows_kq with the 32-thread layout of the matching
convert.cu kernels. supports_op gates the k-quant and i-quant path on
ne0 being a multiple of QK_K, which iq4_nl does not guarantee on its
own (QK4_NL sub-blocks). mxfp4 is left as a TODO.

* cuda: add mxfp4 support to GET_ROWS

Moves the mxfp4 dequantizer into the shared super-block helpers and
reuses it from k_get_rows_kq with the 32-thread layout of the matching
convert.cu kernel. mxfp4 joins the ne0 % QK_K gate in supports_op since
its 32-value sub-blocks do not guarantee QK_K-aligned rows on their own.
This closes GET_ROWS type coverage on CUDA: every quantized GGML type
now takes the direct device path.

* cuda: gate the GET_ROWS row size only for 32-value sub-block types

Address review from @pwilkin: the i-quant commit replaced the return
shared by the whole supported type cascade, so f16/f32/bf16/i32 and the
legacy quants also inherited the ne0 % QK_K == 0 gate and any row size
that is not a multiple of 256 fell back to the scheduler. Split the
cascade: unconditional support is restored everywhere, the gate stays
only on iq4_nl and mxfp4 whose 32-value sub-blocks do not guarantee the
QK_K super-blocks the kernel iterates on.
 		2026-07-30 16:31:47 +03:00
..
template-instances CUDA: fix Gemma E4B MTP FlashAttention (llama/25148) 2026-07-10 13:06:42 +03:00
vendors CUDA: add cublasSgemmBatched mapping for HIP/MUSA vendor headers (llama/25033) 2026-06-26 16:03:57 +03:00
CMakeLists.txt Programmatic Dependent Launch (PDL) for more performance on newer NVIDIA GPUs (Hopper+) (llama/22522) 2026-05-25 12:26:07 +03:00
acc.cu llama/ggml: add LLM training support (llama/10544) 2025-05-13 13:59:21 +03:00
acc.cuh whisper : reorganize source code + improve CMake (#2256) 2024-06-26 19:34:09 +03:00
add-id.cu musa: fix build warnings (llama/15258) 2025-09-20 13:42:38 +03:00
add-id.cuh llama : add gpt-oss (llama/15091) 2025-08-18 20:30:45 +03:00
allreduce.cu Fix for issue #22974. Cast intermediate results to float before adding and casting the result to the destination type. Avoids half+half operator ambiguity. (llama/22994) 2026-05-14 21:26:48 +03:00
allreduce.cuh internal AllReduce kernel for CUDA provider (llama/22299) 2026-05-14 21:26:48 +03:00
arange.cu whisper : reorganize source code + improve CMake (#2256) 2024-06-26 19:34:09 +03:00
arange.cuh whisper : reorganize source code + improve CMake (#2256) 2024-06-26 19:34:09 +03:00
argmax.cu ggml : use WARP_SIZE/2 for argmax reduction offset (llama/18092) 2025-12-18 08:20:56 +02:00
argmax.cuh ggml/ex: calculate accuracy in graph, adapt MNIST (ggml/980) 2024-10-05 15:23:51 +03:00
argsort.cu ggml : process data in smaller chunks in CUDA ggml_top_k() and ggml_argsort() to reduce temporary buffers memory usage (llama/24776) 2026-07-30 16:31:36 +03:00
argsort.cuh ggml : process data in smaller chunks in CUDA ggml_top_k() and ggml_argsort() to reduce temporary buffers memory usage (llama/24776) 2026-07-30 16:31:36 +03:00
binbcast.cu ggml : address integer overflows in binary ops CUDA implementation (llama/24706) 2026-06-26 16:03:57 +03:00
binbcast.cuh CUDA: fuse muls (llama/21665) 2026-04-30 11:29:05 +03:00
clamp.cu cuda: unary ops as float + de-duplicate (ggml/1130) 2025-03-08 15:13:01 +02:00
clamp.cuh whisper : reorganize source code + improve CMake (#2256) 2024-06-26 19:34:09 +03:00
col2im-1d.cu Ggml/cuda col2im 1d (llama/24417) 2026-06-19 12:53:43 +03:00
col2im-1d.cuh Ggml/cuda col2im 1d (llama/24417) 2026-06-19 12:53:43 +03:00
common.cuh cuda: GET_ROWS quants (llama/25962) 2026-07-30 16:31:47 +03:00
concat.cu cuda : relax tensor contiguity requirements for quantized concat (llama/25678) 2026-07-30 16:31:41 +03:00
concat.cuh whisper : reorganize source code + improve CMake (#2256) 2024-06-26 19:34:09 +03:00
conv-transpose-1d.cu ggml-cuda: optimize conv_transpose_1d indexing (llama/25310) 2026-07-10 13:06:42 +03:00
conv-transpose-1d.cuh feat: cuda implementation for `ggml_conv_transpose_1d` (ggml/854) 2024-07-08 14:53:55 +03:00
conv2d-dw.cu CUDA: add conv_2d_dw (llama/14265) 2025-06-21 07:34:17 +03:00
conv2d-dw.cuh CUDA: add conv_2d_dw (llama/14265) 2025-06-21 07:34:17 +03:00
conv2d-transpose.cu CUDA & CPU: support F32 kernel type for `CONV_TRANSPOSE_2D` (llama/17094) 2026-03-29 15:04:36 +03:00
conv2d-transpose.cuh CUDA & CPU: support F32 kernel type for `CONV_TRANSPOSE_2D` (llama/17094) 2026-03-29 15:04:36 +03:00
conv2d.cu CUDA: fix build error from ambiguous __half conversions in conv2d (llama/15690) 2025-09-20 13:42:46 +03:00
conv2d.cuh CUDA: add conv2d (llama/15635) 2025-09-20 13:42:44 +03:00
convert.cu cuda: GET_ROWS quants (llama/25962) 2026-07-30 16:31:47 +03:00
convert.cuh CUDA: fix BF16 FA compilation (llama/20865) 2026-03-29 15:04:36 +03:00
count-equal.cu ggml: fix zero division in ‘dne’ calculation in CUDA COUNT_EQUAL operator when ‘ne’ is small (#10213) 2024-11-15 15:21:04 +02:00
count-equal.cuh ggml/ex: calculate accuracy in graph, adapt MNIST (ggml/980) 2024-10-05 15:23:51 +03:00
cp-async.cuh CUDA: FA support for Deepseek (Ampere or newer) (llama/13306) 2025-05-13 13:59:21 +03:00
cpy-utils.cuh cuda : support non-contiguous i32 to i32 copy (llama/17326) 2025-12-12 17:53:06 +02:00
cpy.cu Added a cudaMemcpy2DAsync fast path to ggml_cuda_cpy (llama/25057) 2026-07-10 13:06:42 +03:00
cpy.cuh cuda : remove legacy copy-op pointer indirection code (llama/16485) 2025-10-15 09:29:17 +03:00
cross-entropy-loss.cu CUDA: add dynamic shared mem to softmax, refactor general usage (llama/14497) 2025-07-12 19:23:56 +03:00
cross-entropy-loss.cuh ggml/examples: add backend support for numerical optimization (ggml/949) 2024-09-24 19:45:08 +03:00
cumsum.cu sampling : add support for backend sampling (llama/17004) 2026-01-14 09:11:59 +02:00
cumsum.cuh Add support for CUMSUM and TRI for CUDA. (llama/17584) 2025-12-12 17:53:17 +02:00
dequantize.cuh cuda: GET_ROWS quants (llama/25962) 2026-07-30 16:31:47 +03:00
diag.cu Add DIAG for CUDA (llama/17873) 2025-12-12 17:53:23 +02:00
diag.cuh Add DIAG for CUDA (llama/17873) 2025-12-12 17:53:23 +02:00
diagmask.cu whisper : reorganize source code + improve CMake (#2256) 2024-06-26 19:34:09 +03:00
diagmask.cuh whisper : reorganize source code + improve CMake (#2256) 2024-06-26 19:34:09 +03:00
dsv4-hc.cu DeepseekV4: Add fused hyper-connection ops (llama/25585) 2026-07-30 16:31:43 +03:00
dsv4-hc.cuh DeepseekV4: Add fused hyper-connection ops (llama/25585) 2026-07-30 16:31:43 +03:00
fattn-common.cuh CUDA: consistent use of __restrict__ + PDL for FA (llama/25185) 2026-07-10 13:06:42 +03:00
fattn-mma-f16.cuh CUDA: fix Gemma E4B MTP FlashAttention (llama/25148) 2026-07-10 13:06:42 +03:00
fattn-tile.cu Add flash attention MMA / Tiles to support MiMo-V2.5 (llama/22812) 2026-05-14 21:26:48 +03:00
fattn-tile.cuh CUDA: fix Gemma E4B MTP FlashAttention (llama/25148) 2026-07-10 13:06:42 +03:00
fattn-vec.cuh Avoid PDL race conditions by disabling __restrict__ when PDL is used (llama/24030) 2026-06-08 14:36:36 +03:00
fattn-wmma-f16.cu Avoid PDL race conditions by disabling __restrict__ when PDL is used (llama/24030) 2026-06-08 14:36:36 +03:00
fattn-wmma-f16.cuh chore : correct typos [no ci] (llama/20041) 2026-03-16 13:10:15 +02:00
fattn.cu CUDA: extend K-type validation to V-types for flash attention (llama/24403) 2026-07-10 13:06:42 +03:00
fattn.cuh cuda: reserve space for quantize kv-cache at startup (llama/23907) 2026-06-08 14:36:36 +03:00
fill.cu ggml : allow fill node alloc inplace (llama/17870) 2025-12-12 17:53:22 +02:00
fill.cuh cuda : add FILL op support (llama/17851) 2025-12-12 17:53:22 +02:00
fwht.cu CUDA: missing PDL sync for FWHT, better fallback (llama/23690) 2026-05-29 09:47:30 +03:00
fwht.cuh CUDA: missing PDL sync for FWHT, better fallback (llama/23690) 2026-05-29 09:47:30 +03:00
gated_delta_net.cu Remove redundant CUDA copies after gated_delta_net. (llama/23940) 2026-07-10 13:06:42 +03:00
gated_delta_net.cuh Remove redundant CUDA copies after gated_delta_net. (llama/23940) 2026-07-10 13:06:42 +03:00
getrows.cu cuda: GET_ROWS quants (llama/25962) 2026-07-30 16:31:47 +03:00
getrows.cuh CUDA: batched+noncont MMQ, refactor bs>1 MoE code (llama/13199) 2025-05-01 13:29:02 +03:00
ggml-cuda.cu cuda: GET_ROWS quants (llama/25962) 2026-07-30 16:31:47 +03:00
gla.cu llama: add support for QRWKV6 model architecture (llama/11001) 2025-01-14 10:38:01 +02:00
gla.cuh llama: add support for QRWKV6 model architecture (llama/11001) 2025-01-14 10:38:01 +02:00
im2col.cu CUDA: handle OW > 65535 in im2col (2D and 3D) (llama/22944) 2026-05-14 21:26:48 +03:00
im2col.cuh ggml: add ops for WAN video model (cuda && cpu) (llama/15669) 2025-09-20 13:42:49 +03:00
lightning-indexer.cu cuda : CUDA GGML_OP_LIGHTNING_INDEXER implementation (generic vector kernel + wmma kernel) (llama/25545) 2026-07-30 16:31:42 +03:00
lightning-indexer.cuh cuda : CUDA GGML_OP_LIGHTNING_INDEXER implementation (generic vector kernel + wmma kernel) (llama/25545) 2026-07-30 16:31:42 +03:00
mean.cu Programmatic Dependent Launch (PDL) for more performance on newer NVIDIA GPUs (Hopper+) (llama/22522) 2026-05-25 12:26:07 +03:00
mean.cuh CUDA: add mean operation (llama/14313) 2025-07-01 17:54:53 +03:00
mma.cuh HIP: RDNA3 mma FA, faster AMD transpose, tune AMD (llama/22880) 2026-05-25 12:26:07 +03:00
mmf.cu CUDA: dedup MoE gate/up activation quantization (llama/25441) 2026-07-30 16:31:42 +03:00
mmf.cuh clean up unused variables warnings (llama/23975) 2026-06-08 14:36:36 +03:00
mmid.cu CUDA: dedup MoE gate/up activation quantization (llama/25441) 2026-07-30 16:31:42 +03:00
mmid.cuh CUDA: dedup MoE gate/up activation quantization (llama/25441) 2026-07-30 16:31:42 +03:00
mmq-config-ampere.cuh CUDA: refactor MMQ kernel configuration (llama/24127) 2026-07-30 16:31:38 +03:00
mmq-config-blackwell.cuh CUDA: refactor MMQ kernel configuration (llama/24127) 2026-07-30 16:31:38 +03:00
mmq-config-cdna.cuh CUDA: refactor MMQ kernel configuration (llama/24127) 2026-07-30 16:31:38 +03:00
mmq-config-pascal.cuh CUDA: refactor MMQ kernel configuration (llama/24127) 2026-07-30 16:31:38 +03:00
mmq-config-rdna2.cuh CUDA: refactor MMQ kernel configuration (llama/24127) 2026-07-30 16:31:38 +03:00
mmq-config-rdna4.cuh CUDA: refactor MMQ kernel configuration (llama/24127) 2026-07-30 16:31:38 +03:00
mmq-load-tiles.cuh cuda: extract Q1_0 elements via __byte_perm (llama/25628) 2026-07-30 16:31:42 +03:00
mmq-vec-dot.cuh CUDA: refactor MMQ kernel configuration (llama/24127) 2026-07-30 16:31:38 +03:00
mmq.cu CUDA: dedup MoE gate/up activation quantization (llama/25441) 2026-07-30 16:31:42 +03:00
mmq.cuh CUDA: tighter MMQ src1 buffer size for native fp4 (llama/25613) 2026-07-30 16:31:41 +03:00
mmvf.cu Avoid PDL race conditions by disabling __restrict__ when PDL is used (llama/24030) 2026-06-08 14:36:36 +03:00
mmvf.cuh CUDA: use mmvq for mul-mat-id for small batch sizes (llama/18958) 2026-02-08 09:29:10 +02:00
mmvq.cu Only index by compile times + always multiply/add (llama/25445) 2026-07-10 13:06:42 +03:00
mmvq.cuh CUDA: route batch>=4 quantized matmul to MMQ on AMD MFMA hardware (llama/23227) 2026-05-29 09:47:30 +03:00
norm.cu Programmatic Dependent Launch (PDL) for more performance on newer NVIDIA GPUs (Hopper+) (llama/22522) 2026-05-25 12:26:07 +03:00
norm.cuh CUDA: fuse adds, fuse add with rms norm (llama/15631) 2025-09-20 13:42:44 +03:00
opt-step-adamw.cu ggml: new optimization interface (ggml/988) 2024-11-20 21:00:08 +02:00
opt-step-adamw.cuh ggml/examples: add backend support for numerical optimization (ggml/949) 2024-09-24 19:45:08 +03:00
opt-step-sgd.cu finetune: SGD optimizer, more CLI args (llama/13873) 2025-08-18 20:30:45 +03:00
opt-step-sgd.cuh finetune: SGD optimizer, more CLI args (llama/13873) 2025-08-18 20:30:45 +03:00
out-prod.cu CUDA: batch out_prod broadcast (dps2>1) path with cublasSgemmBatched (llama/24426) 2026-06-26 16:03:57 +03:00
out-prod.cuh ggml/examples: add backend support for numerical optimization (ggml/949) 2024-09-24 19:45:08 +03:00
pad.cu cuda : extend GGML_OP_PAD to work with non-cont src0 (llama/19429) 2026-02-15 21:44:37 +02:00
pad.cuh whisper : reorganize source code + improve CMake (#2256) 2024-06-26 19:34:09 +03:00
pad_reflect_1d.cu musa: fix build warnings (llama/15611) 2025-09-29 15:18:10 +03:00
pad_reflect_1d.cuh cuda : add Pad Reflect 1D support (llama/14659) 2025-09-20 13:42:39 +03:00
pool2d.cu whisper : reorganize source code + improve CMake (#2256) 2024-06-26 19:34:09 +03:00
pool2d.cuh whisper : reorganize source code + improve CMake (#2256) 2024-06-26 19:34:09 +03:00
quantize.cu CUDA: dedup MoE gate/up activation quantization (llama/25441) 2026-07-30 16:31:42 +03:00
quantize.cuh CUDA: dedup MoE gate/up activation quantization (llama/25441) 2026-07-30 16:31:42 +03:00
reduce_rows.cuh Avoid PDL race conditions by disabling __restrict__ when PDL is used (llama/24030) 2026-06-08 14:36:36 +03:00
roll.cu CUDA: add roll (llama/14919) 2025-08-18 20:30:45 +03:00
roll.cuh CUDA: add roll (llama/14919) 2025-08-18 20:30:45 +03:00
rope.cu Programmatic Dependent Launch (PDL) for more performance on newer NVIDIA GPUs (Hopper+) (llama/22522) 2026-05-25 12:26:07 +03:00
rope.cuh CUDA: fuse rope + set_rows (llama/16884) 2025-11-17 21:05:46 +02:00
scale.cu Programmatic Dependent Launch (PDL) for more performance on newer NVIDIA GPUs (Hopper+) (llama/22522) 2026-05-25 12:26:07 +03:00
scale.cuh whisper : reorganize source code + improve CMake (#2256) 2024-06-26 19:34:09 +03:00
set-rows.cu cuda : add support for f16->f16 GGML_OP_SET_ROWS (llama/25367) 2026-07-10 13:06:42 +03:00
set-rows.cuh CUDA: add set rows for f32 and f16 (llama/14551) 2025-07-20 00:23:50 +03:00
set.cu cuda: add SET operation support (llama/16804) 2025-11-09 23:38:03 +02:00
set.cuh cuda: add SET operation support (llama/16804) 2025-11-09 23:38:03 +02:00
snake.cu cuda: fuse snake activation (mul, sin, sqr, mul, add) (llama/22667) 2026-05-14 21:26:48 +03:00
snake.cuh cuda: fuse snake activation (mul, sin, sqr, mul, add) (llama/22667) 2026-05-14 21:26:48 +03:00
softcap.cu Programmatic Dependent Launch (PDL) for more performance on newer NVIDIA GPUs (Hopper+) (llama/22522) 2026-05-25 12:26:07 +03:00
softcap.cuh cuda : add softcap fusion (llama/14907) 2025-08-18 20:30:45 +03:00
softmax.cu chore : correct typos [no ci] (llama/20041) 2026-03-16 13:10:15 +02:00
softmax.cuh CUDA: backwards pass for misc. ops, add tests (llama/11257) 2025-02-03 22:00:57 +02:00
solve_tri.cu chore : correct typos [no ci] (llama/20041) 2026-03-16 13:10:15 +02:00
solve_tri.cuh SOLVE_TRI CUDA kernel for small matrices (llama/17457) 2025-12-12 17:53:11 +02:00
ssm-conv.cu Avoid PDL race conditions by disabling __restrict__ when PDL is used (llama/24030) 2026-06-08 14:36:36 +03:00
ssm-conv.cuh CUDA: fuse SSM_CONV + ADD(bias) + SILU (llama/22478) 2026-05-01 13:07:35 +03:00
ssm-scan.cu CUDA: Fix ssm_scan_f32 data-races (llama/24360) 2026-06-15 10:33:53 +03:00
ssm-scan.cuh ggml : faster ssm scan (llama/10558) 2025-04-02 15:51:57 +03:00
sum.cu CUDA: Optimize `reduce_rows_f32` kernel, leading up to 25x perf improvement on kernel-level and 10% perf increase for Gemma3n (llama/15132) 2025-08-18 20:30:45 +03:00
sum.cuh tests: add gradient tests for all backends (ggml/932) 2024-09-24 19:45:08 +03:00
sumrows.cu Programmatic Dependent Launch (PDL) for more performance on newer NVIDIA GPUs (Hopper+) (llama/22522) 2026-05-25 12:26:07 +03:00
sumrows.cuh CUDA: add mean operation (llama/14313) 2025-07-01 17:54:53 +03:00
top-k.cu ggml : process data in smaller chunks in CUDA ggml_top_k() and ggml_argsort() to reduce temporary buffers memory usage (llama/24776) 2026-07-30 16:31:36 +03:00
top-k.cuh sampling : add support for backend sampling (llama/17004) 2026-01-14 09:11:59 +02:00
topk-moe.cu cuda: add sqrt_softplus in topk-moe for dsv4 (llama/25896) 2026-07-30 16:31:46 +03:00
topk-moe.cuh cuda: add sqrt_softplus in topk-moe for dsv4 (llama/25896) 2026-07-30 16:31:46 +03:00
tri.cu Add support for CUMSUM and TRI for CUDA. (llama/17584) 2025-12-12 17:53:17 +02:00
tri.cuh Add support for CUMSUM and TRI for CUDA. (llama/17584) 2025-12-12 17:53:17 +02:00
tsembd.cu ggml : fix padding in timestep embedding kernels (llama/15932) 2025-09-20 13:45:30 +03:00
tsembd.cuh whisper : reorganize source code + improve CMake (#2256) 2024-06-26 19:34:09 +03:00
unary.cu Programmatic Dependent Launch (PDL) for more performance on newer NVIDIA GPUs (Hopper+) (llama/22522) 2026-05-25 12:26:07 +03:00
unary.cuh CUDA: fuse relu + sqr (llama/22249) 2026-04-30 11:29:16 +03:00
upscale.cu model: LFM2-VL fixes (llama/17577) 2025-12-12 17:53:14 +02:00
upscale.cuh whisper : reorganize source code + improve CMake (#2256) 2024-06-26 19:34:09 +03:00
vecdotq.cuh cuda: extract Q1_0 elements via __byte_perm (llama/25628) 2026-07-30 16:31:42 +03:00
wkv.cu llama: Add support for RWKV v7 architecture (llama/12412) 2025-03-27 11:06:03 +02:00
wkv.cuh llama: Add support for RWKV v7 architecture (llama/12412) 2025-03-27 11:06:03 +02:00