Eval bug: Abort is called in a thread from a custom thread pool during a llama_decode call

[danilabagroff]

Name and Version

./llama-cli --version
ggml_cuda_init: GGML_CUDA_FORCE_MMQ: no
ggml_cuda_init: GGML_CUDA_FORCE_CUBLAS: no
ggml_cuda_init: found 2 CUDA devices:
Device 0: NVIDIA A2, compute capability 8.6, VMM: yes
Device 1: NVIDIA A2, compute capability 8.6, VMM: yes
register_backend: registered backend CUDA (2 devices)
register_device: registered device CUDA0 (NVIDIA A2)
register_device: registered device CUDA1 (NVIDIA A2)
register_backend: registered backend CPU (1 devices)
register_device: registered device CPU (Intel Xeon Processor (Icelake))
version: 5572 (7675c55)
built with cc (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0 for x86_64-linux-gnu

Operating systems

Linux

GGML backends

CUDA

Hardware

2 x A2 or 3 x A100

Models

SmolLM2-360M-Instruct-BF16

Problem description & steps to reproduce

We are testing inference with 15 threads in the worker pool when an abort is called from one of those threads during a llama_decode call.

• main thread
• 3 threads are started by llama.cpp: a host thread and an extra thread for every single GPU(as far as I understand):
• 15 threads are started by our worker pool

(gdb) info thread
  Id   Target Id                                          Frame 
  1    Thread 0x7ffff7629000 (LWP 4059) "bai-test"        0x00007ffff5898d71 in __futex_abstimed_wait_common64 (private=128, cancel=true, abstime=0x0, op=265, expected=4071, futex_word=0x7fff9cdde2d0) at ./nptl/futex-internal.c:57
  2    Thread 0x7fffb5dff000 (LWP 4062) "cuda00001400006" 0x00007ffff591b4cd in __GI___poll (fds=0x555555bda010, nfds=3, timeout=-1) at ../sysdeps/unix/sysv/linux/poll.c:29
  3    Thread 0x7fffa8dde000 (LWP 4069) "cuda-EvtHandlr"  0x00007ffff591b4cd in __GI___poll (fds=0x7fffa4000c20, nfds=10, timeout=100) at ../sysdeps/unix/sysv/linux/poll.c:29
  4    Thread 0x7fffa2dde000 (LWP 4070) "cuda-EvtHandlr"  0x00007ffff591b4cd in __GI___poll (fds=0x7fff98000c20, nfds=10, timeout=100) at ../sysdeps/unix/sysv/linux/poll.c:29
  5    Thread 0x7fff9cdde000 (LWP 4071) "bai-test"        __futex_abstimed_wait_common (cancel=false, private=<optimized out>, abstime=0x0, clockid=0, expected=3, futex_word=0x555555babc0c) at ./nptl/futex-internal.c:103
  6    Thread 0x7fff93ff0000 (LWP 4072) "bai-test"        __futex_abstimed_wait_common (cancel=false, private=<optimized out>, abstime=0x0, clockid=0, expected=3, futex_word=0x555555babc0c) at ./nptl/futex-internal.c:103
  7    Thread 0x7fff937ef000 (LWP 4073) "bai-test"        __futex_abstimed_wait_common (cancel=false, private=<optimized out>, abstime=0x0, clockid=0, expected=3, futex_word=0x555555bac7ec) at ./nptl/futex-internal.c:103
  8    Thread 0x7fff92fee000 (LWP 4074) "bai-test"        __futex_abstimed_wait_common (cancel=false, private=<optimized out>, abstime=0x0, clockid=0, expected=3, futex_word=0x555555bac7e8) at ./nptl/futex-internal.c:103
  9    Thread 0x7fff927ed000 (LWP 4075) "bai-test"        __futex_abstimed_wait_common (cancel=false, private=<optimized out>, abstime=0x0, clockid=0, expected=3, futex_word=0x555555bac7ec) at ./nptl/futex-internal.c:103
  10   Thread 0x7fff91fec000 (LWP 4076) "bai-test"        0x00007fffea04d771 in ?? () from /lib/x86_64-linux-gnu/libcuda.so.1
  11   Thread 0x7fff917eb000 (LWP 4077) "bai-test"        __futex_abstimed_wait_common (cancel=false, private=<optimized out>, abstime=0x0, clockid=0, expected=3, futex_word=0x555555babc0c) at ./nptl/futex-internal.c:103
  12   Thread 0x7fff90fea000 (LWP 4078) "bai-test"        0x00007ffff7fc3e36 in ?? ()
  13   Thread 0x7fff8bfff000 (LWP 4079) "bai-test"        __futex_abstimed_wait_common (cancel=false, private=0, abstime=0x0, clockid=0, expected=2, futex_word=0x555555babc08) at ./nptl/futex-internal.c:103
  14   Thread 0x7fff8b7fe000 (LWP 4080) "bai-test"        0x00007fffea94e6d0 in ?? () from /lib/x86_64-linux-gnu/libcuda.so.1
  15   Thread 0x7fff8affd000 (LWP 4081) "bai-test"        __futex_abstimed_wait_common (cancel=false, private=<optimized out>, abstime=0x0, clockid=0, expected=3, futex_word=0x555555bac7ec) at ./nptl/futex-internal.c:103
* 16   Thread 0x7fff8a7fc000 (LWP 4082) "bai-test"        __pthread_kill_implementation (no_tid=0, signo=6, threadid=<optimized out>) at ./nptl/pthread_kill.c:44
  17   Thread 0x7fff89ffb000 (LWP 4083) "bai-test"        __futex_abstimed_wait_common (cancel=false, private=<optimized out>, abstime=0x0, clockid=0, expected=3, futex_word=0x555555bac7e8) at ./nptl/futex-internal.c:103
  18   Thread 0x7fff897fa000 (LWP 4084) "bai-test"        __futex_abstimed_wait_common (cancel=false, private=<optimized out>, abstime=0x0, clockid=0, expected=3, futex_word=0x555555babc0c) at ./nptl/futex-internal.c:103
  19   Thread 0x7fff88ff9000 (LWP 4085) "bai-test"        __futex_abstimed_wait_common (cancel=false, private=<optimized out>, abstime=0x0, clockid=0, expected=3, futex_word=0x555555bac7e8) at ./nptl/futex-internal.c:103

Callstack:

Thread 16 "bai-test" received signal SIGABRT, Aborted.
[Switching to Thread 0x7fff8a7fc000 (LWP 4082)]
Download failed: Invalid argument.  Continuing without source file ./nptl/./nptl/pthread_kill.c.
__pthread_kill_implementation (no_tid=0, signo=6, threadid=<optimized out>) at ./nptl/pthread_kill.c:44
warning: 44	./nptl/pthread_kill.c: No such file or directory
(gdb) bt
#0  __pthread_kill_implementation (no_tid=0, signo=6, threadid=<optimized out>) at ./nptl/pthread_kill.c:44
#1  __pthread_kill_internal (signo=6, threadid=<optimized out>) at ./nptl/pthread_kill.c:78
#2  __GI___pthread_kill (threadid=<optimized out>, signo=signo@entry=6) at ./nptl/pthread_kill.c:89
#3  0x00007ffff584527e in __GI_raise (sig=sig@entry=6) at ../sysdeps/posix/raise.c:26
#4  0x00007ffff58288ff in __GI_abort () at ./stdlib/abort.c:79
#5  0x00007ffff5f0f836 in ggml_abort (file=0x7ffff6613268 "/home/ubuntu/sources/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu", line=75, fmt=0x7ffff661325d "CUDA error")
    at /home/ubuntu/sources/llama.cpp/ggml/src/ggml.c:221
#6  0x00007ffff60e6c77 in ggml_cuda_error (stmt=0x7ffff66147e2 "cudaGetLastError()", func=0x7ffff66146fa "ggml_cuda_op_mul_mat", 
    file=0x7ffff6613268 "/home/ubuntu/sources/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu", line=1676, 
    msg=0x7ffff5493f90 "operation failed due to a previous error during capture") at /home/ubuntu/sources/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:75
