Repository navigation
fish_drama3: device_sampling (default since #1063) aborts on HIP with ARGSORT invalid configuration #1071
Description
Summary
On HIP, Fish Audio Drama 3 aborts on the first non-greedy request when fish_drama3.device_sampling is enabled, which has been the default since #1063:
ggml_cuda_compute_forward: ARGSORT failed
ROCm error: invalid configuration argument
current device: 0, in function ggml_cuda_compute_forward at external/ggml/src/ggml-cuda/ggml-cuda.cu:3372
err
external/ggml/src/ggml-cuda/ggml-cuda.cu:106: ROCm error
The server process exits after this, and every request fails until it is restarted. With "fish_drama3.device_sampling": false everything works.
Environment
- AMD Radeon 8060S (Strix Halo, gfx1151), 96 GiB VRAM carve-out (UMA), ROCm 10.0.0 (HIP 7.15.26333, AMD clang 23.0.0git), Debian 13 LXC.
- audio.cpp
mainatdd92e32b, built withscripts/build_linux.sh --backend hip --gpu-targets gfx1151 --model-set custom --models fish_drama3 --target audiocpp_server --build-type Release. Reproduced with bothGGML_HIP_HIPBLASLTON and OFF. - On this iGPU the model only loads with Stage HIP device weight uploads through a bounded host buffer #1070 applied; without it the weight upload stalls first. The crash itself happens later, during generation.
- Model:
drama-3-backbone-4bit-experts-Q4_K_S.gguf(Jmica/Drama-3-GGUF).
Reproduce
Server config entry (backend: hip):
{"id": "drama-3", "family": "fish_drama3", "task": "tts", "mode": "offline",
"path": ".../drama-3-backbone-4bit-experts-Q4_K_S.gguf",
"session_options": {"fish_drama3.cpu_expert_layers": 0, "fish_drama3.context_length": 4096}}curl -s http://127.0.0.1:8002/v1/audio/speech -H 'Content-Type: application/json' \
-d '{"model":"drama-3","input":"Hello, this is a quick test.","voice":"<preset>","response_format":"wav"}'The same happens in mode: streaming with stream_format: sse. Adding "fish_drama3.device_sampling": false to session_options avoids it.
Likely cause (from reading the code)
- With device sampling, the tower samples each codebook in the backend graph and sorts the full codebook:
src/community_models/fish_drama3/tower.cpp,ggml_argsort(ctx, logits, GGML_SORT_ORDER_DESC)overcodebook_size= 4096 logits. - On HIP,
GGML_CUDA_USE_CUBis never defined (external/ggml/src/ggml-cuda/common.cuh,#if !defined(GGML_USE_HIP) && ...). As a result,ggml_cuda_op_argsortalways takes the bitonic path. argsort_f32_i32_cuda_bitoniclaunches one block ofnext_power_of_2(ncols)threads per row (block_dims(ncols_pad, 1, 1)), so 4096 threads here. HIP allows at most 1024 threads per block, so the launch fails withinvalid configuration argument.
On CUDA the same row length goes through CUB (ncols > 1024), which would explain why the problem only shows up on HIP. Greedy decoding (temperature 0 or top_k 1) uses ggml_argmax instead of ggml_argsort and should not be affected; I have not tested that.
Possible fixes (untested suggestions)
- Default
fish_drama3.device_samplingtofalseon HIP until a large-row argsort exists there. - Or give ggml's HIP backend an argsort path for rows longer than 1024, such as hipCUB/rocPRIM or a multi-block bitonic sort, or make it fail with a clear error before the launch.
I can test a patch on gfx1151 if that helps.
Activity
0xShug0 commentedon Oct 24, 2026
@Ramon-Balaguer Thanks for reporting the issue! I'll revert the default for better compatibility as a workaround, but I still plan to add the missing path. For larger models, device-side sampling can significantly improve performance and should be the standard approach going forward.
0xShug0 commentedon Oct 24, 2026
@Ramon-Balaguer Could you test #1076?
ZeroDot1 commentedon Oct 25, 2026
I tested PR #1076 on an AMD Radeon RX 9070 XT (gfx1201, 16 GiB VRAM), ROCm 7.2.53211, Linux x86_64. I built the same upstream base in isolated /tmp worktrees, once without the PR and once with only the PR changes, using Release + HIP/gfx1201.
Results
fish_drama3.device_sampling=true. It generated a 60.511 s mono WAV at 44.1 kHz in 65.539 s (RTF 1.08), without aborting in the large-row sampler.drama3.sampling_msobservations: 7.68 ms median, ranging from 6.43 to 17.56 ms. For context,drama3.backbone_step_mshad a 3,212.51 ms median anddrama3.tower_msa 1,046.56 ms median. These are logged model timing windows, not isolated per-token measurements; in this run, ARGSORT was not the dominant measured cost.ARGSORT(ne=[65000,16,1,1])completed 4,130 repetitions at 440.76–442.85 µs/run across three clean runs (median 442.67 µs/run, approximately 17.5 GB/s). The baseline reports this large HIP ARGSORT as “not supported”, so this measures the newly enabled path; it is not a before/after speedup comparison.Conclusion
The fix works on gfx1201 and enables the missing HIP path. I found no evidence of a material performance regression in these tests. The results support the correctness fix, but they do not establish a speedup over the baseline because the baseline cannot run these large rows.
Possible follow-up optimization to measure
The HIP path currently uses segmented radix sort even for a single row. For
nrows == 1, a non-segmentedhipcub::DeviceRadixSortpath could avoid allocating and initializing the segment-offset array. This is an optimization hypothesis only; it needs profiling and correctness tests before adoption.Sketch:
A useful follow-up benchmark would separate hipCUB setup, temporary-buffer acquisition, and sort execution with HIP events, then compare one-row and multi-row cases around 1,024/1,025, 4,096/4,097, and 32,768 values. I would keep the current bitonic path at or below 1,024 unless measurements show a benefit from changing that threshold.