[BUG]: CUDA errors with two GPUs (multiple parallel requests)

[aropb]

Description

I run several requests (3-4) at the same time, which are executed sequentially by LLamaEmbedder.GetEmbeddings() and StatelessExecutor.InferAsync().

The models for these commands are different.
For Infer (one instance for all users): Qwen2.5-14B-1M-Q5-K-M
For Embedding (one instance for all users): Qwen2.5-1.5B-Q5-K-M

There is always enough memory for queries with a margin.

1. One GPU
-- First there was the CUDA errors:
CUDA error: operation failed due to a previous error during capture
CUDA error: operation not permitted when stream is capturing
ggml_cuda_compute_forward: ADD failed

-- The errors went away when I added thread blocking to GetEmbeddings() and CreateContext/Destroy to InferAsync()

Why did I have to do this, is it right?

Questions:
what are the general limitations of multithreading for LLamaSharp? What should be considered in this case? Does anyone have experience implementing a multi-threaded web application?

2. Two GPUs

GPUSplitMode = GPUSplitMode.Layer;

Despite the fixes for one GPU, errors still occur on two GPUs:

2025-02-09 16:44:06.2064 LLama.Native.SafeLLamaContextHandle.llama_decode Error: CUDA error: operation failed due to a previous error during capture
SafeLLamaContextHandle.Decode => SafeLLamaContextHandle.llama_decode => SafeLLamaContextHandle.llama_decode
2025-02-09 16:44:06.2064 LLama.Native.NativeApi.llama_kv_cache_clear Error: CUDA error: operation not permitted when stream is capturing
SafeLLamaContextHandle.KvCacheClear => NativeApi.llama_kv_cache_clear => NativeApi.llama_kv_cache_clear
2025-02-09 16:44:06.2427 LLama.Native.SafeLLamaContextHandle.llama_decode Error: current device: 1, in function ggml_cuda_op_mul_mat at D:\a\LLamaSharp\LLamaSharp\ggml\src\ggml-cuda\ggml-cuda.cu:1615
SafeLLamaContextHandle.Decode => SafeLLamaContextHandle.llama_decode => SafeLLamaContextHandle.llama_decode
2025-02-09 16:44:06.2427 LLama.Native.NativeApi.llama_kv_cache_clear Error: current device: 1, in function ggml_backend_cuda_buffer_clear at D:\a\LLamaSharp\LLamaSharp\ggml\src\ggml-cuda\ggml-cuda.cu:605
SafeLLamaContextHandle.KvCacheClear => NativeApi.llama_kv_cache_clear => NativeApi.llama_kv_cache_clear
2025-02-09 16:44:06.2427 LLama.Native.NativeApi.llama_kv_cache_clear Error: cudaDeviceSynchronize()
SafeLLamaContextHandle.KvCacheClear => NativeApi.llama_kv_cache_clear => NativeApi.llama_kv_cache_clear
2025-02-09 16:44:06.2427 LLama.Native.SafeLLamaContextHandle.llama_decode Error: cudaGetLastError()
SafeLLamaContextHandle.Decode => SafeLLamaContextHandle.llama_decode => SafeLLamaContextHandle.llama_decode