#7  0x00007ffff60ed62d in ggml_cuda_op_mul_mat (ctx=..., src0=0x5555580d3930, src1=0x7fff0d04ab00, dst=0x7fff0d04ade0, 
    op=0x7ffff6108641 <ggml_cuda_op_mul_mat_vec(ggml_backend_cuda_context&, ggml_tensor const*, ggml_tensor const*, ggml_tensor*, char const*, float const*, char const*, float*, long, long, long, long, CUstream_st*)>, quantize_src1=0x0) at /home/ubuntu/sources/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:1676
#8  0x00007ffff60ef1e7 in ggml_cuda_mul_mat (ctx=..., src0=0x5555580d3930, src1=0x7fff0d04ab00, dst=0x7fff0d04ade0)
    at /home/ubuntu/sources/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:1976
#9  0x00007ffff60f0ac6 in ggml_cuda_compute_forward (ctx=..., dst=0x7fff0d04ade0) at /home/ubuntu/sources/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:2264
#10 0x00007ffff60f221f in evaluate_and_capture_cuda_graph (cuda_ctx=0x7fff14001410, cgraph=0x7fff140fd1c8, graph_evaluated_or_captured=@0x7fff8a7d75db: false, 
    use_cuda_graph=@0x7fff8a7d75d9: true, cuda_graph_update_required=@0x7fff8a7d75da: true) at /home/ubuntu/sources/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:2673
