Оптимизация MUSA и ускорение инференса на Moore Threads MTT S5000 в llama.cpp b11178¶
4.0/10
В релизе llama.cpp b11178 проведена глубокая оптимизация бэкенда MUSA под ускорители Moore Threads MTT S5000 (архитектура mp_31, SDK 5.2.0), устраняющая аппаратные сбои и вычислительные узкие места. Разработчики перевели FlashAttention на 16-байтовые инструкции загрузки K/V-кэша, подключили реализацию CUB для операций ARGSORT и TOP_K, а также перенесли выполнение 48 узлов графа GATED_DELTA_NET с CPU на GPU. На модели Qwen3.8-27B в квантовании Q4_K_M перенос GATED_DELTA_NET увеличил скорость prefill (pp512, FA off) с 964,51 до 2119,26 токенов/с, а генерации (tg64) — с 10,15 до 15,50 токенов/с. Дополнительно исправлена блокировка потоков в fused-ядре TOPK_MOE и отключен нестабильный путь MMQ на PH1, что устранило появление NaN и обеспечило прохождение всех 22 237 тестов backend-ops.
Механизм оптимизации¶
Из-за отсутствия макроса __CUDA_ARCH__ в компиляторе mcc MUSA ошибочно использовала 8-байтовую ширину копирования памяти вместо 16-байтовой; явное задание 16 байт в ggml_cuda_get_max_cpy_bytes() удвоило объем передаваемых данных на поток за инструкцию при стейджинге K/V в shared memory FlashAttention. В ядре TOPK_MOE реализован клампинг индекса строк к последней строке для неполных блоков, что предотвратило ранний выход граничных варпов и зависание барьера синхронизации __syncthreads(). Кроме того, активация совместимых с SDK 5.2.0 ядер GATED_DELTA_NET позволила полностью исключить fallback на CPU для рекуррентных слоев современных гибридных архитектур.
Прирост производительности¶
Аппаратный оффлоад GATED_DELTA_NET на MTT S5000 ускорил стадию prefill для контекста в 512 токенов в 2,2 раза (с 964,51 до 2119,26 ток/с) и поднял темп генерации на 52,7% (с 10,15 до 15,50 ток/с). Расширение единицы копирования памяти в FlashAttention добавило еще 5,8% к скорости prefill на Qwen3.8-27B (рост с 751,15 до 794,73 ток/с) без деградации перплексии.