Repository navigation
musa: fix PH1 (MTT S5000) operator failures and build issues - #29193
Merged
Merged
Conversation
|
Hi @yeahdongcn, thanks for your contribution! Per our contribution guidelines, the automated PR checker found the following issue(s) that need your attention:
Please note that maintainers reserve the right to make final decisions on PRs. If you believe there is a mistake, please comment below. |
yeahdongcn
force-pushed
the
musa-mp31-op-parity
branch
3 times, most recently
from
September 22, 2026 08:57
f2d9f20 to
de5e960
Compare
yeahdongcn
marked this pull request as ready for review
September 22, 2026 10:16
JohannesGaessler
requested changes
Sep 22, 2026
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.
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.
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).
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).
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.
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).
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
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.
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.
yeahdongcn
force-pushed
the
musa-mp31-op-parity
branch
from
September 23, 2026 00:49
de5e960 to
b92019b
Compare
Collaborator
Author
JohannesGaessler
approved these changes
Sep 23, 2026
feal87
added a commit
to feal87/myllama.cpp
that referenced
this pull request
Sep 25, 2026
Merge upstream commits: - llama: add llama_prec_policy + model-driven W4A4 path (ggml-org#24364) - llama: fix tensor split for fused qkv with uneven K/V head sizes (ggml-org#29294) - metal: split fa kernels into per-dtype libraries (ggml-org#29329) - metal: FWHT kernels for block widths above 512 (ggml-org#29095) - CUDA: fuse RMS_NORM + SCALE into one kernel (ggml-org#29393) - common: extract shared unicode path/string helpers (ggml-org#29415) - common,rpc: simplify fs_create_directory_with_parents() (ggml-org#29432) - rpc: include nb in the get_alloc_size cache key (ggml-org#29283) - [SYCL] support sparse FA (ggml-org#28796) - musa: fix PH1 operator failures and build issues (ggml-org#29193) - HIP: bump HIP_VERSION required for fp8 (ggml-org#29231) - opencl: add q5_k bin kernel (ggml-org#29401) - hexagon: add q5_k quant type support (ggml-org#29123) - hexagon: use DMA for contiguous dim1 CONCAT (ggml-org#29404) - mtmd: fix mel preprocessor in LFM2 audio (ggml-org#29403) - vulkan: fix legacy GLSLC without cooperativeMatrix (ggml-org#29409) - gguf-py: ByteLevel processing defaults bos/eos to False (ggml-org#29422) - gguf-py: TemplateProcessing has final word on add_special_token (ggml-org#29417) Assisted-by: Pi
Member
|
@yeahdongcn Dockerfile needs to be updated too: |
Collaborator
Author
Sorry I missed that. I'm working on a new PR to fix it. |
Collaborator
Author
|
Please see: #29481 |
1 task
edwardyoon
pushed a commit
to edwardyoon/focus-llama
that referenced
this pull request
Oct 7, 2026
…g#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. (cherry picked from commit f9af9be)
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Overview
PH1 (MTT S5000, MUSA architecture
31) is the current MUSA target, but parts of the CUDA backend still assumed the first generation (MTT S80,21) or carried compile-time workarounds for older MUSA toolkits. This series lets those paths use the upstream code and fixes the operator failures thattest-backend-opsreports on an MTT S5000 with MUSA SDK 5.2.0:sm_70and up, the CUB based paths, the cooperative launch query and the device warp size,CUDA_SET_SHARED_MEMORY_LIMIT,cudaFuncSetAttribute, and MUSA-only guards that a runtime architecture check already coversq8_0dequantization kernel was compiled out on MUSA (moved to cuda: fix the q8_0 and iq4_nl dequantization kernels #29292), the fusedTOPK_MOEkernel could hang in__syncthreads(),MMQcomputes wrong values on PH1,GATED_DELTA_NETis enabledIQ4_NLis supported as a KV cache type (moved to cuda: support IQ4_NL in the flash attention KV cache #29293)-DGGML_MUSA_GRAPHS=ONbuild compiles31and run in the PH1 MUSA SDK devel imageTesting Done
Measured on one MTT S5000 (MUSA architecture
31,mp_31) with MUSA SDK 5.2.0, single device,-DGGML_MUSA=ON -DMUSA_ARCHITECTURES=31.OPS
./build/bin/test-backend-ops:E2E
End to end with
-ngl 999(FA enabled):Additional information
Requirements