2025-02-09 16:48:54.9660 LLama.Native.SafeLLamaContextHandle.llama_decode Error: ggml_cuda_compute_forward: ADD failed
SafeLLamaContextHandle.Decode => SafeLLamaContextHandle.llama_decode => SafeLLamaContextHandle.llama_decode
2025-02-09 16:48:54.9660 LLama.Native.NativeApi.llama_kv_cache_clear Error: CUDA error: operation not permitted when stream is capturing
SafeLLamaContextHandle.KvCacheClear => NativeApi.llama_kv_cache_clear => NativeApi.llama_kv_cache_clear
2025-02-09 16:48:54.9864 LLama.Native.SafeLLamaContextHandle.llama_decode Error: CUDA error: operation failed due to a previous error during capture
SafeLLamaContextHandle.Decode => SafeLLamaContextHandle.llama_decode => SafeLLamaContextHandle.llama_decode
2025-02-09 16:48:54.9864 LLama.Native.NativeApi.llama_kv_cache_clear Error: current device: 1, in function ggml_backend_cuda_buffer_clear at D:\a\LLamaSharp\LLamaSharp\ggml\src\ggml-cuda\ggml-cuda.cu:607
SafeLLamaContextHandle.KvCacheClear => NativeApi.llama_kv_cache_clear => NativeApi.llama_kv_cache_clear
2025-02-09 16:48:54.9864 LLama.Native.SafeLLamaContextHandle.llama_decode Error: current device: 1, in function ggml_cuda_compute_forward at D:\a\LLamaSharp\LLamaSharp\ggml\src\ggml-cuda\ggml-cuda.cu:2313
SafeLLamaContextHandle.Decode => SafeLLamaContextHandle.llama_decode => SafeLLamaContextHandle.llama_decode
2025-02-09 16:48:54.9864 LLama.Native.NativeApi.llama_kv_cache_clear Error: cudaDeviceSynchronize()
SafeLLamaContextHandle.KvCacheClear => NativeApi.llama_kv_cache_clear => NativeApi.llama_kv_cache_clear
2025-02-09 16:48:54.9864 LLama.Native.SafeLLamaContextHandle.llama_decode Error: err
SafeLLamaContextHandle.Decode => SafeLLamaContextHandle.llama_decode => SafeLLamaContextHandle.llama_decode

Questions:
what to do? What should I pay attention to?

If each subsequent request is sent after 2-3 seconds, then everything works!

As a result of many hours of experimentation, I think that creating and deleting a context (where VRAM memory is allocated) should be performed in thread-safe mode (inside lock).

It may also need to be taken into account in other places where the GPU resource is used.

Thanks.

Reproduction Steps

Multiple parallel requests

Environment & Configuration

• Operating system: Windows Server 2019
• .NET runtime version: 9.0.1
• LLamaSharp version: 0.21.0
• CUDA version (if you are using cuda backend): 12.8
• CPU & GPU device: 2 x RTX 4090 24Gb

Known Workarounds

No response

[martindevans]

the threadsafety story in LLamaSharp is basically just: whatever the thread safety of llama.cpp is. I'm not sure what the state of that is at the moment, so I'd recommend asking about this upstream. Many of the error messages you're getting are straight from llama.cpp, so they should be understandable to the upstream maintainers.

In the past I added a lock (see here) which ensures only one inference call is ever happening (even across multiple contexts). The issues mentioned in that comment have since been resolved, which is interesting, potentially this can be removed!

It sounds like the issues you're getting are mostly around creation and destruction of contexts, so potentially we should add a similar locking system inside those methods.

Would you be interested in working on a PR in this area? It would probably just involve a static lock (like the current one) wrapped around the internals of Create and ReleaseHandle (and of course some testing to see if it actually resolves your issue!).

[aropb]

Thanks for your reply. I'm not ready to make changes at the PR level yet.

[aropb]

#596
ggml-org/llama.cpp#3960

I looked at it, but I didn't understand much :)

[aropb]

@martindevans

I noticed that errors always start:
SafeLLamaContextHandle.llama_decode
or
SafeLLamaContextHandle.llama_init_from_model
or
NativeApi.llama_kv_cache_clear

Maybe they're interfering with each other, too.

At the same time, I see in decode() this:
lock (GlobalInferenceLock)

Apparently, needed to protect all operations with handle using lock()

[martindevans]

Since #3960 is resolved I was hoping we could remove the GlobalInferenceLock entirely, but yeah it's possible it actually needs to be widened.

[aropb]

@martindevans

Very good news, I managed to fix a bug for 2 GPUs. I added to the above a lock to the tokenizer (there is a decoder there) and locked the decoder in InferAsync.

This means that if you add a lock wherever you need it, everything will work.

So far it looks very ugly, but it works! :)

InferAsync:
...
using (var lockerResult = await ContextLocker.LockAsync(cancellationToken: token))
{
result = await Context.DecodeAsync(tokens, LLamaSeqId.Zero, _batch, n_past);
}

...

                using (var lockerResult = await ContextLocker.LockAsync(cancellationToken: token))
                {
                    id = InferenceParams.SamplingPipeline.Sample(...);
                 }

