cuBLAS: refactor and optimize f16 mat mul performance by slaren · Pull Request #1259 · ggml-org/llama.cpp

slaren · 2023-04-30T22:06:50Z

Moves all the cuBLAS specific code from ggml.c to ggml-cuda.cu. This also makes ggml-cuda.h much simpler, since fewer definitions have to exposed now.

Additionally, improves mat mul performance by using multiple stream where possible (when multiplying 3 or 4-dimensional tensors), and by choosing between doing f16 x f32 mat muls either as f16 x f16 or as f32 x f32, depending on what requires less data to be transferred to the GPU.

Overall, improves perplexity times with cuBLAS by ~15%.

`🤖 Generated by Copilot at 4e54943`

Summary

🚀🧹🛠️

This pull request improves the performance, compatibility, and readability of the GGML library and the llama model loader. It refactors the CUDA and BLAS code, simplifies the error checking and memory management, and exposes some useful functions and macros. The main files affected are ggml-cuda.h, ggml.c, ggml.h, llama-util.h, and llama.cpp.

ggml refactored
CUDA and BLAS streamlined
Winter of llama

Walkthrough

Refactored the code for using cuBLAS for matrix multiplication in GGML, by moving the CUDA-related functions and macros to ggml-cuda.h and calling them from ggml.c with conditional compilation (link, link, link, link, link, link, link, link, link, link, link, link, link, link, link, link, link)
Exposed the functions for converting between half-precision and single-precision floating-point numbers as part of the GGML API, by adding their declarations to ggml.h and removing them from ggml.c (link, link)
Moved the macro for asserting conditions from ggml.c to ggml.h, to make it available for other source files that use the GGML library (link, link)
Improved the code style and quality in ggml.c, by removing unused variables, empty lines, and redundant conditional compilation (link, link, link, link, link, link)

From #1233:

Enhanced the llama_buffer and llama_ctx_buffer structs in llama-util.h, by adding default constructors and disabling copy and move constructors and assignment operators, to prevent memory leaks or errors (link, link, link)
Optimized the initialization of temporary buffers in the llama_model_loader struct in llama.cpp, by using the constructor of the std::vector instead of the resize method (link)

slaren · 2023-04-30T22:34:07Z

Exposed the functions for converting between half-precision and single-precision floating-point numbers as part of the GGML API, by adding their declarations to ggml.h and removing them from ggml.c (link, link)

Specifically, this adds vector versions of ggml_fp16_to_fp32 and ggml_fp32_to_fp16. ggml_fp32_to_fp16_row is vectorized with F16C. This was necessary as GGML_FP32_TO_FP16 isn't visible from ggml-cuda.cu, and ggml_fp32_to_fp16 is too slow without inlining.

Moved the macro for asserting conditions from ggml.c to ggml.h, to make it available for other source files that use the GGML library (link, link)

GGML_ASSERT is now exposed in ggml.h, I did this to be able to use it from ggml-cuda.cu, but if this is not desirable I can remove it.

…l_mat_f16

ggerganov · 2023-05-01T12:00:49Z

-    __half m;               // min
+    half d;                 // delta
+    half m;                 // min
    uint32_t qh;            // 5-th bit of quants


At some point, should sync the CUDA block_q5_1 with the CPU one:

https://github.com/ggerganov/llama.cpp/blob/c0335b51f959ddf8c7b58bf497d10e5dc4730267/ggml.c#L736

I am not entirely sure why this isn't the case already, did you have any problems with alignment or anything else?

I updated it in the same way as q5_0 and didn't notice any issues.

For Q5_1 it works both ways.
For Q5_0, the uint32_t way does not work due to alignment issues, so we changed Q5_1 to uint8_t[4] for consistency

* cuBLAS: refactor, convert fp16 to fp32 on device * cuBLAS: use multiple streams, choose smartly between mul_mat_q and mul_mat_f16 * fix build * cuBLAS: update block_q5_1

hrx: name the two backend-ops suite failures that are not the backend's, and refuse the router launch the kernel cannot take (engine ggml-org#1259)

