ggml-org/llama.cpp b11178

on GitHub

musa: fix PH1 (MTT S5000) operator failures and build issues (#29193)

• musa: use 16-byte copies for MUSA like sm_70+

ggml_cuda_get_max_cpy_bytes() derives the copy width from CUDA_ARCH. mcc
never defines it, so MUSA fell into the generic branch and returned 8 bytes
instead of the 16 bytes that every sm_70+ target gets. The value sizes the
per-thread copy unit of the FlashAttention K/V staging code (fattn-common,
fattn-vec, fattn-tile, fattn-mma-f16 shared-memory loads) and of mmq-vec-dot,
so every MUSA FlashAttention kernel moved half as many bytes per instruction.

On an MTT S5000 (mp_31, MUSA SDK 5.2.0) with Qwen3.8-27B-UD-Q4_K_M, -ngl 999,
-p 512 -n 64, -fa on: 751.15 -> 794.73 t/s prefill and 15.59 -> 15.69 t/s
decode. -fa off is unchanged (1050.05 -> 1052.86 t/s prefill), FLASH_ATTN_EXT
is unchanged (3984 ok / 0 fail / 1323 unsupported) and perplexity is
unchanged.

• musa: enable the CUB paths on MUSA

GGML_CUDA_USE_CUB and USE_CUB are selected by "CUDART_VERSION >= 11070", which
the MUSA SDK never satisfies: CUDART_VERSION is not defined anywhere under
/usr/local/musa/include, so the condition is always false and every CUB-based
path stayed compiled out on MUSA even though the SDK ships CUB and the kernels
build for mp_31. Select them from GGML_USE_MUSA as well. The device-wide
algorithms are usable too: cub::DeviceSegmentedSort compiles and produces
correct results on mp_31.

This lifts the ne[0] <= 1024 limit that ggml_backend_cuda_device_supports_op
applied to ARGSORT and TOP_K on MUSA. On an MTT S5000 (S5000, mcc 5.2.0):
ARGSORT 48 ok / 52 not supported -> 100 ok / 0 (CUDA parity), TOP_K 0 ok /
354 not supported -> 527 ok / 0. The other 20 per-op suites are unchanged, the
Qwen3-0.6B f16 (14.4679) and Qwen3.8-27B iq4_nl (5.1724) perplexities are
unchanged, and the 0.6B graph keeps the same nodes and splits (18 CPU + 18
MUSA0, SET_ROWS 1008) as before.

• musa: take the upstream code path where the toolkit supports it

Several guards were written for an older MUSA toolkit. Verified against MUSA SDK
5.2.0 and on an MTT S5000 (mp_31):

• device init: query cudaDevAttrCooperativeLaunch instead of hardcoding false.
  The device reports cooperativeLaunch=1 and musaLaunchCooperativeKernel works
  (verified with a kernel whose result was checked).
• device init: keep prop.warpSize instead of overriding it with 32. The device
  reports 32 anyway, so this only removes the divergence.
• CUDA_SET_SHARED_MEMORY_LIMIT and the FA shared-memory raise: musaFuncSetAttribute
  returns success and sharedMemPerBlockOptin is 192 KiB, so the kernels can use
  more than the default 48 KiB.
• vendors/musa.h: add the cudaDeviceGetAttribute and cudaDevAttrCooperativeLaunch
  mappings the device-init change needs.

Measured on one S5000 with Qwen3.8-27B Q4_K_M (-ngl 999, -r 3): pp512 968.27 ->
957.09 t/s, tg64 10.09 -> 10.23 t/s, FLASH_ATTN_EXT sweep identical (3975/3982
both), perplexity identical (80.2841 +/- 7.26772 both).

• musa: drop compile-time guards that MUSA's runtime gates already cover

mcc never defines CUDA_ARCH, so the arch-gated fallbacks in this group
were already taken on MUSA and the GGML_USE_MUSA guards on top of them only
kept the upstream text from being compiled:

• wkv.cu: the "#pragma unroll" suppression has no effect on the generated
  code that is not already covered by the surrounding guards
• common.cuh: the MUSA-only __builtin_unreachable() in no_device_code() is
  not needed to silence the compiler
• ssm-scan.cu: the SSD (Mamba-2 prefill) block and its dispatch are gated at
  runtime by GGML_CUDA_CC_IS_NVIDIA(cc) and turing_mma_available(cc), which
  are both false for PH1 (cc 0x100310), so compiling them changes nothing
• common.cuh: warp_reduce_max(half2) is guarded the same way as
  warp_reduce_sum(half2) (FP16_AVAILABLE); the MUSA-only guard left the
  function with no return statement. It has no caller today.

MTT S5000 (mp_31, MUSA SDK 5.2.0), MUSA_ARCHITECTURES=31: build rc=0. Against
an unmodified build of the same tree on the same card, FLASH_ATTN_EXT
(3984 ok / 0 fail / 1323 unsupported), SSM_SCAN (15/0), RWKV_WKV6 (6/0),
GATED_DELTA_NET (38/0) and MUL_MAT (1299/0/385 unsupported) are identical, and
perplexity with -fa on is bit-identical (5.1639 +/- 0.36673, 4 chunks).

• musa: do not use MMQ on PH1

test-backend-ops on an MTT S5000 (mp_31, MUSA SDK 5.2.0) fails 260 cases and every
one of them goes through the MMQ path:

• MUL_MAT with a batched src1 (any bs/nr != [1,1]): 109 cases across all
  quantized types, e.g. 12 of 13 cases at n=16, while the plain [1,1] layout
  passes
• every quantized MUL_MAT_ID: 147 cases, while the f16/f32 variants of the same
  shapes pass
• MUL_MAT with more than ~512 tokens: 4 cases (n=509..4096); the small-n cases pass

The cuBLAS/dequant path is correct for all of them and the MMVQ path used for
small batches is unaffected, so quantized matmuls now take that path on PH1
instead of returning wrong values. 27B perplexity with default flags goes from
nan to finite, and the full suite reports 0 failures out of 22237 cases.

The MMQ defect itself (fastdiv, __umulhi, uint3 kernel parameters and
CUDA_ARCH-based MMA availability were all checked and are correct on this
part) is not addressed here.

• musa: keep the block barrier of the fused TOPK_MOE kernel reachable

topk_moe_cuda returns early for the rows past the end of the graph, but one block
covers TOPK_MOE_ROWS_PER_BLOCK (8) rows, so the last block is only partially filled
whenever n_rows is not a multiple of 8. On MUSA a warp that has already returned
blocks the block wide __syncthreads() below, which makes the kernel hang and the
launch time out. CUDA tolerates the exited warps, which is why the CUDA numbers
never showed it.

For MUSA, clamp the row index of those warps to the last row so that every warp of
the block reaches the barrier; they recompute the last row and write the same
values. The CUDA code path is unchanged.

On an MTT S5000 (mp_31) the fused TOPK_MOE cases change from a launch timeout with
no completed case to 418 ok / 0 not supported / 0 failed, i.e. the CUDA result, and
the other 101 per op suites are unchanged (0 failed, no count changes).

• musa: enable GATED_DELTA_NET

The op was turned off for every MUSA target because mcc could not build the kernel
at the time. The current toolkit builds it: with mp_31 and MUSA SDK 5.2.0 the file
compiles with zero errors and all 36 test-backend-ops GATED_DELTA_NET cases pass
against the CPU reference. 27B perplexity is unchanged.

While the op is refused, the scheduler has no choice but to run it on the CPU: 48
GATED_DELTA_NET nodes per forward pass. On an MTT S5000 (Qwen3.8-27B Q4_K_M, -ngl
999, one container, -r 3):

pp512 (FA off)   964.51 -> 2119.26 t/s
tg64  (FA off)    10.15 ->   15.50 t/s

• musa: name the stream capture query API for the graph aware kernels

argsort.cu and mean.cu call cudaStreamCaptureStatus, cudaStreamIsCapturing and
cudaStreamCaptureStatusNone inside their USE_CUDA_GRAPH blocks, but the MUSA
compatibility headers do not alias those names, so building with the experimental
GGML_MUSA_GRAPHS option fails with 7 errors in those two files. Map the three
names to their musa* counterparts, under the same guard that enables the graph
code, so the default build is untouched.

The option stays off by default: on an MTT S5000 the captured path measured
slower (pp512 693 vs 772 t/s, tg128 15.20 vs 15.39 t/s over two sessions) and the
borderline MUL_MAT cases are not reproducible between runs.

• musa: build the CI and docs for PH1 (MTT S5000)

The MUSA CI job and the documented default still targeted the first generation
(MTT S80, MUSA_ARCHITECTURES=21) while the current MUSA SDK targets PH1
(MTT S5000, 31). Move the job, ci/run.sh's default and the build docs to 31,
and run the job in the PH1 MUSA SDK devel image:

registry.mthreads.com/mcconline/inference/pytorch:2.9.1.post1-py3.10-musa5.2.0-mp31-devel-ubuntu22.04-amd64

That image needs two things the previous one did not: python3-venv for the
ccache-buckets step, which builds a virtual environment for the Hugging Face
CLI, and no time prefix on the build command, because container jobs run their
steps with sh and the image ships no time binary.