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

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

Индекс выпуска

  • Ключевое за сегодня: ThreadSync повторно использует частичные барьеры между доменами памяти, восстанавливая производительность после исправления корректности в #3643 (#3761); удалена ветка устаревшего парсера TVM на 3384 строки (#3736); подача входных данных для автотюнинга переведена на текущее устройство (#3546); план синхронизации TileFoundry официально доведён до CUDA lowering (#285); за 24 часа влито 8, обновлено на рассмотрении 48, добавлено 7 issue
    1. Ход работ по основным проектам
      • 1.1 Основная линия: ThreadSync разрешает повторное использование частичных барьеров между доменами памяти, занятость барьеров на B300 возвращается с 12 до 16 блоков/SM (#3761) (11-04)
      • 1.2 Основная линия: удалены устаревшие ветки TVMScript AST и парсера, одним разом вырезано 3384 строки (#3736) (11-03)
      • 1.3 Основная линия: подача входных данных для автотюнинга переведена на текущее устройство, тюнинг на Ascend и MPS снова работает (#3546) (11-03)
      • 1.4 На рассмотрении: скалярный cast в условиях больше не сжимает ширину вектора данных, продолжение вчерашнего #3734 (#3759) (11-03)
      • 1.5 На рассмотрении: две пользовательские ошибки заменены на явные сообщения об ошибках (#3765 / #3767) (11-03)
      • 1.6 На рассмотрении: откат, когда общий tile с padding не может быть скопирован целиком; байты TMA для 2-CTA учитываются по парному лидеру (#3738 / #3766) (11-03)
      • 1.7 На рассмотрении: адаптация sm_103a для B300 в двух ветках, выбор int8 в tcgen05 и архитектурный гейт TMA в CuTeDSL (#3770 / #3772) (11-03)
    1. Адаптация под множество бэкендов
      • 2.1 Ascend: влито пакетное копирование ND2NZ из 3D GM в L1, один дескриптор покрывает весь пакет (#3689) (11-03)
      • 2.2 Ascend на рассмотрении: весь репозиторий TileLang-Ascend перенесён в основную многобэкендовую архитектуру, 2887 файлов (#1975) (11-03)
      • 2.3 Ascend на рассмотрении: аннотации синхронизации copy_pipe_to_l1 разделены по движкам исполнения, исправлен неверный результат, воспроизводившийся 10/10 на A2/A3 (#1981) (11-03)
      • 2.4 Ascend на рассмотрении: широковещание по первой оси Parallel через скалярное чтение, перенос прямого прохода Chunk O, пример MHA (#1979 / #1866 / #1879) (11-03)
      • 2.5 mlir-ascend: рабочий процесс Codex и передача задач влиты в ветку harness (#208) (11-03)
      • 2.6 Hygon: обновлён PR Distributed на рассмотрении, новых коммитов в основной ветке нет (#19) (11-03)
      • 2.7 MetaX и Moore Threads: в окне мониторинга пушей не было (11-03)
    1. Экосистема и стороны внедрения
      • 3.1 TileFoundry: план синхронизации реализован в CUDA lowering с обратной проверкой, влит в основную линию (#285) (11-04)
      • 3.2 TileFoundry на рассмотрении: план синхронизации wgmma, семантика обучения FSDP и HSDP в вычислителе HIR (#286 / #284) (11-03)
      • 3.3 TileOPs: объединение дисперсий считается с первой частичной суммы, дисперсия константных входов больше не даёт NaN (#2821) (11-03)
      • 3.4 TileOPs на рассмотрении: GEMM с опциональным bias, поддержка непрерывной раскладки в FP8 групповом GEMM (#2823 / #2824) (11-03)
      • 3.5 TileOPs на рассмотрении: селекторы W4A16 и плотного GEMM переведены на единую модель стоимости с повторной подгонкой (#2825 / #2826) (11-03)
      • 3.6 TileOPs на рассмотрении: упаковка коротких varlen по KV-группам, L2 flush переведён в режим только чтения (#2819 / #2822) (11-03)
      • 3.7 Стороны внедрения: у FlashQLA и TileKernels пушей нет, но #3761 исправляет fused-ядро TileKernels (11-04)
    1. Сообщество, руководства и мероприятия
      • 4.1 7 новых issue за сутки: одна проблема production-уровня с CUDA graph, четыре от одного и того же внешнего репортёра (11-03)
      • 4.2 Семантика насыщения vcvt на Ascend 950 и ежедневный тест mlir-ascend (#3762 / #1980) (11-03)
      • 4.3 Три обновления сайта документации, расширение VS Code переименовано в vsix (11-03)
      • 4.4 Новостная сторона: по трём каналам попаданий в окне нет, сообщения зарубежных СМИ об обучении на Ascend расходятся (11-03)
    1. Наблюдения за трендами
  • Приложение. Материалы и пояснения к проверке

Ключевое за сегодня

Основная линия этого дня — исправление, которое «возвращает производительность». Слияние #3643 из предыдущего окна решило реальную проблему корректности: раздельные вызовы ThreadSync("shared") и ThreadSync("shared.dyn") могли назначать один и тот же именованный ID барьера непересекающимся группам потоков. Исправление состоит в том, чтобы последующий вызов сохранял уже использованный предыдущим вызовом ID, но это правило сохранения одновременно блокировало другое — одна и та же группа потоков могла бы повторно использовать один и тот же ID между двумя областями видимости. Цена легла напрямую на ядро mhc_pre_big_fuse из TileKernels: это ядро имеет 96 потоков на блок, потоки с 32 по 95 выполняют двухступенчатый конвейер cp.async, синхронизация после асинхронного ожидания в первом проходе получает ID 3, зависимая от буфера синхронизация тех же 64 потоков во втором проходе получает ID 4; лишний ID заставляет объявленные в скомпилированном артефакте ресурсы барьеров вырасти с 4 до 5; при запуске на B300 со скрытой размерностью 1280 предел занятости барьеров из-за этого падает с 16 блоков/SM до 12 блоков/SM, при этом число регистров, разделяемой памяти и инструкций синхронизации не меняется. Решение в #3761 — заставить ThreadSync принимать упорядоченный список областей видимости, конвейер CUDA переписан на ThreadSync(["shared", "shared.dyn"]), синхронизация по-прежнему планируется и вставляется в исходном порядке областей видимости, но ID части барьеров назначаются единообразно только после обработки обеих областей; перед совместным использованием одного ID между областями видимости новый анализатор доказывает, что предикат участия выполняется на всём диапазоне кандидатных потоков, контексты ветвлений и Bind предоставляются через ConstrVisitor, предикаты, связанные с потоками, трактуются как обязательства доказательства, а не считаются истинными по умолчанию; для доказанно плотных групп напрямую используется объём ограничивающего параллелепипеда с пропуском перечисления частичного числа потоков, для недоказанных групп область видимости остаётся в ключе и счётчик участников также входит в ключ, явно пронумерованные барьеры (включая ID, назначенные более ранним вызовом ThreadSync) продолжают сохраняться. Изменение затрагивает 401 строку и 4 файла.

В описании автор провёл очень чистый эксперимент по атрибуции: изменение только объявленного в скомпилированном артефакте числа ресурсов барьеров без изменения каких-либо исполняемых инструкций воспроизводит то же замедление. Такие эксперименты стоит отмечать — они превращают «регрессию производительности» из ощущения в указание на единственную переменную и делают критерий приёмки исправления явным: не общее ускорение бенчмарка, а возврат объявления ресурсов к 4 и занятости к 16 блокам/SM.

Второй пункт — сокращение со стороны предложения и унификация меры. #3736 удаляет tilelang/language/ast/, tilelang/language/parser/ и tilelang/language/tir/entry.py — это модули, оставшиеся от раннего форка локального парсера TVMScript, около 3380 строк Python, вызовов в репозитории больше нет: языковой слой теперь использует интерфейс вышестоящего tvm.tirx.script.parser через tilelang.language.tir.common, регистрация переопределений AST и операторов находится в tilelang.language.overrides, а T.prim_func и T.macro экспортируются из tilelang.language.eager. Фактическое изменение — добавлена 1 строка, удалено 3384 строки, 10 файлов, при этом удалено правило игнорирования *language/ast* в docs/conf.py; публичные экспорты языкового слоя покрываются тестами диалектов и тестами фронтенда. #3546 исправляет меру устройства в автотюнинге: замыкание поставки, устанавливаемое set_autotune_inputs, определяло эталонное устройство из параметра, который ни один вызов никогда не передаёт; фактически передавался список параметров ядра, поэтому резервная ветка CUDA или CPU на машине без CUDA переносила каждый захваченный тензор на CPU, и каждый раунд тюнинга завершался ошибкой несовпадения устройства лаунчера; теперь замыкание использует get_current_device() с той же стратегией, что и поставка тензоров по умолчанию: поведение CUDA не меняется, Ascend и MPS переходят на поставку на текущем устройстве, остальные бэкенды по-прежнему откатываются на CPU. Приёмка на двухкарточной Ascend950: три регрессии автотюнинга Ascend до исправления полностью падали, после — полностью проходят; в основном репозитории testing/python/autotune/ 12 пунктов проходят, 28 пропущены, ещё 3 падают из-за отсутствия CUDA на этой машине, что автор явно оставил на проверку CUDA CI.

Третий пункт на стороне Ascend: копирование из GM в L1 расширено с «один дескриптор — одна матрица» до целого пакета. #3689 позволяет записывать область [B, M, N] одним T.copy, lowering генерирует пакетные дескрипторы ND2NZ, сохраняя пакет и построчный stride на стороне источника, межпакетное расстояние на стороне назначения выводится из физической раскладки, транспонированные пакеты идут через DN2NZ; исходные пакеты, выходящие за границу GM, просто пропускаются, а целевая память остаётся нетронутой, дополнение нулями хвостов M/N применяется только к допустимым пакетам. Существующее поведение для двумерных и одиночных матриц не меняется, реализация ориентирована только на бэкенд AscendC и требует, чтобы целевая сторона L1 имела каноническую матричную раскладку; PTO в этот раз вне области. Проверка на Ascend950DT и CANN 9.2: проходят 50 примерных тестов AscendC и 1733 юнит-теста Ascend, среди них 28 устройственных случаев покрывают пакеты и хвостовые блоки M/N, выход всего пакета за границы, ненепрерывный stride источника, больший объём целевой памяти и прямоугольное транспонирование; случаи CUDA и PTO исключены, CANN 9.3 из образа CI локально не тестировался, и автор написал об этом в описании. Изменение затрагивает 579 строк и 10 файлов.

Четвёртый пункт — это скрытая реализация TileFoundry. В прошлом отчёте он ещё висел в статусе «на рассмотрении», в этом окне мониторинга он был слит: #285 заставляет коммиты, ожидания и барьеры, запланированные анализом sync, действительно генерироваться CUDA lowering, с обратной проверкой каждого пункта на соответствие фактической генерации плану, а при несоответствии выбрасывается LoweringError. Результат на корпусе: в стандартном процессе 14 программ плюс синхронизация, 31 отклонена; у этих 14 TIR отличается от основной ветки только добавленными операторами, ещё 30 полностью совпадают с основной веткой; на одном sm_89 прогнано 7 собираемых программ с синхронизацией, 6 полностью совпадают с эталоном, synccheck и memcheck полностью чистые, тогда как те же 7 программ, lowering которых выполнен способом основной ветки, дают 4 ошибочных результата. Прямой симптом, приведённый автором в мотивации, — максимальное отклонение E4M3 GEMM на sm_89 составляет 659. Этот коммит — 788 добавленных строк, 214 удалённых, 25 файлов — продвигает «анализ может доказать» к «генерация должна соответствовать доказательству».

1. Прогресс основных проектов

1.1 Основная ветка: ThreadSync разрешает повторное использование части барьеров между доменами памяти, занятость барьеров на B300 возвращается с 12 к 16 блокам/SM (#3761)

Слито 11-04 00:50, добавлено 401 строка, удалено 72 строки, 4 файла. Предыстория — #3643 из прошлого окна мониторинга: при раздельных вызовах ThreadSync("shared") и ThreadSync("shared.dyn") непересекающимся группам потоков мог быть назначен один и тот же именованный ID барьера, это реальная проблема корректности, исправление — сохранять ID, уже использованные более ранними вызовами. Побочный эффект — одна и та же группа потоков не может повторно использовать барьер между областями видимости; на ядре mhc_pre_big_fuse из TileKernels это выливалось в лишнее выделение на каждый блок из 96 потоков: двухэтапный конвейер cp.async для потоков с 32 по 95 получал ID 3 на первом проходе и ID 4 на втором, объявленные в скомпилированном продукте ресурсы барьеров росли с 4 до 5, предел занятости барьеров для скрытого измерения 1280 на B300 падал с 16 блоков/SM до 12 блоков/SM, тогда как регистры, разделяемая память и число инструкций синхронизации не менялись вовсе. Этот коммит заставляет ThreadSync принимать упорядоченный список областей видимости, конвейер CUDA переведён на ThreadSync(["shared", "shared.dyn"]): синхронизация по-прежнему планируется и вставляется в исходном порядке, частичные ID барьеров выделяются единообразно после обработки обеих областей видимости; перед совместным использованием одного ID между областями видимости новый анализатор доказывает, что предикат участия тождественно выполняется в диапазоне кандидатных потоков, контексты ветвлений и Bind предоставляет ConstrVisitor, предикаты, зависящие от потоков, рассматриваются как обязательства доказательства, а не как допущения, уже доказанные плотные группы вычисляются по объёму ограничивающего параллелепипеда с пропуском частичного перебора числа потоков, для недоказанных групп область видимости остаётся в ключе и число участников входит в ключ, явно пронумерованные барьеры (включая ID, выделенные более ранними вызовами) продолжают сохраняться, непересекающиеся группы потоков по-прежнему изолированы — исправление #3643 не откатано. Старые вызовы в строковой форме остаются совместимыми, CuTeDSL также проходит через тот же конвейер pass-ов CUDA.

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

1.2 Основная ветка: удаление унаследованного TVMScript AST и ветки парсера, разом удалено 3384 строки (#3736)

Слито 11-03 13:05, добавлена 1 строка, удалено 3384 строки, 10 файлов. Удалены tilelang/language/ast/, tilelang/language/parser/ и tilelang/language/tir/entry.py — остатки раннего локального форка парсера TVMScript, у которых в репозитории нет ни одного вызывающего: языковой слой через tilelang.language.tir.common использует интерфейс вышестоящего tvm.tirx.script.parser, собственные AST и переопределения операторов TileLang регистрируются в tilelang.language.overrides, T.prim_func и T.macro экспортируются из tilelang.language.eager, а в tilelang/language/tir/ потребляются только четыре модуля: common, exports, ir, op. Риск сохранения этого изолированного форка парсера описан вполне конкретно: если кто-то напрямую импортирует tilelang.language.parser.operation, он зарегистрирует конфликтующие операторы на tirx.PrimExpr и tirx.IterVar. Этот коммит также удаляет бесполезный повторный экспорт prim_func в tilelang/language/tir/__init__.py и правило игнорирования для *language/ast* в docs/conf.py (*parser* сохраняется, так как tilelang/language/overrides/parser.py всё ещё используется). Критерий приёмки — публичные экспорты не меняются, что покрывается test_tilelang_language_dialect.py и фронтенд-тестами языка.

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

1.3 Основная ветка: подача входных данных для автоматической настройки переведена на текущее устройство, настройка на Ascend и MPS снова доступна (#3546)

Объединено 11-03 13:19, добавлено 153 строки, удалено 12 строк, 3 файла, исправляет #3545. Замыкание замороженной поставки, устанавливаемое set_autotune_inputs, определяло базовое устройство из некоторого параметра, который «никогда не передаётся из точки вызова» — по контракту процедуры поставки фактически передаётся список параметров ядра — поэтому на машине без CUDA ветка запасного варианта CUDA или CPU переносила каждый захваченный тензор на CPU, и каждый раунд настройки завершался неудачей из-за несовпадения устройства с лаунчером. Теперь замыкание использует get_current_device(), придерживаясь той же стратегии, что и поставка тензоров по умолчанию: поведение на CUDA не меняется, многокарточные benchmark-воркеры по-прежнему копируют по картам; на Ascend и MPS захваченные тензоры поставляются на соответствующих текущих устройствах, благодаря чему настройка становится работоспособной; остальные бэкенды по-прежнему откатываются на CPU. Новый тест-кейс Ascend размещён в testing/ascend/autotune/, охватывает путь декоратора, устройство поставки базы AutoTuner и три случая захвата входов между устройствами; на двухкарточной Ascend950 до исправления все терпели неудачу, после исправления все три проходят; в основном репозитории testing/python/autotune/ 12 проходят, 28 пропущены, ещё 3 (режим eager, скалярный вход, с входом) терпят неудачу из-за отсутствия CUDA на этой машине, автор отмечает, что они оставлены на проверку CUDA runner в CI.

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

1.4 На рассмотрении: скалярный cast в условии больше не сжимает ширину вектора данных, продолжение вчерашнего #3734 (#3759)

Обновлено 11-03 14:09, добавлено 148 строк, удалено 29 строк, 2 файла. Это завершение вчерашнего влитого #3734: #3734 после DecoupleTypeCast корректно ужесточил ограничение ширины для cast данных, но также заставил «cast в скалярном условии» начать ограничивать ширину вектора данных — на b5cb02633 копирование BF16 под охраной могло из-за наличия int64 cast в условии сжаться с 16 lane до 4, даже если это условие полностью согласовано на 16 lane. Исправление состоит в том, чтобы повторно использовать уже имеющийся у планировщика анализ согласованности условий для освобождения от ограничения ширины cast, выделив инструкции if, tirx.if_then_else и условия Select скалярных предикатов из их подвыражений для отдельного обхода; cast в условии больше не накладывает ограничение на ширину типа данных, cast в подвыражениях по-прежнему накладывает, и на целях с поlane-векторными предикатами сохраняется ограничение на условный cast; поведение защитных проверок памяти #3579 и исправленное в #3734 ограничение охраняемого преобразования FP32 в BF16 остаются неизменными. Регрессия охватывает согласованные охраны на 8 lane и 16 lane, три формы условий, cast FP32 внутри подвыражения, векторные предикаты и группировку битов округления; девять тест-кейсов с охраной int64 до исправления терпели неудачу, в итоге два набора — финальная легализация и округление cast — в сумме 74 проходят.

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

1.5 На рассмотрении: две пользовательские ошибки превращены в явные сообщения об ошибке (#3765 / #3767)

Оба обновлены около 11-03 17:10, представлены одним и тем же репортером. #3765 соответствует #3653: частичный барьер, охраняемый границей индексации одного потока, когда эта граница кратна размеру warp, потоки всего warp единодушны в оценке условия; но общее арифметическое доказательство упускало thread_binding < 448, когда внешняя ветка использует побитовое &, и ThreadSync просто прерывался. Этот PR явно принимает такой случай, когда измерение потоков кратно целому числу warp (одновременно охватывая случаи, когда граница равна этому измерению или превышает его), а остальные условия по-прежнему идут по существующему пути доказательства. #3767 соответствует #3005: распределение fragment между потоками, когда внутри одного тела T.Parallel один и тот же fragment читается и записывается по двум структурно разным индексам (например, buf[i] и buf[(i + 1) % N] — такой межlane-шаблон), невозможно понизить без межпоточного обмена; ранее ParallelOpNode::RecordBufferAccess наталкивался на ICHECK StructuralEqual, что проявлялось как внутреннее утверждение компилятора; теперь это заменено на ValueError, оба индексных выражения записываются в сообщение об ошибке с пояснением, что fragment нельзя так обращать. Оба PR относятся к категории «замена внутренних утверждений на диагностику, понятную пользователю», влияние на восприятие пользователей больше, чем масштаб изменений.

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

1.6 На рассмотрении: откат при невозможности целостного копирования разделяемого tile с padding, учёт байтов TMA 2-CTA по парному leader (#3738 / #3766)

Оба коммита обновлены между 13:47 и 13:31 03.11. #3738 исправляет #3704: T.copy из глобальной области меньшего размера, чем общий tile после padding (например, K=100 копируется в BK=112), ошибочно классифицируется ClassifyWarpSpecializedCopy и SelectCopyInstForLowering как цельная TMA-копия, после чего lowering завершается жёсткой ошибкой s_extent == g_extent (112 vs. 100). Причина в том, что логика выбора проверяла только выравнивание последнего измерения, соответствие dtype и глобальный stride, но не проверяла, образует ли диапазон после padding одну цельную копию; данный коммит добавляет это условие и при его невыполнении откатывается к обычному пути копирования. #3766 исправляет подсчёт байтов TMA в сценарии cluster: TMA load с cta_group::2 очищает бит равноправия адреса барьера, байты пары CTA попадают на барьер чётного CTA, тогда как tma_copy с cluster-барьером ранее ожидал cluster_size × tile_bytes только для rank 0 — на (2,1,1) и (1,2,1) это случайно совпадало, а на (4,1,1) и (2,2,1) rank 0 получал только 2 порции вместо ожидаемых 4, а остальные байты попадали на барьер rank 2. Оба коммита не являются новыми функциями и относятся к классу дефектов «при нарушении граничного условия выдаётся внутренняя ошибка».

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

1.7 На рассмотрении: две линии адаптации sm_103a для B300, выбор int8 в tcgen05 и архитектурный гейт TMA в CuTeDSL (#3770 / #3772)

B300 сообщает целевую архитектуру sm_103a, и два PR на рассмотрении обрабатывают две связанные с ней блокирующие проблемы. #3770 соответствует #3769: tcgen05.mma.kind::i8 существует только в архитектурно-специфичных sm_100a и sm_110a, отсутствует как в sm_103a (B300), так и в семействе переносимых sm_10xf, тогда как AllowTcgen5Mma для TargetIsSm100 (arch 100 до 110) всегда принимает операнды int8/uint8, и ядро, успешно проходя lowering, на этапе ptxas выдаёт Feature '.kind::i8' not supported on .target 'sm_103a'; данный коммит заставляет TileLang не выбирать kind::i8 на этих целях. #3772 обрабатывает архитектурный гейт TMA в CuTeDSL: sm_103a для B300 отвергается ещё до входа в DSL с ошибкой, перечисляющей sm_90/90a/100a/110/120/120a/121/121a, но не содержащей 103a; в комбинации основной ветки fa69f7fa с nvidia-cutlass-dsl 4.7.0 и CUDA 13.3 на B300 падают шесть тестов CuTeDSL, включая несколько JIT-случаев tma_load. В той же партии #3768 исправляет пример обратного прохода flash attention для SM100: T.tcgen05_gemm теперь делает mbar обязательным keyword-only параметром и больше не ожидает неявно, а пример по-прежнему вызывается в стиле синхронного GEMM, и оба скрипта выдают missing 1 required keyword-only argument ещё до компиляции; автор также отмечает, что в каталоге examples/flash_attention_sm100/ нет test_*.py, что и стало причиной, по которой проблема так долго не обнаруживалась.

Источники: https://github.com/tile-ai/tilelang/pull/3770, https://github.com/tile-ai/tilelang/pull/3772, https://github.com/tile-ai/tilelang/pull/3768

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

2.1 Ascend: принято пакетное копирование ND2NZ из 3D GM в L1, один дескриптор покрывает весь пакет (#3689)

Объединено 11-03 17:07, добавлено 579 строк, удалено 39 строк, 10 файлов. Ранее копирование GM в L1 на Ascend обрабатывало один матричный DMA-дескриптор за раз, и для копирования пакета приходилось писать for b in T.serial(B) с вызовом для каждого пакета; этот коммит позволяет записать область [B, M, N] одним T.copy, при lowering генерируется пакетный дескриптор ND2NZ, сохраняются пакетные и построчные stride источника, межпакетное расстояние на стороне назначения выводится из физической раскладки, транспонированный пакет идёт через DN2NZ, и переиспользуется существующая обработка границ матричных осей: пакеты источника, выходящие за границу GM, пропускаются, хранилище назначения остаётся нетронутым, а дополнение нулями хвостов M/N применяется только к допустимым пакетам. Поведение существующих двумерных и одноматричных случаев не меняется; реализация затрагивает только бэкенд AscendC и требует, чтобы сторона назначения L1 имела каноническую матричную раскладку, поддержка PTO в это изменение не входит. Среда проверки — Ascend950DT плюс CANN 9.2: пройдены 50 примерных тестов AscendC и 1733 юнит-теста Ascend, из них 28 устройственных кейсов BF16/FP32 покрывают пакетные и хвостовые блоки M/N, выход всего пакета за границы, несмежный stride источника, хранилище назначения больше источника, прямоугольное транспонированное копирование. Автор отмечает, что кейсы CUDA и PTO исключены, а CANN 9.3, используемый в образе CI, локально не проверялся. Это изменение идёт в том же направлении, что и влитое в прошлом выпуске ядро FlashMLA для Ascend 950: и для внимания, и для GEMM пакетная пересылка требует, чтобы один дескриптор покрывал весь пакет, а вызовы по одному пакету на производственных формах дают заметные накладные расходы на запуск.

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

2.2 Ascend на рассмотрении: перенос всего репозитория TileLang-Ascend в многобэкендовую архитектуру основной ветки, 2887 файлов (#1975)

Обновлено 11-03 16:34, масштаб изменений достигает 607317 добавленных строк, 97595 удалённых строк, 2887 файлов — это самый крупный по объёму коммит выпуска и самый достойный отслеживания PR на рассмотрении во всей организации. Цель — согласно предложению #1911 перенести TileLang-Ascend на многобэкендовую архитектуру основной ветки: существующие компоненты Ascend перемещаются в соответствующие каталоги основной ветки под ascend/a2a3 (tilelang/ соответствует tilelang/ascend/a2a3, src/ соответствует src/ascend/a2a3), исходные бэкенды Ascend C и PTO A2A3 независимо регистрируются в основной ветке и ортогональны Ascend и PTO A5 основной ветки, pass и codegen не мешают друг другу и переиспользуют лишь небольшую часть инфраструктуры; обе стороны используют общие target="ascend" и "pto", но бэкенд A2A3 несёт метку arch="dav-2201" (автоматически определяется по имени NPU) и по ней разделяется; примитивы фронтенда A2A3 переходят на import tilelang.ascend.a2a3.language as T и сосуществуют с import tilelang.ascend.language as T для A5, способ написания операторов относительно исходного репозитория не меняется; pass и codegen напрямую используют версии, перенесённые из TileLang-Ascend, чтобы сохранить исходный процесс компиляции, исполнительный бэкенд переиспользует cython основной ветки с добавлением специфичных для Ascend адаптаций вроде индекса workspace, examples и testing размещаются в examples/ascend_a2a3 и testing/ascend_a2a3 соответственно, CI пока использует ci_cd.yml исходного репозитория, и подтверждено, что после адаптации к новым путям он по-прежнему работает. В описании автор предельно честно перечисляет незавершённые задачи: различия файлов в трёх каталогах с исходными файлами ещё предстоит свести, статические проверки и CI форматирования ещё предстоит восстановить, документацию и связанное с Agent ещё предстоит перенести, и отмечает, что «затронуто много файлов, при пакетных операциях Agent неизбежны упущения», прося помощи в проверке. Такой подход — вынести неопределённости миграции прямо в описание PR — гораздо полезнее для рецензирования, чем ждать заключения «готово».

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

2.3 Ascend на рассмотрении: аннотации синхронизации copy_pipe_to_l1 разделены по исполнительным движкам, исправлен неверный результат, воспроизводимый 10/10 на A2/A3 (#1981)

Обновлено 11-03 в 12:36, соответствует #1780. Корневая причина описана чётко: междъядерная коммуникация V→C через PTO в A2/A3 идёт транзитом через GM, три сегмента пересылки относятся к трём движкам — MTE3 в AIV выполняет отправку UB→GM, MTE2 в AIC выполняет приём GM→L1 (TLOAD внутри TPOP), MTE1 в AIC выполняет L1→L0A; при этом copy_pipe_to_l1 в operation_config.h помечен как PIPE_MTE3, поэтому проход автоматической синхронизации генерирует между ним и copy_l1_to_l0a set_flag(PIPE_MTE3, PIPE_MTE1, …) плюс соответствующий wait_flag — но ждёт он MTE3 собственного ядра AIC, а очередь MTE3 в AIC пуста, событие устанавливается немедленно, и чтение L1 движком MTE1 может быть разрешено до завершения записи MTE2; L0A получает старые данные, результат MMA ошибочен. Ошибочная аннотация была внесена вместе с #1238; в той же таблице copy_pipe_to_ub и copy_l0c_to_pipe аннотированы по движку исполнительной стороны, и только эта запись сохранила точку зрения отправляющей стороны. Два доказательства воспроизведения: автор issue заменил PIPE_MTE3 в сгенерированном исходном коде на PIPE_MTE2, после чего 10/10 восстановились до точного результата; локально на последней основной ветке 10/10 запусков ошибочны, максимум за один запуск 16384/16384 несовпадающих элементов. Исправление по предложению из issue — разделить имена, а не просто заменить: аннотацию copy_pipe_to_l1 изменить на PIPE_MTE2 (загрузка GM→L1 для A2/A3), добавить copy_pipe_to_l1_bind с аннотацией PIPE_MTE1 (на прямом пути A5 ожидание готовности L1 и привязка адреса), следуя существующему прецеденту copy_pipe_to_ub и copy_pipe_to_ub_V, чтобы не починить A2/A3 и не сломать A5. Проблема того же происхождения охватывает и направление WAR: пока MTE1 ещё читает, TLOAD движка MTE2 может перезаписать тот же буфер L1.

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

2.4 Ascend на рассмотрении: широковещание по первой оси Parallel через скалярное чтение, миграция прямого прохода Chunk O, пример MHA (#1979 / #1866 / #1879)

Все три обновлены утром 11-03 и относятся к вехам операторов Ascend. #1979 соответствует #1885: семейство сбоев сводится к одному шаблону — одномерный операнд UB, индексируемый первой осью T.Parallel (или внешней serial-переменной), оказывается слева от бинарного выражения; шаблон сокращения из issue dk_i = sum_j v[j] * A[j, i] — ровно эта форма, все 5 FAIL в матрице бисекции из 21 точки принадлежат этому семейству, независимо от направления редукции и способа записи широковещания; корневая причина в том, что AscendLowerParallelToVector при понижении двумерного параллельного вложения в строковый цикл на входе HandleSimpleCase проверяет скалярный доступ только для правого операнда, левый идёт по векторному пути и попадает в ловушку выравнивания UB; в этом коммите широковещательный операнд первой оси также переводится на скалярное чтение. #1866 переносит прямой проход Chunk O (прямой проход O для основанного на chunk гейтированного линейного внимания, GDN) на NPU Ascend, используя режим разработчика плюс четыре конфигурации pass (автоматическое объединение C/V, автоматическая синхронизация C/V, автоматическая синхронизация, планирование памяти), одно ядро @tilelang.jit плюс цикл grid-stride с фиксированными 20 ядрами, обеспечивая внутричиповое слияние, нулевой GM workspace, нулевые хостовые операторы ATen. #1879 добавляет новый пример прямого прохода MHA: раскладка BSHD, трёхдиспетчерный конвейер MMA в режиме Expert, разделённый на 5 Python-файлов по wrapper, общим слоям и трём формам модулей; wrapper не содержит никакого кода T.*, выполняет только проверку метаданных и JIT-диспетчеризацию; is_causal использует маску с выравниванием по правому нижнему углу, полностью замаскированные строки через защиту sumexp + 1e-30 дают выход 0; все 20 тестов точности cann-bench пройдены, среднее ускорение относительно базовой линии torch-npu — 0,97x, совокупный балл 73,13 (по локальным замерам на 910B2). 0,97x означает, что чистого выигрыша пока нет, но как пример он даёт скелет того, «как написать полный оператор внимания на Ascend».

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

2.5 mlir-ascend: рабочий процесс Codex и передача задач слиты в ветку harness (#208)

Объединено 11-03 17:50, добавлено 2715 строк, удалено 39 строк, 66 файлов, целевая ветка слияния — harness, а не main (последний коммит основной линии по-прежнему остановился на 10-31 ci: remove internal cache-service references), при цитировании не следует писать «слито в основную линию». Содержание — добавление первого набора рабочих процессов Codex для harness операторов: инициализация проекта с помощью итеративных вопросов, подключение существующего репозитория операторов, а также восстановление разработки, оптимизации или анализа из персистентного контекста; три набора навыков совместно используют редактируемые записи Markdown в .workspace/, при передаче задач различаются «текущий кандидат» и «проверенный лучший кандидат», сохраняются неудачные и откатанные эксперименты, чтобы суб-агенты получали сфокусированный контекст, а главный агент поддерживал общее состояние; локальный установщик отвечает за регистрацию навыков и добавление записей проекта в AGENTS.md (сохранение существующего содержимого, игнорирование .workspace/, разрешение Git-проекта до корня worktree). README, зависимости, описание повторного использования и дорожная карта описывают установку с приоритетом TileLang и областью охвата Ascend; рабочий процесс ссылается на GSD и сохраняет указание исходной ревизии и декларацию MIT, адаптер исполнения для NPU и распространение через npm всё ещё в списке задач. Этот коммит и связанные с Agent каталоги в основном репозитории и репозитории Ascend — три стороны одного и того же: процесс «написания операторов интеллектуальным агентом» закрепляется как переиспользуемый актив внутри репозитория.

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

2.6 Hygon: обновлён PR Distributed на рассмотрении, в основной ветке новых коммитов нет (#19)

В tilelang-hygon в окне мониторинга была запись о пуше (11-03 11:06 пекинское время), но в основной ветке новых коммитов нет, последний по-прежнему #32 от 11-02 (обработка gfx948 как совместимого с gfx946, чтобы сначала можно было тестировать существующие инструкции). Пуш пришёл из #19 на рассмотрении (заголовок только Distributed, описание пустое), обновлён 11-03 11:06; судя по времени и заголовку, идёт продвижение адаптации, связанной с распределённостью, но описание PR не содержит содержимого, и настоящий отчёт не делает выводов о деталях реализации. Обычный ритм этого репозитория в последнее время — одна правка со стороны HCU (вычислительного блока Hygon) каждые два-три дня: #31 от 10-30 изменил границы MLS и персистентную синхронизацию WASP, #30 дополнил LDS swizzle для обычного копирования, #29 от 10-27 выполнил разделение банков для буфера scale.

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

2.7 MetaX и Moore Threads: в окне мониторинга пушей нет (11-03)

Последний пуш в tilelang-musa (Moore Threads) остановился на 09-30 09:50, в tilelang-metax (MetaX) — на 10-26 07:08; в обоих репозиториях в этом окне не было коммитов и обновлений PR на рассмотрении. Таким образом, динамика бэкендов этого выпуска сосредоточена на сторонах Ascend и Hygon: из 8 слитых коммитов 1 относится к Ascend (#3689), 1 — к mlir-ascend (#208, слит в ветку harness); из 48 обновлений на рассмотрении репозиторий Ascend занимает 15, основной репозиторий — 23.

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

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

3.1 TileFoundry: план синхронизации реализован в CUDA lowering с обратной проверкой, слито в основную линию (#285)

Слито 11-04 00:19, добавлено 788 строк, удалено 214 строк, 25 файлов. Анализ sync планирует коммиты, ожидания и барьеры, необходимые для доступа к разделяемой памяти, но finalize ранее не выпускал ни одного из них, поэтому расписания, буферизованные через разделяемую память, имели гонки, что на карте проявлялось как максимальное отклонение E4M3 GEMM на sm_89 от torch до 659. Этот коммит привязывает шаги анализа к узлам HIR, которым они предшествуют, индексируя по меткам через SyncMetadata: метка собственного оператора узла пуста, копии, добавленные InsertSlice, используют base и update, копии, переносимые циклом, используют ("entry", слот) и ("backedge", слот); ядро lowering перед lowering операторов некоторого узла вызывает новый хук before(owner, ctx, tag), и CudaLowering выпускает там шаги этого узла; schedule finalize отвечает за запуск анализа sync. Контракт записан в четырёх нормативных документах, где спецификация анализа §1.2.5 описывает SyncMetadata, метки, правила и случаи, которые анализ по-прежнему отвергает. Автор также отметил рискованную область: все программы в корпусе, которые могли finalize на основной ветке, по-прежнему могут finalize, лишь добавляются операторы, исключение — hir/staged_dynamic_inner: его lane читают и пишут один и тот же разделяемый tile в цикле, где «число итераций равно lane», и нет барьера, способного их свести; кроме того, сохраняемые вызовы, затрагивающие разделяемую память, трактуются как «одно обращение к неизвестным байтам», и два таких вызова в области без сетки потоков требуют барьера, который TIR не может разместить, поэтому анализ отвергает их, даже если байты не пересекаются; CopyAsyncTensor, CopyReduceAsyncTensor и wgmma пока не имеют эмиттеров.

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

3.2 TileFoundry на рассмотрении: план синхронизации wgmma, семантика обучения FSDP и HSDP в вычислителе HIR (#286 / #284)

#286 обновлён 11-04 02:28, добавлено 1907 строк, удалено 310 строк, 52 файла, это прямое продолжение #285: WGMMA после выпуска может продолжать читать разделяемую память и регистры, и план синхронизации CUDA должен ставить последующие обращения после этих эффектов; этот коммит объявляет семантику завершения для каждого вида эффекта операнда, привязывает барьеры, коммиты и ожидания WGMMA к узлам HIR, которым они нужны, lowering потребляет эти шаги в точке выпуска инструкции, а если последующая область не вмещает коллективное ожидание, операция завершается в области выпуска; одновременно исправляется псевдоним Reshard регистров, чтобы элемент оставался в одном и том же регистре на каждом потоке, держащем копию, и план отслеживает зависимости регистров по этой идентичности хранения. Автор явно отметил этот риск: эмиттер матричных инструкций всё ещё недоступен, этот коммит не устанавливает корректность исполнения или производительность WGMMA, семантика завершения TMA по-прежнему обрабатывается отдельно. #284 обновлён 11-03 14:34, посвящён семантике обучения: прежний вычислитель трактовал reshard как глобальный тензор и упускал вклады разных участников; этот коммит вычисляет локальные значения участников, координаты GPU, вызовы и циклы, gather/scatter, а также разрешение Partial(sum) по осям, и делает доказательство псевдонимов Reshard с семантикой владения, расширяя структурированное жизненное время из #260 и #170; в комплекте добавлено исследование двухблочного GQA, вдохновлённое modded-nanogpt, с явным прямым и обратным проходом, momentum SGD и независимой ссылкой PyTorch, покрывающее FSDP и две стратегии размещения HSDP 2×2 и 4×4, с отчётом о группах коллективной коммуникации, полезной нагрузке и хранении. Критерии приёмки — 775 пунктов основного набора после рефакторинга, 778 пунктов набора бэкенда (13 пропущено на двух картах sm90), десять конфигураций HIR проходят на CPU, обе стратегии на четырёх картах sm90 с реальным PyTorch HSDP 2×2 выполняют два шага обновления, локальные градиенты, параметры и momentum полностью совпадают.

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

3.3 TileOPs: объединение дисперсий начинается с первой частичной величины, дисперсия константного входа больше не даёт NaN (#2821)

Слито 11-03 18:04, добавлено 43 строки, удалено 4 строки, 3 файла. Дефект состоял в том, что объединение дисперсий начиналось с пустого аккумулятора: при объединении конечного большого среднего в пустой аккумулятор квадрат разности переполнялся ещё до умножения на нулевой вес, поэтому дисперсия константного входа становилась NaN; затрагивались вызовы var, std и var_mean, когда до и после сохраняемой оси имеется по одному участку редукции. Исправление состоит в том, чтобы инициализировать объединение с первой реальной частичной величины, а последующие частичные величины продолжают использовать существующую рекурсию Chan; построчная редукция, диспетчеризация, выходной layout и публичный API не меняются. Приёмка: три новых регрессионных теста падают на исходном ядре и проходят после исправления, кандидат прошёл 155 регрессионных тестов и 24 дополнительных численных проверки, tests/ops/test_reduce.py вырос с 155 до 158; три теста автотюнинга исключены, поскольку предыдущая попытка базовой линии превысила 60-секундный лимит теста. Подобные численные ловушки «пустой аккумулятор + большое смещение» в семействе редукций нередки, и стоит запомнить, что они проявляются только на формах с одновременной редукцией до и после сохраняемой оси, обычная редукция по всей оси их не обнаруживает.

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

3.4 TileOPs на рассмотрении: GEMM с опциональным bias, поддержка непрерывной раскладки для FP8 сгруппированного GEMM (#2823 / #2824)

#2823 добавляет опциональный bias[N] в GEMMFwdOp.forward: bias прибавляется к результату накопления в float32, а выход округляется только один раз, что согласуется с семантикой torch.nn.functional.linear; каждая плотная структура добавляет bias в своём эпилоге — warp-специализированный (TMA и построчное хранение), coop2 (включая ширины 176/192 и узкие tile), coop2_swap, ping-pong, coop2s, simple, swap_ab, GEMV с ядрами cp.async, а путь split-K добавляет его только один раз в редукции. #2824 позволяет MoEGroupedGEMMFP8FwdOp помимо маскированных блоков принимать также непрерывные строки экспертов (компактная упаковка или выровненная упаковка, со смещением конца сегмента physical_psum или построчным id эксперта), семантика согласуется с deep_gemm.m_grouped_fp8_gemm_nt_contiguous; ядро SM90 обходит рабочие сегменты каждого эксперта по концу сегмента, компактные построчные id преобразуются в концы сегментов двоичным поиском до разделения на producer и consumer, выровненные построчные id обходятся поблочно по строчным tile с чтением эксперта этого tile из первой строки, компактные эксперты при нехватке последнего блока до block_m строк сохраняются построчно и никогда не записываются в строки следующего эксперта, SM89 отвергает непрерывную раскладку; маскированные вызовы по-прежнему строят ту же самую программу, что и раньше, 35 маскированных построений TIR побайтово совпадают с основной веткой. На стороне бенчмарка добавлены три строки форм EP8 для Qwen3-30B-A3B, использующие измеренные на реальной маршрутизации размеры групп (у модели 128 экспертов, top-8, взятый 44-й слой из 48 со средним по скосу смещением, 16 экспертов того rank в EP8, которому досталось больше всего строк), строки prefill маршрутизируют 4096 token (gate/up и down), строки decode берут последний token 128 последовательностей.

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

3.5 TileOPs на рассмотрении: селекторы W4A16 и плотного GEMM переведены на единую модель стоимости с повторной подгонкой (#2825 / #2826)

Оба PR относятся к подзадаче 2 задачи #2133, обновлены 11-03 в 20:54 и 20:54 соответственно, один отвечает за W4A16, другой за плотный GEMM, причём последний меняет только документацию. #2825 меняет _select_config на argmin _config_cost по _legal_configs, удаляя три исключающих правила: приоритет tile, окно равенства по самому широкому K и joint_switch_margin; в _config_cost добавлен член stage_fill — более глубокий цикл способен скрыть больше задержки на каждый K tile, но само заполнение медленнее, без этого члена на K=4096 модель ставила бы 256-ширинные четыре стадии перед 512-ширинными двумя стадиями, тогда как последняя на практике быстрее до 11%; mma_big, iter_ws и launch_ms удалены, поскольку данных недостаточно, чтобы их зафиксировать. Константы повторно подогнаны по занятому времени на тактовой частоте SM бенчмарка, формы для подгонки берутся вне списка, а сам список оставлен для проверки на отложенной выборке; в строке документации _Calibration описаны три шага — сканирование, подгонка и проверка, — а повторный запуск с нейтральных констант воспроизводит заводские значения. На стороне ядра есть и реальная экономия памяти: CTA split-K буферизует только свои группы scale и zero для своего среза, а не весь K, при выравнивании среза использует векторизованную загрузку; _smem_bytes и _legal_configs считают разделяемую память на каждый split, тогда как в основной ветке каждый CTA split-K буферизует группы всего K. Ещё одна деталь: _ConfigSpace.smem_bytes размещён на 1 KB ниже необязательного верхнего предела — ровно на пределе 17 сканируемых конфигураций не запускались, поскольку аллокатор округляет динамическую разделяемую память вверх, а барьер статический. #2826 же описывает в строке документации, как повторно подгоняются константы селектора плотного GEMM (три шага — сканирование, подгонка и проверка, включая мягкий отжиг и проверку на отложенной выборке), не меняя код.

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

3.6 TileOPs на рассмотрении: упаковка коротких varlen по группам KV, L2 flush переведён в режим только для чтения (#2819 / #2822)

#2819 обновлён 11-03 в 23:37, является частью #2496 и добавляет упакованную warp-специализированную программу для GQAVarlenFwdOp на SM90: когда запрос с группировкой по KV-головам побитово заполняет один tile и это экономит tile, ключи и значения читаются только один раз на KV-голову; логика отказа упакованного ядра использует модель стоимости, оценивающую пакет вызовов, суммарное число упакованных query и размер группы, и принимает работу только тогда, когда гарантируемая формой экономия превышает измеренные 9% накладных расходов на однотипный tile на одну голову; широкие и узкие элементы поочерёдно забираются в одном и том же запуске, без обратного чтения на хосте и без дополнительной синхронизации; когда число неразделённых элементов, посчитанное на устройстве, меньше числа SM, запрос на упаковку одного tile разбивает сканирование ключей на несколько сегментов, объединяемых через рабочую область fp32. На стороне окна также есть два измеренных выигрыша: строки окна экономят от 1% до 9%, строки с чередующимся RoPE — от 3% до 6%. Одновременно исправлены три проблемы в основной ветке: строки без видимых ключей при наличии sink и масштабе score больше примерно 2.4 возвращали NaN (в обеих программах — warp-специализированной для одной головы и warp-специализированной для FP8); резервное вращение query читало таблицу RoPE на строках перед первым ключом запроса; а отклонённый негативный контрольный пример оставлял кадр вызывающей стороны в цикле ссылок, из-за чего CUDA graph одного теста уничтожался во время захвата в другом тесте. Тестовые узлы выросли с 173 до 197. #2822 обновлён 11-03 в 19:01 и меняет L2 flush в bench_kernel с «zero_() по буферу int8» на «чтение приватного буфера int32, размер которого вдвое превышает сообщённый размер L2», буфер один раз извлекается генератором на CPU с фиксированным seed; на CUDA это чтение выполняется ядром Triton (одна загрузка на элемент, один atomic_max на блок в 16 KB), а не собственной редукцией torch, потому что при использовании её в качестве flush следующий вызов Marlin w4a16 замедляется на 1.3–1.5 микросекунды. Выигрыш на H200, 1500 MHz, по времени ядра CUPTI (медиана 30 вызовов в каждом из 3 прогонов): Sum для 2048×11008 f32 — с 27.81 до 21.38 микросекунды, Ge (torch-compile) для 1024×11008 f32 — с 31.31 до 25.67, для 1024×4096 f32 — с 6.30 до 5.47; старый flush оставлял около 120 MB грязных строк, запись которых попадала в следующий измеряемый вызов, а новый flush сам пишет только от 71 до 445 секторов L2 (2–14 KB). Попутно исправлен баг основной ветки: при неизвестном размере L2 журнал писал 256 MB, но выделялось 512 MB.

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

3.7 Адоптеры: FlashQLA и TileKernels без пушей, но #3761 исправляет fused-ядро TileKernels (11-04)

Последний пуш QwenLM/FlashQLA остановился на 10-30 03:02, deepseek-ai/TileKernels — на 10-14 05:49, в обоих репозиториях в этом окне нет коммитов, на стороне адоптеров нового кода нет. Однако главная новость выпуска напрямую связана с адоптерами: performance-регрессия, исправленная в #3761, была вызвана именно ядром mhc_pre_big_fuse из TileKernels — то же самое ядро на B300 при скрытой размерности 1280 потеряло 25% бюджета занятости барьеров. Это показывает, что ядра адоптеров уже стали регрессионным зондом для стороны компилятора: в те недели, когда TileKernels не пушит код, его ядра всё равно пассивно «тестируются» изменениями основной ветки. Для тех, кто следит за экосистемой, это обратный сигнал, который стоит учитывать: в период молчания адоптеров изменения компилятора всё ещё могут ломаться на их формах.

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

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

4.1 7 новых issue за один день: одна проблема production-уровня с CUDA graph, четыре от одного внешнего репортёра (11-03)

В окне мониторинга в основном репозитории добавлено 7 issue (номера #3762–#3775, плюс один автоматический ежедневный отчёт в репозитории Ascend). Наиболее технически содержательным является #3775: сообщивший — команда с производственной нагрузкой: окружение — tilelang 0.1.12, установленный через pip, Python 3.12.3, две H200 с NVLink, драйвер 580.178.04, контейнер CUDA 13, torch 2.14.1+cu130, на стороне сервиса — sglang с одним процессом TP2 (по одному рабочему потоку на устройство, у каждого свой контекст CUDA), модель — GLM-семейства MoE с вниманием DSA, по пути задействованы ядра TileLang mhc_pre_big_fuse_with_norm, mhc_post и ядро prefill DSA. После принудительного переключения бэкенда исполнения на TILELANG_EXECUTION_BACKEND=nvrtc (по умолчанию auto разрешается в tvm_ffi) работа JIT-ядер tilelang во время захвата CUDA graph приводит к повреждению захвата, причём повреждение проявляется в нижележащих несвязанных ядрах после запуска tilelang, и три конфигурации, побайтово идентичные, воспроизводят три разные первые точки отказа: в одном случае cudaError700 (недопустимый дескриптор TMA) выбрасывается при инициализации дескриптора TMA и затем отдаётся эхом в несвязанной точке запуска, в другом — cublasGemmEx сообщает CUBLAS_STATUS_EXECUTION_FAILED (чистый BF16 на входе и выходе), в третьем — недопустимое обращение к памяти в поэлементном ядре, скомпилированном JIT хост-фреймворком; все три происходят во время захвата CUDA graph на decode (корзина размера батча 8), примерно через 100 секунд после запуска сервиса, и разница между /proc/environ и чистым контролем состоит только в одной этой переменной окружения. Это относится к наиболее трудно локализуемому классу: случайный симптом, дрейфующая точка отказа, единственная причина — одна переменная окружения; сообщивший полностью описал условия сравнения для всех трёх запусков.

Ещё четыре поступили от того же внешнего сообщившего; окружение единообразно: основная ветка или pip-пакет 0.1.15, Linux под WSL, Python 3.13.5, RTX 4060 Laptop (sm_89), драйвер 566.26, среда выполнения CUDA 12.6, и ко всем приложены воспроизводящие скрипты с полным выводом: #3774 — перестановка в разделяемой памяти 16×16 выполняет только сдвиг по строке и теряет смещение по столбцу (223 из 256 элементов неверны, и результат полностью совпадает с «потерей сдвига по столбцу»), тогда как та же запись для 32×32 корректна; #3773 — три копирования в разделяемую память плюс GEMM внутри однородного цикла while дают нестабильный результат (последовательный цикл-контроль: 0 ошибок из 100; версия с while: 21 ошибка из 100, 10 различных выводов); #3763 — вложенное целочисленное взятие нижней границы (x // -2) // 4 для int16, 64 элементов, 32 потоков выдаёт Let cannot bind the same var to two different values. Эти три вместе с #3775 образуют концентрированную зону отказов комбинации «потребительская карта + WSL + установка через pip». Оставшиеся два: #3769 — int8 GEMM для tcgen05 на B300 отклоняется ptxas (есть уже рассматриваемый фикс #3770), #3762 — на Ascend 950 T.simd.vcvt при sat=False всё равно насыщает; подробности в следующем пункте.

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

4.2 Семантика насыщения vcvt на Ascend 950 и ежедневный тест mlir-ascend (#3762 / #1980)

#3762 подан пользователем со стороны Ascend; окружение — Ascend950PR (NpuArch 3510 / dav-3510), CANN 9.2.0-beta.2, драйвер 25.7.rc1, torch_npu 2.12.0, TileLang установлен из исходников (e84ce3de от 10-21). Проблема в том, что T.simd.vcvt(src, "uint8", sat=False) для вектора int32 молча игнорирует sat=False и всегда насыщает: отрицательные значения становятся 0, значения больше 255 становятся 255, без ошибок и предупреждений, тогда как вызывающая сторона считает, что использует семантику переноса. Документация для этого случая написана однозначно — при sat=False значения вне диапазона переносятся — ожидаемые значения: -1 → 255, -128 → 128, 300 → 44; фактически: -1 → 0, -128 → 0, 300 → 255. Влияние сказывается на путях вида f32 → int8 в завершающей стадии int8-квантования: как только появляется отрицательное промежуточное значение, семантика меняется с переноса на усечение к нулю. #1980 — автоматический ежедневный тест в репозитории Ascend; в прогоне от 11-03 09:35 все 3030 пунктов пройдены, проходимость 100%, номер запуска workflow 41553942797. Интересно сопоставить эти две цифры: ежедневный тест в репозитории Ascend полностью зелёный, но семантические расхождения, не покрываемые ежедневным тестом (например, фактическое поведение параметра sat), по-прежнему требуют ручных сообщений.

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

4.3 Три обновления сайта документации, расширение VS Code переименовано в vsix (11-03)

Сайт документации tilelang.github.io в пределах окна мониторинга получил три коммита Update docs (11-03 13:11, 13:28, 11-04 00:59 пекинское время) — это регулярные обновления, автоматически собираемые при слиянии в основную ветку; их содержимое соответствует нескольким изменениям, вошедшим в основную ветку в этом окне: послойному отображению #line, списку областей действия ThreadSync и соответствующим разделам документации по пакетному копированию для Ascend. TileOPs.github.io в этом окне не обновлялся, последняя версия по-прежнему от 11-01 — «Следование основной ветке TileOPs: гарантии перекрытия и выравнивания, инструкция по сборке». В репозитории расширения VS Code 11-04 в 02:19 появился коммит «Обновить документацию и переименовать в vsix», следующий за коммитом от 11-02 «Выровнять отладку CUDA по построчному отображению TileLang»: построчная привязка, вошедшая в основную ветку в прошлом выпуске, изменила источник номеров строк в сгенерированном коде, поэтому конфигурацию отладки на стороне расширения пришлось менять вслед за этим — этот коммит скорректировал одновременно имя артефакта и документацию.

Источники: https://github.com/tile-ai/tilelang.github.io, https://github.com/tile-ai/tilelang-vscode-edition

4.4 Новостная сторона: три канала без попаданий в окне, зарубежные СМИ расходятся в оценке обучения на Ascend (11-03)

Три канала — комбинированные запросы Google News на китайском и английском, Hacker News и arXiv — в этом окне не дали существенных попаданий, связанных с TileLang: из 15 пунктов списка кандидатов 12 получены при обнаружении репозиториев GitHub, 2 — не относящиеся к теме Show HN (детективная игра и инструмент подсказок для проектирования аппаратуры), 1 — материал Wccftech о DeepSeek и Ascend. Последнее попадание по all:tilelang на arXiv по-прежнему относится к TileSight от 2026-07-24 (модель профилирования производительности GPU с центром на tile); совместный запрос с TileLang в пределах окна не дал новых статей. Единственная отраслевая новость, заслуживающая упоминания, содержит конфликт формулировок: Wccftech утверждает, что DeepSeek будет обучать V5 с нуля на Atlas 950 SuperPoD, построенном на Ascend 950DT, тогда как опубликованный в тот же период материал The Information говорит о «начале предобучения на 950DT не ранее конца 2026 года», а ранние публикации Bloomberg и Financial Times сообщали, что заказ на 160 тысяч чипов 950DT предназначен в основном для инференса, а после неудач с обучением на Ascend компания вернулась к Nvidia. Три версии расходятся в том, «будет ли и когда 950DT использоваться для обучения», и ни одна не подтверждена официально Huawei или DeepSeek. Приводим это здесь лишь как фон: независимо от того, состоится ли обучение, вложения в компилятор и операторы на стороне Ascend (в этом выпуске #3689, #1975, #1981) закладывают фундамент для этого пути.

Источники: https://wccftech.com/china-going-for-domestic-full-stack-ai-dominance-deepseek-will-train-its-upcoming-v5-model-from-scratch-on-huaweis-ascend-950dt-based-atlas-superpods, https://arxiv.org/

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

Во-первых, регрессии производительности начинают восприниматься как полноценные объекты первого класса, и методология их исправления заслуживает внимания больше, чем сами выводы. Правило сохранения ID, введённое в #3643 ради исправления корректности, само по себе необходимо и правильно; #3761 не откатил его, а позволил сохранению и повторному использованию сосуществовать на разных уровнях: прежде чем совместно использовать один ID в разных областях действия, сначала доказывается, что участвующий предикат выполняется всегда, а если доказать не удаётся — откат к счётчику с областями действия. В сочетании с изолирующим экспериментом автора «изменить только объявленное число ресурсов барьера» эта цепочка исправлений разделила две вещи — «корректность» и «объявление ресурсов»: последнее стало переменной, которую можно отдельно воспроизвести и отдельно принять. Именно такая способность к разделению и является признаком перехода проекта компилятора от «работает» к «поддерживаемо».

Во-вторых, инженерия Ascend продвигается одновременно по двум направлениям. Одно — архитектура: #1975, добавив 2887 файлов и 600 тысяч строк, переносит весь TileLang-Ascend в многофреймворковую основную ветку с разделением бэкендов по таким идентификаторам, как arch="dav-2201", с целью ортогонального сосуществования A2A3 и A5 в одном репозитории; другое — конвейер: #208 в mlir-ascend записывает в .workspace/ три набора рабочих процессов Codex — инициализацию, подключение и восстановление — а также передачу задач, при этом основной репозиторий и репозиторий Ascend имеют собственные каталоги Agent. Глядя на «репозиторий миграции» и «миграцию репозитория с помощью агентов» в один день, можно увидеть, что узкое место на стороне Ascend сместилось с «нехватки операторов» на «нехватку согласованного инженерного каркаса и переиспользуемого процесса миграции».

В-третьих, культура самодоказательства в библиотеках операторов углубляется. В нескольких PR TileOPs, находящихся на рассмотрении в этом выпуске, «доказательства» почти всегда записаны прямо в PR: размер группы EP8 для Qwen3-30B-A3B, измеренный на реальной маршрутизации (взят 44-й слой из 48 — со смещением, близким к медианному), TIR, построенный по 35 маскам, побайтово совпадает с основной веткой; изменение L2 flush сопровождается таблицей сравнения с 46-строчной реализацией и пояснением, что старый flush оставлял 120 МБ грязных строк; модель стоимости W4A16 явно указывает, «какие элементы удалены, потому что данные не удаётся закрепить». TileFoundry идёт другим путём — делает синхронизацию доказуемым анализом, затем заставляет lowering выполнять обратную проверку, а случаи, которые анализ отвергает, прямо перечисляются (семантика завершения TMA и wgmma ещё не объявлена, циклы с неравными границами между потоками). Оба пути вместе демонстрируют одну и ту же ориентацию: лучше признать «не доказано», чем принимать «запустилось» за «правильно».

В-четвёртых, поверхность рисков смещается к форме установки. Из 7 новых issue за день четыре поступили от одного и того же внешнего репортёра, использующего pip-пакет, работающего в WSL и применяющего потребительскую карту sm_89, — все с воспроизводимыми скриптами; #3775 — это комбинация «pip-пакет + несколько контекстов CUDA + бэкенд nvrtc + захват CUDA graph» в производственной среде. Это не глубокие дефекты компилятора, но они определяют опыт нового пользователя при первой попытке. Если сравнить с репозиторием Ascend, где ежедневные 3030 тестов полностью зелёные, а семантическое расхождение sat=False по-прежнему выявляется только ручными отчётами, видно, что между формой покрытия тестами и формой реального использования ещё есть дистанция.

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

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

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

Источник Результат проверки
Организация tile-ai (29 репозиториев) За окно новые пуши в 9 репозиториях
Основной репозиторий tilelang 5 слияний в основную ветку; 6 новых issue; 23 обновления на рассмотрении; 3 автоматические сборки сайта документации
TileOPs 1 слияние в основную ветку; 7 обновлений на рассмотрении
TileFoundry 1 слияние в основную ветку; 2 обновления на рассмотрении
tilelang-ascend В ветку по умолчанию ascendc_pto за окно новых коммитов нет; 15 обновлений на рассмотрении; все 3030 ежедневных тестов пройдены
tilelang-mlir-ascend 1 слияние в ветку harness (не основная линия)
tilelang-hygon В основную ветку новых коммитов нет; #19 на рассмотрении имеет обновления
tilelang-vscode-edition 1 слияние (документация и переименование vsix)
tilelang-metax / tilelang-musa / TileRT / tilescale / DeepStack / tilelang-puzzles За окно пушей нет
Страницы релизов (tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-hygon) За окно новых версий нет; последняя у tilelang по-прежнему v0.1.15 от 09-30
FlashQLA / TileKernels За окно пушей нет
Google News (шесть запросов на китайском и английском) / Hacker News / arXiv Существенных попаданий за окно нет

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

[23] Перенос прямого прохода Chunk O в tilelang-ascend (#1866) — https://github.com/tile-ai/tilelang-ascend/pull/1866