#11 0x00007ffff60f28c8 in ggml_backend_cuda_graph_compute (backend=0x7fff14001950, cgraph=0x7fff140fd1c8)
    at /home/ubuntu/sources/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:2780
#12 0x00007ffff5f27a02 in ggml_backend_graph_compute_async (backend=0x7fff14001950, cgraph=0x7fff140fd1c8) at /home/ubuntu/sources/llama.cpp/ggml/src/ggml-backend.cpp:334
#13 0x00007ffff5f2bb6d in ggml_backend_sched_compute_splits (sched=0x7fff1403ab70) at /home/ubuntu/sources/llama.cpp/ggml/src/ggml-backend.cpp:1404
#14 0x00007ffff5f2c809 in ggml_backend_sched_graph_compute_async (sched=0x7fff1403ab70, graph=0x7fff0cdf6030)
    at /home/ubuntu/sources/llama.cpp/ggml/src/ggml-backend.cpp:1596
#15 0x00007ffff7b88691 in llama_context::graph_compute (this=0x7fff14000b70, gf=0x7fff0cdf6030, batched=false) at /home/ubuntu/sources/llama.cpp/src/llama-context.cpp:1381
#16 0x00007ffff7b84f30 in llama_context::process_ubatch (this=0x7fff14000b70, ubatch=..., gtype=LLM_GRAPH_TYPE_DECODER, mstate=0x7fff14178a00, 
    ret=@0x7fff8a7d7848: -369628841) at /home/ubuntu/sources/llama.cpp/src/llama-context.cpp:683
#17 0x00007ffff7b868bf in llama_context::decode (this=0x7fff14000b70, inp_batch=...) at /home/ubuntu/sources/llama.cpp/src/llama-context.cpp:1018
#18 0x00007ffff7b8d301 in llama_decode (ctx=0x7fff14000b70, batch=...) at /home/ubuntu/sources/llama.cpp/src/llama-context.cpp:2681

All threads are sharing same model and vocabulary but each thread has its own llama_context and llama_sampler:
llama_model_params.n_gpu_layers = 99
llama_context_params.ctx = 1000

• I understand that the bottleneck for inference is 1+2 threads in llama.cpp, but we also need our thread pool for regular business logic, where multithreading fits our needs.
• Same test works steadily on Metal

First Bad Commit

No response

Relevant log output

ggml_gallocr_needs_realloc: node inp_embd is not valid
ggml_gallocr_alloc_graph: cannot reallocate multi buffer graph automatically, call reserve
ggml_backend_sched_alloc_splits: failed to allocate graph, reserving (backend_ids_changed = 0)
check_node_graph_compatibility_and_refresh_copy_ops: disabling CUDA graphs due to batch size > 1 [ffn_inp-0] [960 128 1 1]
check_node_graph_compatibility_and_refresh_copy_ops: disabling CUDA graphs due to batch size > 1 [ffn_inp-17] [960 128 1 1]

CUDA error: operation failed due to a previous error during capture
  current device: 1, in function ggml_cuda_op_mul_mat at /home/ubuntu/sources/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:1676
  cudaGetLastError()