…ng (#220) f148781 (engine ggml-org#681) turned the ctx_other refusal into a warning and passed `model.arch_name()` to LLAMA_LOG_WARN's printf-style format. `llama_model::arch_name()` returns `std::string` (src/llama-model.h:742), so the variadic call passes a non-POD object where the `%s` expects `char *`: src/llama-context.cpp:160:42: error: cannot pass non-trivial object of type 'std::string' to variadic function; expected type from format string was 'char *' [-Wnon-pod-varargs] FAILED: src/CMakeFiles/llama.dir/llama-context.cpp.o That is a hard build failure on the branch tip (clang 24 / gfx1151 build of test-backend-ops and llama-server), not a warning: nothing on 1bit/hrx-vulkan-patched compiles without this line. `.c_str()` is valid here because the temporary lives for the whole full-expression, which is the duration of the log call. Found while building the suite for engine ggml-org#1259 (ask 1: re-run the full suite at pin 710c7c2) - the run could not start until the pin built.

…T when no dispatch covers it (engine ggml-org#1259) (#222) Two bugs of the same shape as the 32-expert router reject that #221 fixed. 1. The matchers accepted the kernel's own upper bound. Every kernel in kernel-corpus/kernels/loom-libs/ops/small_rows_f32.loom carries `index.assume %column_count [range(%column_count, 1, LIMIT)]`, and a Loom range is exclusive of its upper bound, so the kernel refuses LIMIT itself: common.softmax_rows_f32 matcher `<= 4096` kernel range [1, 4096) common.sum_rows_f32 matcher `<= 4096` kernel range [1, 4096) common.argsort_rows_f32 matcher `<= 1024` kernel range [1, 1024), and range(%rank, 0, 1023) common.norm_rows_f32 matcher `<= 65536` kernel range [1, 65536) common.get_rows_small_f32 matcher `<= 65536` kernel range [1, 65536) x4 A shape at exactly the limit was therefore claimed and then rejected at prepare time instead of falling back. The bounds are now exclusive, and get_rows_small also bounds id_stride, which its kernel assumes is < 1048576 and the matcher did not bound at all. 2. ARGSORT was claimed unconditionally by device_supports_op, but the only registration that produces ids is common.argsort_rows_f32. So any ARGSORT shape that matcher declines - and, in particular, the one switched off with GGML_HRX_DISABLE_DISPATCH for a measurement - stayed claimed with no dispatch and failed the WHOLE graph: graph_compute: unsupported HRX node 2: ARGSORT output=3:i32[31,22,1,1] <-ARGSORT inputs=[1:f32[31,22,1,1]<-SOFT_MAX] instead of splitting to the CPU. device_supports_op now asks small_rows_argsort_tensor_supported(), the same bounded-claim pattern TOP_K already uses (hip_top_k_tensor_supported). This also repairs every ARGSORT with more than 1024 columns, which was claimed, declined and failing the graph before. Effect on the suite (pin ced49c6, full `test-backend-ops test -b HRX0`, the TOPK_MOE waiver switched off so the ids are actually compared, one claim): before fix, nothing disabled : 4279/4599, [VIEW] 320, every other node OK after fix, nothing disabled : 4275/4595, [VIEW] 320, every other node OK (no regression) after fix, argsort disabled : 4233/4233 PASS, [VIEW] 0, 0 mismatch dumps The totals move by 4 because four ARGSORT cases with ne[0] >= 1024 used to be claimed and are now honestly declined, so the harness reports them as not supported instead of running them. The second line is the point: with the argsort registration switched off, the ids now fall back to the CPU and the whole suite passes, where before the same arm died on "unsupported HRX node ARGSORT" and only looked like 40 failures - which is what made the debug switch unusable for measurements like the one in engine ggml-org#1259. The kernels themselves are untouched: the ids they produce on device are correct (the routed weights gathered FROM those ids match the CPU in all 320 cases), so nothing here "fixes" a store. The remaining TOPK_MOE difference is the host readback of a non-contiguous VIEW of the argsort output, which is what #221's waiver names.

…ngine ggml-org#1259) gpu-kernels PR #60 fixes hrx_hc_post_comb: it indexed comb as dst*hc + src while ggml stores it as dst + hc*src, so identity=0 multiplied the residual by the transposed comb. With that kernel built, all six DSV4_HC_POST cases pass (identity=0 ERR 0.078-0.217 -> OK). Remove the named waiver; the TOPK_MOE view waiver remains. This must ride the gpu-kernels pin bump. Assisted-by: DeepSeek Harness

…K_MOE waiver (engine ggml-org#1259) (#228) The backend-ops suite compares two values of TOPK_MOE: the routed weights and the expert ids (`selected_experts`, a non-contiguous VIEW of the ARGSORT output). The weights matched everywhere; the ids came back as the ARGSORT **input**'s float bit patterns (`0.202668f` -> 1.045e9 read as int32). The device had the right ids - the weights are gathered from them on device and match - so the value was being read from storage a later dispatch had already reused. HRX already had this concept for one caller: `zha-readback.` + GGML_TENSOR_FLAG_OUTPUT made a fused match decline a covered node (engine ggml-org#1078) so the value exists outside the fused program. This generalises it: - dispatch-registry: accept a second, caller-neutral prefix `hrx-readback.` next to `zha-readback.` (same rule: flag set and prefix match; nothing else changes). - test-backend-ops: declare every value the test compares (`fusion_test_nodes()` / the whole-graph output) as a readback value, generically, by prefixing its name and setting the flag. That is exactly what the suite means by comparing a node, and it applies to every op, not only this one. - test-topk-moe: the `waive_node` override that skipped the ids on HRX is deleted; they are compared like every other backend compares them. Measured at pin 765fa1c with the add-on pin fd69590c, full `test-backend-ops test -b HRX0`: before (waiver deleted, no marker): 4281/4601, rc=1, `[VIEW]` 320, twice after (marker + waiver deleted): 4601/4601, rc=0, `[VIEW]` 0, 0 waivers, twice Fixes ggml-org#1259

…ot turn on a last bit (engine ggml-org#1259) `test_topk_moe` compares the expert ids it selected with the CPU reference's, as an exact index set. With random logits in [-1,1] and 256 experts over 6 routes, the 6th and 7th probabilities are within a few ULPs often enough to matter, and then the two producers legitimately disagree: the CPU gathers `probs` from one softmax and sorts it, the HRX path computes the same quantity in another place, and a last-bit difference flips their order. The ids differ, the weights differ by ~1e-7 (inside this test's nmse tolerance), so the report is "[VIEW] ids only, tiny ERR" - which is exactly what the row at pin a3e8d42 looked like: `TOPK_MOE(ne=[256,22,1,1],n_expert_used=6,with_norm=1,bias_probs=0,gating_func=3)`, `[VIEW] ERR = 2.5e-4`, ~2 full-suite runs in 9, and 12/12 clean when the case is run alone. Measured, not asserted (simulation of this test's own data, numpy float32, 4000 draws per point): exact tie at the k-th boundary, gating_func=3 0.05 % per case -> 0.16 % per full-suite run (cannot explain it) adjacent gap <= 4e-7 relative, gating_func=3 0.12 % per case -> 33 % per full-suite run (matches 2/9) adjacent gap <= 4e-7 relative, softmax 0.10 % per case -> 27 % per full-suite run So the first explanation I wrote for this row - "exact ties, and the reference's tie order is undefined because `ggml_compute_forward_argsort_f32` sorts with an unstable, value-only comparator" - is arithmetically too rare to be the cause (it stays true, it is just not what happens here), and this commit does not rest on it. The fill keeps every other tensor as the base class fills it and gives the logits exactly representable values 2^-9 apart, shuffled per row so the order is not monotone. Adjacent probabilities then differ by ~1e-3, three orders above the few-ULP window, so an id can only differ if a producer is genuinely wrong - which is what the comparison is for. Nothing is skipped, waived or tolerance-widened. Refs ggml-org#1259

…ot turn on a last bit (engine ggml-org#1259) (#238) `test_topk_moe` compares the expert ids it selected with the CPU reference's, as an exact index set. With random logits in [-1,1] and 256 experts over 6 routes, the 6th and 7th probabilities are within a few ULPs often enough to matter, and then the two producers legitimately disagree: the CPU gathers `probs` from one softmax and sorts it, the HRX path computes the same quantity in another place, and a last-bit difference flips their order. The ids differ, the weights differ by ~1e-7 (inside this test's nmse tolerance), so the report is "[VIEW] ids only, tiny ERR" - which is exactly what the row at pin a3e8d42 looked like: `TOPK_MOE(ne=[256,22,1,1],n_expert_used=6,with_norm=1,bias_probs=0,gating_func=3)`, `[VIEW] ERR = 2.5e-4`, ~2 full-suite runs in 9, and 12/12 clean when the case is run alone. Measured, not asserted (simulation of this test's own data, numpy float32, 4000 draws per point): exact tie at the k-th boundary, gating_func=3 0.05 % per case -> 0.16 % per full-suite run (cannot explain it) adjacent gap <= 4e-7 relative, gating_func=3 0.12 % per case -> 33 % per full-suite run (matches 2/9) adjacent gap <= 4e-7 relative, softmax 0.10 % per case -> 27 % per full-suite run So the first explanation I wrote for this row - "exact ties, and the reference's tie order is undefined because `ggml_compute_forward_argsort_f32` sorts with an unstable, value-only comparator" - is arithmetically too rare to be the cause (it stays true, it is just not what happens here), and this commit does not rest on it. The fill keeps every other tensor as the base class fills it and gives the logits exactly representable values 2^-9 apart, shuffled per row so the order is not monotone. Adjacent probabilities then differ by ~1e-3, three orders above the few-ULP window, so an id can only differ if a producer is genuinely wrong - which is what the comparison is for. Nothing is skipped, waived or tolerance-widened. Refs ggml-org#1259

slaren marked this pull request as draft April 30, 2023 22:13

slaren added 3 commits May 1, 2023 13:33

cuBLAS: refactor, convert fp16 to fp32 on device

cf93fdc

cuBLAS: use multiple streams, choose smartly between mul_mat_q and mu…

a9ad140

…l_mat_f16

fix build

4cd0a48

slaren force-pushed the cuda-mat-mul branch from 4fd9d4a to 4cd0a48 Compare May 1, 2023 11:35

slaren marked this pull request as ready for review May 1, 2023 11:38

ggerganov reviewed May 1, 2023

View reviewed changes

Comment thread ggml.c Resolved

slaren mentioned this pull request May 1, 2023

Generalize quantize_fns for simpler FP16 handling #1237

Merged

ggerganov approved these changes May 1, 2023

View reviewed changes

cuBLAS: update block_q5_1

a79756b

slaren merged commit 58b367c into ggml-org:master May 1, 2023

slaren deleted the cuda-mat-mul branch May 1, 2023 16:11

Bearsaerker mentioned this pull request Mar 12, 2025

Eval bug: Gemma 3 extremly slow prompt processing when using quantized kv cache. #12352

Closed

bong-water-water-bong mentioned this pull request Nov 6, 2026

hrx: keep the small-rows matcher bounds exclusive, and decline ARGSORT when no dispatch covers it (engine #1259) 1bit-MONSTER/llama.cpp#222

Merged

bong-water-water-bong mentioned this pull request Nov 7, 2026

hrx: let a caller declare a value it will read back, and drop the TOPK_MOE waiver (engine #1259) 1bit-MONSTER/llama.cpp#228

Merged

bong-water-water-bong mentioned this pull request Nov 8, 2026

test-backend-ops: space TOPK_MOE's logits, so the ids comparison cannot turn on a last bit (engine #1259) 1bit-MONSTER/llama.cpp#238

Merged

Navigation Menu

Repository navigation

cuBLAS: refactor and optimize f16 mat mul performance - #1259

cuBLAS: refactor and optimize f16 mat mul performance#1259
slaren merged 4 commits into
ggml-org:masterfrom
slaren:cuda-mat-mul

slaren commented Apr 30, 2023 •

edited

Loading

Uh oh!

slaren commented Apr 30, 2023 •

edited

Loading

Uh oh!

Uh oh!

ggerganov May 1, 2023

Uh oh!

slaren May 1, 2023

Uh oh!

slaren May 1, 2023

Uh oh!

ggerganov May 1, 2023

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Conversation

slaren commented Apr 30, 2023 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

🤖 Generated by Copilot at 4e54943

Summary

Walkthrough

Uh oh!

slaren commented Apr 30, 2023 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Uh oh!

ggerganov May 1, 2023

Choose a reason for hiding this comment

Uh oh!

slaren May 1, 2023

Choose a reason for hiding this comment

Uh oh!

slaren May 1, 2023

Choose a reason for hiding this comment

Uh oh!

ggerganov May 1, 2023

Choose a reason for hiding this comment

Uh oh!

Reviewers

Uh oh!

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

slaren commented Apr 30, 2023 •

edited

Loading

`🤖 Generated by Copilot at 4e54943`

slaren commented Apr 30, 2023 •

edited

Loading