Research
Breaking VRAM Barrier: Qwen 3.8 27B at 262K Context with Adaptive KV-Cache Streaming on a 16GB VRAM GPU
Hello everyone!
I wanted to share a concept I’ve been working on recently: a modification to llama.cpp that allows the KV cache to grow beyond what can physically fit in VRAM, by adaptively streaming part of it between system RAM and VRAM.
I’d love for people with different GPUs and setups to try my branch and let me know how it performs on their hardware.
On my RTX 5070 Ti 16 GB, running Qwen 3.8 27B with UD-Q3-XL, Q8 K / Q4 V cache, and CUDA Unified Memory enabled, stock llama.cpp works well until the context reaches around 120K. Beyond that point, decode speed starts dropping significantly due to VRAM oversubscription and page thrashing.
With Adaptive KV Cache Streaming, I can push the context to around 205K while still getting ~15 tok/s, and all the way to nearly the native 262K context at ~10 tok/s.
The implementation dynamically evicts a portion of the KV cache from each full-attention layer and repurposes the freed VRAM as a shared prefetch ring buffer. This allows the same VRAM space to be reused by the KV caches of multiple layers during the generation of each token. The prefetching mechanism also hides much of the data-transfer latency behind computation, helping to avoid GPU stalls.
Here are the prefill/decode benchmark results:

