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
Gaurav Garg 98915f889a Improve CUDA graph capture (llama/19754) 
* Improve CUDA graph capture

Currently, CUDA graphs are eagerly enabled on the first call to ggml_backend_cuda_graph_compute. If the graph properties keep changing (4+ consecutive updates), the graph is permanently disabled. This is suboptimal because:

- The first call always incurs CUDA graph capture overhead even if the graph is unstable
- Once permanently disabled, CUDA graphs never re-enable even after the graph stabilizes (e.g., switching from prompt processing to decode)

The new approach delays CUDA graph activation until warmup completes: the same cgraph must be called at least twice with matching properties before CUDA graph capture begins. This avoids wasted capture overhead on volatile graphs and allows graphs to become eligible once they stabilize.
This also fixes issues such as https://github.com/ggml-org/llama.cpp/discussions/19708

* Update ggml/src/ggml-cuda/ggml-cuda.cu

Co-authored-by: Johannes Gäßler <johannesg@5d6.de>

* Remove EM dashes

* Update ggml/src/ggml-cuda/ggml-cuda.cu

Co-authored-by: Aman Gupta <amangupta052@gmail.com>

---------

Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
Co-authored-by: Aman Gupta <amangupta052@gmail.com>
 		2026-02-27 20:57:58 +02:00
