Skip to content

Eval bug: Memory leak? using ROCm #19979

Closed

Description

@MartinEmrich

Name and Version

build/bin/llama-cli --version
ggml_cuda_init: found 1 ROCm devices:
  Device 0: AMD Radeon RX 7800 XT, gfx1101 (0x1101), VMM: no, Wave Size: 32
ggml_vulkan: Found 1 Vulkan devices:
ggml_vulkan: 0 = AMD Radeon RX 7800 XT (RADV NAVI32) (radv) | uma: 0 | fp16: 1 | bf16: 0 | warp size: 64 | shared memory: 65536 | int dot: 1 | matrix cores: KHR_coopmat
version: 8179 (ecbcb7ea9)
built with GNU 14.2.0 for Linux x86_64

Operating systems

Linux

GGML backends

HIP

Hardware

Radeon RX 7800 XT 16GB
Ryzen 2700X on X470, 16GB RAM

Models

bartowski/Qwen_Qwen3.5-27B-GGUF:IQ3_XS

Problem description & steps to reproduce

This model comes in at ~10.8GB.

Running with ROCm and large prompt/context, VRAM usage slowly creeps up while prompt processing until it crashes. Running with Vulkan, VRAM usage stays low and it completes.

Command:

CUDA_VISIBLE_DEVICES=0 GGML_VISIBLE_DEVICES=""  build/bin/llama-bench -v -m ~/.cache/llama.cpp/bartowski_Qwen_Qwen3.5-27B-GGUF_Qwen_Qwen3.5-27B-IQ3_XS.gguf -ngl 99 -fa 1 -ctk q4_0 -ctv q4_0 -p 65536

Watching with amdgpu_top --smi, VRAM usage starts out at 12600MB (10MB GTT), and VRAM usage slowly creeps up. When the VRAM size is reached, it crashes with this error:

Error + Stacktrace ``` ROCm error: out of memory current device: 0, in function alloc at /home/martin/upstream/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:393 ggml_cuda_device_malloc(&ptr, look_ahead_size, device) /home/martin/upstream/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:97: ROCm error [New LWP 773974] [New LWP 773973] [New LWP 773971] [New LWP 773970] [Thread debugging using libthread_db enabled] Using host libthread_db library "/lib/x86_64-linux-gnu/libthread_db.so.1". __syscall_cancel_arch () at ../sysdeps/unix/sysv/linux/x86_64/syscall_cancel.S:56 warning: 56 ../sysdeps/unix/sysv/linux/x86_64/syscall_cancel.S: Datei oder Verzeichnis nicht gefunden #0 __syscall_cancel_arch () at ../sysdeps/unix/sysv/linux/x86_64/syscall_cancel.S:56 56 in ../sysdeps/unix/sysv/linux/x86_64/syscall_cancel.S #1 0x00007fb9a4a9a668 in __internal_syscall_cancel (a1=, a2=, a3=, a4=, a5=a5@entry=0, a6=a6@entry=0, nr=61) at ./nptl/cancellation.c:49 warning: 49 ./nptl/cancellation.c: Datei oder Verzeichnis nicht gefunden #2 0x00007fb9a4a9a6ad in __syscall_cancel (a1=, a2=, a3=, a4=, a5=a5@entry=0, a6=a6@entry=0, nr=61) at ./nptl/cancellation.c:75 75 in ./nptl/cancellation.c #3 0x00007fb9a4b05787 in __GI___wait4 (pid=, stat_loc=, options=, usage=) at ../sysdeps/unix/sysv/linux/wait4.c:30 warning: 30 ../sysdeps/unix/sysv/linux/wait4.c: Datei oder Verzeichnis nicht gefunden #4 0x00007fb9a5334deb in ggml_print_backtrace () from /home/martin/upstream/llama.cpp/build/bin/libggml-base.so.0 #5 0x00007fb9a5334f3e in ggml_abort () from /home/martin/upstream/llama.cpp/build/bin/libggml-base.so.0 #6 0x00007fb9a408dd82 in ggml_cuda_error(char const*, char const*, char const*, int, char const*) () from /home/martin/upstream/llama.cpp/build/bin/libggml-hip.so.0 #7 0x00007fb9a40a16f7 in ggml_cuda_pool_leg::alloc(unsigned long, unsigned long*) () from /home/martin/upstream/llama.cpp/build/bin/libggml-hip.so.0 #8 0x00007fb9a41ff3d1 in void launch_fattn<256, 16, 2>(ggml_backend_cuda_context&, ggml_tensor*, void (*)(char const*, char const*, char const*, char const*, char const*, int const*, float*, HIP_vector_type*, float, float, float, float, unsigned int, float, int, HIP_vector_type, int, int, int, int, int, int, int, int, int, int, int, long, int, int, long, int, int, int, int, int, long), int, unsigned long, int, bool, bool, bool, int) () from /home/martin/upstream/llama.cpp/build/bin/libggml-hip.so.0 #9 0x00007fb9a41f47d1 in void ggml_cuda_flash_attn_ext_tile_case<256, 256>(ggml_backend_cuda_context&, ggml_tensor*) () from /home/martin/upstream/llama.cpp/build/bin/libggml-hip.so.0 #10 0x00007fb9a4095637 in ggml_cuda_graph_evaluate_and_capture(ggml_backend_cuda_context*, ggml_cgraph*, bool, bool, void const*) () from /home/martin/upstream/llama.cpp/build/bin/libggml-hip.so.0 #11 0x00007fb9a4093311 in ggml_backend_cuda_graph_compute(ggml_backend*, ggml_cgraph*) () from /home/martin/upstream/llama.cpp/build/bin/libggml-hip.so.0 #12 0x00007fb9a53505f7 in ggml_backend_sched_graph_compute_async () from /home/martin/upstream/llama.cpp/build/bin/libggml-base.so.0 #13 0x00007fb9a50ac871 in llama_context::graph_compute(ggml_cgraph*, bool) () from /home/martin/upstream/llama.cpp/build/bin/libllama.so.0 #14 0x00007fb9a50aee24 in llama_context::process_ubatch(llama_ubatch const&, llm_graph_type, llama_memory_context_i*, ggml_status&) () from /home/martin/upstream/llama.cpp/build/bin/libllama.so.0 #15 0x00007fb9a50b428f in llama_context::decode(llama_batch const&) () from /home/martin/upstream/llama.cpp/build/bin/libllama.so.0 #16 0x00007fb9a50b5c9b in llama_decode () from /home/martin/upstream/llama.cpp/build/bin/libllama.so.0 #17 0x0000555c314e12ce in test_prompt(llama_context*, int, int, int) () #18 0x0000555c314ddb42 in main () [Inferior 1 (process 773967) detached] ```

