Lumpiasty/llama.cpp
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity
Files
da2888cbab47de58ab13c01e97658655b324cefa
llama.cpp/ggml/src/ggml-cuda
T
History
Piotr Wilkin (ilintar) 305ba519ab CUDA: vectorize same-type get_rows with int4 copy (#25929) 
k_get_rows_float did a scalar one-element-per-thread copy and recomputed the
row-invariant work (index load, fast_div_modulo, src/dst row pointers) for
every element. Hoist that out of the per-element loop, and add a vectorized
path (k_get_rows_float_vec) that copies one int4 (16 B) per thread for the
contiguous same-type (no-cast) case.

The vectorized path is gated at compile time (is_same<src0_t, dst_t>) and at
runtime on 16-byte alignment of the base pointers and all row strides and on
ne00 % VEC == 0. Vectorizing divides the block count by VEC, so a small
single-row gather can drop below the device CU count and regress; an
occupancy gate keeps those on the block-rich scalar path.

On Strix Halo (gfx1151) the DeltaNet recurrent-state gather (ne00=524288)
drops 18.6us -> 13.0us (rocprofv3 HW timestamps), faster than the Vulkan
backend, with no regression on the small conv-state gather; total get_rows
-27%. test-backend-ops GET_ROWS passes (47/47).

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