[slaren]

Does it work if you build with CUDA graphs disabled (GGML_CUDA_GRAPHS=OFF)?

[slaren]

Please check if #14033 fixes this issue.

[danilabagroff]

Does it work if you build with CUDA graphs disabled (GGML_CUDA_GRAPHS=OFF)?

Yes, works like a charm without CUDA graphs.

[danilabagroff]

Please check if #14033 fixes this issue.

Does not help.

[danilabagroff]

We ran a ton of tests in the cloud across different GPUs and nailed down that the issue only crops up just on A2 (1 or 2 units) and A100 (2 or 4 units) were rock-solid on T4 (2 units).

Where to dig further?

[slaren]

I was able to reproduce this issue using the test in #14035. #14033 fixes the issue for me, so it is surprising that it doesn't for you. In my case finding the cause was issue because sometimes would crash in the call to cudaDeviceSynchronize instead of during graph capture. You can try to do the same and see if it crashes in other places (outside of a call to ggml_backend_cuda_graph_compute).

[slaren]

@agray3 tagging you here in case you have any suggestions. It seems that graph capture fails occasionally when running other CUDA code in a different thread. I tracked an instance of this issue to a call to cudaDeviceSynchronize, but maybe there are other CUDA calls that cause issues during graph capture in a different thread?

[danilabagroff]

I was able to reproduce this issue using the test in #14035. #14033 fixes the issue for me, so it is surprising that it doesn't for you. In my case finding the cause was issue because sometimes would crash in the call to cudaDeviceSynchronize instead of during graph capture. You can try to do the same and see if it crashes in other places (outside of a call to ggml_backend_cuda_graph_compute).

I have fetched relevant PR and the test completed successfully:

./test-thread-safety -m ~/models/SmolLM2-360M-Instruct-BF16.gguf -p "Hello, my name is"
...
llama_context: graph splits = 356 (with bs=512), 1 (with bs=1)
Model 2/2, Context 1/1: Result: 'Hello, my name is'
Model 1/2, Context 1/1: Result: 'Hello, my name is'
All threads completed successfully.

...but nothing has changed with our test:

#0  0x00007ffff5a9eb2c in pthread_kill () from /lib/x86_64-linux-gnu/libc.so.6
No symbol table info available.
#1  0x00007ffff5a4527e in raise () from /lib/x86_64-linux-gnu/libc.so.6
No symbol table info available.
#2  0x00007ffff5a288ff in abort () from /lib/x86_64-linux-gnu/libc.so.6
No symbol table info available.
#3  0x00007ffff610f836 in ggml_abort (file=0x7ffff6813728 "/home/ubuntu/sources/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu", line=75, fmt=0x7ffff681371d "CUDA error") at /home/ubuntu/sources/llama.cpp/ggml/src/ggml.c:221
        args = {{gp_offset = 24, fp_offset = 48, overflow_arg_area = 0x7fffa8fd94c0, reg_save_area = 0x7fffa8fd9400}}
#4  0x00007ffff62e7e37 in ggml_cuda_error (stmt=0x7ffff681573c "err", func=0x7ffff6815713 "ggml_cuda_compute_forward", file=0x7ffff6813728 "/home/ubuntu/sources/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu", line=2366, 
    msg=0x7ffff5697910 "operation failed due to a previous error during capture") at /home/ubuntu/sources/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:75
        id = 0
#5  0x00007ffff62f1fbd in ggml_cuda_compute_forward (ctx=..., dst=0x7fff36932660) at /home/ubuntu/sources/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:2366
        err_ = cudaErrorStreamCaptureInvalidated
        err = cudaErrorStreamCaptureInvalidated
        __func__ = "ggml_cuda_compute_forward"
#6  0x00007ffff62f33df in evaluate_and_capture_cuda_graph (cuda_ctx=0x7fff88000e20, cgraph=0x7fff881d17b0, graph_evaluated_or_captured=@0x7fffa8fd95db: false, use_cuda_graph=@0x7fffa8fd95d9: true, cuda_graph_update_required=@0x7fffa8fd95da: true)
    at /home/ubuntu/sources/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:2673
        node = 0x7fff36932660
        ok = true
        i = 567
        integrated = false
        __PRETTY_FUNCTION__ = "void evaluate_and_capture_cuda_graph(ggml_backend_cuda_context*, ggml_cgraph*, bool&, bool&, bool&)"
        __func__ = "evaluate_and_capture_cuda_graph"
