whisper.cpp/ggml/src/ggml-cuda
Johannes Gäßler 8d05cf479c
CUDA: fix overflow in MMA kernel without stream-k (llama/17939)
2025-12-17 15:19:48 +02:00
..
template-instances ggml: CUDA: add head size 72 for flash-attn (llama/16962) 2025-11-09 23:38:03 +02:00
vendors SOLVE_TRI extension to more dimensions (llama/17793) 2025-12-17 15:19:45 +02:00
CMakeLists.txt CUDA: skip fusion for repeating adds in bias (llama/17080) 2025-11-09 23:38:03 +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 cuda : optimize argmax (llama/10441) 2024-12-08 20:14:35 +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 : add error checking for cudaMemcpyAsync in argsort (llama/17599) 2025-12-12 17:53:13 +02:00
argsort.cuh whisper : reorganize source code + improve CMake (#2256) 2024-06-26 19:34:09 +03:00
binbcast.cu ggml: fix CUDA grid launch condition for large block_nums.y in binbcast (llama/16742) 2025-11-09 23:38:03 +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 HIP: enable mmf for RDNA3 (llama/17879) 2025-12-17 15:19:47 +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 musa: add GGML_UNUSED_VARS (llama/15446) 2025-09-20 13:42:38 +03: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 enable fp16/fast_fp16/bf16_mma on PH1 (llama/17551) 2025-12-12 17:53:12 +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 Add support for CUMSUM and TRI for CUDA. (llama/17584) 2025-12-12 17:53:17 +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 overflow in MMA kernel without stream-k (llama/17939) 2025-12-17 15:19:48 +02:00
fattn-mma-f16.cuh CUDA: fix overflow in MMA kernel without stream-k (llama/17939) 2025-12-17 15:19:48 +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 FP16 overflow in tile FA kernel (llama/17875) 2025-12-12 17:53:22 +02:00
fattn-vec.cuh CUDA: fix FA VKQ accumulator overflow (llama/17746) 2025-12-12 17:53:17 +02:00
fattn-wmma-f16.cu CUDA: fix FA VKQ accumulator overflow (llama/17746) 2025-12-12 17:53:17 +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 unpadded strides in MMA FA kernel (llama/17891) 2025-12-12 17:53:23 +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 SOLVE_TRI extension to more dimensions (llama/17793) 2025-12-17 15:19:45 +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 cuda : fix GGML_CUDA_GRAPHS=OFF (llama/15300) 2025-08-18 20:30:45 +03:00
mean.cuh CUDA: add mean operation (llama/14313) 2025-07-01 17:54:53 +03:00
mma.cuh HIP: enable mmf for RDNA3 (llama/17879) 2025-12-17 15:19:47 +02:00
mmf.cu HIP: enable mmf for RDNA3 (llama/17879) 2025-12-17 15:19:47 +02:00
mmf.cuh CUDA: generalized (mma) FA, add Volta support (llama/17505) 2025-12-12 17:53:16 +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: enable WMMA-MMQ INT kernels for RDNA 3 (llama/17576) 2025-12-12 17:53:17 +02:00
mmq.cuh HIP: enable WMMA-MMQ INT kernels for RDNA 3 (llama/17576) 2025-12-12 17:53:17 +02:00
mmvf.cu HIP: enable mmf for RDNA3 (llama/17879) 2025-12-17 15:19:47 +02:00
mmvf.cuh CUDA: fix crash on uneven context without FA (llama/16988) 2025-11-09 23:38:03 +02:00
mmvq.cu CUDA: Remove unneded bias/gate dims in fused mmvq (llama/16858) 2025-11-09 23:38:03 +02:00
mmvq.cuh CUDA: General GEMV fusion (llama/16715) 2025-11-09 23:38:03 +02:00
norm.cu CUDA: Optimize `rms_norm_f32` kernel and its fused variants, giving 1-6% perf E2E (llama/15715) 2025-09-20 13:42:48 +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 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 ggml : add circular tiling support to pad, for Vulkan, CUDA, and CPU (used for making seamless textures) (llama/16985) 2025-12-12 17:53:20 +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: fastdiv, launch bounds for mmvq + q8_1 quant (llama/15802) 2025-09-20 13:42:50 +03:00
quantize.cuh CUDA: batched+noncont MMQ, refactor bs>1 MoE code (llama/13199) 2025-05-01 13:29:02 +03:00
reduce_rows.cuh musa: fix build warnings (llama/15258) 2025-09-20 13:42:38 +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 CUDA: fuse rope + set_rows (llama/16884) 2025-11-17 21:05:46 +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 llama : add gpt-oss (llama/15091) 2025-08-18 20:30:45 +03: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-17 15:19:45 +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 : support LiquidAI LFM2 hybrid family (llama/14620) 2025-07-12 19:23:56 +03:00
ssm-conv.cuh ggml : faster ssm scan (llama/10558) 2025-04-02 15:51:57 +03:00
ssm-scan.cu ggml : fix SSM_SCAN for n_groups > 1 (llama/15625) 2025-09-20 13:42:43 +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 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
topk-moe.cu CUDA: support for weight clamp in top-k norm (llama/16702) 2025-11-09 23:38:03 +02:00
topk-moe.cuh CUDA: support for weight clamp in top-k norm (llama/16702) 2025-11-09 23:38:03 +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: Accelerate MXFP4 table lookup using `__byte_perm` (llama/15451) 2025-09-20 13:42:41 +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