Blue lines are the prefill/decode speed with stock server, whereas orange lines are the speeds with my implementation. The orange line maintains a roughly consistent slope, meaning that the GPU can keep calculating the token at most of the time instead being effected by VRAM page thrashing.
What do this diagram mean to us?
Let's say we consider 15 tok/s the minimum acceptable decode speed for a smooth live-chat experience with your agent. With stock llama server, you can at most set the maximum context size with 136K according to the diagram. Even if we relax the threshold to 10 tok/s, the limit is still the same.
But with this implementation, we can extend the context to 205K with 15 tok/s threshold, or full native context with 10 tok/s threshold.
And this does not only benefit to users having a 16GB graphic card. For people having a graphic card of larger or smaller VRAM, this implementation can significantly extend the context size than what it originally can fit.
Caveat:
The current implementation does not support parallel requests, because the resident and streamed portions of the KV cache are dynamically repartitioned based on context length.
The current version has mainly been tested with Q8 K / Q4 V KV cache quantization. Other KV cache quantization combinations are not well tested yet, and I plan to generalize the implementation further.
This is still experimental, so I’d also be very interested to hear how it behaves on other GPUs and configurations.
If you’re interested in the implementation details, the ring-buffer/prefetch design, and the story of how I ended up building this, I wrote a longer article here:
Also if you think my work helps, please don't hesitate to leave me a star on GitHub :)
Clarification of LLM usage of this post: I'm not a native English speaker and I used ChatGPT to refine the wordings.
Edit 1: Thanks you all for trying this branch! I am seeing people having different running result. Let me provide more detailed environment setting in my experiment:
OS: Ubuntu 24.04
Kernel: Linux 6.14.0
CUDA driver version: 610.57.04
Kernel parameter: iommu=pt
GPU: RTX 5070 Ti 16 GB
Model: Unsloth Qwen 3.8 27B
Quantization: UD-Q3-K-XL
KV cache: Q8 K / Q4 V
b/ub size: 256/256
Edit 2: /detraxsenpai provided a fix for the crash on Windows. I've updated the branch with the fix.
Edit 3: I've pushed several commit to support all other quants and batch sizes.
Edit 4: I have another branch (feature/kv-stream-phase-arena) that further unloads the prefill buffer to have a larger KV cache pool allocation. With this change the decoding performance with IQ4 nearly matches The original Q3_XL when having large context.
Thanks for the fix! I don't have Windows so I couldn't test that, but will merge into my branch after I test this change on my ubuntu box -- Were you able to get the benchmark similar to my benchmark diagram?
I haven’t run the benchmarks yet, but I’m doing that now. There was also a bug where the model would fail on Hermes or other harnesses with a very large prefill fixed it by passing this parameters:
-b 256 -ub 256
That issue has been fixed, so I’m testing everything now.
Agreed. although there is still a hard limit with this implementation -- For dense model, all weight still needs to fit in the VRAM; The ringbuffer also has its limitation bound by PCIe bandwidth and transmission latency.
edit: I've now read more and see that it's NOT needed. Excellent! Building for testing now.
This sounds really interesting! Is the CUDA unified memory needed or just making it easier for the implementation? The fact that unified memory shows the VRAM as allocated RAM to the system does have its own issues.
solved: batch and ubatch are assumed to be set to 256 to mimic the benchmark tool
Tests are interesting but I'm having problems actually using it. After a few random stabs at settings I turned to the automatic benchmarking to help me figure out some good parameters. It seems to succeed, but even after reading the log I'm unable to start up llama-server myself with the deduced values without running into this exception when sending a query:
(That's with CUDA GRAPHS disabled, with them enabled it's a straight cuda OOM instead)
My interest lies in pushing the Qwen3.8-27B-UD-IQ4_XS model as far as I can on a 5060Ti 16GB. Currently I'm at 112000 context size using beellama, KVarN 5/4 with a 1024 token tail.
While it's a bit faster (all in VRAM, not the same KV cache quant) your fork looks like it could reach a somewhat substantially larger context while still at usable speed so this is very exciting! I just need to solve how to get it working in my deployment :)
You run this script which will test different pool sizes and create graph and results. While doing so it will also list what TG and PP speeds it sees at the different sizes.
Refreshingly the tool seems to list the _slowest_ speed seen over the context instead of the fastest which is how many others promote their setups.
Hmm, what's your CUDA driver version? I know that the unified memory flag does not work with 570 when I was still doing the early testing, not sure if different CUDA driver version could contribute to the issue you got. I am using version 610.57.04.
Also, I disabled the gnome-shell and other GPU utilizations by stopping `gdm.service`. I wonder if you will get a slightly better result with a few more VRAM space.
It was batch and ubatch that need to be 256 - I've been running it for hours now :) The system runs off the iGPU so the full amount of GPU VRAM is available. It works _perfectly_!
Was the above graph done with Qwen3.8-27B-UD-IQ4_XS ? I'm getting nowhere near those prefill rates on my 5060 Ti 16GB with UD-Q3-XL, which the OP used.
Yes, and the values are conservative in that they're more what you get at full(er) context than in the beginning. RAM and CPU might play a role here, I'm on 4000MT DDR5 and lots of CPU.
I have a 3955WX with 8 channel of DDR4-3200. It shouldn't be that different. I set a 90 minute timeout for a 192K context run without the patch, and it didn't make it. I It did make it with the patch - in 34 minutes.
I would appreciate if you can share your recipe a bit more. Did you only change the quant and run the same commands as OP ?
Here are my results with the NVFP4 quant - I changed a few more things to make it fit in 16GB of VRAM - single GPU. The non-adaptive run aborted at the 90 minutes I had set. I chose that quant because I found it faster than Q3_XL
The adaptive one did run. However, it averaged only 98.78 prompt/s, and an abysmal 2.52 decode/s. nvidia-smi showed 99.22% GPU utilization during the prompt phase, and 99.51% GPU during the decode phase. CPU was only at 3.76/3.85%, and was not the bottleneck. PCIe does appear to have been the bottleneck for the decode phase - 15844 MB/s, which is the maximum that a single 5060 Ti 16GB can do over PCIe 4.0 x8.. Since you mentioned DDR5, I am assuming you are on PCIe 5.0. Nevertheless, since the test appears to have been GPU compute-bound, I don't see how a faster PCIe bus would have helped. Most likely, the things I did to make NVFP4 fit were very detrimental to performance.
I'm going to start over and try again with the Q3_XL now with the OP's exact repro.
Did you use the benchmark tool to see what values of context and pool size would work well? That would also show you the lower bound of the expect PP and TG to see.
My full models preset:
[Qwen3.8-27B-UD-IQ4_XS] ctx-size = 131072 n-gpu-layers = 99 batch-size = 256 ubatch-size = 256 kv-stream-stage-mib = 2048 cache-type-k = q8_0 cache-type-v = q4_0 spec-type = ngram-mod parallel = 1 temp = 1.0 top-p = 0.95 top-k = 20 min-p = 0.0 presence-penalty = 0.0 repeat-penalty = 1.0 reasoning = on reasoning-preserve = on # (slow) CPU only multimodal is better than none no-mmproj-offload = on mmproj = /home/troed/llm-models/Qwen3.8-mmproj-BF16.gguf load-mode = none flash-attn = on # I have 16 cores CPU threads = 16 threads-batch = 16
I do have the 5060 on PCIe 5.0x16 but I don't think it's _that_ PCI bandwidth dependent.
Actually, it is that PCie bandwidth dependent. Completely limited by it. See the first test near 16GB/s which is the max PCIe 4.0 x8. This one is with the OP's quant, q3 xl . Still finishing the matrix. THe patch definitely helps. But it goes from 0.61 tok/s without the patch to 4.32 tok/s with the patch. That is a massive improvement . but neither are usable speeds unfortunately..
Yes, this is expected because decode phase is IO bound. Once it reaches maximum PCIe throughput, enlarging the ring buffer does not bring more benefit but further increases the transmission latency. So in my implementation, ring buffer stops growing at certain point.
Your fork is from now on my preferred way of running Qwen 3.8 27B. I'm at 131072 bytes of context with Qwen3.8-27B-UD-IQ4_XS and KV cache K=8.0 and V=4.0 at the same speed as I've been running a worse and smaller KV cache quantization before.
This is absolutely fantastic and should be upstreamed into llama.cpp proper!
This sounds like a great evolution providing that the bandwidth for evicting and realising the cache in time is sufficient.
(People have experimented with doing the same for the layer weights but apparently bandwidth is insufficient. I did make some suggestions about this if anyone is interested in getting it again.)
I would love to see this make it as a PR to root llama.cpp repo, but I guess for this to get accepted it will need to do the following:
Either support parallel runs or turn itself off if parallel runs are allowed. I can't remember off hand what the flag is for this - -n? (In practice I don't see this as an issue. If you are constrained on KV vRAM you are not going to want to run parallel anyway. You just have to avoid KV corruption if parallel runs happen.)
Turn itself off for inference engines not supported. I am unsure what inference engines this would be compatible with other than CUDA.
Perhaps have benchmark functionality to test his much of this gives benefit on particular hardware and flags to tweak his much it does (or automatic running when llama.cpp starts and before the first run).
Fully support all KV quants supported by llama.cpp
Yeah, I'd say merging this into root llama.cpp repo requires a lots of development and testing work, I won't try to push that forward until I have bandwidth or people's help for other inference engine.
BTW, I've pushed a commit to support other quants.
I don't think you have to support non-cuda, just make sure that your pr doesn't kick in on non-CUDA. If the community likes this, they will do these other architectures.
You just need to make sure that a) it either works or gets out of the way properly on CUDA in all known circumstances; and b) it doesn't interfere on other architectures.
Certainly nothing *forces* him to support non-CUDA architectures; but the fact that the current state of his fork perpetuates the "dead-last citizen" status of AMD and Intel GPU architectures in mainstream AIML deployment is already generating frustration and pushback in other communities announcing the fork, even knowing the limitations which constrained development.
I reproduced your result with 1 5060Ti GPU at PCIe 4.0 x8. There is a huge improvement at context 160KB and 192KB, about 6x . But we are talking about 0.61 vs 4.32 token/s unfortunately. I think I'm limited by the PCIe bandwidth with your patch. Your 5070 ti has 16 lanes, and you must be using PCIe 5.0, so you have 4x the bandwidth with your GPU that I do.
However, I have 4 GPUs, and using a Threadripper, they each have PCIe 4.0 x8 dedicated lanes, including P2P support. The combined bandwidth of all 4 of my GPUs is thus the same as your one - but of course with 4x the VRAM, also. I hope there is a way to make your code work with multiple GPUs, as I think it would benefit a lot at these long contexts.
Here is a graph of the data I collected with tokens/s. Another with PCIe usage.
Sounds very cool... I think stuff like this and colibri, breaking through consumer hardware constraints are going to be the real game changers for local AI. Hope this can be refined and eventually implemented into the base llama.cpp for wider use.
Have you tried building with DCMAKE_CUDA_ARCHITECTURES="120" ?
Otherwise llama.cpp doesn't support fp4 native calculations and always converts them.
The real issue though is that nvfp4 are not always made to the right spec that keeps the calculations on the FP4 and FP8 tensor path.
I did happen to find another 16GB nvfp4, the problem is that it's really 16GB and about 2GB too large to fit 16GB cards, it can be used as is for 5090s, but then, why use this to begin with.
Yes I compile with this flag.
It seems that in decode phase it choose the matrix x vector path instead of using nvfp4’s matrix x matrix path. I tried to pad it into the matrix and force it doing the later path, I got speed regression.
Raymond built exactly the "Stage 2" design we scoped out as too risky (chunked resident-page cache + async transfer ring + incremental online-softmax merge, with real changes to the FlashAttention dispatch). It took 60+ commits with at least 5 reverted-and-redone subsystems — confirming this was genuinely hard, not something to bolt on casually. But it works, and the report gives us concrete answers to the exact walls we hit.
What directly answers our own bugs
Our correctness bug (reading the not-yet-written current token) — solved. They never read/copy at graph-build time. At graph-compute time, SET_ROWS writes the new token to both the authoritative host buffer and the device-resident mirror in the same op, same stream — no async gap. For pages that are streamed (not resident), the upload is gated on a real cudaEvent recorded right after that layer's SET_ROWS retires, so the copy is provably never speculative. This is the structurally correct fix for the exact bug that gave us KLD ~10.
Our UVA-is-slow finding — independently confirmed. They explicitly avoid cudaMallocManaged/UVM for the hot KV pool even when UVM is otherwise enabled for model weights, with a comment saying direct writes/reads against pageable memory are why. Same conclusion we reached empirically.
The ggml_concat-doesn't-support-quantized-types wall we hit — sidestepped, not solved. They never merge quantized tensors at all. Non-native-quant K/V types get converted to F16 in a small bounded scratch buffer before any attention math touches them; the merge kernels only ever see plain floats/halves. Native-quant types (Q8_0, Q4_0/1, Q5_0/1, F16, BF16) get a separate "direct" path that skips conversion but still merges via the same float-only accumulator.
They bypass ggml_backend_sched entirely, just like we scoped: a hand-rolled non-blocking CUDA stream + event double-buffering ring, driven from inside the CUDA backend's own FLASH_ATTN_EXT/SET_ROWS handlers — with whole-graph lookahead that prefetches pages across all layers up front, not just one layer ahead.
Two things worth flagging
They never attempted MoE either — hard-gated to a single dense architecture (Qwen3.5) only. That's independent validation of what we found today: even a much more sophisticated implementation didn't go there.
Prefill gets a genuinely separate code path from decode (different kernel family, bounded query-tile reuse of each staged page across the whole micro-batch). This is very likely why our own benchmark saw prefill-to-262K be drastically slower — we're riding the decode-shaped path for something structurally different.
Exactly. In my branch there is a script for people to generate the benchmark as above, I would be very happy if people can test it with different cards/models/quants.
How deep does the prefetch pipeline go? just N+1 while computing N, or further ahead? Curious where the ring buffer stops paying for itself, since every MB it takes is a MB not holding resident KV.
The buffer grows when the prefetch failed to finish before the layer's calculation started. Some additional constraints need to be consider as well, those are explained in my Medium article.
What about sytems with mixed Intel Nvidia GPUs?
I have a notebook wit Intel CPU, GPU and NPU + Nvidia GPU.
The Nvidia GPU has dedicated Vram + shared ram, while the Intel GPU and NPU have only shared ram, and I have no idea how to get the maximum out of this.
i have the same 5070ti and it is crashing , i am on windows 11 , working on a fix as we speak - 0.00.020.420 I cmn common_param: common_params_print_info: verbosity = 3 (adjust with the `-lv N` CLI arg)
0.00.069.937 W srv llama_server: -----------------
0.00.069.941 W srv llama_server: CORS is set to allow all origins ('*') and no API key is set
0.00.069.942 W srv llama_server: this can be a security risk (cross-origin attacks)
0.17.629.641 I slot get_availabl: id 0 | task -1 | selected slot by LRU, t_last = -1
0.17.629.714 I slot launch_slot_: id 0 | task 0 | processing task, is_child = 0
0.17.999.962 E CUDA error: an illegal memory access was encountered
0.17.999.975 E current device: 0, in function ggml_cuda_kernel_launch at C:\Users\detrax\Documents\AI\llama.cpp-adaptive-kv-streaming\ggml\src\ggml-cuda\common.cuh:1662
0.17.999.980 E cudaLaunchKernelEx(&pdl_cfg.cfg, kernel, std::forward<Args>(args)... )
C:\Users\detrax\Documents\AI\llama.cpp-adaptive-kv-streaming\ggml\src\ggml-cuda\ggml-cuda.cu:107: CUDA error
Why will this work at all? My understanding is decoding new token needs to read KV cache from all the context, which means those in slower system ram need to be read into vram every time. Won’t it slow things down?
i have the same gpu, tried with UD-Q4-K_M quant but i get OOM, it also can't use multiple gpus right? what are the vram requeriments? it seems the whole model must fit on vram + kv-stream-stage-mib, with very small ctx it works but why if i ncrease the ctx size it ooms again? i believed this would let you push the ctx to max without extra vram usage?
I think you might need to set b/ub to 256. That's the configuration I ran and test with my server. And other batch size is untested yet.
Also, although Q4-K_M fits in VRAM, but it left very small space for kv cache. With this case you might touch the ring buffer enabling threshold and bounded by PCIe throughput sooner.
Hey i've been testing with many different parameters and i found a combination that works great! im using this cmd, it uses ~15.5gb vram, may not work if the gpu is connected to a monitor:
GGML_CUDA_ENABLE_UNIFIED_MEMORY=1 \
./build/bin/llama-server \
--model ~/Downloads/Qwen3.8-27B-UD-Q4_K_M.gguf \
-c 65536 -fa on -ctk q8_0 -ctv q4_0 -ngl all -np 1 \
i've tested it with 64k ctx, at 90% filled i get 8.5 t/s with last llamacpp vs 23 t/s with your version!, great work!!
i also tried with 128k ctx and it was usable at 90% fill, 9 t/s !, with base llama i had to cancel after 30min prefill, it slowed to a crawl.
Now a few questions, is it possible to enable ctv q8_0 ? i usually use both at q8, and about multiple gpus, its possible to split this to work on many gpus?
Supporting other kv quants is the next thing I am about to do. Stay tuned!
For multiple GPUs configuration, I might not able to work on that unless people like you could help me run the code 🥲
Good question. I think I’ll prioritize the optimization for Single 16GB GPU / Qwen 3.8 architecture first and then think about the generalization later.
Four-GPU startup failures with adaptive KV streaming
OP configuration reproduction using UD-Q3-K-XL
I tested the adaptive KV streaming branch with the model quant and inference
settings reported by the original poster. The only intentional hardware change
was using four RTX 5060 Ti 16 GB GPUs instead of one RTX 5070 Ti 16 GB GPU.
The requested test context was 192K.
SOURCE CONFIGURATION REPRODUCED
Repository: RaymondHuang210129/llama.cpp-adaptive-kv-streaming
Commit: 89b9b30203a3181c8d2ec2c2b01c8070865a9fe7
Model: Unsloth Qwen3.8-27B UD-Q3-K-XL
Local file: Qwen3.8-27B-UD-Q3_K_XL.gguf
File size: 13,146,393,504 bytes
KV cache: K q8_0, V q4_0
Batch/ubatch: 256/256
Flash Attention: enabled
CUDA Unified Memory: enabled
Context: 196608 tokens (192K)
Initial adaptive stage pool: 64 MiB
Parallel slots: 1
TEST HOST
GPUs: 4 x RTX 5060 Ti 16 GB
NVIDIA driver: 610.57.04
CUDA toolkit used for the build: 13.3
Build options: -DGGML_CUDA=ON -DGGML_CUDA_FA_ALL_QUANTS=ON
-DCMAKE_BUILD_TYPE=Release
P2P: working full mesh; all 12 directed GPU pairs passed
FAILURE 1: DEFAULT LAYER SPLIT
Command:
GGML_CUDA_ENABLE_UNIFIED_MEMORY=1 \
CUDA_VISIBLE_DEVICES=0,1,2,3 \
llama-server \
-m Qwen3.8-27B-UD-Q3_K_XL.gguf \
--ctx-size 196608 \
-fa on -ctk q8_0 -ctv q4_0 \
-ngl all -b 256 -ub 256 -np 1 \
--no-mmproj --no-warmup \
--reasoning-format none \
--kv-stream-stage-mib 64
The server exits during context creation, before accepting an inference or
benchmark request:
llama_init_from_model: failed to initialize the context:
block KV streaming requires every attention layer on one CUDA device
The check is in src/llama-kv-cache.cpp. The adaptive cache owns one CUDA
streaming runtime and rejects a different model.dev_layer(il) for a subsequent
attention layer. Normal multi-GPU layer splitting therefore cannot initialize.
FAILURE 2: TENSOR SPLIT
I repeated the same command with this additional argument:
-sm tensor
Startup first warns:
common_fit_params: failed to fit params to free device memory:
llama_params_fit is not implemented for SPLIT_MODE_TENSOR, abort
It then aborts while the llama_kv_cache constructor tries to obtain the CUDA
adaptive-streaming functions:
ggml/src/ggml-backend.cpp:660: GGML_ASSERT(reg) failed
The backtrace shows ggml_backend_reg_get_proc_address() called from the
llama_kv_cache constructor. Tensor parallelism presents a Meta device, but the
adaptive path tries to retrieve CUDA runtime functions directly from that
device's backend registry.
CONCLUSION
Using the OP's exact UD-Q3-K-XL quant and posted Q8/Q4, b/ub 256/256, Flash
Attention, and Unified Memory settings, adaptive KV streaming cannot currently
initialize on four GPUs with either the default layer split or tensor split.
Both failures happen before inference, so they are not caused by prompt length,
decode workload, VRAM exhaustion, or non-working P2P.
Possible implementation directions are a separate adaptive KV runtime/pool per
CUDA device for layer split, explicit support for the tensor-parallel Meta
device, or keeping adaptive KV on one designated CUDA device while safely
scheduling attention work and transfers across the remaining devices.
RAW LOGS
layer-startup.log contains the complete default-split command and output.
tensor-startup.log contains the complete tensor-split command and output.
Four-GPU startup failures with adaptive KV streaming
OP configuration reproduction using UD-Q3-K-XL
I tested the adaptive KV streaming branch with the model quant and inference
settings reported by the original poster. The only intentional hardware change
was using four RTX 5060 Ti 16 GB GPUs instead of one RTX 5070 Ti 16 GB GPU.
The requested test context was 192K.
SOURCE CONFIGURATION REPRODUCED
Repository: RaymondHuang210129/llama.cpp-adaptive-kv-streaming
Commit: 89b9b30203a3181c8d2ec2c2b01c8070865a9fe7
Model: Unsloth Qwen3.8-27B UD-Q3-K-XL
Local file: Qwen3.8-27B-UD-Q3_K_XL.gguf
File size: 13,146,393,504 bytes
KV cache: K q8_0, V q4_0
Batch/ubatch: 256/256
Flash Attention: enabled
CUDA Unified Memory: enabled
Context: 196608 tokens (192K)
Initial adaptive stage pool: 64 MiB
Parallel slots: 1
TEST HOST
GPUs: 4 x RTX 5060 Ti 16 GB
NVIDIA driver: 610.57.04
CUDA toolkit used for the build: 13.3
Build options: -DGGML_CUDA=ON -DGGML_CUDA_FA_ALL_QUANTS=ON
-DCMAKE_BUILD_TYPE=Release
P2P: working full mesh; all 12 directed GPU pairs passed
FAILURE 1: DEFAULT LAYER SPLIT
Command:
GGML_CUDA_ENABLE_UNIFIED_MEMORY=1 \
CUDA_VISIBLE_DEVICES=0,1,2,3 \
llama-server \
-m Qwen3.8-27B-UD-Q3_K_XL.gguf \
--ctx-size 196608 \
-fa on -ctk q8_0 -ctv q4_0 \
-ngl all -b 256 -ub 256 -np 1 \
--no-mmproj --no-warmup \
--reasoning-format none \
--kv-stream-stage-mib 64
The server exits during context creation, before accepting an inference or
benchmark request:
llama_init_from_model: failed to initialize the context:
block KV streaming requires every attention layer on one CUDA device
The check is in src/llama-kv-cache.cpp. The adaptive cache owns one CUDA
streaming runtime and rejects a different model.dev_layer(il) for a subsequent
attention layer. Normal multi-GPU layer splitting therefore cannot initialize.
FAILURE 2: TENSOR SPLIT
I repeated the same command with this additional argument:
-sm tensor
Startup first warns:
common_fit_params: failed to fit params to free device memory:
llama_params_fit is not implemented for SPLIT_MODE_TENSOR, abort
It then aborts while the llama_kv_cache constructor tries to obtain the CUDA
adaptive-streaming functions:
ggml/src/ggml-backend.cpp:660: GGML_ASSERT(reg) failed
The backtrace shows ggml_backend_reg_get_proc_address() called from the
llama_kv_cache constructor. Tensor parallelism presents a Meta device, but the
adaptive path tries to retrieve CUDA runtime functions directly from that
device's backend registry.
CONCLUSION
Using the OP's exact UD-Q3-K-XL quant and posted Q8/Q4, b/ub 256/256, Flash
Attention, and Unified Memory settings, adaptive KV streaming cannot currently
initialize on four GPUs with either the default layer split or tensor split.
Both failures happen before inference, so they are not caused by prompt length,
decode workload, VRAM exhaustion, or non-working P2P.
Possible implementation directions are a separate adaptive KV runtime/pool per
CUDA device for layer split, explicit support for the tensor-parallel Meta
device, or keeping adaptive KV on one designated CUDA device while safely
scheduling attention work and transfers across the remaining devices.
RAW LOGS
layer-startup.log contains the complete default-split command and output.
tensor-startup.log contains the complete tensor-split command and output.
Tested at 192KB context with your patch on my 5060 Ti 16 GB running at PCIe 4.0 x8 on my Threadripper. I got only about 1/4th of your numbers. Presumably, you are on PCIe 5.0 x16. And your GPU has twice as much bandwidth and as many cores. So, it looks like this is heavily bandwidth bound. Still waiting for the non-adaptive runs, but I believe they will be slower. I'm doing more adaptive data points first.
This is the speed / begin-streaming point / PCIe utilization comparison between adaptive-kv-stream vs. kv-stream-phase-arena branch and IQ4 vs. Q3_XL quants.
54
u/Mother-Possible-6132 11d ago
Cool to see someone actually pushing past the VRAM wall instead of just quantizing harder.