#7  0x00007ffff62f3a88 in ggml_backend_cuda_graph_compute (backend=0x7fff88001360, cgraph=0x7fff881d17b0) at /home/ubuntu/sources/llama.cpp/ggml/src/ggml-cuda/ggml-cuda.cu:2780
        cuda_ctx = 0x7fff88000e20
        disable_cuda_graphs_due_to_env = false
        use_cuda_graph = true
        cuda_graph_update_required = true
        __func__ = "ggml_backend_cuda_graph_compute"
        graph_evaluated_or_captured = false
#8  0x00007ffff6127a02 in ggml_backend_graph_compute_async (backend=0x7fff88001360, cgraph=0x7fff881d17b0) at /home/ubuntu/sources/llama.cpp/ggml/src/ggml-backend.cpp:334
No locals.
#9  0x00007ffff612bb6d in ggml_backend_sched_compute_splits (sched=0x7fff8800f1f0) at /home/ubuntu/sources/llama.cpp/ggml/src/ggml-backend.cpp:1404
        ec = GGML_STATUS_SUCCESS
        split = 0x7fff881d1748
        split_backend_id = 0
        split_backend = 0x7fff88001360
        i = 1
        splits = 0x7fff881d1690
#10 0x00007ffff612c809 in ggml_backend_sched_graph_compute_async (sched=0x7fff8800f1f0, graph=0x7fff366fb030) at /home/ubuntu/sources/llama.cpp/ggml/src/ggml-backend.cpp:1596
No locals.
#11 0x00007ffff7b8ad95 in llama_context::graph_compute (this=0x7fff88000b70, gf=0x7fff366fb030, batched=false) at /home/ubuntu/sources/llama.cpp/src/llama-context.cpp:1412
        n_threads = 4
        tp = 0x0
        status = 32767
        __func__ = "graph_compute"
#12 0x00007ffff7b875c0 in llama_context::process_ubatch (this=0x7fff88000b70, ubatch=..., gtype=LLM_GRAPH_TYPE_DECODER, mstate=0x7fff881d4330, ret=@0x7fffa8fd984c: 32767) at /home/ubuntu/sources/llama.cpp/src/llama-context.cpp:710
        __func__ = "process_ubatch"
        gf = 0x7fff366fb030
        res = std::unique_ptr<llm_graph_result_i> = {get() = 0x7fff88ab7ac0}
        status = GGML_STATUS_SUCCESS
