VRAM leak during prompt processing

[Artefact2]

I am using llama.cpp a836c8f, compiled on Arch Linux using make LLAMA_HIPBLAS=1 AMDGPU_TARGETS=gfx1030. My GPU is a RX 6750 XT.

I can reproduce a CUDA out of memory error when ingesting a long prompt with the following command:

./main -t 8 -f <(head -n420 wiki.train.raw) -c 24576 -n 1 -m ~/KoboldCpp/models/mixtral-instruct-8x7b-q4k-small.gguf

GGML_ASSERT: ggml-cuda.cu:231: !"CUDA error"
[New LWP 15195]
[New LWP 17341]
[New LWP 17342]
[New LWP 17343]
[New LWP 17344]
[New LWP 17345]
[New LWP 17346]
[New LWP 17347]

This GDB supports auto-downloading debuginfo from the following URLs:
  <https://debuginfod.harting.dev>
  <https://debuginfod.archlinux.org>
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 "/usr/lib/libthread_db.so.1".
0x00007f6feb7181c7 in wait4 () from /usr/lib/libc.so.6
#0  0x00007f6feb7181c7 in wait4 () from /usr/lib/libc.so.6
#1  0x0000557f13db782b in ggml_print_backtrace ()
#2  0x0000557f13e7ff55 in ggml_cuda_error(char const*, char const*, char const*, int, char const*) ()
#3  0x0000557f13e8d6ff in ggml_cuda_pool_malloc_leg(int, unsigned long, unsigned long*) ()
#4  0x0000557f13eab3bc in cuda_pool_alloc<float>::alloc(unsigned long) ()
#5  0x0000557f13ea7d33 in ggml_cuda_op_mul_mat(ggml_tensor const*, ggml_tensor const*, ggml_tensor*, void (*)(ggml_tensor const*, ggml_tensor const*, ggml_tensor*, char const*, float const*, char const*, float*, long, long, long, long, ihipStream_t*), bool) [clone .181] ()
#6  0x0000557f13e80713 in ggml_cuda_compute_forward ()
#7  0x0000557f13de74e6 in ggml_graph_compute_thread ()
#8  0x0000557f13deb62d in ggml_graph_compute ()
#9  0x0000557f13ead75b in ggml_backend_cpu_graph_compute ()
#10 0x0000557f13eb1727 in ggml_backend_sched_graph_compute ()
#11 0x0000557f13e17df4 in llama_decode_internal(llama_context&, llama_batch) ()
#12 0x0000557f13e18ac6 in llama_decode ()
#13 0x0000557f13dabc49 in main ()

As I understand it, this shouldn't happen since the only VRAM that's being used is the BLAS scratch buffer, which is pre-allocated and around a gigabyte in size. However, VRAM use constantly goes up as the prompt is processed, up until it uses up all the VRAM of my card and crashes.

main.log

[BrickBee]

Similar behavior on Windows with a RTX card and a CUDA-enabled build.
When running long prompt processing without any layer offloading, the VRAM usage increases for example to 12 GB before the token generation starts. When now resuming with the same prompt and using --prompt-cache, the VRAM usage is only about 8 GB when tokens are getting generated. So, about 4 GB seem to have leaked here (or been lost to fragmentation) in the process of processing the prompt.

My current process for ingesting large prompts is to split them into 16K token chunks and to process them incrementally using --prompt-cache between process restarts to keep the VRAM usage from increasing too much (and the process from slowing down too much). This works, but isn't ideal. Maybe there's something behind that behavior that can be fixed, so that the VRAM usage and general performance stays on the same level as when using the chunked approach.

[slaren]

This is happening because the attention matrix multiplications are being offloaded automatically, and without the VMM allocator this is very inefficient because each batch will require a larger buffer than the previous because more context is used, it will not be able to reuse the previous allocations and it will allocate more memory. It's not really a leak, it's just a very bad case for the legacy allocator.

The VMM allocator does not have this problem, but currently it is not enabled with HIP. I believe that HIP actually supports the APIs required for the VMM allocator, but someone with an AMD GPU would need to test that.

I am also not sure if it is really intended to offload the KV mat muls automatically. Tagging @JohannesGaessler.

[Artefact2]

Similar behavior on Windows with a RTX card and a CUDA-enabled build.

What GPU do you use? Does it use VMM?

Here's also an easier way to reproduce:

./llama-bench -m Mistral-7B-v0.1-Q2_K.gguf -ngl 0 -p 32768
ggml_init_cublas: GGML_CUDA_FORCE_MMQ:   no
ggml_init_cublas: CUDA_USE_TENSOR_CORES: yes
ggml_init_cublas: found 1 ROCm devices:
  Device 0: AMD Radeon RX 6750 XT, compute capability 10.3, VMM: no
| model                          |       size |     params | backend    | ngl | test       |              t/s |
| ------------------------------ | ---------: | ---------: | ---------- | --: | ---------- | ---------------: |
CUDA error: out of memory
  current device: 0, in function ggml_cuda_pool_malloc_leg at ggml-cuda.cu:7756
  hipMalloc((void **) &ptr, look_ahead_size)
GGML_ASSERT: ggml-cuda.cu:241: !"CUDA error"

each batch will require a larger buffer than the previous because more context is used

I don't really understand the explanation. Why not allocate the biggest batch for max context from the start, surely it's better to OOM early rather than randomly mid-pp?

[BrickBee]

What GPU do you use? Does it use VMM?

RTX 3060 Ti. Yes, it uses VMM.

Either the VMM allocator does not solve the fragmentation issue, or it's a real leak and not fragmentation.

[github-actions]

This issue is stale because it has been open for 30 days with no activity.

[github-actions]

This issue was closed because it has been inactive for 14 days since being marked as stale.