It is directly related to prompt size: With small enough -p it finishes. Also llama-server with same setting runs fine, until the first request with too) big context comes in, then it crashes the same way. -ngl auto or --fit 1 does not take this additional VRAM usage into account either.

Bonus question: why does it not page/spillover to system RAM as Vulkan does?
Also uses 200% CPU cores constantly during the whole run.

With Vulkan:

CUDA_VISIBLE_DEVICES="" GGML_VISIBLE_DEVICES=0  build/bin/llama-bench -v -m ~/.cache/llama.cpp/bartowski_Qwen_Qwen3.5-27B-GGUF_Qwen_Qwen3.5-27B-IQ3_XS.gguf -ngl 99 -fa 1 -ctk q4_0 -ctv q4_0 -p 65536

Uses constant 12200MB VRAM plus 1000MB GTT during the run. CPU usage is only 50% max.

The only issue remotely related I found was this: #4946 but it's old and closed.

First Bad Commit

No response

Relevant log output

(see above)

Activity

  1. buridiaditya commented on Mar 1, 2026

    @buridiaditya

    Running into the same issue with ROCM. the fit-on tells me everything is mint but after a while (50k+ tokens later) runs into out of memory. I'm running llama-server -m local_ai/models/Qwen3.5-35B-A3B-UD-Q4_K_XL.gguf -dev ROCm0 --host 0.0.0.0 --port 8081 -ctk q8_0 -ctv q8_0 --threads 12 -fa 1 --no-mmap -c 200000 on RX 7900XTX (24GB VRAM) (+ 64GB DDR5 DRAM)

    Initialization and Crash Logs

    srv    load_model: loading model 'local_ai/models/Qwen3.5-35B-A3B-UD-Q4_K_XL.gguf'
    common_init_result: fitting params to device memory, for bugs during this step try to reproduce them with -fit off, or provide --verbose logs if the bug only occurs with -fit on
    llama_params_fit_impl: projected to use 22170 MiB of device memory vs. 24124 MiB of free device memory
    llama_params_fit_impl: will leave 1953 >= 1024 MiB of free device memory, no changes needed
    llama_params_fit: successfully fit params to free device memory
    llama_params_fit: fitting params to free memory took 0.39 seconds
    llama_model_load_from_file_impl: using device ROCm0 (Radeon RX 7900 XTX) (0000:03:00.0) - 24376 MiB free
    llama_model_loader: loaded meta data with 52 key-value pairs and 733 tensors from local_ai/models/Qwen3.5-35B-A3B-UD-Q4_K_XL.gguf (version GGUF V3 (latest))
    .....
    
    load_tensors: loading model tensors, this can take a while... (mmap = false, direct_io = false)
    load_tensors: offloading output layer to GPU
    load_tensors: offloading 39 repeating layers to GPU
    load_tensors: offloaded 41/41 layers to GPU
    load_tensors:          CPU model buffer size =   397.85 MiB
    load_tensors:        ROCm0 model buffer size = 19219.42 MiB
    
    .....
    llama_context: constructing llama_context
    llama_context: n_seq_max     = 4
    llama_context: n_ctx         = 200192
    llama_context: n_ctx_seq     = 200192
    llama_context: n_batch       = 2048
    llama_context: n_ubatch      = 512
    llama_context: causal_attn   = 1
    llama_context: flash_attn    = enabled
    llama_context: kv_unified    = true
    llama_context: freq_base     = 10000000.0
    llama_context: freq_scale    = 1
    llama_context: n_ctx_seq (200192) < n_ctx_train (262144) -- the full capacity of the model will not be utilized
    llama_context:  ROCm_Host  output buffer size =     3.79 MiB
    llama_kv_cache:      ROCm0 KV buffer size =  2077.19 MiB
    llama_kv_cache: size = 2077.19 MiB (200192 cells,  10 layers,  4/1 seqs), K (q8_0): 1038.59 MiB, V (q8_0): 1038.59 MiB
    llama_memory_recurrent:      ROCm0 RS buffer size =   251.25 MiB
    llama_memory_recurrent: size =  251.25 MiB (     4 cells,  40 layers,  4 seqs), R (f32):   11.25 MiB, S (f32):  240.00 MiB
    sched_reserve: reserving ...
    sched_reserve:      ROCm0 compute buffer size =   622.52 MiB
    sched_reserve:  ROCm_Host compute buffer size =   399.02 MiB
    sched_reserve: graph nodes  = 6699 (with bs=512), 4389 (with bs=1)
    sched_reserve: graph splits = 2
    sched_reserve: reserve took 192.36 ms, sched copies = 1
    .....
    .....
    /home/buridiaditya/local_ai/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:97: ROCm error
    ROCm error: out of memory
      current device: 0, in function alloc at /home/buridiaditya/local_ai/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:393
      ggml_cuda_device_malloc(&ptr, look_ahead_size, device)
    [New LWP 31311]
    [New LWP 31310]
    [New LWP 31309]
    [New LWP 31308]
    [New LWP 31307]
    [New LWP 31306]
    [New LWP 31305]
    [New LWP 31304]
    [New LWP 31303]
    [New LWP 31302]
    [New LWP 31301]
    [New LWP 31300]
    [New LWP 31299]
    [New LWP 31298]
    [New LWP 31297]
    [New LWP 31296]
    [New LWP 31295]
    [New LWP 31291]
    [New LWP 31289]
    
    This GDB supports auto-downloading debuginfo from the following URLs:
      <https://debuginfod.ubuntu.com>
    Enable debuginfod for this session? (y or [n]) [answered N; input not from terminal]
    Debuginfod has been disabled.
    To make this setting permanent, add 'set debuginfod enabled off' to .gdbinit.
    [Thread debugging using libthread_db enabled]
    Using host libthread_db library "/lib/x86_64-linux-gnu/libthread_db.so.1".
    0x000072be7d310813 in __GI___wait4 (pid=36583, stat_loc=0x0, options=0, usage=0x0) at ../sysdeps/unix/sysv/linux/wait4.c:30
    warning: 30	../sysdeps/unix/sysv/linux/wait4.c: No such file or directory
    #0  0x000072be7d310813 in __GI___wait4 (pid=36583, stat_loc=0x0, options=0, usage=0x0) at ../sysdeps/unix/sysv/linux/wait4.c:30
    30	in ../sysdeps/unix/sysv/linux/wait4.c
    #1  0x000072be7d96e763 in ggml_print_backtrace () from /home/buridiaditya/local_ai/.local/lib/libggml-base.so.0
    #2  0x000072be7d96e90b in ggml_abort () from /home/buridiaditya/local_ai/.local/lib/libggml-base.so.0
    #3  0x000072be7c927582 in ggml_cuda_error(char const*, char const*, char const*, int, char const*) () from /home/buridiaditya/local_ai/.local/lib/libggml-hip.so.0
    #4  0x000072be7c93acc7 in ggml_cuda_pool_leg::alloc(unsigned long, unsigned long*) () from /home/buridiaditya/local_ai/.local/lib/libggml-hip.so.0
    #5  0x000072be7c90ed69 in void launch_fattn<256, 16, 1>(ggml_backend_cuda_context&, ggml_tensor*, void (*)(char const*, char const*, char const*, char const*, char const*, int const*, float*, HIP_vector_type<float, 2u>*, float, float, float, float, unsigned int, float, int, HIP_vector_type<unsigned int, 3u>, int, int, int, int, int, int, int, int, int, int, int, long, int, int, long, int, int, int, int, int, long), int, unsigned long, int, bool, bool, bool, int) () from /home/buridiaditya/local_ai/.local/lib/libggml-hip.so.0
    #6  0x000072be7c8fd50d in ggml_cuda_flash_attn_ext_wmma_f16(ggml_backend_cuda_context&, ggml_tensor*) () from /home/buridiaditya/local_ai/.local/lib/libggml-hip.so.0
    #7  0x000072be7c92ea7b in ggml_cuda_graph_evaluate_and_capture(ggml_backend_cuda_context*, ggml_cgraph*, bool, bool, void const*) () from /home/buridiaditya/local_ai/.local/lib/libggml-hip.so.0
    #8  0x000072be7c92caa1 in ggml_backend_cuda_graph_compute(ggml_backend*, ggml_cgraph*) () from /home/buridiaditya/local_ai/.local/lib/libggml-hip.so.0
    #9  0x000072be7d98b037 in ggml_backend_sched_graph_compute_async () from /home/buridiaditya/local_ai/.local/lib/libggml-base.so.0
    #10 0x000072be7dac0e71 in llama_context::graph_compute(ggml_cgraph*, bool) () from /home/buridiaditya/local_ai/.local/lib/libllama.so.0
    #11 0x000072be7dac2f84 in llama_context::process_ubatch(llama_ubatch const&, llm_graph_type, llama_memory_context_i*, ggml_status&) () from /home/buridiaditya/local_ai/.local/lib/libllama.so.0
    #12 0x000072be7dac9b46 in llama_context::decode(llama_batch const&) () from /home/buridiaditya/local_ai/.local/lib/libllama.so.0
    #13 0x000072be7dacb5df in llama_decode () from /home/buridiaditya/local_ai/.local/lib/libllama.so.0
    #14 0x00006512fc112b88 in server_context_impl::update_slots() ()
    #15 0x00006512fc15ac5e in server_queue::start_loop(long) ()
    #16 0x00006512fc06d0e3 in main ()
    [Inferior 1 (process 31266) detached]
    ./llama-server-launch: line 12: 31266 Aborted                 (core dumped) llama-server -m local_ai/models/Qwen3.5-35B-A3B-UD-Q4_K_XL.gguf -dev ROCm0 --host 0.0.0.0 --port 8081 -ctk q8_0 -ctv q8_0 --threads 12 -fa 1 --no-mmap -c 20000
    

  2. IMbackK commented on Mar 1, 2026

    @IMbackK
    Contributor

    @buridiaditya

    For GFX11 we use rocblas quite heavily in llamacpp. Rocblas allocates some workspace (can be ~600MB) for some of its solutions which means that the vram use rises beyond the value that fit-on can account for. Similarly the wmma fattn kernel allocates some buffers that wont be accounted for by the fit.

    This behavior is normal for GFX11.
    You can force MMQ and turn off ROCWMMA at compile time to avoid most of this at the cost of some performance, otherwise you simply have to increase the free ram target for the fit algo.

  3. MartinEmrich commented on Mar 2, 2026

    @MartinEmrich
    ContributorAuthor

    @IMbackK I assume this does not apply to me? My memory usage makes at no point a 600MB jump, the creep is very slow (over multiple minutes), linear and in the end exceets 600MB.

    If you tell me compilation parameters or environment variables to try, I'd try this out nevertheless just to gather some info for this issue.

    But apart from that, In the meantime I found that Vulkan is in all cases faster (f16/Q8 KV-Cache) or equally as fast as ROCm:

    Vulkan: 
    | qwen35 ?B IQ3_XS - 3.3 bpw     |  10.82 GiB |    26.90 B | ROCm,Vulkan |  99 |           pp512 |        388.15 ± 0.79 |
    | qwen35 ?B IQ3_XS - 3.3 bpw     |  10.82 GiB |    26.90 B | ROCm,Vulkan |  99 |           tg128 |         17.62 ± 0.01 |
    | qwen35 ?B IQ3_XS - 3.3 bpw     |  10.82 GiB |    26.90 B | ROCm,Vulkan |  99 |  1 |           pp512 |        389.21 ± 0.40 |
    | qwen35 ?B IQ3_XS - 3.3 bpw     |  10.82 GiB |    26.90 B | ROCm,Vulkan |  99 |  1 |           tg128 |         17.67 ± 0.00 |
    | qwen35 ?B IQ3_XS - 3.3 bpw     |  10.82 GiB |    26.90 B | ROCm,Vulkan |  99 |   q8_0 |   q8_0 |  1 |           pp512 |        387.65 ± 0.60 |
    | qwen35 ?B IQ3_XS - 3.3 bpw     |  10.82 GiB |    26.90 B | ROCm,Vulkan |  99 |   q8_0 |   q8_0 |  1 |           tg128 |         17.66 ± 0.01 |
    | qwen35 ?B IQ3_XS - 3.3 bpw     |  10.82 GiB |    26.90 B | ROCm,Vulkan |  99 |   q5_1 |   q5_1 |  1 |           pp512 |        107.03 ± 0.61 |
    | qwen35 ?B IQ3_XS - 3.3 bpw     |  10.82 GiB |    26.90 B | ROCm,Vulkan |  99 |   q5_1 |   q5_1 |  1 |           tg128 |         16.78 ± 0.03 |
    | qwen35 ?B IQ3_XS - 3.3 bpw     |  10.82 GiB |    26.90 B | ROCm,Vulkan |  99 |   q4_0 |   q4_0 |  1 |           pp512 |        206.33 ± 0.06 |
    | qwen35 ?B IQ3_XS - 3.3 bpw     |  10.82 GiB |    26.90 B | ROCm,Vulkan |  99 |   q4_0 |   q4_0 |  1 |           tg128 |         17.54 ± 0.01 |
    
    | qwen35 ?B IQ3_XS - 3.3 bpw     |  10.82 GiB |    26.90 B | ROCm,Vulkan |  99 |   q8_0 |   q8_0 |  1 |    1 |   1 |         pp65536 |        250.61 ± 0.15 |
    | qwen35 ?B IQ3_XS - 3.3 bpw     |  10.82 GiB |    26.90 B | ROCm,Vulkan |  99 |   q8_0 |   q8_0 |  1 |    1 |   1 |           tg128 |         17.67 ± 0.01 |
    
    
     CUDA_VISIBLE_DEVICES=0 GGML_VISIBLE_DEVICES=""
    ROCm:
    | qwen35 ?B IQ3_XS - 3.3 bpw     |  10.82 GiB |    26.90 B | ROCm,Vulkan |  99 |           pp512 |        208.42 ± 0.26 |
    | qwen35 ?B IQ3_XS - 3.3 bpw     |  10.82 GiB |    26.90 B | ROCm,Vulkan |  99 |           tg128 |         17.62 ± 0.09 |
    | qwen35 ?B IQ3_XS - 3.3 bpw     |  10.82 GiB |    26.90 B | ROCm,Vulkan |  99 |  1 |           pp512 |        206.32 ± 0.05 |
    | qwen35 ?B IQ3_XS - 3.3 bpw     |  10.82 GiB |    26.90 B | ROCm,Vulkan |  99 |  1 |           tg128 |         17.67 ± 0.09 |
    | qwen35 ?B IQ3_XS - 3.3 bpw     |  10.82 GiB |    26.90 B | ROCm,Vulkan |  99 |   q8_0 |   q8_0 |  1 |           pp512 |        206.29 ± 0.05 |
    | qwen35 ?B IQ3_XS - 3.3 bpw     |  10.82 GiB |    26.90 B | ROCm,Vulkan |  99 |   q8_0 |   q8_0 |  1 |           tg128 |         17.56 ± 0.09 |
    | qwen35 ?B IQ3_XS - 3.3 bpw     |  10.82 GiB |    26.90 B | ROCm,Vulkan |  99 |   q5_1 |   q5_1 |  1 |           pp512 |        107.05 ± 0.17 |
    | qwen35 ?B IQ3_XS - 3.3 bpw     |  10.82 GiB |    26.90 B | ROCm,Vulkan |  99 |   q5_1 |   q5_1 |  1 |           tg128 |         16.68 ± 0.08 |
    | qwen35 ?B IQ3_XS - 3.3 bpw     |  10.82 GiB |    26.90 B | ROCm,Vulkan |  99 |   q4_0 |   q4_0 |  1 |           pp512 |        206.30 ± 0.09 |
    | qwen35 ?B IQ3_XS - 3.3 bpw     |  10.82 GiB |    26.90 B | ROCm,Vulkan |  99 |   q4_0 |   q4_0 |  1 |           tg128 |         17.46 ± 0.09 |
    

    (For a non-expert this seems surprising that a hardware-specific manufacturer-supported library/framework is slower than an open, generic API that cannot do any vendor-specific optimizations?)

  4. IMbackK commented on Mar 2, 2026

    @IMbackK
    Contributor

    On gfx11 some growing buffers are allocated, so the increasing ram with filling context is probubly normal but i would have to check if the magnitute is normal.

    Vulkan being faster is normal, the rocm backend is a port of the cuda backend that gets fairly little optimization work for gfx11 specifically. In the mean time the vulkan backed was written from the start with rdna in mind, and actually dose a lot of vendor specific optimizations.

    Additionally wave64 mode, which is quite beneficial on gfx11 is not available via hip while no relevant feature is avialble via hip that is not available via vulkan so vulkan actually lets you do more vendor specific optimization for rdna3.

    the main benefit of hip is ease of porting cuda code and its much easier to write in general, performance much less of a primary benefit.

  5. MartinEmrich commented on Mar 2, 2026

    @MartinEmrich
    ContributorAuthor

    Thanks, looks like I'll stick to Vulkan then.

  6. Am0g-us commented on Mar 17, 2026

    @Am0g-us

    Same issue (rx 6900 xt). Latest rocm. Tried qwen 3.5 and nemotron 3 super (and nano). No memory leaks on vulkan but very poor input tok/s after ~50k ctx size (~5 input tok/s)

  7. Am0g-us commented on Mar 18, 2026

    @Am0g-us

    Catch it!

    --flash-attn off --no-kv-offload  \         
          --cache-type-k q5_1 \              No VRAM leaks
          --cache-type-v f16 \                         
    

    But if kv cache on GPU

          --flash-attn on  \
          --cache-type-k q5_1 \      OOM after ~ 300k ctx size 
          --cache-type-v q5_1 \
    

    May be flash attention related issue

  8. Am0g-us commented on Mar 19, 2026

    @Am0g-us

    Memory leak without flash attention still happens. I need to test whether CUDA works instead of HIP on AMD graphics cards, and whether a memory leak would occur in this case

  9. IMbackK commented on Mar 19, 2026

    @IMbackK
    Contributor

    gfx1030 should not really have much vram usage increase at runtime, i also cant reproduce any unexpected runtime increase on gfx1100. Could you give us some more information on your environment? Do you see this with regular non-recurrent models too? are you using checkpointing? os?

  10. Am0g-us commented on Mar 19, 2026

    @Am0g-us

    Os: arch linux
    Which models should i test? Checkpoints are enabled by default. For now, i’m trying to run the CUDA version on AMD. I ran nemotron 3 super (with and without ablation—memory leaks occur in both cases) and qwen3.5 (memory leaks occur in the same way). I initially thought the memory leak might be caused by my clumsy model decensoring, but the same thing happens when using unsloth weights

  11. connorbell1231ho-source commented on Mar 19, 2026

    @connorbell1231ho-source
  12. viitana commented on Mar 25, 2026

    @viitana

    Also experiencing this on ROCm with a 32GB MI50 (gfx906) card, build 8505, Linux/NixOS. Creeps up significantly higher than just 600MB.

    Example case:

    I continue a long Opencode session on Qwen3.5 27B with about 60k context size with a fresh start of llama.cpp & cold kv cache. Model takes up about 24GB VRAM (unsloth Q6_K_XL), 200k context takes about 6GB with --cache-type-k q8_0 --cache-type-v q8_0 .

    As the model starts taking in the existing 60k tokens of conversation, VRAM usage slowly climbs from about 90-95% usage to 100% and OOMs. I've tested lowering context size and different model quantizations with around 80% VRAM usage at the start but it still eventually OOMs,, usually well before 100k tokens in the context.

    Does not happen with Vulkan backend with the exact same setups. Worth noting that I am running this through llama-swap.


    Vulkan is unfortunately slower for me and also sometimes gets kv cache misses for some reason, recomputing the whole prompt every few requests.

  13. Am0g-us commented on Mar 27, 2026

    @Am0g-us

    Here are my observations:

    1. Vulkan is MUCH faster (if the entire model fits into VRAM)
    2. Memory leaks with hip always occur when the kv cache is on the GPU. Whether FA is off or on makes no difference. I’ll run tests with the cache in RAM and share my observations
    3. I was unable to test the CUDA version (due to a known zluda bug)
  14. Orcazephyr commented on Apr 3, 2026

    @Orcazephyr

    I also have this issue and it seems to only happen with KV quantization. if i use f16, i can actually fit more context than q8_0 because q8_0 creeps up during PP and eventually OOMs.

  15. 10 remaining items

  16. stragulus commented on Apr 13, 2026

    @stragulus

    I had been working on a fix for this in branches related to the new upcoming turboquant features. However since the fix is independent of turboquants, I propose the PR here: #21830

    For some reason, against the master branch there seems to be a recent regression in prompt processing benchmark speeds that I don't understand just yet (I only updated just now where I was on an earlier april commit before). I did patch in the latest commit that also touched common cuda flash attention logic. Once I figure out why that is, I'll turn it into a regular PR. But, feel free to test, it fixes the VRAM usage creep and should on do so for at least RDNA3 ROCm backend users.

  17. Zgate735 commented on Apr 19, 2026

    @Zgate735

    Hi everyone, I ran into the exact same OOM issue on my Radeon Pro VII (gfx906, VMM: no). #22107

    Instead of changing the Flash Attention kernels, I applied a small patch to the legacy memory pool itself.

    File: ggml/src/ggml-cuda/ggml-cuda.cu (inside ggml_cuda_pool_leg::alloc)
    What it does: If the pool can't find a buffer of the right size (ibest < 0), it proactively flushes the old cached buffers before calling cudaMalloc. This prevents the VRAM hoarding when context grows.

    Here is the patch I used:

        if (ibest >= 0) {
            ggml_cuda_buffer& b = buffer_pool[ibest];
            void * ptr = b.ptr;
            *actual_size = b.size;
            b.ptr = nullptr;
            b.size = 0;
            return ptr;
        }
    
        // Proactive cache flush to prevent fragmentation
        for (int i = 0; i < MAX_BUFFERS; ++i) {
            if (buffer_pool[i].ptr != nullptr) {
                CUDA_CHECK(cudaFree(buffer_pool[i].ptr));
                pool_size -= buffer_pool[i].size;
                buffer_pool[i].ptr = nullptr;
                buffer_pool[i].size = 0;
            }
        }
    
        void * ptr;
        size_t look_ahead_size = (size_t) (1.05 * size);
        look_ahead_size = 256 * ((look_ahead_size + 255)/256);
        ggml_cuda_set_device(device);
        CUDA_CHECK(ggml_cuda_device_malloc(&ptr, look_ahead_size, device));
        *actual_size = look_ahead_size;
        pool_size += look_ahead_size;
    
        return ptr;
    }
    

    It worked perfectly for me. VRAM stays completely flat during prompt processing, and there is no performance drop since the FA math remains untouched. Hope this helps!

  18. Orcazephyr commented on Apr 19, 2026

    @Orcazephyr

    That looks like a good solution for quantized KV but for bf16 kv cache the correct implementation is to perform the entire attention path in bf16 instead of f16.

  19. Orcazephyr commented on Apr 19, 2026

    @Orcazephyr

    @Zgate735 Thank you so much! Your fix is so concise and works perfectly for me.

    Please consider creating it as a PR because many others would benefit from it.

    Now all we need is that fix and ROCM bfloat16 flash attention to be completely implemented and then no more memory overflows

  20. IMbackK commented on Apr 21, 2026

    @IMbackK
    Contributor

    fixed by #22155

  21. williamtwomey commented on Apr 22, 2026

    @williamtwomey

    @IMbackK This doesn't fix the issue, vRAM leak remains due to k/v fragmentation. This should be re-opened. @Zgate735 probably worth getting a PR in still.

  22. stragulus commented on Apr 22, 2026

    @stragulus

    Can confirm, this does not fix the memory leak, it just reduces the severity of the symptoms. The legacy memory pool used keeps getting requests for allocations that scale with kv size. Before this change, it would accumulate many unused fragments, now it deletes the unused ones but it still gets increasing allocation requests, thus total vram usage grows. Just less dramatically.

    I made a small change to llama-bench to demonstrate this easily: https://github.com/ggml-org/llama.cpp/compare/master...stragulus:llama.cpp:debug/master-with-vram-info?expand=1

    Output for latest master revisiion:

    $ build/bin/llama-bench -m /spare/llama.cpp/Qwen3.5-27B.Q4_K_M.gguf -ngl 100 -fa 1 --cache-type-k q8_0 --cache-type-v q8_0 -p 8000 -r 1                                                                                                                                                
    ggml_cuda_init: found 1 ROCm devices (Total VRAM: 24560 MiB):                                                                                                                                                                                                                          
      Device 0: AMD Radeon RX 7900 XTX, gfx1100 (0x1100), VMM: no, Wave Size: 32, VRAM: 24560 MiB                                                                                                                                                                                          
    | model                          |       size |     params | backend    | ngl | type_k | type_v | fa |            test |   VRAM start |     VRAM end |                  t/s |                                                                                                          
    | ------------------------------ | ---------: | ---------: | ---------- | --: | -----: | -----: | -: | --------------: | -----------: | -----------: | -------------------: |                                                                                                          
    | qwen35 27B Q4_K - Medium       |  15.39 GiB |    26.90 B | ROCm       | 100 |   q8_0 |   q8_0 |  1 |          pp8000 |     15.9 GiB |     16.2 GiB |        855.09 ± 0.00 |                                                                                                          
    | qwen35 27B Q4_K - Medium       |  15.39 GiB |    26.90 B | ROCm       | 100 |   q8_0 |   q8_0 |  1 |           tg128 |     15.3 GiB |     15.4 GiB |         28.15 ± 0.00 |                                                                                                          
                                                                                                                                                                                                                                                                                           
    build: 0d0764dfd (8892)                                                                                                                                                                                                                                                                
    $ build/bin/llama-bench -m /spare/llama.cpp/Qwen3.5-27B.Q4_K_M.gguf -ngl 100 -fa 1 --cache-type-k q8_0 --cache-type-v q8_0 -p 32000 -r 1                                                                                                                                               
    ggml_cuda_init: found 1 ROCm devices (Total VRAM: 24560 MiB):                                                                                                                                                                                                                          
      Device 0: AMD Radeon RX 7900 XTX, gfx1100 (0x1100), VMM: no, Wave Size: 32, VRAM: 24560 MiB                                                                                                                                                                                          
    | model                          |       size |     params | backend    | ngl | type_k | type_v | fa |            test |   VRAM start |     VRAM end |                  t/s |                                                                                                          
    | ------------------------------ | ---------: | ---------: | ---------- | --: | -----: | -----: | -: | --------------: | -----------: | -----------: | -------------------: |                                                                                                          
    | qwen35 27B Q4_K - Medium       |  15.39 GiB |    26.90 B | ROCm       | 100 |   q8_0 |   q8_0 |  1 |         pp32000 |     16.6 GiB |     17.7 GiB |        688.59 ± 0.00 |                                                                                                          
    | qwen35 27B Q4_K - Medium       |  15.39 GiB |    26.90 B | ROCm       | 100 |   q8_0 |   q8_0 |  1 |           tg128 |     15.3 GiB |     15.4 GiB |         27.56 ± 0.00 |                                                                                                          
                                                                                                                                                                                                                                                                                           
    build: 0d0764dfd (8892)                                                                                                                                                                                                                                                                
    $ build/bin/llama-bench -m /spare/llama.cpp/Qwen3.5-27B.Q4_K_M.gguf -ngl 100 -fa 1 --cache-type-k q8_0 --cache-type-v q8_0 -p 64000 -r 1                                                                                                                                               
    ggml_cuda_init: found 1 ROCm devices (Total VRAM: 24560 MiB):                                                                                                                                                                                                                          
      Device 0: AMD Radeon RX 7900 XTX, gfx1100 (0x1100), VMM: no, Wave Size: 32, VRAM: 24560 MiB                                                                                                                                                                                          
    | model                          |       size |     params | backend    | ngl | type_k | type_v | fa |            test |   VRAM start |     VRAM end |                  t/s |                                                                                                          
    | ------------------------------ | ---------: | ---------: | ---------- | --: | -----: | -----: | -: | --------------: | -----------: | -----------: | -------------------: |                                                                                                          
    | qwen35 27B Q4_K - Medium       |  15.39 GiB |    26.90 B | ROCm       | 100 |   q8_0 |   q8_0 |  1 |         pp64000 |     17.7 GiB |     20.0 GiB |        545.35 ± 0.00 |                                                                                                          
    | qwen35 27B Q4_K - Medium       |  15.39 GiB |    26.90 B | ROCm       | 100 |   q8_0 |   q8_0 |  1 |           tg128 |     15.3 GiB |     15.4 GiB |         27.82 ± 0.00 |                                                                                                          
                                                                                                                                                                                                                                                                                           
    build: 0d0764dfd (8892)                                                                                                                                                                                                                                                                
    

    For small prompts it will probably indeed not show much of a vram increase due to there being previously allocated fragments still available. But as you see, it quickly grows into multiple GB for large prompts still.

  23. IMbackK commented on Apr 22, 2026

    @IMbackK
    Contributor

    Its not a leak, the fattn kernel allocs space for the dequant that grows with the number of elements but thats necessary and by design and dosent constitute a mem leak.

    The kernel is not as space efficient as possible and you could avoid this by rearchitectureing the fattn kernel to not require dequanted weights but that dosent constituent a bug.

    The allocator was actually leaking memory by accumulating the small chunks.

  24. IMbackK commented on Apr 22, 2026

    @IMbackK
    Contributor

    Now all we need is that fix and ROCM bfloat16 flash attention to be completely implemented and then no more memory overflows

    Unless you are hitting range issues and getting nan or inf in attention, computing the attention in bf16 is pointless and would just decrease precision and performance. Having bf16 kernels available would be fairly trivial, but costs significant compile time and binary size at zero benefit unless range issues are actually hit.

  25. Orcazephyr commented on Apr 23, 2026

    @Orcazephyr

    Now all we need is that fix and ROCM bfloat16 flash attention to be completely implemented and then no more memory overflows

    Unless you are hitting range issues and getting nan or inf in attention, computing the attention in bf16 is pointless and would just decrease precision and performance. Having bf16 kernels available would be fairly trivial, but costs significant compile time and binary size at zero benefit unless range issues are actually hit.

    As I understand it the current kernels cast bf16 to f16 then use f32 accumulation. If a number exceeded 65535 then it would overflow before it is accumulated. I think the whole point of bf16 k/v cache is to prevent any out of range errors.

    But I will defer whether that is worth the cost of maintenance to the maintainers and contributors.

    I've noticed some looping on the qwen models in f16 which went away with bf16 (after I did a short custom patch to the kernels) and also away with q8_0 with rot. I suspect qwen 3.5/3.6 can overflow f16 (in a single cell not just in accumulation which should always be f32) but I really am not sure.

  26. andj1210 commented on May 14, 2026

    @andj1210

    I want to add some clarifyctions based on my own testing.

    The TL/DR:

    • I can confirm technically there is no "leak"
    • but clearly an excessive VRAM usage, which defeats all purpose to use quantized KV cache for ROCM backend.

    I was not quite sure to append it here or open another ticket, since the technicallity of a "Memory leak" is not given.

    Please let me know if I can contribute anything by testing (I have a compiling enviroment setup and can test individual builds/branches etc.) or you want me to open a sepparate ticket for this case.

    I think the testresults speak for themselfes:

    Testcase

    • model is loaded and set to 262144 context token size (Qwen3.6-35B-A3B-UD-IQ2_M.gguf)
    • Fill loaded model with KV cache (prompt = random GUIDs), total Context ~170k tokens, via server Endpoint API call.
    • Start over with new GUIDs.
    • read RAM consumption from AMD Adrenalin tool during testing, (-> it is total VRAM usage, not llama.cpp exclusive).
    Backend KV cache After start (MB) After 1 run (MB) After 2+ runs (MB) delta start -> stable
    Vulkan F16 17650 17632 17640 8 (~0)
    ROCm F16 17930 17945 17945 0
    ROCm Q8_0 15565 21708 21708 +6143
    ROCm Q4_0 14324 20417 20417 +6093

    Vulkan F16 and ROCm F16 are stable across runs.
    ROCm with quantized KV (Q8_0, Q4_0) gradualy grows to additional ~6 GB by the end of first context fill, capped by max context watermark. Since the actual context used was ~170k tokens, results would even be worse when hitting the 256k tokens.

    For completeness:
    I am using llama-swap as llama-server proxy and the Test script was written with AI.
    I monitored the llama-swap model activity and confirmed the server was actually processing runs 2+ and not take it from cache.

    server params (KV cache setting obviously varying):
    --port 8088 -ngl 999 -ctk q8_0 -ctv q8_0 -fa on b--mlock -np 1 -cram 20480 --metrics --batch-size 2048 --temp 0.6 --top-p 0.95 --top-k 20 --min-p 0.0 --repeat_penalty 1.0 --presence_penalty 0.0 --ctx-size 262144 --n-cpu-moe 0 --chat-template-kwargs "{"enable_thinking":false}"

    (Sorry some settings are just noise to this test, and are the artifact of C&P from an existing config)

    System:
    Win10, R9700 AI Pro, llama.cpp/hip: b9133 official

  27. Orcazephyr commented on May 14, 2026

    @Orcazephyr

    @andj1210 can you try @Zgate735 's fix? It worked for me

  28. andj1210 commented on May 14, 2026

    @andj1210

    @Orcazephyr
    It did something, but still a higher ceiling than F16 (which is 18409 MB under same conditions):

    Backend KV cache After start (MB) After 1 run (MB)
    ROCm Q4_0 14727 20916
    ROCm+patch Q4_0 14725 19061

    rev 9156, build from src (hip 7.1)

    I build using a modified cmake preset and toolchain file, catered to my sytem, just for reference:

    Details Cmakepresets:
    ...
        {
            "name": "x64-windows-rocm-release",
            "inherits": [ "base", "release" ],
            "cacheVariables": {
                "CMAKE_TOOLCHAIN_FILE": "${sourceDir}/cmake/x64-rocm-llvm.cmake",
                "CMAKE_CXX_FLAGS": "-Ik:/ROCm/rocWMMA/library/include",
                "GGML_HIP": "ON",
                "GGML_HIP_ROCWMMA_FATTN": "ON",
                "AMDGPU_TARGETS": "gfx1201",
                "LLAMA_CURL": "OFF",
                "hip_DIR": "k:/ROCm/7.1/lib/cmake/hip",
                "hipblas_DIR": "k:/ROCm/7.1/lib/cmake/hipblas",
                "hipblas-common_DIR": "k:/ROCm/7.1/lib/cmake/hipblas-common",
                "rocblas_DIR": "k:/ROCm/7.1/lib/cmake/rocblas",
                "rocblas-common_DIR": "k:/ROCm/7.1/lib/cmake/rocblas-common"
            }
        }
    ...
    

    x64-rocm-llvm.cmake:

    set( CMAKE_SYSTEM_NAME Windows )
    set( CMAKE_SYSTEM_PROCESSOR x86_64 )
    
    set(ROCM_CLANG_HOME $ENV{HIP_PATH_71}/bin)
    
    set( HIP_PLATFORM amd )
    set( CMAKE_C_COMPILER    ${ROCM_CLANG_HOME}/clang.exe )
    set( CMAKE_CXX_COMPILER  ${ROCM_CLANG_HOME}/clang++.exe )
    set( CMAKE_FORCE_C_COMPILER ${ROCM_CLANG_HOME}/clang.exe )
    set( CMAKE_FORCE_CXX_COMPILER ${ROCM_CLANG_HOME}/clang++.exe )
    
    add_compile_options(-Wno-missing-noreturn -Wno-ignored-attributes -Wno-nested-anon-types)
    
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Metadata

Metadata

Assignees

No one assigned

    Type

    No type

    Projects

    No projects

      Milestone

      No milestone

      Relationships

      None yet

      Development

      No branches or pull requests

      Issue actions