HIP/ROCm build: speaker-state graph hangs in hipStreamSynchronize on audio longer than ~2 min because ARGSORT (>1024 cols) falls back to CPU; llama.cpp pin predates the upstream fix

Read the llama.cpp submodule pin and compare its ggml argsort changes with upstream fix #28957; relevant paths mentioned are ggml/src/ggml-cuda/argsort.cu, ggml/src/ggml-cuda/ggml-cuda.cu, and src/speechlm/pee/pee_model.cpp. First build the HIP configuration and reproduce with a 160-second clip, then test the updated pin or carried patch. Done means long clips and repeated server requests complete without hanging on the HIP path.

Summary

On a HIP (ROCm) build of NeMo-Speech.cpp, nemo-speech transcribe of a SALM model hangs forever on audio longer than about 120-160 s: the GPU is idle and a CPU thread spins inside hipStreamSynchronize. nemo-speech serve wedges the same way on a later request. Shorter clips are fine. The cause is in the pinned ggml (no CUB on HIP, so ARGSORT with more than 1024 columns is unsupported on the GPU and falls back to the CPU). This is already fixed in llama.cpp master (ggml-org/llama.cpp#28957). Opening this so the submodule pin can be bumped (or the fix carried in patches/) and so HIP users who hit this can find it.

Environment

• NeMo-Speech.cpp: 20999cc (NEMO_SPEECH_GGML_PATCHED=OFF, stock ggml)
• llama.cpp submodule: bd4f514 (2026-09-23)
• GPU: AMD Radeon AI PRO R9700 (gfx1201, RDNA4)
• ROCm: pip rocm-sdk-core 7.x runtime (hipcc 7.15) at run time; Ubuntu ROCm 7.1 dev packages + clang-21 for compiling
• Platform: ROCm on WSL2 (/dev/dxg + librocdxg) inside a Docker container
• Model: Nemotron-3.5-Transcribe-30B-A4B (HF nvidia/NVIDIA-Nemotron-3.5-Transcribe-30B-A4B, converted to GGUF: AD model IQ4_NL + F16 SALM encoder)

Build flags

cmake --preset cpu-server \
  -DCMAKE_BUILD_TYPE=Release \
  -DNEMO_SPEECH_GGML_PATCHED=OFF \
  -DNEMO_SPEECH_BUILD_SPEECHLM=ON -DNEMO_SPEECH_BUILD_HTTP=ON \
  -DNEMO_SPEECH_BUILD_TTS=OFF -DNEMO_SPEECH_BUILD_NMT=OFF \
  -DGGML_HIP=ON -DGPU_TARGETS=gfx1201 -DAMDGPU_TARGETS=gfx1201 -DCMAKE_HIP_ARCHITECTURES=gfx1201 \
  -DGGML_HIP_GRAPHS=OFF \
  "-DCMAKE_HIP_FLAGS=--rocm-path=/usr --rocm-device-lib-path=/usr/lib/llvm-21/lib/clang/21/amdgcn/bitcode"
cmake --build --preset cpu-server -j"$(nproc)"

(Also needed one local edit so ggml-hip is listed in GGML_DEPENDENCIES in the top-level CMakeLists.txt; unrelated to the hang.)

Repro

export GGML_CUDA_DISABLE_GRAPHS=1
nemo-speech transcribe /audio/input.wav \
  --model /models/<AD model>.gguf --salm-encoder /models/<encoder>.gguf \
  --device auto -l <lang> --word-times \
  --asr.salm.max_sequences 4 --asr.salm.n_ctx_per_sequence 18432 --asr.salm.mtp_draft 0

/audio/input.wav is any >= 160 s, 16 kHz mono wav (speech). Expected: completes (about 22 s for 160 s of audio once fixed). Actual: never finishes, GPU idle, 3 threads at 100 percent CPU. It happens with or without --word-times / --diarize, since the speaker expert runs for SALM regardless. Clips up to about 120 s complete.

Backtrace (engine thread, gdb, trimmed)

#0  rocr::core::BusyWaitSignal::WaitRelaxed
#2  rocr::HSA::hsa_signal_wait_scacquire
#7  hipStreamSynchronize
#8  ggml_backend_cuda_synchronize
#9  ggml_backend_sched_graph_compute_async
#10 ggml_runtime::Session::run_impl
#12 nemo_speech::speechlm::PeeModel::speaker_steps_device
#13 nemo_speech::speechlm::SalmEngine::Impl::encode
#14 nemo_speech::speechlm::SalmEngine::Impl::serve

Other threads are idle in cond_wait; the ROCr DXG queue thread is in D3DKMTWaitForSynchronizationObjectFromCpu.

A HIP API log (AMD_LOG_LEVEL=3 AMD_SERIALIZE_KERNEL=3) of the hung run shows 6 k_argsort<DESC> launches, all with <= 1024 columns, no ascending argsort, then a k_get_rows_float that never returns.

Root cause

ggml/src/ggml-cuda/ggml-cuda.cu, in ggml_backend_cuda_device_supports_op:

case GGML_OP_ARGSORT:
#ifndef GGML_CUDA_USE_CUB
    return op->src[0]->ne[0] <= 1024;

GGML_CUDA_USE_CUB is not defined for HIP, so ARGSORT with more than 1024 columns is unsupported on the GPU. In PeeModel::speaker_steps_device (src/speechlm/pee/pee_model.cpp), select_top and the kept-index ggml_argsort sort rows of (n+1)*num_speakers columns once the speaker cache is compressed (n > kCacheLen = 188, i.e. after the first ~24 s update window of a longer clip). The scheduler therefore moves the argsort to the CPU backend and splits the device-resident speaker-state graph, adding device-to-host-to-device copies inside it. On HIP via WSL2/librocdxg that split graph never completes and hipStreamSynchronize busy-waits forever.

Whether the hang with a CPU-split graph is specific to ROCm-on-WSL2 or would also appear on native Linux ROCm is not known (not tested).

Fix

Upstream ggml fixed the kernel and the supports_op limit in ggml-org/llama.cpp#28957 (merged 2026-09-29): each thread owns several columns, the block is capped at 1024 threads, and supports_op checks shared memory instead of 1024. bd4f514 does not contain it. I have not yet re-tested NeMo-Speech.cpp against a llama.cpp that includes #28957. A local equivalent that I did test (against bd4f514; same 160 s clip completes, 600 s clips complete at RTF ~0.13, 12 sequential server requests with no hang):

diff --git a/ggml/src/ggml-cuda/argsort.cu b/ggml/src/ggml-cuda/argsort.cu
index 24115da..b8c3dbd 100644
--- a/ggml/src/ggml-cuda/argsort.cu
+++ b/ggml/src/ggml-cuda/argsort.cu
@@ -168,41 +168,38 @@ static inline __device__ void ggml_cuda_swap(T & a, T & b) {
 
 template<ggml_sort_order order>
 static __global__ void k_argsort_f32_i32(const float * x, int * dst, const int ncols, int ncols_pad) {
-    // bitonic sort
-    int col = threadIdx.x;
+    // bitonic sort; each thread handles ncols_pad/blockDim.x columns (block size is capped at 1024)
     int row = blockIdx.x;
 
-    if (col >= ncols_pad) {
-        return;
-    }
-
     const float * x_row = x + row * ncols;
     extern __shared__ int dst_row[];
 
-    // initialize indices
-    dst_row[col] = col;
-
+    for (int col = threadIdx.x; col < ncols_pad; col += blockDim.x) {
+        dst_row[col] = col;
+    }
     __syncthreads();
 
     for (int k = 2; k <= ncols_pad; k *= 2) {
         for (int j = k / 2; j > 0; j /= 2) {
-            int ixj = col ^ j;
-            if (ixj > col) {
-                if ((col & k) == 0) {
-                    if (dst_row[col] >= ncols ||
-                        (dst_row[ixj] < ncols && (order == GGML_SORT_ORDER_ASC ?
-                            x_row[dst_row[col]] > x_row[dst_row[ixj]] :
-                            x_row[dst_row[col]] < x_row[dst_row[ixj]]))
-                    ) {
-                        ggml_cuda_swap(dst_row[col], dst_row[ixj]);
-                    }
-                } else {
-                    if (dst_row[ixj] >= ncols ||
-                        (dst_row[col] < ncols && (order == GGML_SORT_ORDER_ASC ?
-                            x_row[dst_row[col]] < x_row[dst_row[ixj]] :
-                            x_row[dst_row[col]] > x_row[dst_row[ixj]]))
-                    ) {
-                        ggml_cuda_swap(dst_row[col], dst_row[ixj]);
+            for (int col = threadIdx.x; col < ncols_pad; col += blockDim.x) {
+                int ixj = col ^ j;
+                if (ixj > col) {
+                    if ((col & k) == 0) {
+                        if (dst_row[col] >= ncols ||
+                            (dst_row[ixj] < ncols && (order == GGML_SORT_ORDER_ASC ?
+                                x_row[dst_row[col]] > x_row[dst_row[ixj]] :
+                                x_row[dst_row[col]] < x_row[dst_row[ixj]]))
+                        ) {
+                            ggml_cuda_swap(dst_row[col], dst_row[ixj]);
+                        }
+                    } else {
+                        if (dst_row[ixj] >= ncols ||
+                            (dst_row[col] < ncols && (order == GGML_SORT_ORDER_ASC ?
+                                x_row[dst_row[col]] < x_row[dst_row[ixj]] :
+                                x_row[dst_row[col]] > x_row[dst_row[ixj]]))
+                        ) {
+                            ggml_cuda_swap(dst_row[col], dst_row[ixj]);
+                        }
                     }
                 }
             }
@@ -210,8 +207,7 @@ static __global__ void k_argsort_f32_i32(const float * x, int * dst, const int n
         }
     }
 
-    // copy the result to dst without the padding
-    if (col < ncols) {
+    for (int col = threadIdx.x; col < ncols; col += blockDim.x) {
         dst[row * ncols + col] = dst_row[col];
     }
 }
@@ -233,7 +229,7 @@ void argsort_f32_i32_cuda_bitonic(const float *   x,
     // bitonic sort requires ncols to be power of 2
     const int ncols_pad = next_power_of_2(ncols);
 
-    const dim3 block_dims(ncols_pad, 1, 1);
+    const dim3 block_dims(ncols_pad < 1024 ? ncols_pad : 1024, 1, 1);
     const dim3 block_nums(nrows, 1, 1);
     const size_t shared_mem = ncols_pad * sizeof(int);
 
diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu
index 27b8350..12379b0 100644
--- a/ggml/src/ggml-cuda/ggml-cuda.cu
+++ b/ggml/src/ggml-cuda/ggml-cuda.cu
@@ -5532,7 +5532,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
 #endif // defined(GGML_USE_HIP) || defined(GGML_CUDA_USE_CUB)
         case GGML_OP_ARGSORT:
 #ifndef GGML_CUDA_USE_CUB
-            return op->src[0]->ne[0] <= 1024;
+            return op->src[0]->ne[0] <= 8192;
 #else
             return true;
 #endif

Suggested: bump the llama.cpp submodule past #28957 (or carry the change in patches/ for the HIP path), and, independent of the sort fix, avoid CPU-backend splits inside the speaker-state graph (or document that ops unsupported on the GPU can stall HIP). I can test a bump on this setup and report back.