..
template-instances CUDA: faster FA for GQA > 1 but not power of 2 (llama/19092) 2026-01-30 15:56:40 +02:00
vendors HIP: add fattn-mma-f16 for RDNA4 (llama/18481) 2026-01-30 15:56:40 +02:00
CMakeLists.txt CUDA : Update CCCL-tag for 3.2 to final release from RC (llama/19486) 2026-02-15 21:44:37 +02: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
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 CUDA: Fix builds for older CCCL versions by ifdefing strided_iterator (llama/18964) 2026-01-30 15:56:40 +02:00
argsort.cuh sampling : add support for backend sampling (llama/17004) 2026-01-14 09:11:59 +02:00
binbcast.cu ggml : extend bin bcast for permuted src1 (llama/19484) 2026-02-15 21:44:37 +02:00
binbcast.cuh CUDA: fuse adds, fuse add with rms norm (llama/15631) 2025-09-20 13:42:44 +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
common.cuh Improve CUDA graph capture (llama/19754) 2026-02-27 20:57:58 +02:00
concat.cu musa: fix all warnings, re-enable `-DLLAMA_FATAL_WARNINGS=ON` in ci and update doc (llama/12611) 2025-03-31 14:56:53 +03:00
concat.cuh whisper : reorganize source code + improve CMake (#2256) 2024-06-26 19:34:09 +03:00
conv-transpose-1d.cu musa: add GGML_UNUSED_VARS (llama/15446) 2025-09-20 13:42:38 +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: add conv_2d_transpose (llama/14287) 2025-06-21 07:34:17 +03:00
conv2d-transpose.cuh CUDA: add conv_2d_transpose (llama/14287) 2025-06-21 07:34:17 +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: loop over ne2*ne3 in case it overflows (llama/19538) 2026-02-15 21:44:37 +02:00
convert.cuh HIP: RDNA4 tensor core support for MMF (llama/17077) 2025-12-12 17:53:06 +02: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 cuda : fix copy of large tensors (ggml_nbytes <= INT_MAX assertion) (llama/18433) 2026-01-14 09:11:59 +02: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: replace GGML_CUDA_F16 with CUDA arch checks (llama/15433) 2025-09-20 13:42:38 +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
fattn-common.cuh cuda : fix "V is K view" check for non-unified KV cache (llama/19145) 2026-01-30 15:56:40 +02:00
fattn-mma-f16.cuh CUDA: fix padding of GQA to power of 2 in FA (llama/19115) 2026-01-30 15:56:40 +02:00
fattn-tile.cu ggml: CUDA: add head size 72 for flash-attn (llama/16962) 2025-11-09 23:38:03 +02:00
fattn-tile.cuh CUDA: fix kernel selection logic for tile FA (llama/19686) 2026-02-27 20:57:58 +02:00
fattn-vec.cuh CUDA: fix allignment on register spill for FA (llama/18815) 2026-01-30 15:56:40 +02:00
fattn-wmma-f16.cu Adjust workaround for ROCWMMA_FATTN/GFX9 to only newer ROCm veresions (llama/19591) 2026-02-27 20:57:58 +02:00
fattn-wmma-f16.cuh CUDA: generalized (mma) FA, add Volta support (llama/17505) 2025-12-12 17:53:16 +02:00
fattn.cu cuda : fix nkvo, offload and cuda graph node properties matching (llama/19165) 2026-01-30 15:56:40 +02:00
fattn.cuh CUDA: refactor FA support/selection code (llama/15454) 2025-09-20 13:42:38 +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
getrows.cu CUDA: fix GET_ROWS for large tensors (llama/15882) 2025-09-20 13:42:52 +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 Improve CUDA graph capture (llama/19754) 2026-02-27 20:57:58 +02: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: fix im2col_3d to respect non-contiguous inputs (views) (llama/15956) 2025-09-20 13:45:30 +03:00
im2col.cuh ggml: add ops for WAN video model (cuda && cpu) (llama/15669) 2025-09-20 13:42:49 +03:00
mean.cu ggml-cuda: enable cuda-graphs for `n-cpu-moe` (llama/18934) 2026-01-30 15:56:40 +02:00
mean.cuh CUDA: add mean operation (llama/14313) 2025-07-01 17:54:53 +03:00
mma.cuh HIP: add mmf for CDNA (llama/18896) 2026-01-30 15:56:40 +02:00
mmf.cu HIP: add mmf for CDNA (llama/18896) 2026-01-30 15:56:40 +02:00
mmf.cuh HIP: add mmf for CDNA (llama/18896) 2026-01-30 15:56:40 +02:00
mmid.cu CUDA: add fp kernel for larger batch size MoE (llama/16512) 2025-10-15 09:29:17 +03:00
mmid.cuh CUDA: add fp kernel for larger batch size MoE (llama/16512) 2025-10-15 09:29:17 +03:00
mmq.cu HIP: adjust RDNA3.5 MMQ kernel selction logic (llama/18666) 2026-01-14 09:11:59 +02:00
mmq.cuh cuda: optimize iq2xxs/iq2xs/iq3xxs dequantization (llama/19624) 2026-02-27 20:57:58 +02:00
mmvf.cu CUDA: use mmvq for mul-mat-id for small batch sizes (llama/18958) 2026-02-08 09:29:10 +02: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 CUDA: use mmvq for mul-mat-id for small batch sizes (llama/18958) 2026-02-08 09:29:10 +02:00
mmvq.cuh cuda : enable CUDA graphs for MMID 1 <= BS <= 4 (llama/19645) 2026-02-27 20:57:58 +02:00
norm.cu CUDA: Factor out and re-use `block_reduce` function (llama/18785) 2026-01-30 15:56:40 +02: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 CPU/CUDA: fix (GQA) mul mat back, add CUDA support (llama/11380) 2025-02-03 22:00:57 +02: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: experimental native mxfp4 support for blackwell (llama/17906) 2025-12-31 17:52:09 +02:00
quantize.cuh CUDA: experimental native mxfp4 support for blackwell (llama/17906) 2025-12-31 17:52:09 +02:00
reduce_rows.cuh CUDA: Factor out and re-use `block_reduce` function (llama/18785) 2026-01-30 15:56:40 +02: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 CUDA: Fix non-contig rope (llama/19338) 2026-02-15 21:44:37 +02:00
rope.cuh CUDA: fuse rope + set_rows (llama/16884) 2025-11-17 21:05:46 +02:00
scale.cu ggml: add ops for WAN video model (cuda && cpu) (llama/15669) 2025-09-20 13:42:49 +03:00
scale.cuh whisper : reorganize source code + improve CMake (#2256) 2024-06-26 19:34:09 +03:00
set-rows.cu CUDA: use fastdiv in set-rows (llama/16834) 2025-11-09 23:38:03 +02: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
softcap.cu cuda : add softcap fusion (llama/14907) 2025-08-18 20:30:45 +03:00
softcap.cuh cuda : add softcap fusion (llama/14907) 2025-08-18 20:30:45 +03:00
softmax.cu CUDA: Factor out and re-use `block_reduce` function (llama/18785) 2026-01-30 15:56:40 +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 SOLVE_TRI extension to more dimensions (llama/17793) 2025-12-18 08:20:56 +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 model : add ASR support for LFM2-Audio-1.5B (conformer) (llama/18106) 2025-12-31 17:52:09 +02:00
ssm-conv.cuh ggml : faster ssm scan (llama/10558) 2025-04-02 15:51:57 +03:00
ssm-scan.cu ggml : optimize cuda ssm_scan using warp-level reduction (llama/18505) 2026-01-14 09:11:59 +02: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 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
sumrows.cuh CUDA: add mean operation (llama/14313) 2025-07-01 17:54:53 +03:00
top-k.cu CUDA: Replace init_offsets kernel with iterators in cub-based argsort (llama/18930) 2026-01-30 15:56:40 +02:00
top-k.cuh sampling : add support for backend sampling (llama/17004) 2026-01-14 09:11:59 +02:00
topk-moe.cu CUDA: refactor topk-moe to enable more models (GLM 4.7, Nemotron etc.) (llama/19126) 2026-01-30 15:56:40 +02:00
topk-moe.cuh CUDA: refactor topk-moe to enable more models (GLM 4.7, Nemotron etc.) (llama/19126) 2026-01-30 15:56:40 +02: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 ggml : add ops SOFTPLUS, EXPM1, TRI, SOLVE_TRI, CUMSUM (llama/17063) 2025-11-17 21:05:46 +02:00
unary.cuh ggml : add ops SOFTPLUS, EXPM1, TRI, SOLVE_TRI, CUMSUM (llama/17063) 2025-11-17 21:05:46 +02: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: optimize iq2xxs/iq2xs/iq3xxs dequantization (llama/19624) 2026-02-27 20:57:58 +02: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