Окно мониторинга: 2026-11-02 07:00 ~ 2026-11-03 07:00 (пекинское время, 24 часа; точка отсечения предыдущего выпуска — 11-02 07:00).

Источники: порепозиторная проверка организации tile-ai на GitHub (29 репозиториев, пуши в 9 за окно) + поиск подтверждённых данных по смерженным и находящимся на рассмотрении PR (смержено 18, обновлено на рассмотрении 44) + полная проверка новых issue (8) + попатентный обзор коммитов в основную ветку (tilelang 5, TileOPs 7, TileFoundry 1, tilelang-ascend 4, tilelang-hygon 1, tilelang-lsp 1, сайт документации 5) + проверка страниц релизов и тегов (5 репозиториев) + два репозитория потребителей + запросы Google News на китайском и английском (через прокси) + Hacker News + arXiv.

Индекс текущего выпуска

  • Главное за сегодня: реализация постоянного ядра FlashMLA для Ascend 950 (#3678, добавлено 1997 строк, 42 коммита); исправление реальной регрессии в векторизованном планировании (#3734); директива #line переведена на построчную перепривязку (#3703); TileOPs — 7 слияний производительности и тестовой инфраструктуры за день; TileFoundry превращает сортировку в разделяемой памяти в доказуемый аналитический слой (#283) и показывает, что lowering в main на sm_89 даёт 4 ошибки из 7 программ; за 24 часа слито 18 изменений, 44 обновления в ревью, 8 новых issue
    1. Ход работ по ключевым проектам
      • 1.1 Основная линия: постоянное ядро FlashMLA для Ascend 950, двойные пути BF16 prefill и компактного FP8 decode (#3678) (11-02)
      • 1.2 Основная линия: остаточный 512-битный вектор после пропуска развязки cast, сбой компиляции ViT forward на SM100 (#3734) (11-02)
      • 1.3 Основная линия: директива #line переведена на построчную перепривязку, отладочное отображение может использовать плотный режим (#3703) (11-02)
      • 1.4 Основная линия: в примере внимания tile выводится из лимита разделяемой памяти устройства, на RTX PRO 6000 с 2 провалов из 7 до 9 полных прохождений (#3702) (11-02)
      • 1.5 Основная линия: битовое преобразование CUDA float32x8 переведено на встроенный bit cast, число инструкций SASS снижено с 19 до 18 (#3748) (11-02)
    1. Адаптация под множество бэкендов
      • 2.1 Ascend: добавлен forward-оператор WY Fast, обратный сегментный перенос wy_fast_bwd_split (#1827 / #1828) (11-02)
      • 2.2 Ascend: слияние проверки фактов документации T.tile.fill и T.tile.clear (#1634 / #1628) (11-02)
      • 2.3 Ascend в ревью: cython wrapper кэширует разбор формы символов, продолжение линии накладных расходов хоста (#1967) (11-02)
      • 2.4 Ascend в ревью: синхронизация версии z3 перенесена из основной линии, продолжается перенос синхронизации всего репозитория в основную линию (#1965 / #1975) (11-02)
      • 2.5 Ascend в ревью: два дефекта оператора редукции, broadcast по первой оси Parallel заменён на скалярное чтение (#1978 / #1968 / #1979) (11-02)
      • 2.6 Hygon: gfx948 обрабатывается как совместимый с gfx946, сначала обеспечена тестируемость существующих инструкций (#32) (11-02)
      • 2.7 Moore Threads, MetaX, mlir-ascend — без пушей в окне мониторинга (11-03)
    1. Экосистема и внедряющие стороны
      • 3.1 TileOPs: 7 слияний за день, 6 по производительности, 1 по тестовой инфраструктуре (11-02)
      • 3.2 TileOPs: softmax выделенного warp decode с перекрытием QK/PV, геометрическое среднее 0,952x на 478 вызовах сканирования (#2817) (11-02)
      • 3.3 TileOPs: MoE собирает токены по предварительной перестановке строк, 387,8 TFLOPS на H200 (#2815) (11-02)
      • 3.4 TileOPs: 16-битные поэлементные операции без деления с проверкой диапазона, 8-байтовые операнды — один вектор на поток (#2814 / #2813) (11-02)
      • 3.5 TileOPs: статистики Norm переведены на вычитание элементов строки, устранено переполнение вблизи 1e36 (#2811) (11-02)
      • 3.6 TileOPs: цель tileops_test позволяет запускать tests/ops без CUDA (#2816) (11-02)
      • 3.7 TileFoundry: сортировка в разделяемой памяти оформлена как доказуемый аналитический слой, из 83 фикстурных планов 52 отклонены, 31 принят (#283) (11-02)
      • 3.8 TileFoundry в ревью: CUDA lowering выдаёт план синхронизации и выполняет обратную проверку, в main 4 ошибки из 7 программ (#285) (11-02)
      • 3.9 Внедряющие стороны: FlashQLA, TileKernels — без пушей в окне мониторинга (11-03)
    1. Сообщество, руководства и мероприятия
      • 4.1 Корректность синхронизации Ascend становится центром сосредоточенных сообщений об ошибках: 8 новых issue за день, два репортёра систематически отправляют их раздельно (11-02)
      • 4.2 SIMD-загрузка многовекторных и транспонированных представлений теряет зависимость DMA, корневая причина — восстановление логической области (#3749) (11-02)
      • 4.3 Новостная сторона: Google News на китайском и английском, Hacker News и arXiv — ни одного попадания в окне мониторинга (11-03)
    1. Наблюдения за трендами
  • Приложение. Материалы и пояснения к проверке

Главное за сегодня

В этот день основная линия была только одна: TileLang перенёс ядро декодирования MLA из DeepSeek на Ascend 950, причём перенёс в форме, планируемой компилятором. #3678 за один раз добавляет 1997 строк, 20 файлов, 42 коммита, охватывая два пути — BF16 prefill и компактный FP8 decode, а также опциональный дополнительный KV-кэш FP8 и FP4, attention sink, недопустимые индексы и независимую длину для каждого query. Способ реализации заслуживает внимания: цикл внимания, онлайн softmax, деквантование и планирование полностью написаны на TileLang, с помощью T.Stage компилятор управляет версиями буферов и синхронизацией, QK и PV используют два разных представления layout одного и того же кольца памяти L0B, размещение в L1 задаётся напрямую через T.annotate_buffer_addresses; prefill в одном выделении аккумулятора разделяется на четыре непересекающихся выходных среза, и после вычисления нормализации каждый срез сразу выгружается. Чтобы такая форма записи стала возможной, на стороне компилятора одновременно добавлены три возможности — копирование из L0C в UB с сохранением вывода NZ, выбор каждого потока AIV перед собственным конвейером, сохранение исходного порядка для явно планируемых стадий при совпадении физической фазы выпуска. Целевая форма задана очень узко: число голов 64, Dqk и Dv оба равны 512, topk кратен 64.

Вторая линия — исправление реальной регрессии компилятора. #3734 касается прямого прохода ViT на SM100: Cannot convert type float32x16 to CUDA type. Путь длинный, но каждый шаг ясен — планировщик векторизации выбрал 16 lane для 256-битного хранилища BF16 и отложил ограничения cast, рассчитывая, что DecoupleTypeCast разделит вычисления и обращение к памяти; а #3579 незадолго до этого исключил «охранные условия чтения памяти» из развязки (дублирование этих охранных условий меняет семантику), в результате на последнем шаге легализации ширина всё ещё выбиралась по отложенной стратегии, остаточное выражение FP32 раздувалось в 512-битный вектор, который бэкенд CUDA не может выразить. Исправление состоит в том, чтобы сделать отложенные ограничения параметром планировщика, а не свойством IR; только последний шаг легализации передаёт false, преобразование FP32 в BF16 для охранных условий снижается до 8 lane, успешно развязанные копии BF16 по-прежнему сохраняют 16 lane; cast адреса также вынесен из ограничений типа вектора данных и определяется легальностью на основе существующего анализа доступа.

Третья линия — опыт отладки. #3703 меняет директиву #line с «привязки только к первой строке каждого оператора TIR» на опциональную многоуровневую генерацию: уровень 0 отключает, 1 сохраняет прежнее поведение, 2 и выше заново привязывает каждую непустую сгенерированную строку, включая операторы, развёрнутые в несколько строк, и закрывающие фигурные скобки. Реализация размещена в общем потоке вывода для языков C, используемом CUDA, HIP, CPU C, C host, Metal, Ascend, и не затрагивает vendored TVM; операторы и функции без исходного диапазона явно переключаются на <tilelang-generated>, исходный пользовательский текст сохраняется дословно, чтобы не нарушить продолжение макросов и строковые литералы. Добавлена переменная окружения TILELANG_EMIT_LINE_DIRECTIVES в качестве значения по умолчанию на стороне Python, которая сводится в конфигурацию pass перед поиском в кэше. Локально 42 пункта пройдено, 6 пропущено, случаи CUDA и HIP пропущены на машине автора, фактическое поведение cuda-gdb ещё не проверено напрямую — автор сам это указал. В тот же день в репозитории LSP последовал коммит, выравнивающий отладку CUDA по новому отображению строк.

Темп на стороне операторов не остановился, но направление изменилось. Из 7 слитых за день коммитов TileOPs 6 посвящены производительности и 1 — тестовой инфраструктуре, все от одного автора; каждый по-прежнему сопровождается таблицей фактических замеров на H200 и порогом регрессии по чек-листу. Наиболее ценным является #2816: добавлена цель tileops_test, существующая только во время тестов, которая на PrivateUse1 вычисляет общий эталонный workload для каждого оператора, чтобы дымовые тесты tests/ops больше не требовали наличия CUDA. Изменение составляет 1281 строку, 59 файлов, автор также чётко обозначил границы — сравнение с общим эталоном может доказать только независимость контракта от бэкенда и переносимость, но не численную корректность реального бэкенда, эта линия по-прежнему висит на #2611.

Наконец, скрытая линия TileFoundry. #283 оформляет «как следует упорядочивать обращения к разделяемой памяти» в слой анализа: семейство синхронизации проходит по каждому функции в порядке выпуска lowering, для каждого обращения к разделяемой памяти берёт инструкцию, меш выпуска, операнды и диапазон байтов, перед конфликтующими обращениями планирует commit/wait и барьеры в области видимости. Из 83 фикстур планирования он планирует 52, отклоняет 31 (24 используют копирование TMA, 6 используют wgmma, семантика завершения обоих ещё не объявлена; 1 требует вставки барьера в цикле с неодинаковыми границами через потоки). Сопутствующий #285 на рассмотрении, превращает план в lowering CUDA и выполняет обратную проверку: из 7 программ, собираемых на одной карте sm_89, 6 полностью совпадают с эталоном, synccheck и memcheck полностью чистые, а при lowering способом main 4 из тех же 7 дают неверный результат; непосредственный симптом, указанный в описании #283, — максимальное отклонение E4M3 GEMM на sm_89 составляет 659.

1. Ход развития основных проектов

1.1 Основная линия: постоянное ядро FlashMLA для Ascend 950, двойной путь BF16 prefill и компактного FP8 decode (#3678)

Объединено 11-02 09:26, добавлено 1997 строк, удалено 17 строк, 20 файлов, 42 коммита — наиболее весомый вклад в этом окне мониторинга. Ядро ориентировано на Ascend 950: prefill идёт через BF16, decode — через компактный FP8, опционально FP8 и FP4 для дополнительного KV-кэша, а также поддерживает attention sink, недопустимые индексы и независимую длину для каждого query. Цикл внимания, онлайн-softmax, деквантование и планирование полностью выражены на TileLang и управляются через T.Stage, а также версии буферов и синхронизацию, управляемые компилятором; QK и PV совместно используют два разных представления layout одного кольца памяти L0B, размещение в L1 задаётся через T.annotate_buffer_addresses; prefill в рамках существующего распределения аккумулятора выделяет четыре непересекающихся выходных среза и комплектует их независимыми событиями синхронизации — по завершении нормализации одного среза он немедленно выгружается; decode использует копирование вывода целой строкой. На стороне компилятора добавлены три улучшения: копирование из L0C в UB сохраняет выход NZ, каждый поток AIV выбирается перед своим конвейером, явные фазы планирования сохраняют порядок исходного кода при выравнивании физических фаз запуска. Цель — число голов 64, Dqk и Dv оба 512, topk кратен 64; pair gather использует уже задокументированный обход генерации кода с нулевым burst; прилагаются раннеры корректности и бенчмарков, регрессии повторного использования буферов и границ query, а также README с объявлением лицензии FlashMLA MIT.

Источник: https://github.com/tile-ai/tilelang/pull/3678

1.2 Основная ветка: остаточный 512-битный вектор после пропуска развязки cast, сбой компиляции прямого прохода ViT на SM100 (#3734)

Объединено 11-02 09:05. На e8225d7f5 прямой проход ViT для SM100 сообщает Cannot convert type float32x16 to CUDA type в месте защищённого сохранения вывода FP32→BF16: планировщик выбирает 16 lane для 256-битного сохранения BF16 и откладывает ограничения cast, ожидая, что DecoupleTypeCast разделит вычисления и обращения к памяти; после того как #3579 заставил развязку пропускать те защиты, которые читают память, последующий LegalizeVectorizedLoop продолжает откладывать по той же стратегии, и остаточное выражение FP32 превращается в 512-битный вектор, не представимый на стороне CUDA. Исправление делает defer_cast_constraints C++-параметром планировщика (а не свойством IR), выбираемым вызывающей стороной по фазе: LayoutInference и LowerTileOp сохраняют true, финальная легализация всегда передаёт false, а вычисленная ширина передаётся вперёд как явный vectorize_hint, чтобы переписыватель не перепланировал по стратегии по умолчанию. Защищённое преобразование FP32→BF16 снижается до 8 lane, успешно развязанное копирование BF16 сохраняет 16 lane, проверка безопасности защиты сохраняется; приведение адреса и параметры указателя доступа исключаются из ограничений типа вектора данных, непрерывный адрес низводится до скалярного базового адреса плюс шаг, легальность определяется существующим анализом доступа, выровненный векторный доступ с индексом int64 и четырёхlane атомарные случаи не затрагиваются.

Источник: https://github.com/tile-ai/tilelang/pull/3734

1.3 Основная ветка: директива #line переходит к построчной перепривязке, отображение отладчика получает опциональный плотный режим (#3703)

Объединено 11-02 04:06, 682 добавленные строки, 11 файлов. Прежняя генерация #line привязывала только первую строку, сгенерированную одним оператором TIR, последующие строки продвигались логическим номером строки компилятора, и cuda-gdb относил сгенерированные загрузки, операции и сохранения к посторонним строкам Python; синтезированные операторы и завершающие фигурные скобки могли наследовать уже продвинутое отображение, а функции без исходного диапазона наследовали исходный файл предыдущей функции. Данный коммит расширяет tl.emit_line_directives до уровней: 0 или False отключает и сохраняет поведение по умолчанию, 1 или True сохраняет уровень операторов, 2 и выше перепривязывает каждую непустую сгенерированную строку (включая многострочно развёрнутые операторы и завершающие фигурные скобки). Реализация находится в выходном потоке C-семейства, охватывает CUDA, HIP, CPU C, C host, Metal и Ascend, не изменяя vendored TVM; одновременно записываются исходные диапазоны операторов и функций, для областей без источника явно переключается на <tilelang-generated>, пользовательский текст сохраняется дословно, чтобы не нарушить продолжение макросов и строковые литералы. Добавлена переменная TILELANG_EMIT_LINE_DIRECTIVES как значение по умолчанию на стороне Python, сворачиваемая в конфигурацию pass перед поиском в кэше, явная конфигурация имеет приоритет (включая случаи недопустимости значения по умолчанию). Локально проверено сборкой Metal плюс LLVM: две группы по 42 элемента — директивы строк и compile-only — прошли, 6 пропущено; после удаления директив два фикстура — векторизованное сложение и условный выбор — дают полностью идентичные артефакты на уровнях 0, 1, 2. Автор отмечает, что случаи CUDA и HIP на его машине были пропущены, поведение cuda-gdb напрямую не тестировалось.

Источник: https://github.com/tile-ai/tilelang/pull/3703

1.4 Основная ветка: пример внимания выводит tile из лимита разделяемой памяти устройства, на RTX PRO 6000 результат меняется с 2 успехов и 7 провалов на 9 успехов из 9 (#3702)

Объединено 11-02 06:59, добавлено 345 строк, 13 файлов. В нетюнинговых путях нескольких примеров внимания размеры tile для Hopper были жёстко зашиты; на sm_86, sm_89, sm_120 лимит динамической разделяемой памяти на блок составляет около 99 KiB, и запуск сразу падал: в исходном тексте ошибки требовалось установить разрешённую динамическую разделяемую память в 163840 байт, тогда как устройство 0 поддерживает максимум 101376 байт на блок. Теперь каждый пример выводит tile из shared_memory_per_block_optin, следуя подходу существующего примера GQA decode (example_gqa_decode.py); конфигурация для дата-центров не меняется, поскольку каждая ветка срабатывает только при потребности в tile ниже предыдущего уровня. Результаты на устройствах с ~99 KiB: четыре примера attention sink — 64×64 двухстадийный со 128 потоками (прямой проход) и 64×32 одностадийный со 128 потоками (обратный), example_gqa_bwd.py — прямой 64×64, обратный 64×32, example_mla_decode.py — BLOCK_N 32, BLOCK_H 16. В последние два добавлен параметр threads: при tile в 64 строки или 16 строк T.gemm требует четыре warp, чтобы каждый warp занимал целые 16 строк; при 256 потоках lowering сразу падает (сообщается о конфликте layout для acc_s и acc_s_cast, а также об отсутствии допустимого разбиения на warp при M=16 N=32). Кроме того, flash-attn 2.x не имеет ядер для sm_12x, и четыре примера varlen изначально падали на этапе эталона; теперь однократно проверяется flash_attn_varlen_supported, и при отсутствии поддержки используется чисто PyTorch-реализация varlen_attention_ref (покрывает GQA, causal с выравниванием по правому нижнему краю и window_size). Проверка на RTX PRO 6000 Blackwell (sm_120): CUDA — все 9 пунктов пройдены (ранее 2 пройдено, 7 провалено), CuTeDSL — также все 9 пройдены; на H100 с доступным лимитом 227 KiB по-прежнему используются исходные tile, также все 9 пунктов пройдены.

Источник: https://github.com/tile-ai/tilelang/pull/3702

1.5 Основная ветка: преобразование битов CUDA float32x8 переведено на встроенный bit cast, число инструкций SASS сократилось с 19 до 18 (#3748)

Объединено 11-02 09:58, добавлено 20 строк, удалено 3 строки, исправляет #3679 и заменяет #3681. Изменение затрагивает только CUDA: переинтерпретация битов float32x8 заменена с последовательности битовых операций на __builtin_bit_cast(uint64_t, make_float2(v, v)), число генерируемых инструкций SASS сократилось с 19 до 18. Прирост небольшой, но это типичный для бэкенда компилятора случай завершающей правки класса «семантика эквивалентна, код короче».

Источник: https://github.com/tile-ai/tilelang/pull/3748

2. Мультибэкендная адаптация (Ascend / Hygon / Moore Threads / MetaX)

2.1 Ascend: добавлен прямой оператор WY Fast, перенесён сегментированный обратный проход wy_fast_bwd_split (#1827 / #1828)

Оба изменения влиты 11-02 02:33, с разницей менее минуты. #1827 способом «перенос и добавление» переносит оператор WY Fast на сторону Ascend NPU, #1828 переносит и дополняет сегментированный обратный проход wy_fast_bwd_split. Представление WY — распространённый приём в моделях класса линейного внимания, позволяющий свести умножение на рекуррентную матрицу к низкоранговой форме; прямой проход вместе с сегментированным обратным замыкают эту цепочку операторов на Ascend до состояния, пригодного для обучения. В описаниях обоих PR есть только шаблонное приветствие; объём изменений определяется по заголовкам и коммитам, и настоящий отчёт не делает выводов о деталях внутренней реализации.

Источники: https://github.com/tile-ai/tilelang-ascend/pull/1827, https://github.com/tile-ai/tilelang-ascend/pull/1828

2.2 Ascend: влита проверка фактов документации T.tile.fill и T.tile.clear (#1634 / #1628)

В том же пакете влиты две проверки фактов со стороны документации, касающиеся поведения API T.tile.fill и T.tile.clear. В обсуждении #1628 оставлена договорённость о поддержке: эти два интерфейса довольно близки, при изменении одного следует проверить, не требуется ли соответствующая правка другого. Это часть непрерывной серии проверок фактов API; в то же время на рассмотрении находятся T.tile.bitwise_not (#1656), T.tile.bitwise_and (#1651), T.tile.bitwise_lshift (#1659) и T.tile.bitwise_or (#1655).

Источники: https://github.com/tile-ai/tilelang-ascend/pull/1634, https://github.com/tile-ai/tilelang-ascend/pull/1628

2.3 Ascend на рассмотрении: кэширование разбора формы символов в cython wrapper, продолжение линии накладных расходов хоста (#1967)

Обновлено 11-02. #1967 кэширует результаты разбора формы символов в cython wrapper, цель — дальнейшее снижение накладных расходов dispatch на каждый вызов, соответствует дефекту #1898. Эта линия и влитый в прошлом выпуске #1948 — две точки одной цели: #1948 заменил безусловное конструирование TVM Analyzer при каждом вызове на ленивое конструирование (на 910B3 с примерно 450 микросекунд до примерно 106 микросекунд), #1967 занимается стоимостью самого повторного разбора на символьном пути. В том же семействе находится на рассмотрении #3757, на стороне основной ветки выравнивается динамический ABI и проверка общих символов: ключ по объекту Var, а не по печатному имени, чтобы одноимённые разные переменные не считались одним символом, а stride по одноэлементным измерениям рассматривается только как запасной вариант и не позволяет ему первым определять значение.

Источники: https://github.com/tile-ai/tilelang-ascend/pull/1967, https://github.com/tile-ai/tilelang/pull/3757

2.4 Ascend на рассмотрении: версия z3 синхронно переносится из основной ветки, перенос всего репозитория на основную ветку продвигается (#1965 / #1975)

#1965 обновлён 11-02 в 12:18, переносит insert sync версии z3 из основной ветки в репозиторий Ascend. #1975 обновлён в тот же день в 03:28, по-прежнему переносит TileLang-Ascend целиком на многобэкендовую архитектуру основной ветки. Это два уровня одного пути сходимости: один меняет реализацию алгоритма синхронизации, другой — скелет репозитория. Сопутствующие действия на стороне основной ветки — #3752 (конфигурация z3 rlimit) на рассмотрении и #3735, предлагающий доказывать непересечение модульных индексов в thread sync — проблема бюджета доказательств, оставленная откатом ThreadSync в прошлом отчёте, решается одновременно с двух концов: «настройка ресурсов решателя» и «сокращение форм, требующих доказательства».

Источники: https://github.com/tile-ai/tilelang-ascend/pull/1965, https://github.com/tile-ai/tilelang-ascend/pull/1975

2.5 Ascend на рассмотрении: два дефекта операторов редукции, броадкаст по первой оси Parallel переведён на скалярное чтение (#1978 / #1968 / #1979)

Все три обновлены 11-02. #1978 исправляет ошибочное уменьшение размера tile-оператора вдвое для целого блока [1,N] UB в цикле разбиения по строкам в AscendVidReduction (соответствует #1800); #1968 вдвое уменьшает счётчики sigmoid и round при threads=2 (соответствует #1897); #1979 переводит операнд броадкаста по первой оси Parallel на скалярное чтение, обходя ловушку выравнивания UB в VEC (соответствует #1885). Все три относятся к одному классу проблем: на стороне Ascend предположения векторизованного пути о форме и выравнивании жёстче, чем на стороне CUDA, и нарушение граничных условий даёт жёсткую ошибку на этапе компиляции или выполнения.

Источники: https://github.com/tile-ai/tilelang-ascend/pull/1978, https://github.com/tile-ai/tilelang-ascend/pull/1968

2.6 Hygon: gfx948 обрабатывается как совместимый с gfx946, сначала обеспечить тестируемость существующих инструкций (#32)

Влито 11-02 в 10:46. Изменение трактует gfx948 как совместимый с gfx946, переиспользуя существующие MLS atoms, MMAC, barrier и опции компиляции, чтобы до подключения gather load сначала прогнать существующие builtins на новой модели. Это типичный двухшаговый подход: сначала подключить номер новой архитектуры к существующему пути и получить тестируемую базовую линию, затем дополнить то, что действительно требует новых инструкций. Предыдущее исправление LDS swizzle обычным копированием в этом репозитории влито в прошлом окне.

Источники: https://github.com/tile-ai/tilelang-hygon/pull/32

2.7 Moore Threads, MetaX, mlir-ascend: в окне нет пушей (11-03)

Последний пуш tilelang-musa остановился на 09-30, tilelang-metax — на 10-26, tilelang-mlir-ascend — на 10-31 08:21 (раньше начала этого окна). Динамика бэкендов в этом выпуске поэтому сосредоточена на стороне Ascend и Hygon, причём Ascend занимает 4 из 18 влитых и большинство из 44 обновлений на рассмотрении.

Источники: https://github.com/tile-ai/tilelang-musa, https://github.com/tile-ai/tilelang-metax

3. Экосистема и стороны внедрения (TileOPs / TileFoundry / TileKernels / FlashQLA)

3.1 TileOPs: 7 влитий за один день, 6 по производительности, 1 по тестовой инфраструктуре (11-02)

Все 7 коммитов были объединены в период с 05:10 до 05:22 11-02, автор тот же — lcy-seso; они охватывают шесть направлений: внимание, MoE, поэлементные операции, нормализация, тестовые цели и CI-контракты. По сравнению с предыдущим выпуском «13 коммитов за день» количество сократилось вдвое, но структура изменилась: в предыдущем выпуске это была хвостовая настройка по каждому shape, в этом выпуске появились два типа непроизводительных изменений — тестовая инфраструктура (#2816) и CI-контракты (#2818), что свидетельствует о том, что библиотека операторов начинает структурную подготовку к мультибэкендности и поддерживаемости. Каждый коммит по-прежнему сопровождается таблицей фактических замеров на H200 и заявлением об отсутствии регрессий уже измеренных строк в manifest.

Источники: https://github.com/tile-ai/TileOPs/commits/main

3.2 TileOPs: softmax для warp-специализированного decode с перекрытием QK/PV, геометрическое среднее по 478 вызовам сканирования — 0.952× (#2817)

Объединён в 05:21 11-02. Изменение перекрывает softmax warp-специализированного GQA decode с матричными умножениями QK и PV: вероятности остаются в регистрах, K и V используют независимые барьеры, совместно используя трёхслотовый кольцевой буфер. Условия выбора: D равно 128, tile из 64 ключей, одна строка tile на каждый KV-заголовок, отсутствие softcap, а сетка «один CTA — один SM» способна удовлетворить границам занятости и длины потока; вызовы, не удовлетворяющие условиям, продолжают идти по последовательной программе. Все девять границ планировщика проверяются на стороне хоста с использованием фиксированных фактов устройства, не требуя реального выделения тензоров внимания и не требуя карты с 132 SM. Среда фактических замеров: H200, CUDA 13.2, PyTorch 2.13.0+cu132, TileLang 0.1.12, с использованием времени занятости устройства CUPTI: для формы (16, 8, 4096, 64, 8, 128) float16 составляет 0.074610 мс против 0.077490 у fa3 (1.039×, 3.654 TB/s), bfloat16 — 0.074260 против 0.077220 (1.040×, 3.671 TB/s); зафиксированное геометрическое среднее задержки по 478 затронутым вызовам сканирования составляет 0.952× от main.

Источники: https://github.com/tile-ai/TileOPs/pull/2817

3.3 TileOPs: MoE собирает по токенам с предварительной перестановкой строк, 387.8 TFLOPS на H200 (#2815)

Объединён в 05:19 11-02. Вход MoE с tight-раскладкой изменён на сбор по токенам, одна и та же строка hidden переиспользуется между несколькими её маршрутами; когда токены не заполняют GPU, маршруты разбиваются по разным блокам. Одновременно удалены неиспользуемые в tight-сканировании выделение и запись индекса источника экспертов. В планировании добавлено правило: для неслитых tight-групп с K-мажорным порядком, если среднее число строк больше 128 и не превышает 192, выбирается ping-pong tile 64×256 на H200; вызовы с явно заданным размером группы по-прежнему идут через модель стоимости. Фактические замеры на (4096, 256, 8, 7168, 2048): bfloat16 — 7.4417 мс против 7.6982 у flashinfer-cutlass (1.034×, 387.844 TFLOPS), float16 — 7.5040 против 7.9869 (1.064×, 384.624 TFLOPS); записанные сканирования маршрутов и чётных групп покрывают оба dtype и режимы смещения, в manifest нет регрессий выше уровня шума.

Источники: https://github.com/tile-ai/TileOPs/pull/2815

3.4 TileOPs: 16-битные поэлементные операции избавляются от деления с проверкой диапазона, 8-байтовые операнды — один вектор на поток (#2814 / #2813)

#2814 объединён в 05:18 11-02: для bfloat16 Sigmoid и 16-битного SiLUAndMul используется обратная величина с масштабом в четверть, что обходит деление с проверкой диапазона, и этот вспомогательный метод сделан общим с SiLU; одновременно убрана ненужная защита от субнормальных чисел в операции экспоненты, выход fused-гейта размещается между потоками смежными 16-байтовыми векторами, а также сохраняется нерундированная активация SiLU с умножением на бесконечность при значениях gate, близких к минус 88. Фактические замеры: bfloat16 Sigmoid на 16777216 элементах — 0.017504 мс против 0.018656 у torch-compile (1.066×, 3.834 TB/s), на 268435456 элементах — 0.248545 против 0.262529 (1.056×, 4.320 TB/s); SiLUAndMul на (2048, 14336) — 0.041856 против 0.045440 (1.086×). #2813 объединён в 05:17: 8-байтовые операнды при длинных запусках обрабатываются как один вектор на поток; перед изменением с помощью roofline-анализа TileFoundry подтверждено, что все такие побитовые операторы ограничены памятью: при пропускной способности HBM 4.8 TB/s получены нижние границы для каждой нагрузки — hidden-state-prefill для int64 — 41.9 мкс, для int32 — 21.0 мкс, llama-7b-ffn для int64 — 56.4, для int32 — 28.2, cnn-feat-broadcast для int64 — 42.8, linear-bias для int64 — 14.0, для int32 — 7.0.

Источники: https://github.com/tile-ai/TileOPs/pull/2814, https://github.com/tile-ai/TileOPs/pull/2813

3.5 TileOPs: статистики Norm переведены на вычитание элемента строки, устранено переполнение вблизи 1e36 (#2811)

Смержен 02.11 в 05:10. Корневая причина описана прямо: все ядра среднего и дисперсии выполняют fp32-суммирование по исходной строке, сумма растёт вместе со смещением строки (примерно N на x ), строки вблизи 1e36 переполняются напрямую, и все выходы превращаются в NaN; даже без переполнения среднее округляется на масштабе смещения, x − mean вносит это округление в каждый элемент, и тогда дисперсия, вычисленная для константной строки, равна квадрату округления, а нормализация даёт ±weight вместо bias. Исправление добавляет инвариант для каждой fp32-величины, образуемой строкой: она ограничена только разбросом строки, а не смещением — на практике берётся «некоторый элемент внутри строки» как база и затем вычитается: в warp-ориентированных ядрах используется первый элемент warp, в шардированных ядрах — первый элемент строки. Это прямое исправление того же класса проблем, что и четыре дефекта операторов Norm, перечисленные в #2812 в прошлом отчёте.

Источник: https://github.com/tile-ai/TileOPs/pull/2811

3.6 TileOPs: цель tileops_test позволяет запускать tests/ops без CUDA (#2816)

Смержен 02.11 в 05:20, это самое крупное изменение за окно мониторинга: 1281 добавленная строка, 524 удалённых, 59 файлов, закрывает #2803. Добавлена цель tileops_test, существующая только во время тестов: на PrivateUse1 для каждого оператора вычисляется общая workload-референция, и дымовые тесты tests/ops выполняются через неё; неподдерживаемые вызовы референции явно перечислены как unvalidated и пропускаются, а исключения, действующие только на CUDA, проверяются в рамках той же цели. Реестр сопоставления операторов workload-ам разделяется через workloads/calls, привязки, специфичные для бенчмарков, учёт копирований и переопределения входов остаются в benchmarks/_cases; дополнительно добавлены генератор тестовых устройств, поддержка свёрток, исправлены метки возможностей тестов, а референсные реализации приведены в соответствие с семантикой scale и пустых строк в обратном проходе GQA. Автор явно указывает ограничения: эта цель проверяет независимый от бэкенда контракт и переносимость тестов, сравнение с общей референцией не доказывает численную корректность реального бэкенда, #2611 остаётся открытым; вызовы без доступной референции по-прежнему видны в форме unvalidated. В той же партии #2818 отделяет контракт dispatch от генерации данных workload: для контрактных用例 создаются только формы, без генерации и проверки неиспользуемого содержимого тензоров, а пять наборов структурных тестов хоста прогоняются в compile-contract-gate только один раз.

Источник: https://github.com/tile-ai/TileOPs/pull/2816

3.7 TileFoundry: упорядочивание обращений к разделяемой памяти оформлено как доказуемый аналитический слой, из 83 фикстур 52 запланированы и 31 отклонена (#283)

Смержен 02.11 в 02:09. Проблема вполне реальна: finalize не выдаёт commit, wait или barrier между обращениями к разделяемой памяти, поэтому планировщики, использующие разделяемую память как промежуточное хранилище, создают гонки, и на карте sm_89 E4M3 GEMM отклоняется от torch максимум на 659. Направление, заданное в #267: «анализ порядка обращений на HIR» плюс «lowering, выдающий порядок согласно цели», причём выдаётся только минимально корректная форма, а при невозможности доказательства — отказ; данный коммит поставляет аналитическую часть. Новое семейство синхронизации проходит по каждой функции в порядке выдачи lowering: для каждого обращения к разделяемой памяти берутся инструкция, mesh выдачи и операнды из analysis/emission.py, область видимости из IterationScope, число байт из размещения, записанного семейством памяти; перед конфликтующим обращением планируются один commit и один wait (для незавершённых записей cp.async) и один барьер, покрывающий «самую внутреннюю область, одновременно содержащую оба», с нумерацией по активным позициям его ранее выполняющихся дочерних узлов. Операторы через новый реестр visitor_registry/completion.py объявляют, как завершается эффект; всё, что ничего не объявляет, трактуется как UNKNOWN; семейство отклоняет обращения после UNKNOWN, обращения UNKNOWN, всё ещё находящиеся в полёте при возврате, вставку барьера в цикле, не дожидающемся на границе потоков, и случаи сохранённого вызова, когда вызываемая сторона владеет разделяемой памятью. По умолчанию ни один вызывающий не запрашивает это семейство, поэтому 75 планировочных фикстур, выполняющих finalize, побайтово совпадают с main; из 83 фикстур запланированы 52, отклонены 31 (24 используют копирование TMA, 6 используют wgmma и ещё не объявили семантику завершения; 1 требует вставки барьера в цикле на границе, зависящей от потока), и用例 корпуса построчно утверждают эти отказы.

Источник: https://github.com/tile-ai/TileFoundry/pull/283

3.8 TileFoundry на рассмотрении: CUDA lowering выдаёт план синхронизации и выполняет обратную проверку, на main 4 из 7 программ ошибочны (#285)

Обновлено 11-02 12:40. Этот коммит — исполнительная часть #283: CUDA lowering выполняет шаги согласно плану синхронизации и построчно проверяет, что фактическое выполнение совпадает с планом; при расхождении выбрасывается LoweringError. Результаты на корпусе: в стандартном процессе 14 программ проходят с синхронизацией, 31 отклонена; у этих 14 TIR отличается от main только добавленными инструкциями, ещё 30 полностью совпадают с main; после прохождения через основной visitor 75 программ, которые main финализирует, побайтово идентичны. На одном sm_89 запущены 7 программ, которые собираются и содержат синхронизацию: 6 полностью совпадают с эталоном и чисты под synccheck и memcheck; те же 7 программ, пониженные способом main, дают 4 ошибочных результата. Автор также сообщает о двух уже существующих в main проблемах: на GEMM и loop_outputs racecheck сообщает, что несколько lane одной и той же cp.async пишут по одному адресу (copy_async в HIR не принимает разбиение цели, поэтому каждый поток копирует весь tile), а loop_outputs к тому же записывает один выход в неправильный элемент.

Источник: https://github.com/tile-ai/TileFoundry/pull/285

3.9 Принимающие стороны: FlashQLA и TileKernels без пушей в окне (11-03)

Последний пуш QwenLM/FlashQLA остановился на 10-30 03:02, deepseek-ai/TileKernels — на 10-14 05:49; в обоих репозиториях в этом окне коммитов не было. TileKernels с 10-14 уже несколько периодов подряд без пушей; вся видимая активность принимающих сторон в этом периоде приходится на два собственных репозитория — TileOPs и TileFoundry.

Источники: https://github.com/QwenLM/FlashQLA, https://github.com/deepseek-ai/TileKernels

4. Сообщество, руководства и мероприятия

4.1 Корректность синхронизации на Ascend стала средоточием сообщений об ошибках: 8 новых issue за один день, два репортёра систематически подают отдельно (11-02)

В окне в основном репозитории появилось 8 новых issue с номерами преимущественно от #3740 до #3749, все касаются генерации кода и вставки синхронизации для двух бэкендов — AscendC и PTO. Из них 7 поданы одним репортёром в окружении TileLang 0.1.15, PyTorch 2.10.0+cpu, torch_npu 2.10.0; тема — расхождение поведения одной и той же программы на двух бэкендах: сгенерированный код по-разному выбирает запускаемые ядра (AscendC выдаёт mix(1,2), PTO — vector), при enable_auto_schedule=True AscendC не вставляет синхронизацию, при выключенном автошедулинге генерация успешна, а при включённом падает с неясным сообщением об ошибке, один и тот же код AscendC генерирует, а PTO — нет. Общая черта этой группы отчётов — воспроизводимый prim_func и построчное сопоставление сгенерированного кода, то есть форма, напрямую преобразуемая в регрессионные тесты. Вместе с находящимися на рассмотрении #3735 (доказательство непересечения модульных индексов в thread sync) и #3580 (синхронизация предикатов и продвинутых индексов) корректность синхронизации на стороне Ascend стала сейчас наиболее средоточенной линией.

Источники: https://github.com/tile-ai/tilelang/issues/3747, https://github.com/tile-ai/tilelang/issues/3745

4.2 SIMD-загрузки с межстрочными и транспонированными представлениями теряют зависимость от DMA, корневая причина — в реконструкции логической области (#3749)

Самая технически содержательная запись из этой партии; окружение репортёра — ветка main e8225d7f5, Ascend950PR, CANN 9.2.0. Проблема в том, что две последовательные SIMD-загрузки теряют зависимость от более раннего копирования GM в UB: во-первых, загрузка по строкам начинается с середины строки и переходит на следующую строку; во-вторых, доступ идёт к транспонированному представлению, и физически непрерывное измерение не является последним логическим. Путь корневой причины описан подробно: сгенерированный tl.access_ptr покрывает физический диапазон загрузки, а логическая область, восстановленная MemoryAccessDetector::LogicalPointerRegion, упускает часть её, поэтому анализ зависимостей считает входной DMA не связанным с загрузкой, AutoSchedule переносит копирование входа после загрузки, а InsertSync не выдаёт set/wait MTE2 к V. В минимальном воспроизведении work.shape равен (2, 128), сначала T.copy(A, work[1, 0:64]), затем в T.SimdVF() T.simd.vld(work[0, 64]) читает 128 последовательных элементов половинной точности: сгенерированный дескриптор чтения — tl.access_ptr(work[0, 64], 128, 1), физически покрывающий [64, 192), где [128, 192) — как раз только что записанная DMA часть, а восстановленная область — work[0:1, 0:128], то есть физически [0, 128), полностью минуя записанный интервал; число строк вычислено как ceil(128/128) = 1 строка, начальный столбец 64 не учтён. Эта запись относится к тому же фундаменту, который в этом окне в основном дорабатывают для FlashMLA и векторизации PTO; при попадании она даёт тихо неверный результат.

Источник: https://github.com/tile-ai/tilelang/issues/3749

4.3 Сторона новостей: ни в Google News (китайский и английский), ни в Hacker News, ни в arXiv нет совпадений в окне мониторинга (11-03)

Шесть групп запросов в Google News (английские TileLang, TileLang GPU kernel, tile-ai TileLang; китайские оператор TileLang, TileLang Ascend, MetaX или Hygon или Moore Threads) вернули в качестве последней релевантной публикации статью SemiAnalysis от 10-21 о Huawei и DeepSeek, остальные сосредоточены в партии от 09-29 до 10-01 об открытых компонентах Ascend от DeepSeek — все они раньше данного окна мониторинга. Последнее совпадение по all:"TileLang" на arXiv — TileSight от 2026-07-24, совместный запрос TileLang и Ascend не дал результатов; среди кандидатов в Hacker News с TileLang связаны только записи о пушах в репозиторий. Из 14 кандидатов 12 происходят из обнаружения репозиториев GitHub, 2 — не относящиеся к теме Show HN и репозиторий LSP; информационная картина этого выпуска полностью сформирована стороной репозиториев.

Источник: https://news.google.com/, https://arxiv.org/

5. Наблюдения за трендами

Самое ценное суждение этого дня: Ascend больше не «цель адаптации», а «цель первого выпуска». FlashMLA — ключевой оператор на собственной ветке декодирования DeepSeek — напрямую записан в основную ветку с целью Ascend 950, причём используются низкоуровневые возможности вроде T.Stage, annotate_buffer_addresses, предназначенные для явного планирования, а не отдельный репозиторий для портирования. В сочетании с тремя изменениями на стороне компилятора (сохранение NZ при передаче из L0C в UB, выбор потока AIV перед конвейером, эмиссия явных стадий в порядке исходного кода) видно, что путь Ascend переходит от «может работать» к «может работать согласно замыслу планирования».

Вторая линия — центр тяжести корректности смещается вниз. В прошлом выпуске основная ветка откатила целый набор улучшений корректности ThreadSync ради миграции TIRX, и в этом выпуске сразу появились две компенсации: в основной ветке #3752 настраивается решатель z3 для ограничений ресурсов, #3735 предлагает использовать непересекаемость модульных индексов для сокращения числа форм, требующих доказательства; при этом #3749 вскрыл дефект «анализ зависимостей не видит физический диапазон загрузки» вплоть до уровня MemoryAccessDetector::LogicalPointerRegion. Если сложить это с той скрытой линией TileFoundry — анализ может доказать упорядочение разделяемой памяти в 52 из 83 тестовых примеров, а lowering по способу main приведёт к ошибке в 4 из 7 программ — один и тот же сигнал появляется трижды: главный риск TileLang сейчас не в недостатке функциональности, а в том, что «сгенерированный код выглядит правильным, но на деле неверен», и команда выбрала подход, превращающий корректность в опровержимое доказательство, а не в латание дыр по одной заплатке за раз.

Третья линия — тестирование и контракты начинают выделяться в отдельный слой. tileops_test в TileOPs отвязывает smoke-тесты от CUDA, compile-contract-gate разделяет контракт dispatch и генерацию данных workload — оба влиты в течение одного часа; автор при этом прямо написал, что «общая ссылка не может доказать корректность реального бэкенда». Такая привычка разделять «что можно протестировать» и «что доказано» говорит о выходе проекта из прототипной стадии больше, чем несколько лишних патчей производительности.

Наконец, стоит отметить контраст: на стороне операторов за день 7 влитых патчей производительности, 6 из них с таблицами измерений на H200, а на стороне новостей три канала за 24 часа — ноль совпадений. Внешняя заметность TileLang по-прежнему сосредоточена на событиях вроде «открытых компонентов Ascend от DeepSeek», повседневный прогресс почти исключительно происходит внутри репозиториев. Для тех, кто следит за проектом, это означает, что основой ежедневного отчёта могут быть только коммиты и PR, а не публикации.

Приложение. Материалы и пояснения к проверке

Таблица проверки источников

Источник Результат проверки
Организация tile-ai (29 репозиториев) В окне мониторинга пуши в 9 репозиториев
Основной репозиторий tilelang 5 вливаний в основную ветку; 8 новых issue; 5 автоматических сборок сайта документации
TileOPs 7 вливаний в основную ветку; 1 для сайта
TileFoundry 1 вливание в основную ветку; 1 новое изменение на рассмотрении
tilelang-ascend 4 вливания в основную ветку; 6 обновлений на рассмотрении
tilelang-hygon 1 вливание
tilelang-lsp 1 вливание (выравнивание отладочного сопоставления строк с CUDA)
TileOPs-nightly 1 автоматический пуш
tilelang-metax / tilelang-musa / tilelang-mlir-ascend / TileRT / tilescale / DeepStack В окне мониторинга пушей нет
Страницы релизов (tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-hygon) В окне мониторинга новых версий нет
FlashQLA / TileKernels В окне мониторинга пушей нет
Google News (шесть групп запросов на китайском и английском) / Hacker News / arXiv Существенных совпадений в окне мониторинга нет

Полный список источников