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

Источники: порепозиторная проверка организации GitHub tile-ai (tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-mlir-ascend / tilelang-hygon / tilelang-metax / tilelang-musa / TileRT / tilescale / DeepStack / TileOPs-nightly / сайт документации) + поиск фактических данных по смерженным и находящимся на рассмотрении изменениям (смержено 35, обновления на рассмотрении 82) + полная проверка новых issue (11) + поштучная проверка коммитов в основную ветку (tilelang 7, TileOPs 15, TileFoundry 3, tilelang-ascend 5, tilelang-hygon 2, tilelang-mlir-ascend 1, сайт документации 4) + проверка страниц релизов и тегов (6 репозиториев) + проверка сторонних экосистемных репозиториев (FlashQLA, TileKernels, три репозитория исследовательской цепочки RISC-V, PTO-ISA/tilelangGym) + запросы Google News на китайском и английском (через прокси) + Hacker News + arXiv.

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

  • Ключевое за сегодня: основная линия — откат точного подсчёта потоков ThreadSync при миграции на TIRX (#3733), а также новый тикет отслеживания, признающий отзыв улучшений корректности; TileOPs — 13 перформанс-мержей за один день; фиксированные накладные расходы хоста Ascend сокращены с примерно 450 мкс до примерно 106 мкс; за 72 часа 35 мержей, 82 обновления на рассмотрении, 11 новых issue
    1. Основные достижения проекта
      • 1.1 Основная линия: откат точного подсчёта потоков ThreadSync ради миграции на TIRX (#3733) (01.11)
      • 1.2 Основная линия: межwarp-ное полное суммирование CUDA объединено в один обмен через разделяемую память, для 256 потоков число барьеров снижено с 6 до 2 (#3715) (30.10)
      • 1.3 Основная линия: первый этап векторизации T.Parallel для бэкенда PTO, поэлементные операторы используют инструкции VMI (#3512) (30.10)
      • 1.4 Основная линия: буфер scale в DeepGEMM для Ascend переведён на многобуферность со счётчиком записей, загрузка данных может быть поднята перед условной загрузкой (#3629) (30.10)
      • 1.5 Основная линия: два исправления смержены — скалярный GEMM на CPU очищает только целевую область, легализация отрицательных индексов пропускает недостижимые ветви (#3709 / #3712) (30.10)
      • 1.6 Основная линия: на уровне языка удалены два неиспользуемых Ref-заглушки в proxy (#3722) (31.10)
      • 1.7 Новые дефекты и на рассмотрении: TMEM молча теряет результаты, в ROCm не находятся символы, TF32 округление к чётному (#3721 / #3711 / #3714) (30.10 ~ 31.10)
    1. Адаптация под бэкенды
      • 2.1 Ascend: Cython wrapper лениво конструирует TVM Analyzer, одиночный вызов на 910B3 сокращён с примерно 450 мкс до примерно 106 мкс (#1948) (30.10)
      • 2.2 Ascend: RFC предлагает заменить реализацию синхронизации на ядро основной линии плюс отдельный сервис доказательств Z3 (#1971) (30.10)
      • 2.3 Ascend на рассмотрении: полный перенос TileLang-Ascend на мультибэкендную архитектуру основной линии (#1975) (31.10)
      • 2.4 Ascend на рассмотрении: пример fp8_mqa_logits для kpool в GLM-5.3, двойной kernel с разделением задач AIV/Cube (#1966) (31.10)
      • 2.5 Ascend: смержены четыре проверки фактов документации и API, clamp реально работает на PTO (#1899 / #1893 / #1620 / #1600) (30.10)
      • 2.6 Ascend: ежедневный тест — все 2970 тестов пройдены, 01.11 был один сбой workflow (#1977 / #1976) (31.10 ~ 02.11)
      • 2.7 Hygon: смержены исправления границ MLS и постоянной синхронизации WASP, совместимость с gfx948 на рассмотрении (#31 / #32) (30.10)
      • 2.8 mlir-ascend: CI очищен от ссылок на внутренний кэш-сервис, workflow Codex на рассмотрении (#206 / #208) (31.10)
      • 2.9 Moore Threads и MetaX: за окно мониторинга пушей не было (02.11)
    1. Экосистема и потребители
      • 3.1 TileOPs: 15 мержей за 72 часа, 13 из них пришлись на 01.11, почти каждый с таблицей измерений и порогом регрессии по чек-листу (01.11)
      • 3.2 TileOPs: плотный forward GQA выбирает программу decode или prefill по одному измеренному правилу, сожаление выбора снижено с 1.818 до 1.017 (#2808) (01.11)
      • 3.3 TileOPs: пейджинговый MHA decode расширен на каждый head dim, укладывающийся в разделяемую память, пейджинговый forward GQA избавлен от старого ядра decode (#2807 / #2797) (01.11)
      • 3.4 TileOPs: coop2-производитель сведён к одному warp, добавлен транспонированный блок ширины 192, часть W4A8 переведена с одной волны на один CTA на tile (#2806 / #2805) (01.11)
      • 3.5 TileOPs: кэш компиляции ключуется по цели сборки, слой операторов единообразно проверяет перекрытие выхода и входа (#2802 / #2801) (01.11)
      • 3.6 TileOPs, исправляемые дефекты: четыре оператора Norm теряют точность на строках вдали от нуля, переполнение около 1e36 превращается в NaN (#2812) (01.11)
      • 3.7 TileFoundry: lowering сводит issuing mesh и разбиение на окна в одно решение, 83 фикстуры планирования побайтово идентичны (#280 / #282) (31.10 ~ 01.11)
      • 3.8 TileFoundry: на сайте опубликована статья о планировании версии 0.3 «Scheduling, Designed for Agents» (#4 / #281) (31.10 ~ 01.11)
      • 3.9 Потребители: FlashQLA смержил Dual-CP forward и backward для SM120, на 2×RTX 5090 сравнение с Inter-CP для последовательностей от 32K до 128K (#51) (30.10)
      • 3.10 Потребители: TileKernels без пушей с 14.10 (02.11)
    1. Сообщество, учебные материалы и мероприятия
      • 4.1 Формируется исследовательская цепочка RISC-V: три репозитория преобразуют Numba CUDA Python в TileLang, затем ограниченный PrimFunc опускается в RVV C для одного hart (31.10)
      • 4.2 PTO-ISA/til

elangGym выделен из PTO-ISA/tilelang в отдельный репозиторий, основным регрессионным шлюзом является путь Ascend SimtVF (10-30)

  • 4.3 Новостная сторона: в Google News (запросы на китайском и английском), Hacker News и arXiv существенных совпадений в окне мониторинга нет (11-02)
    1. Наблюдения за трендами
  • Приложение: материалы и пояснения по проверке

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

Самое весомое действие этого окна — это вычитание. Основная ветка целиком откатывает (#3733) точный подсчёт потоков (#3291), влитый три недели назад, и причина написана прямо: тот набор изменений переносил дорогостоящий запрос подсчёта перед проверкой барьера полного блока, из-за чего даже при достижении барьера всеми потоками приходилось прогонять доказательство, текущий бюджет решателя TVM исчерпывался и возвращался частичный подсчёт, и в результате изначально легальные вспомогательные kernel remap и sequence просто не компилировались с ошибкой «Cannot prove a thread rectangle». Откат выполнен одновременно для двух реализаций ThreadSync, удалены тесты, введённые #3291, и добавлен случай редукции «цикл пуст для большинства потоков, но все 256 потоков всё равно достигают барьера»; улучшения корректности — коммерческий guard, точная граница участвующих потоков, не прямоугольные множества потоков — также отозваны, а новый #3732 оставлен как тикет для повторной оптимизации после миграции. Это явно помеченная как временная мера уступка ради миграции: TileLang готовит миграцию на TIRX, чтобы избавиться от текущей зависимости от форка TVM, а сопутствующее предложение по оптимизации форка TVM уже закрыто.

Вторая линия — Ascend наращивает фундамент. #1948 исправляет на стороне хоста фиксированные накладные расходы, не зависящие от масштаба оператора: Cython wrapper при каждом вызове безусловно конструировал TVM Analyzer, что на 910B3 измеренно составляло около 0,28 миллисекунды за вызов, тогда как он нужен только когда форма является символьным выражением; после перехода на ленивое конструирование vec add для одного блока снизился с 442–467 микросекунд до 105,1 микросекунды, для 32 блоков — с 441–493 микросекунд до 108,0 микросекунды, сжатие kpool — с 447–450 микросекунд до 111,3 микросекунды, а служащий нижней границей голый вызов ctypes около 12 микросекунд не изменился. В тот же период RFC #1971 даёт маршрут замены реализации синхронизации: портировать ядро синхронизации основной ветки плюс отдельный сервис доказательств Z3, подключить через слой адаптации старого TIR к существующему конвейеру компиляции, сохранить TVM, планирование, планирование памяти, AscendC/PTO codegen и ABI выполнения, а после прохождения приёмки переключить реализацию по умолчанию и удалить старый алгоритм — это две точки одной и той же задачи, что и #1975, переносящий весь репозиторий Ascend на многобэкендовую архитектуру основной ветки.

Третья линия — в библиотеке операторов. TileOPs за один день 11-01 влила 13 коммитов, почти все — патчи производительности, и каждый приложен таблицей измерений на H200 и порогом регрессии manifest. Среди них #2808 сводит выбор программы для плотного прямого прохода GQA к одному правилу, сокращая сожаление о выборе относительно самого быстрого кандидата с 1,818 до 1,017 (2–128 позиций query, 586 вызовов), и исправляет реальную ошибку основной ветки: масштаб score после умножения на log2(e) терялся в FP32, и выход становился NaN или неверным значением.

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

1.1 Основная ветка: откат точного подсчёта потоков в ThreadSync, уступая место миграции на TIRX (#3733)

Влито 11-01 14:50. Откатываемый #3291 переносил точный подсчёт потоков перед проверкой барьера полного блока и добавлял прямоугольную сертификацию, ценой чего становился прогон CountSatisfyingValues даже при достижении барьера всеми потоками; после исчерпания бюджета решателя на стороне TVM возвращался частичный подсчёт, и вспомогательные kernel remap и sequence сразу падали с ошибкой «[ThreadSync] Cannot prove a thread rectangle». Регрессия воспроизводится на основной ветке cc903dcb9 с неизменённым пином TVM 4211874e9e, а новый случай сводит её к «цикл пуст для большинства потоков, но все 256 потоков всё равно достигают барьера». Откат охватывает одновременно обе реализации — общую и Ascend, — восстанавливая быстрый путь полного блока и прежнее поведение частичного подсчёта; в тексте прямо сказано, что это временное смягчение на период миграции на TIRX, предложение по оптимизации на стороне форка TVM tile-ai/tvm#89 уже закрыто, а последующая оптимизация перенесена в #3732.

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

1.2 Основная ветка: полная межуarpовая редукция в CUDA объединена в один обмен через разделяемую память, для 256 потоков число барьеров снижено с 6 до 2 (#3715)

Влито 10-30 20:26. Старая реализация в заголовочном файле редукции CUDA reduce.h рекурсивно шла от threads/2 до scale, и на каждом шаге с offset не меньше 32 выполняла полный обмен записью-чтением через разделяемую память с двумя барьерами: для блока из 256 потоков три шага (128, 64, 32) в сумме давали 3 обмена через разделяемую память и 6 барьеров, для блока из 512 потоков — 4 обмена и 8 барьеров. Изменение сжимает все межуровневые шаги в один обмен (одна операция st.shared на элемент, всего 2 барьера): сначала выполняется участок с scale меньше 32 внутри warp, затем не более cross_radix lane в каждой группе внутри warp совместно сворачивают, при этом давление на регистры ограничено сверху. Изменены оба пути — скалярный и пакетный.

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

1.3 Основная ветка: первый этап векторизации T.Parallel в бэкенде PTO реализован, поэлементные операторы идут через инструкции VMI (#3512)

Совместимо с 10-30 01:29. Цель — позволить пользователю писать в T.SimdVF(lanes=L) скалярно-параллельный поэлементный T.Parallel, который компилятор векторизует в инструкции PTO VMI, а затем через PTOAS исполняет на Ascend 950. Разделение обязанностей написано чётко: сторона TileLang отвечает за вывод логической раскладки lane, конвейер проходов векторизации, генерацию инструкций VMI и логическое отображение Fragment; сторона PTOAS отвечает за физическое разбиение регистров, конкретизацию масок, распределение буферов Fragment и векторных регистров, вставку mem_bar. Дорожная карта разбита на пять этапов; данный коммит — это M1, базовый каркас (поэлементные вычисления, охватывающие два случая: делится ли extent на lanes нацело), далее последовательно идут M2 управляющие потоки, M3 reduce, M4 планирование ёмкости UB и распределение регистров, M5 несмежные обращения и более широкие типы данных PACK.

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

1.4 Основная ветка: буфер scale в Ascend DeepGEMM переведён на многобуферность со счётчиком записей, загрузка данных может быть вынесена до условной загрузки (#3629)

Совместимо с 10-30 05:48. DeepGEMM повторно использует scale factor между несколькими K-блоками; данный коммит добавляет явный режим многобуферности со счётчиком записей, позволяющий версиям буфера scale продвигаться по фактическим записям: в каждой итерации цикла owner версия продвигается только один раз — на первой фактически выполняемой задаче записи, последующие записи используют эту версию, итерации без записи сохраняют старое значение, поддерживаются несколько непересекающихся owner и общие часы. Благодаря этому загрузка данных может быть выстроена перед условной загрузкой scale; одновременно K-цикл DeepGEMM развёрнут, и в L1 и L0 данные сначала загружаются, а затем scale, а также добавлены случаи защищённых обновлений, вложенных обновлений, сохраняемых чтений, flag и повторного использования буферов, а также FP4 с K-мажорным и MN-мажорным порядком с хвостовыми блоками и кольцевым оборачиванием. Этот режим требует явного числа версий и предполагает не более одного обновления на итерацию owner.

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

1.5 Основная ветка: совмещены два исправления — скалярный CPU- GEMM очищает только целевую область, легализация отрицательных индексов пропускает недостижимые ветви (#3709 / #3712)

Оба совмещены 10-30 и относятся к закрытию пунктов, находившихся на рассмотрении в прошлом выпуске. #3709 заставляет скалярный CPU- GEMM очищать только целевую область, тогда как ранее очищались также диапазоны, не относящиеся к текущему выводу; #3712 заставляет LegalizeNegativeIndex пропускать недостижимые ветви, избегая переписывания отрицательных индексов в неисполняемом коде. Эти два пункта относятся к той же партии очистки дефектов, что и #3708, #3710 из прошлого отчёта, причём #3708 (изменение разрядности на int64) и #3710 (обход двух уровней внутрипроцессных кэшей при отключении кэша) всё ещё на рассмотрении.

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

1.6 Основная ветка: на языковом уровне удалены два неиспользуемых заполнителя Ref в proxy (#3722)

Совместимо с 10-31 17:31. В tilelang/language/proxy.py одновременно существовали два Ref: один — заглушка tirx.Var для проверки типов, другой — пустой runtime-класс, и у обоих в дереве не было потребителей; внешний T.Ref фактически происходит из eager.builder, и именно тамошняя реализация является настоящим телом макроса ссылочного параметра. Данный коммит удаляет оба заполнителя, переменную типа _T и ставшие ненужными импорты Generic, TypeVar, сохраняя T.Ref и поведение макроса ссылки, и убирает лишь неиспользуемый путь прямого импорта tilelang.language.proxy.Ref. 179 фронтенд-тестов конструкций прошли, 6 пропущены, проверка идентичности во время выполнения подтверждает, что T.Ref по-прежнему является eager.builder.Ref.

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

1.7 Новые дефекты и находящиеся на рассмотрении: тихая потеря результатов TMEM, ненайденный символ ROCm, округление TF32 к ближайшему чётному (#3721 / #3711 / #3714)

#3721 — наиболее тревожный пункт данного окна: при выполнении T.gemm над TMEM на sm_100a, когда threads/32 не делит N нацело, компиляция проходит, но в PTX нет ни одной tcgen05.mma, T.copy(Ct, Cl) читает TMEM, в который никто не писал, и произведение тихо отбрасывается; в воспроизведении при threads=256 получается одна инструкция, при threads=384 — ноль, автор подозревает, что src/cuda/op/gemm.cc при возврате разбиения {1, num_warps} для TCGEN05 упускает проверку покрытия. #3711 сообщает, что под PyTorch wheel для ROCm 7.14 и 10.0 ядро выдаёт «named symbol not found». #3714 на рассмотрении — добавляет округление к ближайшему чётному для преобразования fp32 в tf32 (cvt.rn.tf32.f32 и насыщенная форма), ненасыщенная требует SM90 и выше, насыщенная — SM100 и выше, закрывая соответствующий запрос функции #3713.

Источники: https://github.com/tile-ai/tilelang/issues/3721, https://github.com/tile-ai/tilelang/issues/3711, https://github.com/tile-ai/tilelang/pull/3714

2. Адаптация под множество бэкендов (Ascend / Hygon / Moore Threads / MetaX)

2.1 Ascend: ленивое создание TVM Analyzer в Cython wrapper, одиночный вызов на 910B3 снижен с примерно 450 микросекунд до примерно 106 микросекунд (#1948)

Смержен 10-30 в 07:44. CythonKernelWrapper.forward() ранее при каждом вызове безусловно создавал объект TVM Analyzer(), что на 910B3 давало фиксированные накладные расходы на хосте около 0,28 миллисекунды, хотя этот объект нужен только тогда, когда форма result или workspace является символьным выражением; для ядер со статической формой он вообще не используется, но расходы всё равно оплачивались. Изменение откладывает создание до ветки tir.PrimExpr и переиспользует объект по мере необходимости. Измерения: vec add с одним блоком снизился с 442–467 микросекунд до 105,1 микросекунды, с 32 блоками — с 441–493 микросекунд до 108,0 микросекунды, сжатие kpool — с 447–450 микросекунд до 111,3 микросекунды, прямой вызов ctypes — с примерно 12,5 микросекунды до примерно 12 микросекунд. Добавлены два регрессионных теста, проверяющих соответственно, что статический путь не создаёт объект, а символьный путь создаёт его по требованию.

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

2.2 Ascend: RFC предлагает заменить синхронную реализацию на ядро из основной ветки плюс отдельный сервис доказательств Z3 (#1971)

Данный RFC не постулирует вывод об ускорении, а лишь фиксирует базу сравнения: репозиторий Ascend на TileLang 43b0b69 и TVM c2921fd, основой IR является старый tvm.tir; основная ветка на TileLang 1c0f402 и TVM 4211874e9e, основой IR уже являются tvm.tirx и tvm.s_tir, объектная и регистрационная системы также заменены на новый FFI и средства рефлексии. Поэтому миграцию нельзя свести к копированию одного файла InsertSync. Предлагаемый маршрут — портировать ядро синхронизации из основной ветки и отдельный сервис доказательств Z3, подключить их к существующему конвейеру компиляции через слой адаптации старого TIR, сохранить существующие TVM, планирование, планирование памяти, codegen AscendC/PTO и исполняемый ABI, а после приёмки по охвату поддержки, корректности, производительности и установке переключить реализацию по умолчанию и удалить старый алгоритм синхронизации. Сопутствующий PR — #1965.

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

2.3 Ascend на рассмотрении: полный перенос TileLang-Ascend на многобэкендовую архитектуру основной ветки (#1975)

Предложен 10-31 и постоянно обновляется; цель — согласно предложению #1911 перенести TileLang-Ascend на многобэкендовую архитектуру основной ветки TileLang; в описании приведены два сопоставления различий: одно с веткой ascendc_pto, другое с основной веткой. Этот пункт и RFC из 2.2 — два уровня одного пути сходимости: один меняет алгоритм синхронизации, другой — каркас репозитория.

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

2.4 Ascend на рассмотрении: пример fp8_mqa_logits для kpool GLM-5.3, два ядра и разделение труда AIV/Cube (#1966)

Сделана версия для Ascend по примеру fp8_mqa_logits для GLM-5.3 из каталога kpool основного репозитория. Семантика оператора — пулинговое историческое оценивание indexer разреженного внимания GLM-5.3-Flash: для каждой строки query в пределах допустимого интервала pool по таблице страниц читается ключ E4M3 из страницы 64-pool (байты в uint8 плюс масштаб fp32), после деквантизации вычисляются оценки MQA для 32 голов indexer, умножаются на масштаб per-pool, проходят через ReLU-гейт и взвешенное по головам сведение в один FP32 logit на каждый pool; вне диапазона или для недопустимых страниц выдаётся минус бесконечность. Реализация состоит из двух ядер: ядро маски по тайлу 64-pool генерирует маску 0 или минус бесконечность для каждой дорожки, заменяя предварительное заполнение на хосте; основное ядро для каждой строки и каждой страницы выделяет один блок, два AIV обрабатывают по половине страницы в 32-pool, Cube читает всю страницу и выполняет один GEMM 64×32×128. Число голов 32, размерность 128, размер страницы 64.

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

2.5 Ascend: смержены четыре проверки фактов документации и API, clamp реально работает на PTO (#1899 / #1893 / #1620 / #1600)

Все четыре изменения влиты с 10-30, относятся к продолжающейся серии фактической проверки API: #1899 проверяет документацию T.tile.bitwise_xor (повторное поднятие #1681), #1893 проверяет T.tile.bitwise_rshift (повторное поднятие #1676), #1620 обеспечивает реальное срабатывание T.tile.clamp на PTO с синхронизацией документации, а также добавляет тесты и документацию API, #1600 обновляет docstring axpy, усиливает тесты и дополняет документацию API. В той же серии на рассмотрении находятся #1968, #1967, #1974 — соответственно сокращение вдвое счётчиков sigmoid и round, кэширование разбора символьных форм, исправление ошибочного сокращения вдвое размеров tile при построчном разбиении целыми блоками.

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

2.6 Ascend: ежедневный тест 2970 случаев — все пройдены, 11-01 был один сбой workflow (#1977 / #1976)

Запись ежедневного теста от 11-02 06:42: 2970 пунктов, пройдено 2970, провалено 0, проходимость 100%; выпуск от 10-31 06:56 также 2970 пунктов — все пройдены. Между ними выпуск от 11-01 05:08 помечен как workflow test job failure, что совпало по времени с миграционными изменениями #1975 и RFC #1971 в данном окне мониторинга — это флуктуация, требующая внимания, а не стабильная регрессия.

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

2.7 Hygon: влиты исправления границ MLS и непрерывной синхронизации WASP, совместимость с gfx948 на рассмотрении (#31 / #32)

#31 влит 10-30 10:56, три вещи: улучшено динамическое выведение границ MLS и обработка stride в MatrixStore; исправлено опускание clear_accum в MMAC и blockscaled GEMM; в непрерывном пути WASP использованы потребитель M-разбиения и независимые барьеры A/B с сохранением ротации slot и phase и добавлением синхронизации для хвостового потребителя. #32 на рассмотрении: gfx948 трактуется как совместимый с gfx946, чтобы существующие инструкции напрямую работали на новой модели. Сообщённый в предыдущем окне мониторинга LDS swizzle для обычного копирования (#30) был реализован в тот же период.

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

2.8 mlir-ascend: CI очищен от ссылок на внутренний кэш-сервис, workflow Codex на рассмотрении (#206 / #208)

#206 влит 10-31 07:24, из CI удалены ссылки на внутренний кэш-сервис — очистка в целях воспроизводимости. #208 на рассмотрении: для Operator Harness добавляется первая партия workflow Codex: инициализация проекта через итеративные вопросы, подключение существующего репозитория операторов и восстановление разработки, тюнинга или анализа из персистентного контекста; три skill совместно используют редактируемые Markdown-записи в .workspace/, при передаче задач различаются «текущий кандидат» и «проверенный лучший кандидат», сохраняются неудачные и откаченные эксперименты. Автор явно указывает, что реальных измерений на NPU не было; адаптация исполнения на NPU и распространение через npm отнесены к последующей работе.

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

2.9 Moore Threads и MetaX: в окне нет пушей (11-02)

Последний пуш tilelang-metax остановился на 10-26, tilelang-musa — на 09-30; в обоих репозиториях за 72 часа данного окна мониторинга нет коммитов. Это означает, что динамика бэкендов в этом выпуске сосредоточена на трёх сторонах: Ascend, Hygon и mlir-ascend.

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

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

3.1 TileOPs: 15 вливаний за 72 часа, 13 из них пришлись на 11-01, почти каждое сопровождается таблицей фактических измерений и порогом регрессии по чек-листу (11-01)

После релиза v0.1.0 TileOPs за один день 11-01 принял 13 коммитов, плюс #2794 от 10-30 и #2795 от 10-31 — всего 15 коммитов за 72 часа. Изменения почти все относятся к категории perf и охватывают внимание, GEMM, поэлементные операции, редукции и MoE; для каждого коммита приводятся две группы времён на H200 — busy и back-to-back, либо сравнение с vLLM и FlashInfer, и заявлено, что в manifest нет регрессий по уже измеренным строкам. Такой способ «принятия с таблицами» показывает, что библиотека операторов уже вошла в стадию построчной калибровки по shape, а не дописывания функциональности.

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

3.2 TileOPs: плотный GQA-форвард выбирает decode- или prefill-программу по одному измеренному правилу, regret выбора снижен с 1.818 до 1.017 (#2808)

Принят 11-01 в 13:04. GQADenseFwdOp и GQADenseWithLSEFwdOp теперь определяют, какой набор программ использовать, через одно и то же правило AttentionCall.dense_decode_program; на входе — длины query и key, batch, группировка голов, head dim, RoPE и число SM: от 2 до 64 позиций query идут в warp-специализированный decode, где каждая позиция внутри группы KV-голов упаковывается в блок из 64 строк; при единственной позиции query — разбиение на 32 сегмента для заполнения одной волны либо batch-1 программа с 1 последовательностью и 1–2 KV-головами. Regret выбора (геометрическое среднее и максимум) относительно самой быстрой доступной программы: для одной позиции 1.068 и 1.87 снижены до 1.018 и 1.69; для 2–128 позиций 1.818 и 25.2 снижены до 1.017 и 1.81. Одновременно исправлены две ошибки основной ветки: умножение score scale на log2(e) исчезало в FP32, вызывая NaN или неверные значения; вызов D256 с несколькими позициями на SM90 падал при запуске из-за нехватки разделяемой памяти.

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

3.3 TileOPs: постраничный MHA decode расширен на каждый head dim, помещающийся в разделяемую память; из постраничного GQA-форварда удалено старое decode-ядро (#2807 / #2797)

#2807 принят 11-01 в 12:30: постраничный MHA decode расширен на каждый head dim при условии, что «блок помещается в разделяемую память»; ранее при большем head dim этот путь просто не использовался. #2797 принят в тот же день в 03:54: из GQAPagedFwdOp удалено старое paged decode kernel, а верхняя граница постраничного MHA decode урезана до head dim, для которого удаётся собрать код, — то есть отозвана возможность выхода за границы, оставленная предыдущей версией. Вместе эти два коммита — завершающий шаг «сначала расширить возможности, затем урезать избыточность».

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

3.4 TileOPs: coop2-производитель сведён к одному warp, добавлен транспонированный блок ширины 192, часть волн W4A8 переведена на схему один CTA — один tile (#2806 / #2805)

#2806 принят 11-01 в 12:02: coop2-производитель Hopper GEMM сведён с нескольких warp к одному загрузочному warp, а при допускающей разделяемой памяти трёхуровневой конвейеризации добавлен равномерный транспонированный блок ширины 192; на H200 wgrad-q-b снижен с 384.9 микросекунд до 382.3, mlp-up fp16 — с 317.9 микросекунд до 311.9 (back-to-back соответственно с 458.5 до 449.6 и с 369.0 до 367.3); автор заявляет отсутствие регрессий по измеренным строкам GEMM; коммит — часть цели #2791, подтверждение на двух картах всё ещё не завершено. #2805 принят в тот же день в 10:36: для части волн W4A8 prefill, заполняющих не менее трёх четвертей SM, сетка переведена со stream-K на схему один CTA — один tile; llama-8b-up-b33 снижен с 23.52 микросекунд до 17.98 (vLLM 20.42), b128 — с 31.33 до 23.39 (vLLM 24.05); во внесписочном сканировании 44 комбинаций N/K для M от 33 до 128 изменённая форма планирования быстрее на 5%–36%, медиана 18%, ни один случай не замедлился.

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

3.5 TileOPs: кэш компиляции ключуется по цели сборки, слой операторов единообразно проверяет перекрытие выхода и входов (#2802 / #2801)

#2802 объединён 01.11 в 11:31, исправляет проблему повторного использования кэша компилятора между устройствами: записи сборщика и Op используют в качестве ключа цель сборки, ленивая компиляция и автотюнинг остаются на ожидаемом устройстве, приёмка — 99 целей изоляции, тесты GEMM поставщиков и групповых GEMM.#2801 объединён в тот же день в 10:33, проверка «выход не должен разделять байты с любым другим вызывающим тензором» поднята на уровень оператора, исключения объявляются по оператору через поле overlaps манифеста, удалены четыре проверки на стороне ядра kernel: MoE unpermute, MoE групповой GEMM, шаблон постоянного GEMM и др.; попутно исправлена гонка в основной ветке: IndexedExpertMLPFwdOp разрешает перекрытие out с topk_weights, тогда как последний всё ещё читается при редукции во время записи. Ведущее время MoEPostPermuteFwdOp с out снизилось с 51,48 микросекунд до 39,33.

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

3.6 TileOPs исправляет дефекты: четыре оператора Norm теряют точность на строках вдали от нуля, около 1e36 переполняется в NaN (#2812)

Открыто 01.11. LayerNorm, AdaLayerNorm (включая -Zero), GroupNorm и BatchNorm в настоящее время напрямую суммируют исходные строки в fp32 для получения среднего и дисперсии (BatchNorm использует E[x²] минус квадрат среднего), это тот же класс дефектов, который был исправлен в #2773 для FusedAddLayerNorm и исправляется в #2811 для тренировочного ядра InstanceNorm. Проявление: строка fp32 с константой 10000.1 после нормализации даёт на выходе примерно плюс-минус weight вместо bias, строка из 10000 плюс случайное число не совпадает с эталоном float64 в допуске fp32; строка вблизи 1e36 переполняется при суммировании, в fp32 и bf16 все выходы NaN. Затронутые пути охватывают fragment, обход через разделяемую память, одну строку на warp, блочные и потоковые реализации и др.

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

3.7 TileFoundry: lowering сводит issuing mesh и разбиение окна в одно решение, 83 фикстуры планирования побайтово идентичны (#280 / #282)

#280 объединён 31.10 в 00:47, решение «что испускать при каждом вызове» сведено в одно место; #282 объединён 01.11 в 11:47, продолжает сводить в одно место «каким mesh испускать» и «как резать окно»: в analysis/emission.py добавлены copy_mesh, issue_owner, issue_mesh, copy_mesh и issue_owner используют общего помощника, их единственное различие (как читать Partial) остаётся в одном месте; три дублирующихся разбиения Slice унифицированы помощником _cut. Приёмка выполняется побайтовым сравнением: из 83 фикстур планирования 75 выводов finalize полностью совпадают с основной веткой, остальные 8 в основной ветке также падают; раздел 7.1 docs/spec/passes.md вносит это множество решений mesh в спецификацию.

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

3.8 TileFoundry: сайт публикует статью о планировании версии 0.3 «Scheduling, Designed for Agents» (#4 / #281)

Репозиторий сайта объединил #4 30.10, опубликована статья версии 0.3 «Scheduling, Designed for Agents», одновременно добавлены структура каталогов статей и иллюстраций, поддержка MathJax для формул, подписи к рисункам и стили таблиц, а также якоря для каждого графика; карточка Blog на главной странице теперь открывает индекс блога, индекс перечисляет две статьи; upstream-ref не изменён, сайт по-прежнему рендерит v0.0.3. На стороне репозитория #281 от 31.10 анонсирует эту статью в README. Для линии TileFoundry это шаг переноса «планирования, ориентированного на agent» из кода и примечаний к релизу в публичное изложение.

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

3.9 Со стороны внедряющих: FlashQLA объединяет прямой и обратный проход Dual-CP для SM120, сравнение с Inter-CP на 2× RTX 5090 для последовательностей от 32K до 128K (#51)

30-10 03:02 объединено. Изменения добавляют для SM120 прямой и обратный проход Dual-CP: прямой проход повторно использует существующее ядро state-scan для агрегации состояния на уровне карты, упорядоченное произведение M вычисляется блоками по 32 столбца, чтобы уложиться в разделяемую память; обратный проход использует транспонированную M для распространения состояния в обратном порядке и повторно использует матрицы M, закэшированные прямым проходом; одновременно устранена гонка за разделяемую память в обратном ядре — синхронизацией 7-й стадии с warp сохранения dV. Производительность на 2 картах RTX 5090 с BF16, chunk 32, длиной последовательности от 32K до 128K сравнивалась с Inter-CP, по 3 измерения на случай по медиане, каждое — 25 прогревов плюс 100 замеров.

Источник: https://github.com/QwenLM/FlashQLA/pull/51

3.10 Сторона внедрения: TileKernels без пушей с 14-10 (02-11)

Последний пуш TileKernels от DeepSeek остановился на 14-10, в текущем окне изменений нет. Все значимые изменения на стороне внедрения в этом выпуске исходят от Tongyi FlashQLA.

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

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

4.1 Формируется исследовательская цепочка RISC-V: три репозитория переводят Numba CUDA Python в TileLang, а затем ограниченный PrimFunc — в RVV C для одного hart (31-10)

Исследователь сообщества zishu0131-png 31-10 в тот же день оформил три репозитория: tilelang-riscv как общую точку входа с описанием цели исследования, двухэтапной цепочки компиляции и комбинаций версий; numba2tilelang отвечает за статический разбор Numba CUDA Python и преобразование в исходный код TileLang; tilelang2rvv преобразует поддерживаемый реальный tvm.tirx.PrimFunc в RVV intrinsic C, затем через Clang собирает в RV64GCV ELF и запускает на QEMU virt в режиме одного hart. Автор чётко обозначил границы: это не официально зарегистрированный бэкенд TileLang и не универсальный симулятор GPU-моделей; адресация, зависящая от threadIdx, межhart-синхронизация и неподдерживаемый IR отклоняются, блоки выполняются последовательно, shared и fragment отображаются в локальные массивы одного hart; исследовательская среда — TileLang 0.1.15, Clang 23.1.3, QEMU 11.1.0, совместимость с другими версиями не гарантируется, и явно указано, что потоковые артефакты вышестоящего уровня нельзя автоматически считать приемлемым входом для этого бэкенда.

Источник: https://github.com/zishu0131-png/tilelang-riscv

4.2 PTO-ISA/tilelangGym выделен из PTO-ISA/tilelang в отдельный репозиторий, главный шлюз регрессии — путь Ascend SimtVF (30-10)

Этот набор тестов чувствительности и регрессии выделен из examples/ascend/t_parallel_st_suite репозитория PTO-ISA/tilelang, причина — дать набору тестов возможность независимо фиксировать версии TileLang и PTOAS. Текущие три основных пути: Ascend SoftLib SimtVF (target=ascend, cython, главный шлюз регрессии, полностью проходится на версии remat), VectorizeParallelToPTO и vload/vstore для VL, а также ручные эталонные реализации SIMD DSL на @pto.jit в каталоге benches; в каталоге каждого тестового случая есть config.toml и VERDICT.md, дополнительно есть недельные и CI-скрипты регрессии.

Источник: https://github.com/PTO-ISA/tilelangGym

4.3 Новостная сторона: запросы Google News на китайском и английском, Hacker News и arXiv не дали существенных попаданий в окне (02-11)

Трёхдневные запросы Google News на китайском и английском не вернули ни одного относящегося к TileLang материала в пределах окна, Hacker News и arXiv также не дали новых записей; все 20 пунктов в списке кандидатов происходят из пушей в репозитории GitHub и представляют собой обычную активность в коде, а не новости. Информационная картина этого выпуска поэтому полностью складывается из стороны репозиториев и репозиториев сообщества.

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

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

Наиболее значимое суждение этого окна исходит от самого вышестоящего проекта: лучше откатить улучшения корректности, чем убрать сложность ThreadSync с критического пути. В описании отката ясно указана цена — три класса улучшений (проверка частного, точные границы участвующих потоков, не прямоугольные наборы потоков) отозваны вместе, оставлен только старый консервативный путь, а #3732 отложен до завершения миграции на TIRX. Готовность проекта компилятора публично написать «бюджета нашего доказателя не хватает, откатываемся к консервативной реализации» говорит о реальных ограничениях текущего этапа больше, чем любое наращивание функциональности: миграция на TIRX — приоритет номер один основной линии, все прочие оптимизации уступают ей дорогу.

В унисон с этим звучат два документа на стороне Ascend, описывающих «схождение к основной ветке»: RFC #1971 предлагает заменить алгоритм синхронизации на ядро основной ветки плюс отдельный сервис доказательств Z3 — сначала исследование, затем замена, а между ними слой адаптации старого TIR; #1975 напрямую переносит весь репозиторий Ascend на многобэкендовую архитектуру основной ветки. Оба пути требуют затрат на миграцию, поэтому видимая работа в последнее время смещается к фундаментальным задачам: #1948 снижает фиксированные накладные расходы хоста на каждый вызов с примерно 450 микросекунд до примерно 106 микросекунд, четыре проверки фактов в документации, очистка ссылок на внутренний кэш в CI — всё это работа такого рода.

Ритм TileOPs сместился с «выпуска релизов» на «настройку». Из 15 коммитов за 72 часа 13 пришлись на один день, почти каждый сопровождается таблицей реальных замеров на H200 и заявлением «в manifest нет регрессий», а для описания качества правил планирования начинают использоваться такие метрики, как regret выбора (#2808 с 1.818 до 1.017). Это означает, что точка конкуренции в библиотеках операторов сместилась с «есть ли такой оператор» на «правильно ли выбран программа для каждой shape», и именно в этом заключается долгосрочная ценность DSL вроде TileLang по сравнению с рукописными kernel.

На стороне экосистемы появились две новые развилки. Одна — аппаратные диалекты: Hygon трактует gfx948 как совместимый с gfx946, Ascend продолжает дополнять документацию PTO и символьные формы. Другая — языковой выход: исследовательская цепочка RISC-V понижает ограниченный PrimFunc до RVV C для одного hart, звёзд очень мало, и автор явно заявляет, что это не официальный бэкенд, но она использует IR TileLang как переносимый актив — это сигнал того же рода, что и адаптация на уровне репозиториев для Ascend и Hygon. Несколько развилок существуют одновременно, и то, сможет ли миграция вышестоящего TIRX дать стабильную границу IR, определит, осядут ли эти нижестоящие вложения или потребуют переработки.

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

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

Источник Результат проверки
Организация tile-ai (29 репозиториев) В окне мониторинга пуш в 10 репозиториев
Основной репозиторий tilelang 7 коммитов в основную ветку; 4 новых issue; 3 коммита в сайт документации
TileOPs 15 коммитов в основную ветку (13 — 11-01); 3 новых issue; 1 коммит в сайт
TileFoundry 3 коммита; 1 коммит в сайт
tilelang-ascend 5 коммитов; 4 новых issue (включая записи ежедневных тестов и RFC)
tilelang-hygon 2 коммита; 1 новый коммит на рассмотрении
tilelang-mlir-ascend 1 коммит; 1 новый коммит на рассмотрении
tilelang-metax / tilelang-musa / TileRT / tilescale / DeepStack В окне мониторинга пуша не было
Страницы релизов (tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-hygon / mlir-ascend) В окне мониторинга новых версий нет
FlashQLA 1 коммит (#51)
TileKernels С 10-14 пуша не было
Три репозитория исследовательской цепочки RISC-V / PTO-ISA/tilelangGym Есть пуши, три репозитория исследовательской цепочки создали цепочку в один день
Google News / Hacker News / arXiv Существенных попаданий в окне мониторинга нет

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

Синхронизация WASP (#31) — https://github.com/tile-ai/tilelang-hygon/pull/31

** Страница релизов TileFoundry — https://github.com/tile-ai/TileFoundry/releases