#13 0x00007ffff7b88fdd in llama_context::decode (this=0x7fff88000b70, inp_batch=...) at /home/ubuntu/sources/llama.cpp/src/llama-context.cpp:1051
        status = 32767
        ubatch = @0x7fff88abd3b0: {equal_seqs = false, n_tokens = 1, n_seq_tokens = 1, n_seqs = 1, token = 0x7fffa8fd9c28, embd = 0x0, pos = 0x7fff883a7100, n_seq_id = 0x7fff88e6d990, seq_id = 0x7fff88e6e0e0, output = 0x7fff8823c7f0 "\001&Sw\370\177"}
        res = std::unique_ptr<llm_graph_result_i> = {get() = 0x7fffa8fd98f0}
        t_logits = 0x20000
        t_embd = 0x555555ad3060
        __func__ = "decode"
        batch_allocr = {batch = {n_tokens = 1, token = 0x7fffa8fd9c28, embd = 0x0, pos = 0x7fff883a7100, n_seq_id = 0x7fff88e6d990, seq_id = 0x7fff88e6e0e0, logits = 0x7fff88aba930 "\0019\036w\370\177"}, seq_id_0 = {_M_elems = {0}}, 
          pos = std::vector of length 1, capacity 1 = {288}, n_seq_id = std::vector of length 1, capacity 1 = {1}, seq_id = std::vector of length 2, capacity 2 = {0x7fffa8fd9998, 0x0}, logits = std::vector of length 1, capacity 1 = {1 '\001'}}
        batch = @0x7fffa8fd9960: {n_tokens = 1, token = 0x7fffa8fd9c28, embd = 0x0, pos = 0x7fff883a7100, n_seq_id = 0x7fff88e6d990, seq_id = 0x7fff88e6e0e0, logits = 0x7fff88aba930 "\0019\036w\370\177"}
        vocab = @0x555555bf34a8: {pimpl = std::unique_ptr<llama_vocab::impl> = {get() = 0x555555b627b0}}
        hparams = @0x555555bf1978: {vocab_only = false, rope_finetuned = false, use_par_res = false, swin_norm = false, n_ctx_train = 8192, n_embd = 960, n_embd_features = 0, n_layer = 32, n_rot = 64, n_embd_head_k = 64, n_embd_head_v = 64, n_expert = 0, 
          n_expert_used = 0, n_rel_attn_bkts = 0, n_embd_head_k_mla = 0, n_embd_head_v_mla = 0, posnet = {n_embd = 0, n_layer = 0}, convnext = {n_embd = 0, n_layer = 0}, n_head_arr = {_M_elems = {15 <repeats 32 times>, 0 <repeats 480 times>}}, n_head_kv_arr = {
            _M_elems = {5 <repeats 32 times>, 0 <repeats 480 times>}}, n_ff_arr = {_M_elems = {2560 <repeats 32 times>, 0 <repeats 480 times>}}, n_layer_dense_lead = 0, n_lora_q = 0, n_lora_kv = 0, n_ff_exp = 0, n_ff_shexp = 0, n_expert_shared = 0, n_norm_groups = 0, 
          expert_weights_scale = 0, expert_weights_norm = false, expert_gating_func = 0, moe_every_n_layers = 0, f_norm_eps = 0, f_norm_rms_eps = 9.99999975e-06, f_norm_group_eps = 0, f_attn_logit_softcapping = 50, f_final_logit_softcapping = 30, 
          rescale_every_n_layers = 0, time_mix_extra_dim = 0, time_decay_extra_dim = 0, wkv_head_size = 0, token_shift_count = 2, n_lora_decay = 0, n_lora_iclr = 0, n_lora_value_res_mix = 0, n_lora_gate = 0, rope_attn_factor = 1, rope_freq_base_train = 100000, 
          rope_freq_base_train_swa = 100000, rope_freq_scale_train = 1, rope_freq_scale_train_swa = 1, n_ctx_orig_yarn = 8192, rope_yarn_log_mul = 0, rope_sections = {_M_elems = {0, 0, 0, 0}}, swa_type = LLAMA_SWA_TYPE_NONE, n_swa = 0, swa_layers = {_M_elems = {
              false <repeats 512 times>}}, ssm_d_conv = 0, ssm_d_inner = 0, ssm_d_state = 0, ssm_dt_rank = 0, ssm_dt_b_c_rms = false, f_clamp_kqv = 0, f_max_alibi_bias = 0, f_logit_scale = 0, f_residual_scale = 0, f_embedding_scale = 0, f_attention_scale = 0, 
          causal_attn = true, use_alibi = false, attn_soft_cap = false, use_kq_norm = true, n_cls_out = 1, n_moe_layer_step = 0, n_no_rope_layer_step = 4, n_attn_temp_floor_scale = 8192, f_attn_temp_scale = 0.100000001, dec_start_token_id = -1, 
          pooling_type = LLAMA_POOLING_TYPE_NONE, rope_type = LLAMA_ROPE_TYPE_NORM, rope_scaling_type_train = LLAMA_ROPE_SCALING_TYPE_LINEAR}
        n_vocab = 49152
        n_tokens_all = 1
        n_embd = 960
        embd_pooled = false
        n_outputs_all = 1
        did_optimize = false
        mstate = std::unique_ptr<llama_memory_state_i> = {get() = 0x7fff881d4330}
        n_outputs_prev = 0
#14 0x00007ffff7b8fbd9 in llama_decode (ctx=0x7fff88000b70, batch=...) at /home/ubuntu/sources/llama.cpp/src/llama-context.cpp:2812
        ret = 32767
        __func__ = "llama_decode"

2 threads by llama.cpp and others are mine:

