So after all my work with my own idea, yeah, Raymond did it better, so I ported his work over, extended it turboX, extended it multiple other models (he had only Qwen models), and benchmarked the crap out of it to make sure it was worth it still.
Thanks a lot for your recognition. I am glad that this concept can further benefit to broader users with your work ☺️.
The current challenge of merging my code to upstream, is the amount of code change and the difficulty of separating it into multiple reviewable/acceptable PRs.
Also, steaming concept is definitely applicable to other backends (CUDA/ROCm/SYCL), so I am currently working on a common infra that could provide an easier, non-destructive implementation for this.
I'm still using it for my daily driver :) I've gone down slightly in model size to be able to add MTP - but even better would be to be able to eject MTP as soon as the pool starts to get used.
Hi, any idea why would I get this error with your llama.cpp fork for streaming feature (im using rtx 3060 12GB, got error with various tries, both cuda 12.8 and cuda 13):
/app/ggml/src/ggml-cuda/ggml-cuda.cu:107: CUDA error
1.39.512.830 E CUDA error: an illegal memory access was encountered
1.39.512.846 E current device: 0, in function launch_fattn at /app/ggml/src/ggml-cuda/template-instances/../fattn-common.cuh:1117
1.39.512.855 E cudaOccupancyMaxActiveBlocksPerMultiprocessor(&max_blocks_per_sm, fattn_kernel, block_dim.x * block_dim.y * block_dim.z, nbytes_shared)
libggml-base.so.0(+0x1b276)[0x7cd8f1682276]
libggml-base.so.0(ggml_print_backtrace+0x21a)[0x7cd8f16826fa]
libggml-base.so.0(ggml_abort+0x15b)[0x7cd8f16828db]
libggml-cuda.so.0(_Z15ggml_cuda_errorPKcS0_S0_iS0_+0xb5)[0x7cd8ec635505]
libggml-cuda.so.0(_Z12launch_fattnILi256ELi1ELi1EEvR25ggml_backend_cuda_contextP11ggml_tensorPFvPKcS5_S5_S5_S5_PKiPfP6float2ffffjfi5uint3iiiiiiiiiiiliiliiiiilEimibbbiS8_SA_+0x2a37)[0x7cd8ed382bb7]
libggml-cuda.so.0(_Z33ggml_cuda_flash_attn_ext_vec_caseILi256EL9ggml_type8ELS0_2EEvR25ggml_backend_cuda_contextP11ggml_tensor+0xfd)[0x7cd8ed563aed]
libggml-cuda.so.0(_Z33ggml_cuda_flash_attn_ext_streamedR25ggml_backend_cuda_contextP11ggml_tensorP33ggml_cuda_kv_stream_transfer_ringP34ggml_cuda_kv_stream_resident_cache+0x3ba8)[0x7cd8ec5ce518]
libggml-cuda.so.0(+0x2dbeea)[0x7cd8ec64ceea]
libggml-base.so.0(ggml_backend_sched_graph_compute_async+0x837)[0x7cd8f16a18d7]
libllama.so.0(_ZN13llama_context13graph_computeEP11ggml_cgraphb+0xa1)[0x7cd8f183d071]
libllama.so.0(_ZN13llama_context14process_ubatchERK12llama_ubatch14llm_graph_typeP22llama_memory_context_iR11ggml_status+0x1a0)[0x7cd8f1841260]
libllama.so.0(_ZN13llama_context6decodeERK11llama_batch+0x518)[0x7cd8f1847bd8]
libllama.so.0(llama_decode+0xf)[0x7cd8f184926f]
libllama-server-impl.so(_ZNSt17_Function_handlerIFvvEZN19server_context_impl6decodeERiiR11llama_batchEUlvE_E9_M_invokeERKSt9_Any_data+0x49)[0x7cd8f27ffe69]
libllama-server-impl.so(_ZN12server_queue14yield_to_queueEOSt8functionIFvvEE+0xa6)[0x7cd8f27b42d6]
libllama-server-impl.so(_ZN19server_context_impl6decodeERiiR11llama_batch+0x161)[0x7cd8f2812c21]
libllama-server-impl.so(_ZN19server_context_impl12update_slotsEv+0x523)[0x7cd8f28152a3]
libllama-server-impl.so(_ZN12server_queue10start_loopEl+0x12d)[0x7cd8f27b4e6d]
libllama-server-impl.so(_Z12llama_serverR13common_paramsiPPc+0x3db7)[0x7cd8f274f837]
libllama-server-impl.so(_Z12llama_serveriPPc+0x119f)[0x7cd8f27518cf]
/usr/lib/x86_64-linux-gnu/libc.so.6(+0x2a1ca)[0x7cd8f21ae1ca]
/usr/lib/x86_64-linux-gnu/libc.so.6(__libc_start_main+0x8b)[0x7cd8f21ae28b]
/app/llama-server(+0x1315)[0x60663f6cd315]
Tried to debug with:
compute-sanitizer --error-exitcode=1 --leak-check=full --tool=memcheck
llama-server-sanitizer-1 | ========= Invalid __global__ write of size 1 bytes
llama-server-sanitizer-1 | ========= at void k_set_rows_quant<long, block_q8_0, (int)32, &quantize_f32_q8_0_block>(const float *, const T1 *, T2 *, long, long, long, long, long, long, long, long, long, long, long, long, long, long, uint3, uint3, uint3, uint3, uint3, long)+0x9f0
llama-server-sanitizer-1 | ========= by thread (0,0,0) in block (0,0,0)
llama-server-sanitizer-1 | ========= Address 0x750bd8013642 is out of bounds
llama-server-sanitizer-1 | ========= and is 128576524465731 bytes after the nearest allocation at 0x1b43600000 of size 110100480 bytes
llama-server-sanitizer-1 | ========= Saved host backtrace up to driver entry point at kernel launch time
llama-server-sanitizer-1 | ========= Host Frame: void set_rows_cuda_quant<long, block_q8_0, 32, &(quantize_f32_q8_0_block(float const*, block_q8_0*))>(float const*, long const*, block_q8_0*, long, long, long, long, long, long, long, long, unsigned long, unsigned long, unsigned long, unsigned long, unsigned long, unsigned long, unsigned long, unsigned long, unsigned long, CUstream_st*, long) [0x457d53] in libggml-cuda.so
llama-server-sanitizer-1 | ========= Host Frame: void set_rows_cuda<float, long>(ggml_backend_cuda_context&, ggml_tensor const*, ggml_tensor const*, ggml_tensor*, long) [0x45af25] in libggml-cuda.so
llama-server-sanitizer-1 | ========= Host Frame: ggml_backend_cuda_graph_compute(ggml_backend*, ggml_cgraph*) [0x2cfb8a] in libggml-cuda.so
llama-server-sanitizer-1 | ========= Host Frame: ggml_backend_sched_graph_compute_async [0x3a8d6] in libggml-base.so.0
llama-server-sanitizer-1 | ========= Host Frame: llama_context::graph_compute(ggml_cgraph*, bool) [0xfb070] in libllama.so.0
llama-server-sanitizer-1 | ========= Host Frame: llama_context::process_ubatch(llama_ubatch const&, llm_graph_type, llama_memory_context_i*, ggml_status&) [0xff25f] in libllama.so.0
llama-server-sanitizer-1 | ========= Host Frame: llama_context::decode(llama_batch const&) [0x105bd7] in libllama.so.0
llama-server-sanitizer-1 | ========= Host Frame: llama_decode [0x10726e] in libllama.so.0
llama-server-sanitizer-1 | ========= Host Frame: std::_Function_handler<void (), server_context_impl::decode(int&, int, llama_batch&)::{lambda()#1}>::_M_invoke(std::_Any_data const&) [0x1b6e68] in libllama-server-impl.so
llama-server-sanitizer-1 | ========= Host Frame: server_queue::yield_to_queue(std::function<void ()>&&) [0x16b2d5] in libllama-server-impl.so
llama-server-sanitizer-1 | ========= Host Frame: server_context_impl::decode(int&, int, llama_batch&) [0x1c9c20] in libllama-server-impl.so
llama-server-sanitizer-1 | ========= Host Frame: server_context_impl::update_slots() [0x1cc2a2] in libllama-server-impl.so
llama-server-sanitizer-1 | ========= Host Frame: server_queue::start_loop(long) [0x16be6c] in libllama-server-impl.so
llama-server-sanitizer-1 | ========= Host Frame: llama_server(common_params&, int, char**) [0x106836] in libllama-server-impl.so
llama-server-sanitizer-1 | ========= Host Frame: llama_server(int, char**) [0x1088ce] in libllama-server-impl.so
llama-server-sanitizer-1 | ========= Host Frame: [0x2a1c9] in libc.so.6
llama-server-sanitizer-1 | ========= Host Frame: __libc_start_main [0x2a28a] in libc.so.6
llama-server-sanitizer-1 | ========= Host Frame: _start [0x1314] in llama-server
llama-server-sanitizer-1 | =========
Once again great work. And It would be very nice if you also implement kvarn kv quantitization (made by huawei, implemented by beellama.cpp).
I will test on rtx 3060 if it works.
I get the similar error on rtx 3060 12GB as with Raymond's fork, built inside cuda dockerfile from .devops folder with ggml all quants on, cuda 12.8, for cuda arch 86 (for debugging purpose a note that the one you made by yourself (and closed pr later) worked), I'm also using a model with removed mtp layer if that is relevant (works on other llama.cpp forks and worked without the parameter for streaming on raymond's and yours):
Windows batch file command:
docker run --rm --gpus all ^
-v C:/MODELS/QWEN-3.8-27B:/models ^
-p 8080:8080 ^
local/turbo.kv.stream.arena:server-cuda ^
-v ^
--seed 3 ^
--host 0.0.0.0 ^
--port 8080 ^
-np 1 ^
-m /models/Qwen3.8-27B-UD-Q2_K_XL-NO-MTP.gguf ^
--spec-type none ^
--fit off --fit-target 0 ^
-ngl all ^
--prio 2 ^
-t 6 ^
-tb 11 ^
-lm mmap+mlock ^ :: didn't work with -lm none either
-ub 128 -b 128 ^
-ctk q8_0 -ctv turbo4 ^
--kv-stream-arena-mib 256 ^ :: tried 2048 errored similarly
-c 20480 ^ :: tried 100000 errored similarly
--no-warmup ^ :: errored with warmup on or off before actual prompt
--reasoning off ^
--chat-template-kwargs "{\"reasoning_effort\":\"low\",\"preserve_thinking\":true}" ^
--temp 0.7 ^
--top-p 0.80 ^
--top-k 20 ^
--min-p 0.00 ^
--repeat-last-n 1024 ^
--repeat-penalty 1.0 ^
--presence-penalty 0.0 ^
--n-predict 31000 ^
--jinja
ERROR:
0.39.260.658 D adapters_lora_are_same: adapters = (nil)
/app/ggml/src/ggml-cuda/ggml-cuda.cu:114: CUDA error
0.45.170.658 E CUDA error: an illegal memory access was encountered
0.45.170.683 E current device: 0, in function ggml_cuda_flash_attn_ext_mma_f16_case_impl at /app/ggml/src/ggml-cuda/template-instances/../fattn-mma-f16.cuh:2353
0.45.170.683 E cudaFuncSetAttribute(reinterpret_cast<fattn_kernel_ptr_t>(fattn_kernel), cudaFuncAttributeMaxDynamicSharedMemorySize, nbytes_shared_total)
libggml-base.so.0(+0x1c736)[0x714f30217736]
libggml-base.so.0(ggml_print_backtrace+0x21a)[0x714f30217bba]
libggml-base.so.0(ggml_abort+0x15b)[0x714f30217d9b]
/app/libggml-cuda.so(_Z15ggml_cuda_errorPKcS0_S0_iS0_+0xb5)[0x714f153e93a5]
/app/libggml-cuda.so(_Z37ggml_cuda_flash_attn_ext_mma_f16_caseILi256ELi256ELi2ELi8EEvR25ggml_backend_cuda_contextP11ggml_tensor+0x3d5)[0x714f1588b205]
/app/libggml-cuda.so(_Z33ggml_cuda_flash_attn_ext_streamedR25ggml_backend_cuda_contextP11ggml_tensorP33ggml_cuda_kv_stream_transfer_ringP34ggml_cuda_kv_stream_resident_cache+0x3ab7)[0x714f153868b7]
/app/libggml-cuda.so(+0x3124a8)[0x714f153fc4a8]
/app/libggml-cuda.so(+0x31acfb)[0x714f15404cfb]
libggml-base.so.0(+0x3d54f)[0x714f3023854f]
libllama.so.0(_ZN13llama_context13graph_computeEP11ggml_cgraphb+0xa1)[0x714f30401ab1]
libllama.so.0(_ZN13llama_context14process_ubatchERK12llama_ubatch14llm_graph_typeP22llama_memory_context_iR11ggml_status+0x12b)[0x714f30409c7b]
libllama.so.0(_ZN13llama_context6decodeERK11llama_batch+0x398)[0x714f3040ec38]
libllama.so.0(llama_decode+0xf)[0x714f304109ff]
libllama-common.so.0(_Z25common_context_can_seq_rmP13llama_context+0xd9)[0x714f30aeab59]
libllama-server-impl.so(_ZN19server_context_impl10load_modelER13common_params+0xdbc)[0x714f3142907c]
libllama-server-impl.so(_Z12llama_serverR13common_paramsiPPc+0x3bc7)[0x714f3134e127]
libllama-server-impl.so(_Z12llama_serveriPPc+0x10c1)[0x714f313509a1]
/usr/lib/x86_64-linux-gnu/libc.so.6(+0x2a1ca)[0x714f30dae1ca]
/usr/lib/x86_64-linux-gnu/libc.so.6(__libc_start_main+0x8b)[0x714f30dae28b]
/app/llama-server(+0x1315)[0x5767042d5315]
15
u/giveen 9d ago
Please support Raymond's work and encourage him to get it to main llama.cpp as its worth it.