Окно мониторинга: 2026-10-21 07:00 ~ 2026-10-22 07:00 (пекинское время, 24 часа).

Источники: поэлементная проверка репозиториев организации GitHub tile-ai (поток коммитов Atom, охват tilelang / TileOPs / TileFoundry / tilelang-ascend〔ascendc_pto〕/ tilelang-hygon〔ветки main и feature〕/ tilelang-metax〔dev〕/ tilelang-musa / tilelang-mlir-ascend / TileRT / tilescale / DeepStack) + поиск достоверных данных о смерженных и находящихся на рассмотрении (две группы is:merged / is:open, статистика с группировкой по репозиториям) + проверка тегов на странице релизов + проверка коммитов сайта документации + проверка репозиториев потребителей (TileKernels / FlashQLA) + поиск новых репозиториев сообщества + множественные запросы Google News на китайском и английском языках (через прокси) + Hacker News + arXiv + сторонние отчёты отраслевого анализа.

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

  • Главное за сегодня: официальный релиз TileFoundry 0.0.3, планирование на уровне инструкций передано agent; TileOPs за один день влил 20 коммитов, обратный проход MLA и маломасштабный путь SM89 дополнены в один день; восемь коммитов в основном репозитории сосредоточены на компиляторных дефектах класса «не выдаёт ошибку, но результат недостоверен»; выравнивание DeepGEMM для Ascend переведено из статуса на рассмотрении во влито (10-21)
    1. Основные результаты по проектам
      • 1.1 Основной репозиторий: пакетный finalize_reducer ранее сворачивал только первый выходной блок (#3070) (10-22)
      • 1.2 Основной репозиторий: векторизованное преобразование BF16→FP32 переведено на inline PTX с упаковкой lane-конверсий (#3595) (10-21)
      • 1.3 Основной репозиторий: pass разворачивания сохраняет пустые stepped-циклы, Select в ROCm переведён на понижение по lane (#2887) (10-21)
      • 1.4 Основной репозиторий: программный конвейер сохраняет транзитивные зависимости копирования, домены потоков участвуют в определении конфликтов областей (#2822) (10-21)
      • 1.5 Основной репозиторий: предварительная проверка трёх комбинаций источник/цель/округление для cast, отказ от молчаливой потери аннотаций (#3350) (10-21)
      • 1.6 Основной репозиторий: CUDA target без arch отклоняется вместо молчаливого заполнения sm_50 (#2860) (10-21)
      • 1.7 На рассмотрении: исправления знаковости сдвига и области действия цикла TIRx парными коммитами в tvm fork и основном репозитории (#3597 / tvm #87 / #86) (10-21)
    1. Адаптация под множество бэкендов (Ascend / MetaX / Hygon / Moore Threads и др.)
      • 2.1 Оценка Ascend: в глубоком отчёте SemiAnalysis выделена отдельная глава о TileLang Ascend, где утверждается, что из шести репозиториев на тот момент работал только TileLang (10-21)
      • 2.2 Основной репозиторий, Ascend: эвристика DeepGEMM и выравнивание epilogue переведены из статуса на рассмотрении во влито (#3577) (10-21)
      • 2.3 Основной репозиторий, Ascend: исправление целочисленной знаковости warp_reduce_max/min на PTO (#3582) (10-21)
      • 2.4 Репозиторий Ascend: суб-32B компактный источник построчного интервала невозможно выразить, вывод keepdim reduce искажается (#1795) (10-21)
      • 2.5 Репозиторий Ascend: JIT-кэш переведён на атомарную публикацию с восстановлением при удалении кэша (#1903) (10-21)
      • 2.6 Репозиторий Ascend: в cann-bench добавлен пример depthwise_conv_2d (#1915) (10-21)
      • 2.7 Ascend на рассмотрении: indexer_topk_reducesum, GLM-5.3 kpool compress и миграция серии операторов WY fast (10-21)
      • 2.8 Hygon: на feature-ветке продвигается FP8-декодирование GEMM, локализовано, что цепочка адресов занимает 13% ядра (10-21)
      • 2.9 MetaX / Moore Threads / MLIR-Ascend: в окне мониторинга нет пушей, три линии остаются на существующих версиях (10-21)
    1. Экосистема и пользователи (TileOPs / TileFoundry / библиотеки операторов)
      • 3.1 TileFoundry: тег 0.0.3 и релиз, планирование на уровне инструкций + рефакторинг анализа + примеры для H200 и Apple AMX (10-21)
      • 3.2 TileOPs: обратный проход префилла MLA дополнен на уровне op и warp-специализированным ядром SM90 (#2678 / #2680 / #2688 / #2686) (10-22)
      • 3.3 TileOPs: dense GQA поддерживает FP32 attention sink, -inf отключает, +inf обнуляет этот head (#2515) (10-21)
      • 3.4 TileOPs: сосредоточенное дополнение маломасштабного пути SM89, плитка FP8 BMM по умолчанию снижена до 64×64 (#2681 / #2672 / #2669 / #2654) (10-21)
      • 3.5 TileOPs: NSA top-k на SM89 расширен до 2048 сжатых блоков и дополнены выходные слоты (#2674 / #2675 / #2673) (10-21)
      • 3.6 TileOPs: прямой проход GDN завершается одним ядром, когда walk заполняет SM; нормализация с квантованием по кластерам режет широкие строки (#2682 / #2679) (10-21)
      • 3.7 TileOPs: обратный проход GQA в warp-специализированном ядре напрямую обслуживает сгруппированные KV head (#2685) (10-21)
      • 3.8 TileOPs на рассмотрении: сворачивание edge-оси одним запуском, сужение столбцовой плитки NSA top-k (#2689 / #2690) (10-22)
      • 3.9 Пользователи: TileKernels и FlashQLA в окне мониторинга без пушей (10-21)
    1. Сообщество, обучающие материалы и мероприятия
      • 4.1 Обучающий материал сообщества: обновлён китайский «TVM / TileLang Cookbook», от jit до cubin (10-21)
      • 4.2 Новый репозиторий сообщества: создан tilelang-wrapper-for-sail, описание пока отсутствует (10-21)
      • 4.3 Инструмент сообщества: репозиторий удалённого agent автотюнинга продолжает пополняться (10-21)
      • 4.4 Документация и версии: сайт документации обновлялся автоматически семь раз, VERSION основного репозитория по-прежнему 0.1.15 (10-22)
    1. Наблюдения за тенденциями

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

В данном окне включено 32 коммита (TileOPs 20, tilelang 8, tilelang-ascend 3, сайт документации TileOPs 1), а также один официальный релиз: TileFoundry 21.10 в 10:26 выпустил 0.0.3. Акценты трёх линий различаются. Восемь коммитов в основном репозитории почти все представляют собой устранение дефектов класса «не выдаёт ошибку, но результат недостоверен» — пакетное сокращение обрабатывает только первый выходной блок, аннотация округления cast молча отбрасывается, CUDA target без arch молча откатывается до sm_50, планирование конвейера теряет транзитивные зависимости; TileOPs за день довёл обратное распространение MLA prefill от уровня оператора до warp-специализированного kernel для SM90, одновременно продвинув четыре небольших пути на SM89 — FP8 BMM, W4A16, paged DSA и NSA top-k, — а очередь на рассмотрении сократилась с 16 до 9; на стороне Ascend вчера ещё находившиеся на рассмотрении эвристика DeepGEMM и выравнивание epilogue (#3577) сегодня включены, а в репозитории Ascend исправлен дефект моделирования доступа к памяти, вызывавший искажение вывода keepdim reduce (#1795). В индустрии SemiAnalysis в глубоком отчёте о Huawei-DeepSeek посвятил отдельный раздел TileLang Ascend, сделав вывод, что из шести репозиториев, получивших поддержку Ascend, на тот момент работал только TileLang.

1. Прогресс ключевых проектов

1.1 Основной репозиторий: пакетный finalize_reducer ранее сокращал только первый выходной блок (#3070)

Включён 22.10 в 01:16, 2 файла +29/-5, исправляет #2623. Понижение пакетного finalize_reducer вызывало run_batch только один раз в начале буфера, остальные последовательные выходные блоки вообще не сокращались: при размерности вывода больше batch последующие сегменты получали не результат сокращения, а неинициализированные значения или значения, оставшиеся от предыдущего раунда. Исправление выпускает пакетный all-reduce с корректным смещением поблочно по ceil(OutputDim / batch) — при 128 выходах и batch=4 выпускается 32 вызова, — одновременно переключая буфер ремарпинга reducer на плоское представление с адресацией каждой партии по константному линейному смещению, убирая прежнее перестроение формы и шага и сохраняя все вызовы chunk. Проверка выполнена на L40S: добавлены числовые регрессии sum/max/min, покрывающие fp16 и fp32, batch 4/8/16, 23 связанных с finalize_reducer теста прошли, внешняя матрица воспроизведения из 16 случаев — ноль расхождений. Это типичный дефект класса «корректно на малых размерах, молчаливо ошибочно на больших».

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

1.2 Основной репозиторий: векторизованное преобразование BF16→FP32 переведено на inline PTX упакованного lane-преобразования (#3595)

Включён 21.10 в 19:28, 4 файла +19/-7.
Добавлен tl::__tl_cvt_bf16x2_to_f32x2, который с помощью inline PTX преобразует пары BF16 lane за один раз в f32x2, заменяя скалярный путь повышения по каждому lane; поддержка всех трёх ширин lane — 2, 4 и 8 — сохранена. Одновременно обновлены ожидаемые значения векторизованного cast и вывода раскладки, а тесты BF16 cast сведены к существующему покрытию без расширения набора утверждений. Изменение не затрагивает семантику, меняется только генерируемая последовательность инструкций, а область влияния — все пути вычислений BF16, идущие к квантованному cast.

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

1.3 Основной репозиторий: pass разворачивания сохраняет пустые stepped-циклы, Select в ROCm переведён на поlane-понижение (#2887)

Включён 21.10 в 19:20. Pass разворачивания циклов ранее удалял stepped-циклы с trip count 0 как пустую оболочку, изменяя структуру программы, видимую последующим pass; исправление сохраняет такие пустые циклы в процессе разворачивания. В том же коммите изменён векторизованный Select на стороне ROCm: он переведён на поlane-понижение, чтобы избежать отказа группового понижения на неподдерживаемых целях. Одно изменение относится к структуре IR, другое — к понижению на бэкенде, и они объединены в одном коммите, что указывает на их происхождение из одного сценария воспроизведения.

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

1.4 Основной репозиторий: программный конвейер сохраняет транзитивные зависимости копирования, домен потоков участвует в определении конфликтов областей (#2822)

Включён 21.10 в 18:13. Планирование конвейера ранее отслеживало только прямые зависимости, косвенные (транзитивные) зависимости копирования терялись на этапе планирования, а вместе с ними удалялось ожидание dependent copy; кроме того, проверка конфликтов физических областей имела два разных набора правил на сторонах планирования и понижения. Исправление использует обратный worklist, отслеживающий одновременно скалярные зависимости и зависимости копирования буферов, заставляет планирование и понижение использовать одну и ту же проверку конфликтов физических областей, сохраняет привязку воспроизведения в области видимости и дополняет покрытие транзитивных производителей, strided alias и пути последовательного отката TMA. Второй коммит вводит домен потоков в определение конфликтов областей: обращения сначала разворачиваются до материализованного домена потоков, а затем сравниваются области, что позволяет конвейеризовать непересекающиеся срезы разделяемой памяти, сохраняя при этом физическое перекрытие межпоточных зависимостей и strided-буферов; один и тот же домен используется для планирования конвейера, проверки зависимостей и версионирования буферов.

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

1.5 Основной репозиторий: предварительная проверка трёх комбинаций источника/цели/округления cast (#3350)

Смержено 21.10 в 13:49. round может понижаться только для явно перечисленных комбинаций: PTX cvt.rs требует источник f32, 32-битный операнд rbits и определённую упакованную цель. Остальные комбинации ранее пропускались при диспетчеризации ветки, и аннотация округления молча исчезала; скалярные циклы такие комбинации и так отвергали, векторизованные циклы не отвергали, поэтому один и тот же cast вёл себя по-разному в двух формах циклов. Типичный пример — f32→E8M0: его helper округляет к плюс бесконечности и вообще не принимает параметр округления, а векторизованная ветка никогда не читала round. Теперь недопустимые комбинации единообразно отвергаются до диспетчеризации, и ни одна ветка больше не может потерять аннотацию; сопутствующий тест передаёт этот параметр только в режиме rs, который действительно принимает rbits.

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

1.6 Основной репозиторий: CUDA target без arch отвергается, а не молча заполняется sm_50 (#2860)

Смержено 21.10 в 13:32, 1 файл, исправляет #1049. На машине с CUDA toolkit, но без доступного GPU, CUDA target без указанного arch вырождается в голый cuda, TVM затем молча подставляет arch=sm_50, и компиляция продолжается для неверного SM без какой-либо диагностики на всём пути — именно эта ловушка возникает при tilelang.compile(func) с target="auto". Исправление находится в _cuda_target_from_arch: единственная точка входа, общая для автодетекта и явного target="cuda", теперь выбрасывает действенную ошибку и приводит в ней два способа явного указания arch ({"kind":"cuda","arch":"sm_90"} или TILELANG_DEFAULT_TARGET). Headless-компиляция по-прежнему работает, если задан конкретный arch, поскольку codegen и nvcc зависят только от toolkit, а доступ к устройству откладывается до запуска kernel. В ходе работы над PR одно время были добавлены три юнит-теста без GPU, но в финальном коммите они были убраны, и изменение свелось только к проверке на входе.

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

1.7 На рассмотрении: исправления знаковости сдвига и области действия цикла TIRx оформлены парными коммитами

Основной репозиторий #3597 исправляет потерю знаковости при смешанных знаковых сдвигах, параллельно в fork tile-ai/tvm отправлен #87 «TIRx сохраняет тип левого операнда сдвига» — это соответствующие изменения одной и той же логики на уровне языка и на уровне TIRx. Кроме того, tvm #86 переносит обратно исправление области действия цикла из upstream apache/tvm#19927. В fork на рассмотрении 25 штук, в основном репозитории — 163; опережающие изменения на стороне TIRx приводят к тому, что fork берёт на себя предварительную проверку части изменений, а основной репозиторий затем подхватывает их — этот путь сотрудничества встречается в данном окне второй раз.

Источник: https://github.com/tile-ai/tilelang/pull/3597, https://github.com/tile-ai/tvm/pull/87, https://github.com/tile-ai/tvm/pull/86

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

2.1 Оценка Ascend: в глубоком отчёте SemiAnalysis выделена отдельная глава о TileLang Ascend

SemiAnalysis 21.10 опубликовал «The Huawei-DeepSeek Symbiosis», в подзаголовке которого прямо перечислены «глубокое погружение и микробенчмарки Huawei Ascend 950PR, анализ DSv4.1 Flash, поддержка Ascend в DeepSeek, AICPU, TileLang Ascend». Три пункта, напрямую связанные с TileLang: во-первых, из шести репозиториев, получивших поддержку Ascend (TileLang, DeepGEMM, DeepEP, TileKernels, FlashMLA, DeepSelect), на момент написания статьи был доступен только TileLang, поскольку CANN 9.2.0 ещё не выпущен, а DeepSeek требует CANN 9.2.0 и компилятор C++20, тогда как 9.2.0-beta.2 последний не предоставляет; во-вторых, бэкенд Ascend 950 уже включён в основной репозиторий TileLang, а не существует только в отдельном репозитории адаптации под Ascend, и обеспечивает нативную компиляцию, планирование, синхронизацию и интеграцию со средой выполнения; AutoSchedule на основе зависимостей по данным и оценочного времени выполнения определяет перекрытие и число слотов он-чип буфера, на 950PR сокращая время до 40% для небольших одиночных операторов, а компилятор также автоматически вставляет синхронизации, удаляет лишние ожидания и проверяет условные ветвления; в-третьих, при сравнении с CANN и Triton-ascend «на некоторых размерах можно сравнить или превзойти CANN», причём оптимальная реализация переключается в зависимости от размера (автор отмечает, что это наивный бенчмарк при горячем кэше). Статья также указывает на пробел в экосистеме: большинство оптимизаций ориентировано на суперузлы, а на x86-хостах vLLM-Ascend 0.26.0rc1 всё ещё имеет множество проблем.

Источник: https://newsletter.semianalysis.com/p/huawei-deepseek

2.2 Основной репозиторий, Ascend: эвристика DeepGEMM и выравнивание epilogue переведены из рассмотрения в мерж (#3577)

Объединено 10-21 10:20, 29 файлов +1392/-1036. В прошлом выпуске отмечался как находящийся на рассмотрении, в этом выпуске завершено слияние: дополнены выбор конфигурации и epilogue для DeepGEMM-Ascend, охватывающие fused bias, Alpha, стохастическое округление BF16 и квантование FP8 для плотных, пакетных и групповых GEMM, а также в собственных коммитах автора удалён избыточный setup и восстановлено тестовое стробирование. Этот коммит взаимно подтверждается с внешней оценкой из 2.1 — на стороне Ascend повторно оттачиваются именно эвристики и числовое завершение пути GEMM, и это как раз и есть разрыв между «работает» и «работает хорошо».

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

2.3 Основной репозиторий Ascend: исправление знаковости целых чисел для warp_reduce_max/min на PTO (#3582)

Объединено 10-21 10:21. pto.redux_max/min в PTO отвергает типизированные целые числа: PTOAS требует носитель без знаковости плюс явную signedness, поэтому T.warp_reduce_max/min для int32/uint32 напрямую не компилируется, выдавая «requires signedness=’signed’ or ‘unsigned’ for integer values». Исправление унифицирует warp_reduce_sum/max/min через обёртки tl.redux_* из contrib.ptodsl (теперь экспортируются), где обёртка предоставляет носитель и знаковость, а также устанавливает ограничение на value dtype по образцу AscendAllReduce; в задаче указаны основной репозиторий #3541 и PTOAS #1579. Верификация двухуровневая: тесты lower-only codegen по сетке op × dtype, а также разрежённые тесты исполнения с двумя целями (ascend/pto) — с использованием lane uint32 ≥2^31, что напрямую выявляет неверную знаковость.

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

2.4 Репозиторий Ascend: источник с компактным межстрочным интервалом менее 32B невыразим, вывод keepdim reduce искажается (#1795)

Объединено 10-21 17:15, 2 файла +302/-3, исправляет #1682, это наиболее технически содержательный коммит данного окна. Корневая причина на уровне шаблона: copy_ub_to_gm безусловно генерирует многобёрстный DataCopyPad, а движок MTE3 реализует системное соглашение «узкие строки UB хранятся по 32B-слотам» — адрес источника на каждый burst продвигается на ceil(blockLen/32) + srcStride блоков по 32B. Для производителей, пишущих строки по границам 32B, это точное совпадение; но вывод (M,1) fp32 из T.reduce_max записывается компактно через V-pipe с межстрочным интервалом всего 4B, и движок не может это выразить. Практически второй burst читает element 8 вместо element 1 (в [127, 1151, 71, 199] число 1151 как раз является значением element 8), сторона записи в порядке, ошибка в последовательности адресов чтения источника DMA. Исправление разделено на две части: когда обе стороны packed и total ≤ 65535, свёртка в один плоский одиночный burst (M×L байт, используя контракт «длина одиночного burst произвольна»); невыразимый межстрочный интервал, не кратный 32B, заменяется на явное прерывание trap(), защищённое compile-time if constexpr, без генерации дополнительного кода для кратных 32B и для раскладки слотов менее 32B. Почему не деградировать: автоматическая межконвейерная синхронизация bisheng покрывает только прямолинейные intrinsics; обёртывание DataCopyPad условным ветвлением или размещение скалярной записи GM в братской ветви заставит bisheng генерировать некорректный код синхронизации и приведёт к зависанию устройства (на практике построчный DMA, обёрнутый рантайм-ветвлением, вызывает исключение aicore). 14 тестовых случаев охватывают пути flat/multi, воспроизведение keepdim reduce, повторное использование цикла UB, легальную одиночную burst-форму невыразимого pitch и явное прерывание; тестовая среда — Ascend910 A3 (dav-c220) / локальная сборка CANN 9.0.0. В описании PR сохранены две границы: компактный межстрочный интервал менее 32B с назначением со страйдом (например, запись (M,1) fp32 в некоторый столбец широкого тензора) по-прежнему идёт пошагово по слотам; blockCount > 4095 (12-битное поле) по-прежнему будет молча усекаться (на практике при 4200 строках копируются только первые 104 строки), планируется отдельный issue.

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

2.5 Репозиторий Ascend: JIT-кэш переведён на атомарную публикацию с восстановлением при удалении кэша (#1903)

Объединено 10-21 16:20, 7 файлов +1026/-86, связано с #1836. Исходная реализация перезаписывала скомпилированный артефакт на месте в kernel_lib.so, и если другие процессы уже загрузили эту разделяемую библиотеку, они читали файл в момент перезаписи. Исправление сначала копирует артефакт в отдельный временный файл в том же каталоге, а затем атомарно публикует его через os.replace(); если перед загрузкой кэшированная библиотека уже удалена, после подтверждения по сообщению об ошибке и состоянию файла возвращается cache miss и выполняется повторная компиляция; если на этапе сохранения каталог исчезает, сохранение пропускается, а уже полученные результаты компиляции или тюнинга сохраняются; ошибки прав доступа, отсутствие зависимых библиотек, недействительный ELF и внутренние исключения сериализации продолжают распространяться вверх и не поглощаются как miss. Охват включает KernelCache, AutoTunerCache и AutotuneResult, cache key, логика вычисления операторов и параллелизм CI не изменяются. 41 host-only тестовый пример пройден (9 на атомарную публикацию, 32 на восстановление после удаления кэша), в описании PR указано, что интеграционная проверка на стороне NPU не проводилась. Часть с атомарной публикацией использует реализацию из #1853 и сохраняет информацию об оригинальном авторе.

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

2.6 Репозиторий Ascend: в cann-bench добавлен пример depthwise_conv_2d (#1915)

Объединено 10-21 16:19. В examples/cann-bench/ добавлен бенчмарк depthwise двумерной свёртки. Прилагаемые детали smoke-теста стоит отметить: критерием PASSED в bench_test.sh является появление в выводе KERNEL OUTPUT MATCH или TEST PASSED!, а скрипт ранее печатал «smoke test passed» и при коде выхода 0 всё равно считался проваленным; исправление выровнено по smoke-тестам LSTM/GRU, чтобы новые тестовые примеры не давали постоянных ложных красных результатов в CI.

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

2.7 Ascend на рассмотрении: indexer_topk_reducesum, kpool compress для GLM-5.3 и серия миграций WY fast

В очереди на рассмотрении репозитория Ascend в этом окне 13 обновлений, сгруппированных в три набора. #1917 реализует оператор indexer_topk_reducesum, соответствующий суммированию после выбора top-k в индексаторе разреженного внимания; #1910 добавляет оператор kpool-сжатия для GLM-5.3, указывая на путь KV-сжатия модели нового поколения; #1827 / #1828 / #1802 / #1803 / #1799 — это группа миграций операторов серии WY fast и chunk (сам WY fast и его обратный проход, обратный проход chunk_o, Chunk Scaled Dot KKT, Chunk Delta H), явно нацеленная на цепочку обучения линейного внимания/гейтовых delta-сетей. Кроме того, #1785 исправляет индексацию многомерного скалярного BufferLoad в бинарных операциях T.tile, #1854 проверяет конфигурацию автоматической межъядерной синхронизации, #1731 исправляет ошибку результата W4A4 int4×int4 в gemm_v0 (#438).

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

2.8 Hygon: ветка feature продвигает FP8-декодирующий GEMM, цепочка адресов определена как занимающая 13% kernel

В ветке main репозитория tilelang-hygon в окне нет новых коммитов (последний — #25, объединённый 10-20), но в ветке feat/hcu-smls-b-vgpr 10-21 в 11:04 и 15:27 было три пуша с переиспользуемыми выводами по профилированию производительности и настройке параметров. Пуш в 11:04 разворачивает k-loop декодирующего GEMM и возвращает tile с M>96 к bn96; ключевое суждение автора — «фиксированные накладные расходы на каждом шаге происходят от цепочки адресов, а не от barrier» — символьная индексация на уровне разделяемой памяти (k+d) % 3 опускается до v_mul_lo_u16 / v_sub_u16 / v_readfirstlane_b32 и попадает на путь адресации matrix_load, занимая 13% kernel, тогда как T.Pipelined и так выдаёт те же два barrier. Пуш в 15:27 приоритетно разбивает средний участок wide-N при M∈(64,128] по числу block, для gate_up при M∈(64,96] берёт bm96/bn128/bk128/ns2 и переносит builder внутрь пакетного decode_cfg; одновременно документация по декодирующему GEMM сокращена с 3825 строк до 588 строк в двухчастном виде «результаты + журнал оптимизаций», дополнительно приложены индекс закрытых рычагов и модель пропускной способности.

Источник: https://github.com/tile-ai/tilelang-hygon/tree/feat/hcu-smls-b-vgpr

2.9 MetaX / Moore Threads / MLIR-Ascend: в окне нет пушей

tilelang-metax (ветка по умолчанию dev, последний push 10-16, релиз всё ещё на v0.1.14 от 09-17), tilelang-musa (последний 09-30, v0.1.15+musa.1), tilelang-mlir-ascend (последний 10-19, tag остановился на v0.1.2.020.sglang.poc от 09-09) — во всех трёх репозиториях в данном окне мониторинга нет коммитов и нет новых tag.

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

3. Экосистема и пользователи (TileOPs / TileFoundry / библиотека операторов)

3.1 TileFoundry: релиз 0.0.3 с tag, планирование на уровне инструкций + рефакторинг анализа + примеры для H200 и Apple AMX

В прошлом выпуске отмечалось, что «анонс 0.0.3 уже в README, tag ещё не проставлен»; в этом периоде 10-21 в 10:26 состоялся официальный релиз, на PyPI доступен tilefoundry==0.0.3 с кодовым названием «Agent-driven scheduling». Ключевое изменение — кому принадлежит право планирования: компилятор перечисляет, какими инструкциями может быть реализована каждая операция HIR и что для этого требуется каждой из них, agent выбирает из кандидатов, а компилятор затем понижает выбранное расписание в TIR. В CLI появились tilefoundry schedule candidates и schedule facts для показа кандидатов и ограничений, tf.schedule фиксирует выбор, tilefoundry schedule finalize завершает понижение; инструкции объявляются как операции AtomSched с явным контрактом операндов TIR, начиная с CUDA MMA atom, SM90 WGMMA и асинхронного копирования. Аналитическая часть переработана в три направления: память (сегментированное распределение, footprint и окна переиспользования, учитывающий циклы жизненный цикл buffer), стоимость (работа и трафик по областям выполнения, каждая связь доступа объявляет, что читает) и структура (определение фактов, контекст оценки и снимки областей унифицированно идут через visit context, переиспользование проекций позволяет анализировать каждый пример корпуса лишь один раз). Новый туториал ведёт через GEMM 8192×5120 на 5120×17408 вплоть до оптимального kernel AtomSched на H200.

По производительности приведены три группы реальных цифр: префилл Qwen3.8-27B на одной H200 — 80.29 / 309.61 / 631.19 ms для 1024 / 4096 / 8192 token соответственно, у SGLang — 89.06 / 317.11 / 638.25 ms; короткие последовательности 128 / 256 / 512 / 1024 token дают 1.88× / 1.41× / 1.10× / 1.03× от SGLang; пример префилла Qwen на Apple M2 Pro реализует большое произведение через внешнее произведение FP16 AMX. Два прогона Qwen3.8-27B на каждой длине дают тот же первый token, что и transformers. Известны две проблемы: wheel пока не упаковывает заголовочные файлы CUDA-рантайма TileFoundry, поэтому из установленного пакета tf.compile не может собрать CUDA kernel (нужен checkout исходников, исправят в следующей версии), а также таймлайн производительности неверно оценивает insert_slice / tf.zeros для позиций в mesh и не показывает недозаполнениеmesh lane (#239).

Источники: https://github.com/tile-ai/TileFoundry/releases/tag/v0.0.3

3.2 TileOPs: обратный проход префилла MLA дополнен уровнем op и warp-специализированными kernel для SM90 (#2678 / #2680 / #2688 / #2686)

Четыре коммита образуют единую линию. #2678 (10-21 11:34) добавляет два уровня op — MLADenseBwdOp и MLAVarlenBwdOp, выравнивая их по интерфейсу _flash_attn_backward из FA3; #2680 (10-21 20:58) добавляет kernel для SM90 — для D_qk=192, D_v=128 реализованы две warp-специализированные версии MLADenseBwdWSKernel и MLAVarlenBwdWSKernel; #2688 (10-21 23:28) меняет порядок выборки данных в резидентном CTA в MLAVarlenPrefillWSFwdKernel — прежде каждому head сначала выделялся один query tile, и после завершения первого tile приходилось выстраиваться в очередь по head по одному, теперь за раз обрабатываются query tile только одного и того же head; #2686 (10-22 03:08) — финальная уборка: константы с единственным читателем и helper перенесены к месту использования без изменения поведения. Сайт документации #78 синхронно вносит два новых op в страницу Attention; причиной стало то, что check_api_pages.py сообщил об отсутствии двух страниц среди 207 экспортируемых op основного репозитория.

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

3.3 TileOPs: dense GQA поддерживает FP32 attention sink (#2515)

Смёржено 21.10 в 22:42. Для префилла и декодирования dense GQA добавлен опциональный посимвольный по head FP32 attention sink, охватывающий FP16/BF16/FP8 на SM90, RoPE, softcap и скользящее окно; на других архитектурах вызов sink напрямую отклоняется. Семантически каждый sink — это logit в знаменателе без масштабирования и без softcap: -inf отключает этот sink, +inf напрямую обнуляет выход этого head; split decode при больших logit также должен сохранять вес sink. Попутно исправлено искажение бенчмарка — бенчмарк блочного префилла ранее измерялся как одиночное query-декодирование; хвостовой кейс с fused RoPE переводит длину KV на 271/1025, чтобы покрыть раздельное выполнение.

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

3.4 TileOPs: точечное закрытие пробелов в маломасштабных путях SM89 (#2681 / #2672 / #2669 / #2654)

Все четыре смёрженных PR закрывают слабые места уровня L40S/4090. #2681 меняет тайл по умолчанию для FP8 BMM на SM89 на 64×64 — при условии, что в результирующем grid не более одного блока на SM; для больших grid сохраняется 128×128; автор приводит замеры на одной карте с PyTorch 2.13.0 + CUDA 13.0 + TileLang 0.1.12: 1×16×64×128 снижается с 5.161 до 3.799 us, 16×128×128×1024 — с 12.943 до 7.414 us, а 4096³, наоборот, меняется с 512.573 на 583.680 us (тайл в этой строке не менялся, это шум). #2672 расширяет верхнюю границу M по умолчанию для 64-колоночного W4A16 с 16 до 64; время replay для M=17 / 32 / 64 снижается на 43.35% / 47.20% / 40.00% соответственно, при этом сохраняются явное переопределение конфигурации и числовой допуск. #2669 реализует на MMA постраничный DSA для SM89, сохраняя 656-байтовую раскладку FP8 latent cache, округление BF16, натуральный логарифм LSE и поведение dead position: batch 4 / 128 head / 256 позиций — 0.269 ms против эталона PyTorch FP32 0.727 ms, 48×64 head×2048 позиций — 2.543 против 2.535 ms; автор прямо пишет, что «на больших кейсах результаты практически равны, это не является универсальным преимуществом». #2654 для малого числа токенов вычисляет и переносит только покрывающий префикс аккумулятора на 8 токенов; на H200 со 132 SM, при N, делящемся на 128, 133–225 выходных tile и K∈{7168, 7680, 8192}, открывает декодирование 40 токенов до M33–40.

Источники: https://github.com/tile-ai/TileOPs/pull/2681, https://github.com/tile-ai/TileOPs/pull/2672, https://github.com/tile-ai/TileOPs/pull/2669, https://github.com/tile-ai/TileOPs/pull/2654

3.5 TileOPs: NSA top-k на SM89 расширен до 2048 сжатых блоков и дополнены выходные слоты (#2674 / #2675 / #2673)

Три PR касаются одного и того же пути SM89. #2674 заставляет group48 переиспользовать существующую программу tiled top-k: один token — один tile (48 строк query head), поколоночная раскладка warp для GEMM; #2675 расширяет выбор top-k на SM89 с прежнего потолка до 1024–2048 сжатых блоков, покрывая head dim 64 / 128 / 256; #2673 исправляет заполнение выхода при недостатке кандидатов — каждый запрошенный слот top-k должен заполняться -1, ранее полоса сортировки на 32 слота оставляла в слотах 32–63 старые значения. Кроме того, #2684 добавляет верхнюю границу ёмкости сортировки NSA для упакованных запросов: по min(total_chunks, ceil(total_tokens / block_size)) с последующим выравниванием по tile скоринга.

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

3.6 TileOPs: однокернельный префилл GDN и разбиение количественной нормализации (#2682 / #2679)

#2682 нацелен на прямой проход GDN линейного внимания: его префилл на 79 строках ночного пакетного прогона медленнее FlashInfer (0.57×–0.98×); исправление заключается в том, чтобы в случае, когда «walk может заполнить SM», упаковать префилл в одно ядро. #2679 занимается FusedAddRMSNormPerTokenQuant: около 25 инструкций SASS на один элемент тратятся на поэлементную математику; вместо этого предлагается упаковать поэлементную математику и разделить широкие строки между кластерами. Оба изменения заменяют «последовательные накладные расходы внутри одного оператора» на «более широкий параллелизм и меньше инструкций», что совпадает по идее с той же группой мелкомасштабной настройки путей для SM89 в тот же день.

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

3.7 TileOPs: обратный проход GQA напрямую обслуживает сгруппированные KV head в warp-специализированном ядре (#2685)

Смержен 10-21 22:58. На SM90 при head dim 128 GQABwdOp ранее передавал сгруппированные KV head конвейеризованному ядру, обходя CTA по одному на каждый query head; теперь warp-специализированное ядро обратного прохода GQA напрямую принимает сгруппированные KV head, сокращая повторные пересылки и накладные расходы планирования. В тот же день #2687 исправляет проблему автозагрузки бэкенда устройства в регрессионном тесте переименования PrivateUse1, что относится к очистке тестовой инфраструктуры.

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

3.8 TileOPs на рассмотрении: сворачивание edge-оси в один запуск, сужение column tile для NSA top-k (#2689 / #2690)

В конце окна мониторинга два изменения всё ещё на рассмотрении. #2689 касается reduction: edge-ось сворачивается в один запуск, а хвостовая загрузка fold отправляется первой; #2690 использует более узкий column tile для NSA top-k на SM89. После слияния 20 изменений очередь на рассмотрении в TileOPs сократилась с 16 в прошлом выпуске до 9 — это единственный репозиторий в организации, где количество изменений на рассмотрении заметно снизилось за это окно мониторинга.

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

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

deepseek-ai/TileKernels (1991★) и QwenLM/FlashQLA (732★) в этом окне мониторинга не имеют коммитов, последние пуши обоих остановились на 10-14; версия FlashQLA по-прежнему v0.1.3 от 09-30. Затишье на стороне принимающих сторон длится уже неделю, динамика основной линии остаётся внутри организации tile-ai.

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

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

4.1 Руководство сообщества: обновлён китайский «TVM / TileLang Cookbook»

RuneFang/tvm_tilelang_cookbook (34★, репозиторий создан 7-21) имел пуш 10-21. Самоописание весьма чёткое: ориентировано на читателей, которые «немного пишут на Python, слышали о CUDA, но не читали код компилятора», и шаг за шагом ведёт от @tilelang.jit вплоть до .cubin. Такие китайские материалы, подробно разбирающие всю цепочку от DSL до продукта компиляции, по-прежнему редки и особенно полезны новичкам на новых бэкендах вроде Ascend и Hygon.

Источник: https://github.com/RuneFang/tvm_tilelang_cookbook

4.2 Новый репозиторий сообщества: создан tilelang-wrapper-for-sail, без описания

flytiger-eco/tilelang-wrapper-for-sail создан 10-21 09:04, без описания и README; название репозитория указывает на предоставление обёртки TileLang для SAIL. Такие небольшие репозитории «внутренний runtime производителя + оболочка TileLang» обычно являются аванпостом подключения нового оборудования, но содержимое пока не опубликовано, в этом выпуске только фиксируется факт.

Источник: https://github.com/flytiger-eco/tilelang-wrapper-for-sail

4.3 Инструмент сообщества: репозиторий удалённого агента автотюнинга продолжает получать пуши

cc-c122/TileLang-Remote-Autotuning-Agent (1★, репозиторий создан 7-01) имел пуш 10-21 22:57. В русле того же направления, что и TileFoundry с передачей планирования агенту, сторона сообщества также пытается сделать автотюнинг агентом с возможностью удалённого управления, а не перебором локальными скриптами.

Источник: https://github.com/cc-c122/TileLang-Remote-Autotuning-Agent

4.4 Документация и версии: семь автоматических обновлений сайта документации, VERSION основного репозитория по-прежнему 0.1.15

tilelang.github.io в окне мониторинга автоматически обновлялся семь раз (10-21 10:18, 11:26, 11:33, 13:54, 18:26, 19:23 и 10-22 01:23), все обновления запускались слиянием в основной репозиторий. VERSION основного репозитория по-прежнему 0.1.15 (этот релиз вышел 09-30), последняя веха в README по-прежнему датируется 09-30 — бэкенд Ascend 950; в TileOPs новых tag нет; TileFoundry сегодня добавил tag v0.0.3 — единственный новый релиз в окне. Хотя pushed_at у TileOPs-nightly изменился, в его ветке по умолчанию с 08-30 нет новых коммитов, изменения пришли из второстепенных веток и не учитываются как активность разработки.

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

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

Характер дефектов компилятора сместился от «неверных вычислений» к «неверному выбору цели». Из восьми слияний в основном репозитории за этот период четыре меняют то, «что компилятор думает, что делает»: CUDA target без arch молча откатывается к sm_50 (#2860), аннотация округления при cast молча отбрасывается в зависимости от ветки (#3350), пакетная редукция обрабатывает только первый выходной блок (#3070), планирование конвейера теряет транзитивные зависимости (#2822). Их общая черта — отсутствие ошибок: они дают недостоверные результаты лишь при определённых размерах или формах циклов — такое сложнее локализовать, чем падение, и это показывает, что проект ставит молчаливую корректность выше расширения функциональности.

Узкое место линии Ascend сместилось к модели памяти и границам рантайма. Плотный межстрочный интервал менее 32B невозможно выразить через DMA, остаётся только явное прерывание (#1795); конкурентность кэша переведена на атомарную публикацию с добавлением восстановления после удаления (#1903); целочисленная знаковость PTO дополняется через слой-обёртку (#3582). Все три изменения — не «добавление операторов», а «обеспечение корректности существующих операторов в граничных условиях»; вместе со слиянием эвристического выравнивания DeepGEMM сторона Ascend переходит от покрытия возможностей к инженерной надёжности.

Аналитические агентства начинают использовать TileLang как критерий доступности Ascend. SemiAnalysis определяет степень завершённости поддержки Ascend в DeepSeek через «из шести репозиториев только TileLang тогда был работоспособен» — это лучше любого отдельного технического блога показывает место TileLang в нарративе отечественных вычислительных мощностей: из экспериментального DSL он стал слоем по умолчанию. Тот же отчёт даёт и обратный сигнал — оптимизация сосредоточена на суперузлах, поддержка x86-хостов остаётся неполной, универсальность отстаёт от производительности.

Центр тяжести стека внимания смещается от прямого прохода инференса к обратному проходу обучения. TileOPs за один день дополнил op-слой обратного прохода префилла MLA, warp-специализированный kernel для SM90 и оптимизацию порядка выборки данных, а вместе с обработкой сгруппированных KV head в обратном проходе GQA это полностью совпадает с направлением дополнения обратных операторов на стороне Ascend за последний месяц (indexer_bwd, обратный SiTU-GLU, обратный WY fast). Покрытие TileLang на стороне обучения систематически дополняется.

SM89 переходит от «работает» к «тюнингу». Тайлы FP8 BMM, верхняя граница M для W4A16, постраничный DSA, top-k NSA — четыре направления в один день подверглись изменениям, и все сопровождаются измерениями на одной карте и оговорками вроде «не является универсальным преимуществом». Потребительские сценарии и инференс на одной карте стали одной из целей производительности, при этом самоограничение авторов в отношении методики бенчмарков заметно ужесточилось.

Разделение труда между планированием и корректностью в TileFoundry и основном репозитории становится ясным. Основной репозиторий продолжает ужесточать семантику и корректность, а TileFoundry 0.0.3 передаёт «выбор инструкции» агенту, оставляя «гарантию семантики» компилятору, и приводит воспроизводимые цифры на двух примерах — H200 и Apple AMX. Если это разделение стабилизируется, в экосистеме TileLang возникнет двухуровневая структура: «языковой слой отвечает за корректность, агентный слой борется за производительность».

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

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

Источники Результаты проверки
tile-ai/tilelang (main) 8 коммитов и 8 слияний в окне мониторинга; VERSION по-прежнему 0.1.15; на рассмотрении 163, из них 21 обновлён в окне
tile-ai/TileOPs (main) 20 коммитов и 20 слияний в окне мониторинга; на рассмотрении 9, 3 обновлены в окне
tile-ai/tilelang-ascend (ascendc_pto) 3 коммита и 3 слияния в окне мониторинга; на рассмотрении 235, 13 обновлены в окне
tile-ai/tilelang-hygon (main / feat-hcu-smls-b-vgpr) в main нет коммитов; в ветке feature 3 пуша (10-21 11:04, 15:27); на рассмотрении 3
tile-ai/TileFoundry (main) новых коммитов нет; v0.0.3 помечен тегом и выпущен 10-21 10:26; на рассмотрении 1
tile-ai/tilelang-metax (dev) / musa / mlir-ascend во всех трёх репозиториях в окне нет коммитов (последние соответственно 10-16, 09-30, 10-19), новых тегов нет
tile-ai/TileRT / tilescale / DeepStack во всех трёх репозиториях в окне нет коммитов (последние соответственно 09-24, 08-25, 10-05)
deepseek-ai/TileKernels / QwenLM/FlashQLA коммитов нет, последние пуши остановились на 10-14
tilelang.github.io / TileOPs.github.io 7 и 1 автоматическое обновление соответственно, все инициированы слияниями в основном репозитории
TileOPs-nightly в ветке по умолчанию нет новых коммитов с 08-30, изменение pushed_at вызвано ветками, отличными от ветки по умолчанию
Google News (несколько наборов на китайском и английском) / Hacker News / arXiv три отраслевых материала в окне посвящены взаимодействию Huawei-DeepSeek, только статья SemiAnalysis содержит отдельную главу о TileLang; последняя релевантная работа на arXiv по-прежнему от 07-24

Дополнительные пояснения: пуши по направлению Hygon находятся не в ветке по умолчанию; метод определения — сначала получить поток коммитов ветки feat/, затем сверить ref и before→head события PushEvent через интерфейс событий репозитория; изменение pushed_at в TileOPs-nightly также исключено этим методом. Количество на рассмотрении взято как общее число результатов поиска is:open по каждому репозиторию, количество обновлённых в окне — как общее число результатов поиска updated:>=начало окна. Коммиты, попавшие в интервал половины дня до начала окна (10-21 07:00), покрываются предыдущим отчётом и в этом выпуске не учитываются повторно.

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

ek-ai/TileKernels>