(gdb) info threads
  Id   Target Id                                           Frame 
  1    Thread 0x7ffff60bb000 (LWP 30568) "bai-test"        0x00007ffff5a98d71 in ?? () from /lib/x86_64-linux-gnu/libc.so.6
  2    Thread 0x7fffb83ff000 (LWP 30571) "cuda00001400006" 0x00007ffff5b1b4cd in poll () from /lib/x86_64-linux-gnu/libc.so.6
  3    Thread 0x7fffb6783000 (LWP 30585) "cuda-EvtHandlr"  0x00007ffff5b1b4cd in poll () from /lib/x86_64-linux-gnu/libc.so.6
  4    Thread 0x7fffb4940000 (LWP 30586) "bai-test"        0x00007ffff5a98c37 in ?? () from /lib/x86_64-linux-gnu/libc.so.6
  5    Thread 0x7fffa9fff000 (LWP 30587) "bai-test"        0x00007fffeac4f804 in ?? () from /lib/x86_64-linux-gnu/libcuda.so.1
  6    Thread 0x7fffa97fe000 (LWP 30588) "bai-test"        0x00007ffff5a98c37 in ?? () from /lib/x86_64-linux-gnu/libc.so.6
* 7    Thread 0x7fffa8ffd000 (LWP 30589) "bai-test"        0x00007ffff5a9eb2c in pthread_kill () from /lib/x86_64-linux-gnu/libc.so.6
  8    Thread 0x7fff9fff0000 (LWP 30590) "bai-test"        0x00007ffff5a98c37 in ?? () from /lib/x86_64-linux-gnu/libc.so.6
  9    Thread 0x7fff9f7ef000 (LWP 30591) "bai-test"        0x00007fffeac4f909 in ?? () from /lib/x86_64-linux-gnu/libcuda.so.1
  10   Thread 0x7fff9efee000 (LWP 30592) "bai-test"        0x00007ffff5a98c37 in ?? () from /lib/x86_64-linux-gnu/libc.so.6
  11   Thread 0x7fff9e7ed000 (LWP 30593) "bai-test"        0x00007fffeac4f83e in ?? () from /lib/x86_64-linux-gnu/libcuda.so.1
  12   Thread 0x7fff9dfec000 (LWP 30594) "bai-test"        0x00007fffeac4f83e in ?? () from /lib/x86_64-linux-gnu/libcuda.so.1
  13   Thread 0x7fff9d7eb000 (LWP 30595) "bai-test"        0x00007fffeaa1a350 in ?? () from /lib/x86_64-linux-gnu/libcuda.so.1
  14   Thread 0x7fff9cfea000 (LWP 30596) "bai-test"        0x00007ffff5a98c37 in ?? () from /lib/x86_64-linux-gnu/libc.so.6
  15   Thread 0x7fff7ffff000 (LWP 30597) "bai-test"        0x00007ffff5a98c37 in ?? () from /lib/x86_64-linux-gnu/libc.so.6

Could it be due to different CC versions? I reproduced the issue with A100 (SM80) and A2 (SM86) but couldn’t reproduce it with T4 (SM75).

[agray3]

@agray3 tagging you here in case you have any suggestions. It seems that graph capture fails occasionally when running other CUDA code in a different thread. I tracked an instance of this issue to a call to cudaDeviceSynchronize, but maybe there are other CUDA calls that cause issues during graph capture in a different thread?

I'll check this with our CUDA Graphs team and report back

[agray3]

@agray3 tagging you here in case you have any suggestions. It seems that graph capture fails occasionally when running other CUDA code in a different thread. I tracked an instance of this issue to a call to cudaDeviceSynchronize, but maybe there are other CUDA calls that cause issues during graph capture in a different thread?

Can you try replacing cudaDeviceSynchronize() with cudaStreamSynchronize()? cudaDeviceSynchronize() isn't allowed during stream capture, even when called by a different thread to the one that is associated with the capture.

[slaren]

I did that already in #14033, but graph capture still fails occasionally when using multiple threads. I even tried using a global mutex so that only one thread is running CUDA functions at the same time, but that also didn't work, so I am at completely at a a loss at this point. The only thing I can say for sure is that when using multiple threads, a random operation during graph capture may fail with "operation failed due to a previous error during capture".