...
using (var lockerResult = await ContextLocker.LockAsync(cancellationToken: token))
{
NativeApi.llama_kv_cache_seq_rm(Context.NativeHandle, LLamaSeqId.Zero, tokensKeep, tokensKeep + n_discard);
NativeApi.llama_kv_cache_seq_add(Context.NativeHandle, LLamaSeqId.Zero, tokensKeep + n_discard, n_past, -n_discard);
}

...

                using (var lockerResult = await ContextLocker.LockAsync(cancellationToken: token))
                {
                    returnCode = await Context.DecodeAsync(_batch, token);
                 }
                
                if (returnCode != 0)
                    throw new LLamaDecodeError(returnCode);

...

ContextLocker - based on SemaphoreSlim.
Please fix this at the LLamaSharp level. This is a problem for many people, I think.

[aropb]

But there is a question about performance so that it does not degrade.

There are no errors after adding the locks. But of course I see delays in responses due to waiting for locks. And it is not clear whether this is correct. Here you need to understand very well what is doing inside llama.cpp.

Maybe the locking solution is not quite right and there is another way to properly solve the problem with CUDA errors for multiple requests.

[aropb]

I checked again, as soon as I removed the locks inside InferAsync(), the error immediately appeared:

2025-02-10 14:21:49.7078 LLama.Native.NativeApi.llama_kv_cache_clear Error: CUDA error: operation not permitted when stream is capturing
SafeLLamaContextHandle.KvCacheClear => NativeApi.llama_kv_cache_clear => NativeApi.llama_kv_cache_clear
2025-02-10 14:21:49.7078 LLama.Native.SafeLLamaContextHandle.llama_decode Error: ggml_cuda_compute_forward: RMS_NORM failed
SafeLLamaContextHandle.Decode => SafeLLamaContextHandle.llama_decode => SafeLLamaContextHandle.llama_decode
2025-02-10 14:21:49.7322 LLama.Native.NativeApi.llama_kv_cache_clear Error: current device: 1, in function ggml_backend_cuda_buffer_clear at D:\a\LLamaSharp\LLamaSharp\ggml\src\ggml-cuda\ggml-cuda.cu:605
SafeLLamaContextHandle.KvCacheClear => NativeApi.llama_kv_cache_clear => NativeApi.llama_kv_cache_clear
2025-02-10 14:21:49.7322 LLama.Native.SafeLLamaContextHandle.llama_decode Error: CUDA error: operation failed due to a previous error during capture
SafeLLamaContextHandle.Decode => SafeLLamaContextHandle.llama_decode => SafeLLamaContextHandle.llama_decode
2025-02-10 14:21:49.7322 LLama.Native.SafeLLamaContextHandle.llama_decode Error: current device: 1, in function ggml_cuda_compute_forward at D:\a\LLamaSharp\LLamaSharp\ggml\src\ggml-cuda\ggml-cuda.cu:2313
SafeLLamaContextHandle.Decode => SafeLLamaContextHandle.llama_decode => SafeLLamaContextHandle.llama_decode
2025-02-10 14:21:49.7322 LLama.Native.SafeLLamaContextHandle.llama_decode Error: err
SafeLLamaContextHandle.Decode => SafeLLamaContextHandle.llama_decode => SafeLLamaContextHandle.llama_decode
2025-02-10 14:21:49.7322 LLama.Native.NativeApi.llama_kv_cache_clear Error: cudaDeviceSynchronize()
SafeLLamaContextHandle.KvCacheClear => NativeApi.llama_kv_cache_clear => NativeApi.llama_kv_cache_clear

[aropb]

Since #3960 is resolved I was hoping we could remove the GlobalInferenceLock entirely, but yeah it's possible it actually needs to be widened.

Or maybe it's the assembly llama.cpp [Adds the build parameter LLAMA_SCHED_MAX_COPIES]?

ggml-org/llama.cpp#6017

[aropb]

ggml-org/llama.cpp#3960 (comment)

But here we are talking about launching from different threads and on different GPUs!? If so, how can this be managed? The more I read, the more questions I have :)

[martindevans]

The more I read, the more questions I have

Exactly my feelings on llama.cpp thread safety! This is why I was a little vague about thread safety in my initial reply.

[aropb]

Maybe you need to ask the llama.cpp team a question?
It's only about CUDA one or multiple GPUs.
But I can't even formulate it correctly :)

They write that llama.cpp thread-safe.