nikola/llama_cpp
1
0
Fork 0
Code Issues Pull Requests Actions 18 Packages Projects Releases Wiki Activity
8,722 Commits 1 Branch 0 Tags
8a132faaa06d904bbeb3ebfbc8f7e2ed921b8ded
Commit Graph
100 Commits
This Branch
This Branch
All Branches
Author SHA1 Message Date
Aman GuptaandGitHub d12cc3d1ca CUDA: also store node->src->data ptrs for equality check (#21635) 
* CUDA: also store node->src->data ptrs for equality check

* address review comments
 2026-04-09 01:01:56 +08:00
Aman GuptaandGitHub c5ce4bc227 CUDA: make cuda graphs props check faster (#21472) 
* CUDA: compute fast hash instead of expensive props check

* use seen node

* use memcp
 2026-04-08 09:05:51 +08:00
Aman GuptaandGitHub de1aa6fa73 CUDA: check for buffer overlap before fusing (#21566) 
* CUDA: check for buffer overlap before fusing

* use ggml_cuda_check_fusion_memory_ranges
 2026-04-08 00:57:04 +08:00
Aman GuptaandGitHub 94ca829b60 llama-bench: add -fitc and -fitt to arguments (#21304) 
* llama-bench: add `-fitc` and `-fitt` to arguments

* update README.md

* address review comments

* update compare-llama-bench.py
 2026-04-06 22:26:02 +08:00
Aman GuptaandGitHub b7ad48ebda llama: add custom newline split for Gemma 4 (#21406) 2026-04-04 15:06:34 +08:00
Aman GuptaandGitHub 278521c33a llama-model-loader: print warning when using overrides with mmap (#20978) 
* llama-model-loader: use pinned memory for tensor overrides

* change to warning
 2026-03-30 17:40:17 +08:00
Aman GuptaandGitHub 9c600bcd4b llama-bench: print -n-cpu-moe when offloaded layers > 1 (#20984) 2026-03-25 21:17:27 +08:00
Aman GuptaandGitHub 3fc6f1aed1 ggml-backend: re-enable graph reuse with pipeline parallelism (#20927) 2026-03-24 20:47:00 +08:00
Aman GuptaGitHubJohannes Gäßler
 		e852eb4901 llama-fit: fix regex pattern for gate_up tensors (#20910) 
* llama-fit: fix regex pattern for gate_up tensors

* Apply suggestions from code review

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

---------

Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
 2026-03-24 12:57:57 +08:00
Aman GuptaandGitHub 3c8521c4f5 llama-graph: replace cont with reshape for alpha in qwen35 (#20640) 2026-03-16 22:07:13 +08:00
Aman GuptaandGitHub 34818ea6c0 CUDA: GDN hide memory latency (#20537) 2026-03-16 11:41:45 +08:00
Aman GuptaandGitHub bd1ec818e9 compare-llama-bench: check remotes as well (#20406) 2026-03-12 00:14:42 +08:00
Aman GuptaandGitHub e8bbc736cb ggml-cuda: disable gdn for musa (#20278) 2026-03-09 16:15:36 +08:00
Aman GuptaandGitHub e2763a6723 contributing: limit open PRs for new contributors to 1 (#20036) 2026-03-09 15:05:34 +08:00
Aman GuptaGitHubGeorgi Gerganov
 		c5a778891b ggml: add GATED_DELTA_NET op (#19504) 
* ggml: add GATED_DELTA_NET op

* remove the transpose

* add KDA

* add qwen35 dense

* llama : check for fused gated delta net backend support

---------

Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
 2026-03-07 15:41:10 +08:00
Aman GuptaGitHubJohannes Gäßler
 		d48e876467 ggml-cuda: add mem check for fusion (#19916) 
* ggml-cuda: add mem check for fusion

* Replace NaNs with -FLT_MAX

* fix typo

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

---------

Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
 2026-03-07 00:05:43 +08:00
Aman GuptaGitHubJohannes Gäßler
 		1e38a7a6fa CUDA: use shared mem for ssm_conv (#20128) 
* CUDA: use shared mem for ssm_conv

* fuse silu + ssm_conv

* fuse unary + mul

* enable for fp16

* formatting

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

---------

Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
 2026-03-06 23:09:59 +08:00
Aman GuptaandGitHub d903f30e25 ggml-cpu: add repack for mxfp4 (#19738) 2026-02-27 18:15:09 +08:00
Aman GuptaGitHubSigbjørn Skjæret
 		b68d75165a llama: Add option to merge gate and exp weights (#19139) 
* llama: Add option to merge gate and exp weights

* Update convert_hf_to_gguf.py

Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com>

* Update convert_hf_to_gguf.py

Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com>

* update constants.py

* add gate_up for the all MoE models

* convert: simplify merge tensor condition

* update constants.py

* reduce number of models, add create_tensor_gate_up helper

---------

Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com>
 2026-02-26 21:01:08 +08:00
Aman GuptaandGitHub 684b36101c ggml-cpu: FA add GEMM microkernel (#19422) 
* ggml-cpu: FA add GEMM microkernel

* add guard for sizeless vector types

* fix case where DV % GGML_F32_EPR !=0

* move memset out of the loop

* move another memset out of the loop

* use RM=4 for arm

* simd_gemm: convert everything to int

* convert everything to size_t to avoid warnings

* fixup

* add pragma for ignoring aggressive loop optimizations
 2026-02-15 11:09:24 +05:30
Aman GuptaandGitHub 5065da554e CUDA: loop over ne2*ne3 in case it overflows (#19538) 
* CUDA: loop over ne2*ne3 in case it overflows

* use fastdiv
 2026-02-13 17:01:40 +05:30
Aman GuptaandGitHub 2ceda3f662 ggml-cpu: use LUT for converting e8->f32 scales on x86 (#19288) 
* ggml-cpu: use LUT for converting e8->f32 scales on x86

* add dispatch based on macro
 2026-02-04 09:43:29 +08:00
Aman GuptaandGitHub 8bece2eb20 CUDA: use mmvq for mul-mat-id for small batch sizes (#18958) 
* CUDA: use mmvq for mul-mat-id for small batch sizes

* add mmvq too

* Fix perf issue on ampere. Use mmvf mm-id only for non-nvidia GPUs

* templatize multi_token_path
 2026-02-03 23:31:23 +08:00
Aman GuptaandGitHub 9f682fb640 ggml-cpu: FA split across kv for faster TG (#19209) 
* ggml-cpu: split across kv for faster TG

* simplify sinks application

* add ref impl
 2026-02-03 01:19:55 +08:00
Aman GuptaandGitHub 15818ac44c ci: add test-backend-ops test for CPU (#19268) 2026-02-02 22:40:28 +08:00
Aman GuptaandGitHub 3bcc990997 CUDA: refactor topk-moe to enable more models (GLM 4.7, Nemotron etc.) (#19126) 2026-01-29 10:31:28 +08:00
Aman GuptaandGitHub bcb43163ae ggml-cpu: Use tiled FA for prompt-processing (#19012) 
* ggml-cpu: Use tiled FA for prompt-processing

the FA performance is gimped on CPU on long contexts because it essentially uses a vector kernel. This PR adds a tiled FA for PP. Perf tuning for tile sizes done on a AMD EPYC single-socket 64-c machine.

* fix out of bounds for mask

* skip rows where there are all masks

* skip tile if mask is inf

* store mask in worksize

* check inf tile earlier
 2026-01-25 23:25:58 +08:00
Aman GuptaandGitHub 81ab64f3c8 ggml-cuda: enable cuda-graphs for n-cpu-moe (#18934) 
* ggml-cuda: add split-wise cuda graph

* add n-cpu-moe compare_llama_bench.py

* fix hip/musa builds
 2026-01-24 14:25:20 +08:00
Aman GuptaandGitHub e34d6d03b2 convert_hf_to_gguf.py: refactor modify_tensors to call super (#18866) 2026-01-23 02:58:07 +08:00
Aman GuptaandGitHub b70d251076 CUDA: add gqa_ratio 4 for GLM 4.7 flash (#18953) 2026-01-22 18:51:53 +08:00
Aman GuptaandGitHub 47f9612492 llama-model: fix unfortunate typo (#18832) 2026-01-14 17:55:15 +08:00
Aman GuptaandGitHub b137718878 test-backend-ops: fix mxfp4 tests on blackwell (#18736) 2026-01-11 01:12:57 +08:00
Aman GuptaandGitHub 090b137e56 ggml-cuda: refactor cuda graph usage (#18637) 
* ggml-cuda: refactor cuda graph usage

* use is_enabled() instead of enabled
 2026-01-06 23:48:45 +08:00
Aman GuptaandGitHub f47edb8c19 ggml-cuda: check for srcs outside the cgraph (#18583) 
* ggml-cuda: check for srcs outside the cgraph

* review: use leafs instead
 2026-01-05 22:46:36 +08:00
Aman GuptaandGitHub 908a9e5a1e CUDA: disable cuda graph when using n-cpu-moe (#18593) 
* CUDA: disable cuda graph when using n-cpu-moe

* call ggml_cuda_set_device
 2026-01-05 01:37:48 +08:00
Aman GuptaandGitHub 5126c41c1c ggml-cuda: remove unused params in ggml_cuda_graph (#18579) 2026-01-05 01:37:09 +08:00
Aman GuptaandGitHub e57f52334b ggml-cuda: fixes for concurrent streams (#18496) 2026-01-03 23:15:01 +08:00
Aman GuptaandGitHub 26831bded9 ggml-cuda: remove unneccesary prints on ggml_cuda_init (#18502) 2026-01-01 19:18:43 +08:00
Aman GuptaandGitHub 7bcaf815c2 sycl: add newline at the end of CMakeLists.txt (#18503) 2025-12-31 14:23:44 +08:00
Aman GuptaandGitHub d77d7c5c06 CUDA: add log line when mxfp4 acceleration is used (#18483) 
* CUDA: add log line when mxfp4 acceleration is used

* add in backend_get_features
 2025-12-30 17:40:46 +08:00
Aman GuptaandGitHub 5fa66c6e67 cuda: fix race condition in cumsum (#18448) 
* ggml-cuda: fix race condition in cumsum

* remove unneccesary sync_threads
 2025-12-29 14:07:17 +08:00
Aman GuptaandGitHub 07a0c4ba92 Revert "ggml-cuda: use CMAKE_CUDA_ARCHITECTURES if set when GGML_NATIVE=ON (#18413)" (#18426) 2025-12-28 20:53:36 +08:00
Aman GuptaandGitHub 06705fdcb3 ggml-cuda: Use same regex for GGML_NATIVE=OFF (#18407) 2025-12-27 19:56:27 +08:00
Aman GuptaandGitHub 85c40c9b02 ggml-cuda: fix regex for arch list (#18371) 
* ggml-cuda: fix regex for arch list

* make regex exact
 2025-12-26 01:35:14 +08:00
Aman GuptaandGitHub 83b3b1c271 cuda: optimize cumsum cub path (#18362) 
* cuda: optimize cumsum cub path

* remove heavy perf test
 2025-12-25 23:55:38 +08:00
Aman GuptaGitHubAman Gupta <aman>
 		b0fb0f0aee ggml-cuda: fix blackwell native builds (#18361) 
* ggml-cuda: fix blackwell native builds

Replace 12x in native architectures by 12xa

* replace for GGML_NATIVE=OFF too

* only replace for native

* remove 120f-virtual for default compilation

---------

Co-authored-by: Aman Gupta <aman>
 2025-12-25 22:12:11 +08:00
Aman GuptaGitHubAman Gupta <aman>
 		c8a2417d7b CUDA: experimental native mxfp4 support for blackwell (#17906) 
* CUDA: experimental native mxfp4 support for blackwell

* optimize load_tiles

* optimize quantize_mxfp4

* cleanup

* first pass review: formatting

* use interleaved layout for mma

* mmq: add assert for size

* use __nv_fp4x4_e2m1

* use iter_k as 512, cleanup

* Use 1200 as blackwell instead of 1000

* address review comments

* mmq: fix stride

* quantize.cu: use reference impl of e8m0 scale

* address review comments

* add 120f-virtual + minor fixes

---------

Co-authored-by: Aman Gupta <aman>
 2025-12-24 22:28:26 +08:00
Aman GuptaandGitHub cc0a04343e server: friendlier error msg when ctx < input (#18174) 
* llama-server: friendlier error msg when ctx < input

This PR adds formatted strings to the server's send_error function

* llama-server: use string_format inline

* fix test
 2025-12-19 12:10:00 +01:00
Aman GuptaandGitHub ed32089927 ggml-cuda: reorder only relevant nodes (#17639) 2025-12-02 12:36:31 +08:00
Aman GuptaandGitHub 6eea666912 llama-graph: avoid expand_forward for fusion (#17633) 2025-12-01 11:12:48 +02:00
Aman GuptaGitHubJohannes GäßlerCarl Philipp Klemm
 		c7af376c29 CUDA: add stream-based concurrency (#16991) 
* CUDA: add stream-based concurrency

* HIP: fix hipStreamWaitEvent define and nodiscard warnings

* ggml-cuda: fix fusion inside stream

* ggml-cuda: fix bug w.r.t first stream launch

* ggml-cuda: format

* ggml-cuda: improve assert message

* ggml-cuda: use lambda instead of duplicating code

* ggml-cuda: add some more comments

* ggml-cuda: add more detailed comments about concurrency

* ggml-cuda: rename + remove unused var

* ggml-cuda: fix condition for stream launch

* ggml-cuda: address review comments, add destructor

* common.cuh: add is_valid for concurrent events

* common.cuh: make comment better

* update comment

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

* update comment

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

* common.cuh: fix lower_bound condition + remove join_node data from write_ranges

* ggml-cuda: fix overlap condition + shadowing parameter

---------

Co-authored-by: Carl Philipp Klemm <carl@uvos.xyz>
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
 2025-11-30 08:17:55 +08:00
Aman GuptaandGitHub 2e7ef98f18 ggml-cuda: add stricter checking for fusion (#17568) 
* ggml-cuda: make conditions for fusion more explicit

* ggml-cuda: remove size check as std::equal already does it
 2025-11-28 20:34:51 +08:00
Aman GuptaandGitHub fd7353d5eb cuda: fix rope fusion for gemma3 (#17378) 2025-11-19 18:25:05 +08:00
Aman GuptaandGitHub a90eb94ca9 CUDA: fuse rope + set_rows (#16884) 
* CUDA: add fused rope

* move k forward_expand up

* create helper function instead of re-using params

* make assert statement more in line with comment

* rope_norm: coalesced writes to global mem
 2025-11-13 08:50:01 +08:00
Aman GuptaandGitHub 64fe17fbb8 Revert "CUDA: add expert reduce kernel (#16857)" (#17100) 2025-11-08 21:05:19 +08:00
Aman GuptaandGitHub c1b187688d CUDA: skip fusion for repeating adds in bias (#17080) 2025-11-08 16:58:05 +08:00
Aman GuptaandGitHub 2759ccdb4a CUDA: avoid mul + bias fusion when doing fusion (#16935) 2025-11-04 10:53:48 +08:00
Aman GuptaGitHubJohannes Gäßler
 		4146d6a1a6 CUDA: add expert reduce kernel (#16857) 
* CUDA: add expert reduce kernel

* contigous checks, better formatting, use std::vector instead of array

* use vector empty instead of size

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

---------

Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
 2025-10-31 20:05:07 +08:00
Aman GuptaandGitHub e41bcce8f0 CUDA: use fastdiv in set-rows (#16834) 
* CUDA: use fastdiv in set-rows

* add assert about value fitting in u32
 2025-10-29 21:11:53 +08:00
Aman GuptaandGitHub 9a3ea685b9 CUDA: Fix bug in topk-moe for gpt-oss (#16821) 
* CUDA: Fix bug in topk-moe for gpt-oss

When using ggml_can_fuse_subgraph, the output nodes which are passed are wrong. This causes `test-backend-ops` to still fuse ndoes (because the nodes are not used elsewhere in the graph),
but it actually doesn't fuse in the actual gpt-oss

* fix for qwen3 too

* change ifndef to ifdef
 2025-10-29 15:55:06 +08:00
Aman GuptaandGitHub 463bbf20bf CUDA: add unused vars to mmvf and mmvq (#16807) 2025-10-28 10:31:21 +08:00
Aman GuptaandGitHub 75cbdd3fce test-backend-ops: print failed tests at the end (#16785) 2025-10-27 09:25:10 +08:00
Aman GuptaandGitHub 75d33b9302 CUDA: support for weight clamp in top-k norm (#16702) 2025-10-27 09:06:16 +08:00
Aman GuptaandGitHub f77c13b91f CUDA: General GEMV fusion (#16715) 2025-10-26 19:28:04 +08:00
Aman GuptaandGitHub 0bcb40b48c CUDA: use CUB for arbitary size argsort (#16754) 2025-10-24 20:46:19 +08:00
Aman GuptaandGitHub 061f0eff02 ggml-cuda: use passed ops instead of hardcoded ops (#16712) 2025-10-23 19:14:06 +08:00
Aman GuptaandGitHub 9285325ce0 CUDA: fix bug in topk-moe softmax (#16711) 2025-10-22 12:33:08 +08:00
Aman GuptaandGitHub 03792ad936 CUDA: topk-moe: add optional parameter for gpt-oss (#16649) 2025-10-21 22:40:38 +08:00
Aman GuptaandGitHub 4926419c4d ggml: add ggml_can_fuse_subgraph (#16662) 
* ggml: add ggml_can_fuse_subgraph

* ggml-cuda: use ggml_can_fuse_subgraph for topk-moe

* format

* 1. remove inputs from signature as they are transient nodes
2. add check for views: view_src should be part of the subgraph

* - combine check into one loop
- check all view_src parents
- other minor review comments

* remove redudant if test

* - rename and other minor review comments

* add assert about count < 32
 2025-10-21 16:43:14 +08:00
Aman GuptaandGitHub 55754bebd5 CODEOWNERS: update for ggml-cuda/mmf (#16660) 2025-10-19 10:37:12 +03:00
Aman GuptaandGitHub 38355c6c8e CUDA: use registers instead of smem in topk-moe (#16647) 
Uses the technique used in the vulkan PR #16641. Neat trick!
 2025-10-18 11:52:53 +02:00
Aman GuptaandGitHub 120bf7046d CUDA + openCL: fix bug in accessing rms_norm->src while doing fusion (#16577) 2025-10-14 07:48:08 -07:00
Aman GuptaandGitHub 1ee9d0b415 CUDA: use fastdiv + ggml_cuda_mad for mmvf (#16557) 
* CUDA: use fastdiv + ggml_cuda_mad for mmvf

* use bf16 directly + fix formatting

* Add exception for HIP code
 2025-10-14 13:16:21 +02:00
Aman GuptaandGitHub 48e2fa9fb7 CUDA: add fp kernel for larger batch size MoE (#16512) 
* CUDA: kernel for larger batch sizes for MoE

* WIP

* WIP

* WIP

* WIP

* WIP

* WIP

* fixup

* tests

* Move mmq_ids_helper to mmid

* cleanup

* Remove redundant checks
 2025-10-14 13:15:15 +02:00
Aman GuptaandGitHub c0bfc57af4 CUDA: mul_mat_id for mmf for bs <= 64 for f16 and bs <= 32 for f32 (#16277) 
* CUDA: mul_mat_id for mmf for bs <= 64 for f16 and bs <= 32 for f32

This commit adds mul_mat_id support for ncols_dst >= 16. It does this by
packing ncols_dst tiles into the blockDim.y.

My tests on a RTX 3090 show that this is faster than the cuBLAS fallback
for f16 till bs=64, and for f32 till bs=32

* Review: refactor if statement
 2025-09-27 18:49:32 +02:00
Aman GuptaandGitHub 077c94d0ca CUDA: add a fused top-K MoE kernel (#16130) 
* CUDA: add a fused top-K MoE kernel

This kernel does the following:
1. softmax over the logits per token [n_experts, n_tokens]
2. argmax reduce over the top-k (n_experts_used) logits
3. write weights + ids to global memory

It is intended as fusion of softmax->top-k->get_rows pipeline for MoE models

* Refactor into ggml_cuda_should_use_topk_moe

* Review: Use better coalescing pattern, use WARP_SIZE, store logits into registers before

* Review: format + micro-optimizations

* Fix bug: fix tie breakers

* Add optional norm + clean-up code

* Use smem for final write

* Add bounds check

* Use better memory pattern for writeback
 2025-09-25 16:35:05 +02:00
Aman GuptaandGitHub 6d758839ff Add LLaDA-7b-MoE diffusion model (#16003) 2025-09-16 10:38:28 +08:00
Aman GuptaandGitHub 106220562a CUDA: some micro-optimizations in mmf.cuh for mul_mat_id (#15926) 2025-09-15 17:35:11 +08:00
Aman GuptaandGitHub a972faebed CUDA: Add mul_mat_id support for the mmf kernel (#15767) 
* CUDA: Add mul_mat_id support the mmf

Add support for mul_mat_id for bs < 16

* Review: use warp_size, fix should_use_mmf condition

* Launch one block per expert, stride along n_expert_used

* templatize mul_mat_id

* Pad shmem to 16 bytes, add helper function mul_mat_f_switch_ids

* Reduce compile times by dividing mmf into f16, bf16 and f32 variants

* Divide mmf by ncols_dst

* Add missing files

* Fix MUSA/HIP builds
 2025-09-09 14:38:02 +08:00
Aman GuptaandGitHub 0a16bf52e6 CUDA: generate_cu_files.py - add missing mxfp4 (#15880) 2025-09-09 01:23:46 +08:00
Aman GuptaandGitHub 81017865ee CUDA: fix bug in rms_norm fusion (#15660) 
* CUDA: fix bug in rms_norm fusion

* Fix bug for OP_REPEAT

* Fix index for add
 2025-08-29 21:30:06 +08:00
Aman GuptaandGitHub 009b709d6e CUDA: fuse adds, fuse add with rms norm (#15631) 
* CUDA: fused add with rms_norm_mul

* Non-broadcast fuse works

* Add fused adds

* format

* Remove n_fuse from template params

* Address review comments

* Move template inside binbcast
 2025-08-29 11:35:58 +08:00
Aman GuptaandGitHub 55042b3692 scripts: add sqlite3 check for compare-commits.sh (#15633) 2025-08-28 19:23:22 +08:00
Aman GuptaandGitHub efe3a90996 CUDA cmake: add -lineinfo for easier debug (#15260) 2025-08-12 17:21:45 +08:00
Aman GuptaandGitHub 34c9d765bf CUDA: add attention sinks for tile and wmma (#15178) 
* CUDA: add attention sinks for tile and wmma

* Review: formatting changes + remove syncthreads from tile + remove warp_reduce_max from wmma
 2025-08-09 20:00:24 +08:00
Aman GuptaandGitHub 784524053d Fix params bug in diffusion example (#14993) 2025-08-01 01:22:58 +08:00
Aman GuptaandGitHub 8a4a856277 Add LLaDA 8b Diffusion model (#14771) 
* Add support for Llada-8b: diffusion model

* Add README

* Fix README and convert_hf_to_gguf

* convert_hf_to_gguf.py: address review comments

* Make everything in a single example

* Remove model-specific sampling

* Remove unused argmax

* Remove braced initializers, improve README.md a bit

* Add diffusion specific gguf params in set_vocab, remove setting rope_theta and rms_norm_eps

* Remove adding the mask token

* Move add_add_bos_token to set_vocab

* use add_bool in gguf_writer.py
 2025-07-31 19:49:09 +08:00
Aman GuptaandGitHub 0a5036bee9 CUDA: add roll (#14919) 
* CUDA: add roll

* Make everything const, use __restrict__
 2025-07-29 14:45:18 +08:00
Aman GuptaandGitHub 446595b9b3 Docs: add instructions for adding backends (#14889) 2025-07-27 09:36:43 +08:00
Aman GuptaandGitHub 8c988fa41d CUDA: add fused rms norm (#14800) 2025-07-23 09:25:42 +08:00
Aman GuptaandGitHub 2be60cbc27 docs : fix link for tools/perplexity in README.md (#14780) 2025-07-20 20:13:47 +02:00
Aman GuptaandGitHub 938b785764 Clang-format: local files first + fix BinPacking (#14779) 2025-07-20 19:42:34 +08:00
Aman GuptaandGitHub f9a31eea06 CUDA: set_rows + cpy.cu refactor (#14712) 2025-07-18 14:54:18 +08:00
Aman GuptaandGitHub ab14019821 Support diffusion models: Add Dream 7B (#14644) 
* Support diffusion models: Add Dream 7B

* Move diffusion to examples

* Move stuff to examples. Add patch to not use kv-cache

* Address review comments

* Make sampling fast

* llama: remove diffusion functions

* Add basic timings + cleanup

* More cleanup

* Review comments: better formating, use LOG instead std::cerr, re-use batch, use ubatch instead of max_length

* fixup!

* Review: move everything to diffusion-cli for now
 2025-07-16 20:03:51 +08:00
Aman GuptaandGitHub 9c9e4fc635 llama-context: add ability to get logits (#14672) 2025-07-14 21:01:41 +08:00
Aman GuptaandGitHub 7de5c7cab6 CUDA: add set rows for f32 and f16 (#14551) 
* CUDA: add set rows for f32 and f16

* Review: change kernel params, use strides from host

* Use 1-d kernel

* Review: use int64_t for blockDim.x, rename nb->s for clarity
 2025-07-12 16:31:38 +03:00
Aman GuptaGitHubGeorgi Gerganov
 		11ee0fea2a Docs: script to auto-generate ggml operations docs (#14598) 
* Docs: script to auto-generate ggml operations docs

* Review: formatting changes + change github action

* Use built-in types instead of typing

* docs : add BLAS and Metal ops

---------

Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
 2025-07-10 23:29:01 +08:00
Aman GuptaandGitHub 75c91de6e9 CUDA: add bilinear interpolation for upscale (#14563) 2025-07-08 10:11:18 +08:00
Aman GuptaandGitHub b9c3eefde1 CUDA: add bf16 and i32 to getrows (#14529) 2025-07-07 21:45:43 +08:00
Aman GuptaandGitHub 55c2646b45 CUDA: add dynamic shared mem to softmax, refactor general usage (#14497) 2025-07-03 07:45:11 +08:00