whisper.cpp/ggml/src/ggml-cuda
Sam/Samuel d8a146b0f9 metal: optimise `GGML_OP_SUM` (llama/16559)
* optimise GGML_OP_SUM

* add non-contiguous tests by permuting the input

* change tests to require full contiguity of OP_SUM

* cuda : add check GGML_OP_SUM

---------

Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
2025-10-22 12:58:11 +03:00
..
template-instances CUDA: faster tile FA, add oob checks, more HSs (llama/16492) 2025-10-15 09:29:17 +03:00
vendors HIP: Disable ROCWMMA fattn on CDNA when compiled against ROCWMMA 2.0.0 (llama/16221) 2025-10-12 11:16:23 +03:00
CMakeLists.txt CUDA: faster tile FA, add oob checks, more HSs (llama/16492) 2025-10-15 09:29:17 +03:00
acc.cu
acc.cuh
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
arange.cuh
argmax.cu
argmax.cuh
argsort.cu
argsort.cuh
binbcast.cu musa: fix build warnings (llama/15611) 2025-09-29 15:18:10 +03:00
binbcast.cuh CUDA: fuse adds, fuse add with rms norm (llama/15631) 2025-09-20 13:42:44 +03:00
clamp.cu
clamp.cuh
common.cuh cuda : remove legacy copy-op pointer indirection code (llama/16485) 2025-10-15 09:29:17 +03:00
concat.cu
concat.cuh
conv-transpose-1d.cu musa: add GGML_UNUSED_VARS (llama/15446) 2025-09-20 13:42:38 +03:00
conv-transpose-1d.cuh
conv2d-dw.cu
conv2d-dw.cuh
conv2d-transpose.cu
conv2d-transpose.cuh
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 ggml: allow casting between f32 and i32 (llama/15783) 2025-09-20 13:42:51 +03:00
count-equal.cu
count-equal.cuh
cp-async.cuh
cpy-utils.cuh HIP: Cleanup hipification header (llama/15285) 2025-08-18 20:30:45 +03:00
cpy.cu cuda : remove legacy copy-op pointer indirection code (llama/16485) 2025-10-15 09:29:17 +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
cross-entropy-loss.cuh
dequantize.cuh CUDA: replace GGML_CUDA_F16 with CUDA arch checks (llama/15433) 2025-09-20 13:42:38 +03:00
diagmask.cu
diagmask.cuh
fattn-common.cuh CUDA: faster tile FA, add oob checks, more HSs (llama/16492) 2025-10-15 09:29:17 +03:00
fattn-mma-f16.cuh musa: add GGML_UNUSED_VARS (llama/15446) 2025-09-20 13:42:38 +03:00
fattn-tile.cu CUDA: faster tile FA, add oob checks, more HSs (llama/16492) 2025-10-15 09:29:17 +03:00
fattn-tile.cuh CUDA: fix numerical issues in tile FA kernel (llama/16540) 2025-10-15 09:29:17 +03:00
fattn-vec.cuh CUDA: enable FA for FP32 KV cache (llama/16546) 2025-10-15 09:29:17 +03:00
fattn-wmma-f16.cu HIP: Disable ROCWMMA fattn on CDNA when compiled against ROCWMMA 2.0.0 (llama/16221) 2025-10-12 11:16:23 +03:00
fattn-wmma-f16.cuh CUDA: faster tile FA, add oob checks, more HSs (llama/16492) 2025-10-15 09:29:17 +03:00
fattn.cu CUDA: enable FA for FP32 KV cache (llama/16546) 2025-10-15 09:29:17 +03:00
fattn.cuh CUDA: refactor FA support/selection code (llama/15454) 2025-09-20 13:42:38 +03:00
getrows.cu CUDA: fix GET_ROWS for large tensors (llama/15882) 2025-09-20 13:42:52 +03:00
getrows.cuh
ggml-cuda.cu metal: optimise `GGML_OP_SUM` (llama/16559) 2025-10-22 12:58:11 +03:00
gla.cu
gla.cuh
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
mma.cuh CUDA: Add mul_mat_id support for the mmf kernel (llama/15767) 2025-09-20 13:42:52 +03:00
mmf.cu CUDA: add fp kernel for larger batch size MoE (llama/16512) 2025-10-15 09:29:17 +03:00
mmf.cuh CUDA: add fp kernel for larger batch size MoE (llama/16512) 2025-10-15 09:29:17 +03: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 CUDA: add fp kernel for larger batch size MoE (llama/16512) 2025-10-15 09:29:17 +03:00
mmq.cuh CUDA: MoE helper in device code, better tile sizes (llama/15525) 2025-09-20 13:42:41 +03:00
mmvf.cu CUDA: use fastdiv + ggml_cuda_mad for mmvf (llama/16557) 2025-10-15 09:29:17 +03:00
mmvf.cuh CUDA: GEMM for FP32/FP16/BF16 and ne11 <= 16 (llama/15131) 2025-08-18 20:30:45 +03:00
mmvq.cu musa: fix build warnings (llama/15611) 2025-09-29 15:18:10 +03:00
mmvq.cuh
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
opt-step-adamw.cuh
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
out-prod.cuh
pad.cu ggml: add ops for WAN video model (cuda && cpu) (llama/15669) 2025-09-20 13:42:49 +03:00
pad.cuh
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
pool2d.cuh
quantize.cu CUDA: fastdiv, launch bounds for mmvq + q8_1 quant (llama/15802) 2025-09-20 13:42:50 +03:00
quantize.cuh
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
rope.cuh
scale.cu ggml: add ops for WAN video model (cuda && cpu) (llama/15669) 2025-09-20 13:42:49 +03:00
scale.cuh
set-rows.cu ggml : implement set_rows with i32 index (llama/16159) 2025-09-29 15:18:09 +03:00
set-rows.cuh CUDA: add set rows for f32 and f16 (llama/14551) 2025-07-20 00:23:50 +03: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
ssm-conv.cu
ssm-conv.cuh
ssm-scan.cu ggml : fix SSM_SCAN for n_groups > 1 (llama/15625) 2025-09-20 13:42:43 +03:00
ssm-scan.cuh
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
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
topk-moe.cu musa: update compile flags (llama/16265) 2025-10-12 11:16:23 +03:00
topk-moe.cuh CUDA: add a fused top-K MoE kernel (llama/16130) 2025-09-29 15:18:10 +03:00
tsembd.cu ggml : fix padding in timestep embedding kernels (llama/15932) 2025-09-20 13:45:30 +03:00
tsembd.cuh
unary.cu model : Apertus model implementation (llama/15852) 2025-10-12 11:16:23 +03:00
unary.cuh model : Apertus model implementation (llama/15852) 2025-10-12 11:16:23 +03:00
upscale.cu
upscale.cuh
vecdotq.cuh CUDA: Accelerate MXFP4 table lookup using `__byte_perm` (llama/15451) 2025-09-20 13:42:41 +03:00
wkv.cu
wkv.cuh