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

Источники: порепозиторная проверка организации GitHub tile-ai (tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-hygon / tilelang-metax / tilelang-musa / tilelang-mlir-ascend / TileRT / tilescale / DeepStack / сайт документации / репозиторий языковых сервисов / репозиторий ночных бенчмарков) + поиск подтверждённых данных по смерженным и находящимся на рассмотрении изменениям (30 смерженных, 49 обновлений на рассмотрении) + полная проверка новых issue (11) + проверка diff на уровне коммитов и страниц релизов + проверка состояния артефактов PyPI + проверка репозиториев внедряющих сторон (TileKernels / FlashQLA) + поиск новых репозиториев сообщества + множественные запросы Google News на китайском и английском (через прокси) + Hacker News + arXiv + отраслевые СМИ.

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

  • Ключевое за сегодня: за 24 часа принято 30 коммитов, 49 обновлений на рассмотрении, открыто 11 новых issue; основная линия — три тихие ошибки в «семантике потоков и барьеров»; первый релиз TileOPs всё ещё застрял на последней ячейке PyPI; два крупных слияния в Hygon
    1. Основные результаты по проекту
      • 1.1 Основной репозиторий: выделение вывода out_idx перенесено в поток вызывающей стороны, пять из трёх повреждений при воспроизведении на H100 сведены к нулю (#3649) (26.10)
      • 1.2 Основной репозиторий: компиляция NVRTC на стороне устройства подключена к пути исполнения TVM-FFI, добавлен новый ключ компиляции (#3285) (26.10)
      • 1.3 Основной репозиторий: в cache hint пакетной загрузки TMA не хватает одного параметра шаблона, путь TMA в DeepSelect не компилируется (#3683) (26.10)
      • 1.4 Основной репозиторий: нефинитные константы FP8/FP4 в редукции выводятся как недопустимые литералы, суббайтовый аккумулятор теперь отклоняется заранее (#3189) (26.10)
      • 1.5 Основной репозиторий: блочно-масштабируемый GEMM для SM120 с явным развёртыванием K-атомов, конфигурация FullCol даёт прирост 51% (#3286) (26.10)
      • 1.6 Основной репозиторий: скалярные параметры в выходной форме JIT разрешаются до выделения (#3602) (26.10)
      • 1.7 Основной репозиторий: dtype операнда метаданных разреженности sm8x больше не проваливается без запасного варианта (#3335) (26.10)
      • 1.8 Основной репозиторий: пример DeepSelect Top-K слит, в 129 наборах конфигураций максимум 8,00x относительно torch.topk (#3304) (26.10)
      • 1.9 На рассмотрении: масштаб рефакторинга TIRx расширен до 1194 файлов, привязка к непатченной upstream-версии (#3666) (26.10)
      • 1.10 Новый дефект: три типа смещённых барьеров ThreadSync, с тремя исправлениями на рассмотрении (#3680 / #3682 / #3681) (26.10)
      • 1.11 Новый дефект: конфликт batch и раскладки редукции на sm_100 и выше, уже есть одно исправление на рассмотрении (#3667 / #3684) (26.10)
      • 1.12 На рассмотрении: проверка целевой формы транспонирования, регионализованная перестановка kill (#3677 / #3670) (26.10)
      • 1.13 Закрыто: шесть падений на этапе компиляции из серии фаззинга (26.10)
    1. Адаптация под множество бэкендов (Ascend / Hygon / MetaX)
      • 2.1 Основной репозиторий Ascend: интерфейс точной адресации с версионированным on-chip буфером (#3607) (26.10)
      • 2.2 Ascend на рассмотрении: FlashMLA с ручным планированием, BF16 prefill и FP8 decode (#3678) (26.10)
      • 2.3 Ascend на рассмотрении: раннее отклонение переполнения on-chip памяти, копирование E4M3 из UB в L1, раскладка представления и единицы указателей (#3655 / #3639 / #3623) (24–26.10)
      • 2.4 Ascend на рассмотрении: опциональность пакетного планирования, требование пакетного режима для межъядерной синхронизации, прямой двойной копии L0C (#3675 / #3614 / #3468) (26.10)
      • 2.5 Репозиторий Ascend: двойное исправление секции SHMEM в установщике, пустой wheel больше не даёт ложного успеха (#1943 / #1946) (26.10)
      • 2.6 Репозиторий Ascend: неверный результат int4-матричного умножения W4A4, корневая причина — упаковка полубайтов воспринимается как 1 байт (#1731) (26.10)
      • 2.7 Репозиторий Ascend: копирование GM в UB более чем двух измерений сворачивает ведущие измерения по числу строк DMA (#1805) (26.10)
      • 2.8 Новый дефект Ascend: обёртка Cython создаёт анализатор при каждом вызове, около 0,28 мс (#1945) (26.10)
      • 2.9 Hygon: конвейер регистров и дивергенция warp слиты, 28 файлов (#18) (26.10)
      • 2.10 Hygon: нативное хранение B16 и динамический GEMM WASP слиты, двойной банк scale-буфера на рассмотрении (#28 / #29) (26.10)
      • 2.11 MetaX: в организационном представлении есть записи о пушах, ни в одной из шести веток нет новых коммитов в окне мониторинга (26.10)
    1. Экосистема и потребители (TileOPs / TileFoundry)
      • 3.1 TileOPs: первый релиз всё ещё на последней ячейке, артефакт по-прежнему dev-версия от 2025-10-09 (#2738) (26.10)
      • 3.2 TileOPs: плотный forward D128 с разделением key, prefill-rope снижен с 31,04 мкс до 17,09 (#2777) (26.10)
      • 3.3 TileOPs: персистентный forward D128 переписан по потокам с выводом LSE (#2774) (26.10)
      • 3.4 TileOPs: построчное разбиение LayerNorm с fused-сложением между CTA, для дефекта точности открыт отдельный тикет (#2768 / #2773) (26.10)
      • 3.5 TileOPs: четыре поэлементных и редукционных, строка bf16 clamp ускорена на 5,9% (#2766 / #2767 / #2769 / #2770) (26.10)
      • 3.6 TileOPs: проверка выравнивания вызываемых тензоров поднята на уровень оператора, три завершающих (#2772 / #2778 / #2780) (26.10)
      • 3.7 TileOPs: три типа исправлений смещённого хранения, групповой GEMM SM90 и выходной ABI (#2761 / #2765 / #2762) (26.10)
      • 3.8 TileOPs на рассмотрении: постраничный prefill FP8-запросов переведён на умножение value по e4m3 (#2775) (26.10)
      • 3.9 TileOPs на рассмотрении: разреженные длинные последовательности GD

Сканирование разбиения предзаполнения N (#2779) (26.10)

  • 3.10 TileFoundry: динамическая разделяемая память размеруется по фактическому представлению ядра (#271) (26.10)
  • 3.11 TileFoundry: целое окно передаётся самим операндом, GEMM для sm_89 снова компилируется (#273) (26.10)
  • 3.12 Новый дефект TileFoundry: сторона анализа завышает пик разделяемой памяти и не передаёт его в TIR (#272) (26.10)
    1. Сообщество, руководства и мероприятия
      • 4.1 Языковой сервис: запущена экспериментальная отладка CUDA с подключением cuda-gdb и Nsight (26.10)
      • 4.2 Сайт документации: восемь обновлений за один день (26.10)
      • 4.3 Четыре улучшения управления процессом внесения вклада: регламент описания PR, пропуск серверных тестов для чисто документационных изменений, подтверждение заголовка Agent, снижение шума бота проверки (#3671 / #3672 / #3673 / #3674) (26.10)
      • 4.4 Новые репозитории сообщества: центральный репозиторий сообщества RISC-V, репозиторий операторов CANN, репозиторий бэкенда Hexagon, репозиторий инструментов отладки (26.10)
      • 4.5 Промышленность: публикация о сокращении разрыва в экосистеме между Ascend и DeepSeek (26.10)
    1. Наблюдения за трендами
  • Приложение. Материалы и пояснения к проверке

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

За 24 часа организация tile-ai слила 30 коммитов (tilelang 14, TileOPs 14, tilelang-hygon 2), на рассмотрении 49 обновлений, открыто 11 issue, закрыто 12. Вчерашние 78 коммитов за 72-часовое окно за один день сократились до 30 — темп нормальный, но содержательный центр явно сместился с «зачистки производительности» на «семантику потоков и барьеров»: слитое в основной репозиторий исправление выделения выходных буферов (#3649) свело к нулю проблему повреждения выходных данных в трёх запусках из пяти на H100, причина в том, что три адаптера принимали stream= вызывающей стороны, но выделяли выходные буферы out_idx в текущем потоке, а кеширующий аллокатор упорядочивает повторное использование по «потоку выделения», из-за чего работа, стоящая в очереди в другом потоке, перезаписывала результаты ядра; в тот же день всплыли три дефекта ThreadSync (#3680), указывающие на то, что барьер попадает в дивергентную ветвь, барьер во ветви ошибочно снимает внешнюю отложенную зависимость, а также на позицию барьера в цикле с подсчётом по потокам, вместе с двумя исправлениями на рассмотрении #3682 и #3681, плюс один коммит в форке tvm переводит побитовые булевы операции в логические, так что прувер больше не сдаётся в пределах бюджета 10000 шагов (#88). По релизам прогресса нет: в стартовом трекинговом тикете TileOPs (#2738) остались только пробный прогон на TestPyPI и первый tag, на PyPI последним по-прежнему остаётся 0.0.1.dev2 от 2025-10-09, в репозитории нет ни tag, ни release. На стороне бэкендов Hygon за один день слил два крупных коммита: конвейер регистров с warp-дивергенцией (#18, 28 файлов +5312/-64) и нативное хранилище B16 с динамическим WASP GEMM (#28); в основную ветку Ascend слит интерфейс точной адресации с версионированным on-chip буфером (#3607, 14 файлов +982/-26). Самым тяжёлым на рассмотрении остаётся #3666, масштаб вырос с чуть более тысячи вчера до 1194 файлов +106648/-104497.

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

1.1 Основной репозиторий: выделение выходных буферов out_idx перенесено в поток вызывающей стороны, воспроизводимое на H100 повреждение «три из пяти» сведено к нулю (#3649)

Слито 26.10 в 22:36, 5 файлов +183/-31, исправляет #3648. Три адаптера cython, nvrtc и cutedsl принимают stream= и запускаются в этом потоке, но выходные буферы out_idx выделялись через torch.empty в текущем потоке torch; кеширующий аллокатор упорядочивает повторное использование только по «потоку выделения», поэтому блок видеопамяти, в котором ещё стояла работа в очереди, возвращался как свободный, и эта работа затем перезаписывала результаты ядра. Исправление добавляет output_allocation_stream(stream, device): если поток вызывающей стороны отличается от текущего, возвращается поток вызывающей стороны, иначе возвращается пустое значение; при нулевом дескрипторе выполняется отображение на поток по умолчанию, поскольку в старых версиях PyTorch ExternalStream(0) даёт пулированный поток, а не поток по умолчанию. Воспроизведение на H100 80GB из повреждений 4/5, 5/5, 5/5 превратилось в 0/5 для каждого из трёх адаптеров, новый тестовый файл — 10 пройденных тестов. Автор явно отклонил вариант с добавлением wait_stream: это заставляло бы каждый вызов ставить поток запуска после всей работы текущего потока.

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

1.2 Основной репозиторий: компиляция на стороне устройства через NVRTC подключена к пути исполнения TVM-FFI, добавлен новый элемент ключа компиляции (#3285)

Смержен 27.10 в 00:02, 12 файлов +133/-133. Добавлен tl.cuda_compiler="nvrtc", работающий совместно с существующим хост-рантаймом TVM-FFI: NVRTC компилирует device-исходники CUDA в cubin, общий модуль рантайма CUDA импортирует эти байты, а хост-сторона — привязки, преобразования скаляров, условия, циклы и подготовка запуска — по-прежнему генерируется TVM-FFI; значение по умолчанию остаётся nvcc, старый адаптер execution_backend="nvrtc" сохраняет прежнюю реализацию. В ключ бинарного кэша добавлены имя компилятора и версии NVCC/NVRTC, которые участвуют в хешировании вместе с деталями цели и эффективными опциями компиляции; архитектурные токены вроде 90a и 100f сохраняются как есть, для несовместимых целей диагностика выдаётся до компиляции; каждая успешно созданная NVRTC program уничтожается в finally. CI этого PR на H20 показал 5506 пройденных, 491 пропущенный, примеры 76 пройденных, канал ROCm 7.2 — 3220 пройденных, канал Ascend — 1797 пройденных.

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

1.3 Основной репозиторий: у cache hint пакетной TMA-загрузки не хватает одного шаблонного параметра, TMA-путь DeepSelect не компилируется (#3683)

Смержен 27.10 в 00:11, 2 файла +20/-11. Пакетная TMA-загрузка с eviction_policy="evict_first" или "evict_last" не компилируется на main: генерация кода выдаёт один шаблонный аргумент CacheHintSm90, тогда как перегрузка tl::tma_load для сырых указателей принимает только тип барьера. Этот дефект напрямую блокирует TMA-путь DeepSelect; воспроизведено на GB200 с CUDA 13.1 на примерах ring и replay из этого примера. Исправление выделяет блок cache hint из #3296, добавляет недостающий шаблонный параметр и передаёт выбранную политику в PTX-инструкцию; часть с выравниванием разделяемой памяти остаётся в #3296. После исправления на той же машине все 144 пункта файла DeepSelect проходят, 22 теста TMA-копирования проходят, а проверка компиляции SM90 с CUDA 12.9 также подтверждает, что PTX для всех трёх политик содержит надлежащие операнды (остальные изменения #3296 ещё не смержены).

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

1.4 Основной репозиторий: нефинитные константы FP8/FP4 в редукции порождают недопустимые литералы, суббайтовые аккумуляторы теперь отклоняются заранее (#3189)

Смержен 26.10 в 15:44, 7 файлов +204/-4, исправляет #2998. T.reduce_max и T.reduce_min при выводе в FP8 E4M3 выдавали недопустимый литерал inff; теперь нефинитные константы E4M3, E8M0 и FP4 идут через существующий конструктор преобразования с передачей допустимых выражений CUDART_INF_F и CUDART_NAN_F, поведение насыщения и NaN сохраняется. То же изменение заставляет обычную и отложенную редукцию заранее отклонять суббайтовые числовые аккумуляторы (FP4, int4 и подобные), выдавая явную диагностику; вход FP4 с аккумулятором FP16 или FP32 по-прежнему поддерживается, булевы редукции освобождены от проверки, поддержка редукции E8M0 по-прежнему вне рамок этого изменения. На B300 с CUDA 13.1 все 28 фокусных тестов проходят, полный файл тестов редукции — 150 пройденных, ещё три сбоя пакетной валидации уже существуют на main.

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

1.5 Основной репозиторий: блочно-масштабируемый GEMM SM120 явно разворачивает K-атомы, конфигурация FullCol даёт прирост 51% (#3286)

Смержен 26.10 в 16:36, 4 файла +308/-2. Это завершающий штрих после введения T.gemm_blockscaled и T.mma_gemm_blockscaled в #3257: когда операнд A берётся из регистрового фрагмента, цикл K-атомов ранее использовал T.serial, и каждое обращение к фрагменту было индексацией времени выполнения; nvcc в конвейеризованных и warp-специализированных ядрах без самостоятельного разворачивания оставлял фрагмент в локальной памяти, то есть вытеснял MMA-путь; после замены на T.unroll(..., explicit=True) каждое обращение становится константой времени компиляции. Второе место — когда и SFA, и SFB находятся в разделяемой памяти, раньше на каждую пару MMA приходилось три чтения разделяемой памяти, теперь на каждую lane на каждый K-атом загружается один компактный пакет выбора, покрываются оба layout: строчно-ориентированный и блочно-K-ориентированный. На RTX PRO 6000 Blackwell 143 пройденных, 4 пропущенных, вывод побитово совпадает с эталоном float32; в бенчмарке регистрового A FullRow даёт прирост 2.7%, Square — 11.1%, FullCol — 51% (с 348.5 до 526.7 TFLOPS), а ptxas сообщает о снижении spill с 664 и 484 байт до 320 и 160. Автор также отмечает: путь с block_K = 128 и строчно-ориентированным scale в разделяемой памяти и на main не проходил валидацию, это не связано с данным изменением и оставлено для отдельной проверки.

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

1.6 Основной репозиторий: скалярные параметры в выходных формах JIT разрешаются до выделения памяти (#3602)

Смёржен 26.10 в 10:06, 7 файлов +242/-3, исправляет #3406. Обёртки Cython и PTO ранее при выделении выходных данных не могли разрешить выходные размерности, выраженные через явные скалярные параметры PrimFunc; теперь адаптер передаёт отображение «выходная размерность → позиция скалярного параметра», и обёртка напрямую использует значение фактического аргумента, не добавляя дублирующих параметров запуска. Регрессионные тесты покрывают оба бэкенда, нульмерные и ненульмерные случаи, скалярные и многомерные выходы, порядок выходов, размещение на устройстве и количество параметров запуска. Ещё два коммита из той же группы дефектов находятся на рассмотрении: #3636 обрабатывает составные выражения размерностей вида (n, 2*n), #3638 — проблему слияния одноимённых, но независимых динамических символов по имени.

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

1.7 Основной репозиторий: dtype операнда для разрежённых метаданных sm8x больше не проваливается без проверки (#3335)

Смёржен 26.10 в 16:54, 1 файл +12/-5. Конструктор раскладки разрежённых метаданных диспетчеризуется по архитектуре и сопоставляет dtype операнда MMA с шириной буфера метаданных; ветка sm90 отклоняет неподдерживаемые dtype, а ветка sm8x была написана в виде двух проверок по белому списку без запасного варианта, поэтому любое другое написание попадало в нижележащую логику размещения, которая выводит group и interweave только из ширины буфера. Следствие: float8_e4m3fn (каноническое написание в torch) принимался как 16-битные метаданные и применял соглашения fp16 — group 32 и interweave 4; любое нераспознанное написание, включая опечатки, также принималось при произвольной ширине. Теперь неподдерживаемый dtype сразу вызывает ошибку.

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

1.8 Основной репозиторий: пример DeepSelect Top-K смёржен, до 8,00× относительно torch.topk на 129 конфигурациях (#3304)

Смёржен 26.10 в 16:53, 4 файла +1265/-0. Пример переписывает DeepSelect от DeepSeek на TileLang без встроенного PTX; автор измерил 129 конфигураций, 10 из которых непригодны, на RTX 5090 (170 SM) с CUDA 12.9 и PyTorch 2.7, прогоняя двумя методами: с повторяющимся вводом и с очисткой кэша. В режиме повторяющегося ввода конфигурация BF16, батч 6, N = 16384, K = 512 показала задержку 5,059 микросекунды против 40,448 у torch.topk — ускорение 8,00×; конфигурация с N = 1048576 — 15,503 против 100,254, 6,47×. Комбинации FP32 и K = 4096 в этой обёртке непригодны, нативный путь для K = 1 имеет ошибку выравнивания адреса; автор помечает в таблице «непригодно» прочерком, а не нулём. Исправление подсказки кэша, от которого зависит этот пример, — это смёрженный в тот же день #3683.

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

1.9 На рассмотрении: масштаб рефакторинга TIRx вырос до 1194 файлов, привязка к непатченной версии upstream (#3666)

Открыт 25.10, трижды обновлён 26.10, текущий объём 1194 файла +106648/-104497. План — перевести TileLang с встроенного TVM на отдельно выпускаемый TIRx и наложить ноль патчей на фиксированный базовый upstream: удалить 3rdparty/tvm и его элементы подмодулей, в 3rdparty/ остаются только CUTLASS, заголовочные файлы HIP и файл игнорирования; неизменённый исходный код apache/tvm-tirx закреплён на 9af5544d, зависимости — apache-tvm-tirx==0.1.0a2 и apache-tvm-ffi==0.1.15.post0, прежние интеграционные патчи с суффиксом +tilelang.1 больше не нужны. Это самый крупный коммит в текущей очереди, и он определяет, как последующие исправления на стороне TVM (например, булевы побитовые операции Z3 в форке tvm того же дня) попадут в основную ветку.

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

1.10 Новый дефект: три класса смещённых барьеров ThreadSync, три исправления на рассмотрении (#3680 / #3682 / #3681)

#3680 открыт 26.10 в 19:56, указывает на неправильно размещённые барьеры tl.ThreadSync в трёх местах: защита для по-поточных значений ошибочно считается согласованной внутри блока, поэтому барьер остаётся внутри ветки и ждёт потоки, пропустившие ветку; барьер в одной из веток if очищает отложенные зависимости внешней последовательности, и при пропуске этой ветки блоком запись перед if конфликтует с чтением после if; когда число итераций цикла является по-поточным (например, по модулю номера потока), барьер остаётся внутри тела цикла, и частичное понижение синхронизации сразу завершается ошибкой «невозможно доказать прямоугольность потоков». #3682 (изменения уровня 3 файла +39/-7, 12 новых тестов) использует однопроходный предварительный анализ для сбора буферов с возможными по-поточными значениями и сокращает сводку синхронизации по достижимости ветвей и циклов; автор заявляет, что изменение только сообщает о гонках, не устраняет их и не выполняет разбиение ветвей. #3681 исправляет другой однородный сбой: начиная с sm_100 циклы float32 векторизуются до 256 бит, и при трансляции в 8 дорожек сгенерированное выражение берёт адрес временного объекта, NVRTC отказывается компилировать с ошибкой «должно быть lvalue», исправление собирает биты напрямую через __float_as_uint со сдвигом. Сопутствующий PR #88 форка tvm переводит булевы побитовые операции в логические, что позволяет доказателю прийти к тому же выводу в пределах бюджета 10000 шагов.

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

1.11 Новый дефект: конфликт batch и раскладки при редукции на sm_100 и выше, есть один рассматриваемый патч (#3667 / #3684)

#3667 открыт 26.10 в 08:31: T.reduce_* на sm_100 и sm_120 не понижается, когда указан batch, равный N, поскольку копирование фрагментов на sm_100 и выше идёт через 256-битные загрузки, число выходных элементов на поток уменьшается, и batch, допустимый на sm_90, вызывает жёсткую проверку. #3684 открыт 27.10 в 05:44 вместе с исправлением: ReduceOp теперь берёт наибольший собственный делитель N, не превышающий batch, и выдаёт предупреждение; finalize_reducer остаётся строгим, так как его тесты фиксируют поведение; новые тесты понижения для sm_100a и sm_120a падают на main и проходят в этой ветке, тестовая машина — RTX 5070 Ti. Другой патч по той же проблеме редукции #3341 (разрешение диапазона потоков для пакетной полной редукции) всё ещё на рассмотрении.

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

1.12 На рассмотрении: проверка целевой формы транспонирования, регионализированная перестановка kill (#3677 / #3670)

#3677 переносит проверку целевой формы tl.transpose в конструктор TransposeNode, выполняемый после NormalizeToAccessRegion, поэтому покрывается и прямое конструирование в обход Python-обёртки; равенство определяется через arith::Analyzer::CanProveEqual вместо сравнения списков измерений Python, (3, n+1) и независимо сконструированное (n+1, 3) по-прежнему принимаются, а при несовпадении констант выдаётся ошибка с номерами осей и длинами обеих сторон. Ранее несовпадающая цель понижалась в выход за границы индексов, давая неверные результаты или запись за границы. 13 тестов CPU-цели проходят. #3670 заменяет оставшийся механизм scope_level в анализе живости общей памяти: это значение увеличивается только внутри тела For, содержащего SeqStmt, и никогда не увеличивается в While, поэтому ошибается в обе стороны — в случае While точка kill не переходит обратное ребро, упаковщик перекрывает взаимоисключающие области, и на sm90 измеренный результат удваивается каждый раунд; в случае соседних циклов kill переходит границу всей области. Новое правило вычисляет по реальным цепочкам внешних For и While и непрерывным интервалам каждого цикла, ограничивая оба направления — выход за границы и преждевременную остановку.

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

1.13 Закрыто: шесть сбоев времени компиляции из серии фаззинга (26.10)

26.10 за один раз закрыто шесть дефектов времени компиляции или неверного кода из фаззинга: T.atomic_add на целях int16 и uint16 напрямую прерывается, так как в nvcc нет 16-битной перегрузки (#2572); небулев элемент буфера, используемый напрямую как условие if или assert, прерывается внутренним утверждением is_bool() (#3009); вещественные операнды, поданные в семейства div и mod, принимающие только целые, также вызывают внутреннее утверждение (#3008); целочисленные ширины вроде int2, uint2, не входящие в ряд 1, 4, 8, 16, 32, 64, прерывают генерацию кода CUDA (#2996); reduce_max и reduce_min при записи в фрагменты FP8 или FP4 падают (#2998); а атомарная редукция с целым буфером в качестве назначения молча отбрасывает все потоки, кроме потока 0 (#2959). Для последнего также влито исправление документации (#3340), меняющее назначение в примере с целого буфера на именованный элемент, чтобы пример сам не продолжал учить неправильному написанию.

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

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

2.1 Основной репозиторий Ascend: интерфейс точной адресации версионированных внутричиповых буферов (#3607)

Объединено 10-26 10:04, 14 файлов +982/-26. Добавлен T.annotate_buffer_addresses, позволяющий в конвейере Ascend после планирования выполнять точную побайтовую адресацию выделений в пулах UB, L1, L0A, L0B, L0C и SS; буферы, помещаемые в адресацию, должны быть предварительно объявлены с фиксированным числом копий через T.annotate_buffer_versions. Проверка выполняется после развёртывания версий и дополнения раскладки и покрывает отсутствующие или неподдерживаемые числа копий, перекрытия, невыравненность, динамические размеры и переполнение ёмкости; аллокатор также должен резервировать контекст UB для SIMT и уступать приватное адресное пространство AIV, а дескрипторы scale factor для L0 следуют за своими выделениями данных. Частичное отображение закрепляет только заданный интервал, остальные буферы размещаются в выровненном свободном пространстве, и повторное использование занятых адресов не допускается; пулы без аннотаций адресации сохраняют прежнее поведение. При проекции на субъядра сохраняются только аннотации на активных копиях хранилища.

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

2.2 Ascend на рассмотрении: вручную спланированный FlashMLA, BF16 prefill и FP8 decode (#3678)

Открыт 10-26 19:31, охватывает BF16 prefill и компактный FP8 decode на Ascend 950, включая опциональный дополнительный KV cache в FP8 и FP4, attention sink, недопустимые индексы и длину по каждому запросу. Цикл внимания, онлайн-softmax, деквантование и планирование выражены на TileLang, лишь небольшой фрагмент intrinsic-заголовка оборачивает одну аппаратную инструкцию. Ручная реализация явно управляет конвейером и синхронизацией AIC и AIV; постоянная версия использует T.Stage вместе с управляемыми компилятором версиями буферов и синхронизацией, QK и PV используют разные представления раскладки одного и того же кольца хранилища L0B.

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

2.3 Ascend на рассмотрении: раннее отклонение переполнения on-chip памяти, копирование E4M3 из UB в L1, раскладка представлений и единицы указателей (#3655 / #3639 / #3623)

#3655 заставляет MergeUBAllocations отклонять переполнение уже на этапе снижения по целевой физической ёмкости: два независимых цикла каждый легально выбирают по две 32-КБ доли L0A, после объединения 64 КБ ровно заполняют объём, ранее выдавался указатель со смещением за конец физического буфера без каких-либо метаданных о размере, теперь выдаётся короткая диагностика.#3639 открывает копирование E4M3 из UB в L1: ND-вход автоматически переупорядочивается через существующий scatter и конвейер после копирования либо копируется напрямую в предварительно упорядоченный NZ-вход, переиспользуя существующие проверки геометрии и добавляя обработку размера элемента FP8 в анализ занимаемого объёма и задержки.#3623 заставляет буферы, создаваемые T.view и T.reshape, каждый сохранять свою раскладку, избегая выбора неверного отображения Cube или адреса DMA при разрешении раскладки через общую переменную данных, и позволяет смещению по «сырому» указателю переотображаться в собственных единицах элемента.

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

2.4 Ascend на рассмотрении: опциональное пакетное планирование, требование пакетного режима для кросс-ядерной синхронизации, прямая двойная копия L0C (#3675 / #3614 / #3468)

#3675 добавляет для Ascend опциональное пакетное планирование (schedmode 1): в обычном режиме планировщик запускает оператор, как только частично доступны необходимые ресурсы ядра, полноядерные синхронизирующие операторы нескольких потоков могут перемежаться, пакетный режим ждёт готовности всех ресурсов и запускает целиком, его обязательно использовать в сценариях полноядерной синхронизации; #3614 описывает симптомы и исправление того же механизма в режиме 0.#3468 обрабатывает ограничение A5, при котором FixPipe не может одновременно включать режим двух целей и quant_pre, заставляя T.dual_copy напрямую выполнять преобразование L0C из fp32 в f16/bf16, что устраняет одно преобразование Vector и буфер UB на fp32, а также полосу fp32 от L0C до UB.

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

2.5 Репозиторий Ascend: двойное исправление секции SHMEM в установщике, пустой wheel больше не даёт ложного успеха (#1943 / #1946)

Два рассматриваемых PR указывают на один и тот же дефект (#1728): install_ascend.sh --enable-shmem после сбоя build.sh не проверяет код возврата, продолжает bdist_wheel, подстановочный шаблон package_data не находит файлы и не сообщает об ошибке, в результате собирается пустой пакет с одним __init__.py и выводится Successfully installed shmem-1.0.0, и лишь более чем двадцать строк спустя на source ../../../install/set_env.sh сообщается «не удалось настроить окружение shmem», что уводит диагностику в сторону конфигурации окружения.#1946 меняет только два места: добавляет проверку кода возврата с немедленным выходом, делает set_env.sh абсолютным путём на основе каталога скрипта и проверяет с помощью контролируемого harness с дословным извлечением секции shmem (до исправления 6 из 10 пунктов падали, 4 проходили, после исправления проходят все 10).#1943 идёт другим путём: перед удалением существующего пакета последовательно проверяет нативную сборку, скрипт окружения, сборку wheel и наличие wheel, добавляет 10 тестов потока управления CPU и поясняет, что исправляет продолжение работы установщика и диагностику, но не сам сбой компиляции вышестоящего SHMEM.

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

2.6 Репозиторий Ascend: неверный результат int4-матричного умножения W4A4, корневая причина — полубайтовая упаковка трактуется как 1 байт (#1731)

#1731 обновлён 10-26, локализация и воспроизведение достаточно полные. Весь стек шаблонов C++ выводит фрактальную геометрию из sizeof(AscendC::int4b_t) == 1, но этот тип имеет полубайтовую упаковку, и правильное число элементов в 32-байтовом блоке C0 — 64, а не 32, поэтому во всех трёх уровнях перемещения возникает двукратное отклонение: параметры LoadData от L1 к L0A/L0B вычисляются по неверной константе, данные на чипе укладываются в неверные позиции; транспонированная загрузка операнда B жёстко отключена в load_cbuf_to_cb_s4 с молчаливым отбрасыванием флага, и B никогда не получает раскладку nZ, ожидаемую mad_s4; длина в байтах между UB и GM считается как 1 байт на элемент, каждая строка копируется вдвое больше и смешивается с соседними строками. Сами вычислительные инструкции верны, поэтому симптом — «масштаб нормальный, значения полностью неверны» (в исходном отчёте 99,6% несовпадений). Исправления все находятся в src/tl_templates/ascend/common.h и двух преобразованиях, добавлены пример W4A4 и 10 регрессионных тестов int4, сквозной тест 1024-куба из 99,6% несовпадений стал нулевым несовпадением, существующий набор gemm_v0 из 48 тестов, 263 теста l1_to_l0 и gm_to_ub, 492 поэлементных теста — все проходят.

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

2.7 Репозиторий Ascend: копирование GM в UB более чем двух измерений сворачивает ведущие измерения в число строк DMA (#1805)

#1805 исправляет #1231. Модель DMA в copy_gm_to_ub и copy_ub_to_gm — двумерный прямоугольник, а AscendCopy::Lower берёт из активных измерений region только последние два как строки и столбцы; при более чем двух активных измерениях ведущие измерения молча отбрасываются: копирование 3D-блока A[t:t+16, :, :] (region 16, 4, 8, всего 512 элементов) понижается до 4 и 8, DMA переносит только 32 элемента, остальные 15 токенов фрагмента UB не инициализированы, последующие вычисления читают мусорные значения. Способ исправления — свернуть ведущие измерения в число строк DMA, изменены оба направления. В тот же день в репозитории Ascend обновлена партия PR проверки фактов документации (sin и cos, axpy, clamp, серия bitwise, clear, fill), причём в коммите clamp попутно исправлена неработающая нижняя граница T.tile.clamp на PTO.

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

2.8 Новый дефект Ascend: обёртка Cython создаёт анализатор при каждом вызове, около 0,28 мс (#1945)

Открыт 10-26 15:52, автор отчёта приводит послойную локализацию и причинно-следственную проверку. CythonKernelWrapper.forward() при каждом вызове безусловно конструирует TVM Analyzer(), медиана на локальной машине около 0,28 мс, а используется он лишь один раз в ветке динамического shape; для ядер со статическим shape dynamic_symbolic_map — пустой словарь, simplify() никогда не выполняется. Причинно-следственная проверка состояла в замене глобального Analyzer в модуле на заглушку, «бросающую ошибку при вызове simplify»; все вызовы по-прежнему успешны, полное время вызова снижается с 442–467 микросекунд до 104–107 микросекунд и не зависит от размера ядра (1 block и 32 block — один порядок). Измерения на Ascend 910B3 с CANN 9.2.0. Для инференс-конвейеров с плотными вызовами мелких операторов это фиксированные накладные расходы, стоит следить за направлением исправления.

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

2.9 Hygon: конвейер регистров и warp-дивергенция слиты, 28 файлов (#18)

Слито 10-26 13:33, 28 файлов +5312/-64, открыто 10-08, сошлось за более чем две недели. Изменение добавляет в tilelang-hygon конвейер регистров, который может работать совместно с существующим конвейером разделяемой памяти, соответствующие инструкции коммита и барьеров вставляются по ступеням конвейера; при числе warp более 4 включается warp-дивергенция. Вывод автора: example_gemm достигает производительности, сопоставимой с rocblas. В том же репозитории в этот день следует обратить внимание на порядок слияния #18: он слит раньше #28 и представляет собой крупное слияние изменений основной ветки в ветку, поэтому число строк намного больше, чем у #28.

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

2.10 Hygon: нативное хранилище B16 и динамический WASP GEMM слиты, двойной банк буфера scale на рассмотрении (#28 / #29)

#28 объединён 26.10 в 16:23, 6 файлов +340/-230, три вещи: сохранить порядок испускания параметров Let в MLS, распространить разбиение warp из GEMM на матричное хранилище, отклонить упакованный выходной алиас B16 в B32; одновременно сделать измерение WASP динамическим, выводить выравнивание по осям, исправить фильтрацию хвоста K в pong и отделить сквозную регрессию WASP от базовых тестов MLS. #29 открыт 26.10 в 17:13 и находится на рассмотрении: выделить для планирования строк буфера scale отдельный двухбанковый раздел, отдельный от банков LDS, чтобы выделение, копирование scale и понижение MMAC использовали одно и то же отображение строк, сохранить жизненный цикл физических интервалов и выравнивание страниц m0, распределить работу простаивающих wave по двум банкам и использовать это для блочного масштабированного GEMM при ручной специализации warp.

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

2.11 MetaX: на странице организации есть запись о push, но ни в одной из шести веток нет новых коммитов в окне мониторинга (26.10)

Страница организации показывает, что у tilelang-metax в окне мониторинга есть отметка времени push (пекинское время 26.10 15:08), однако при проверке по каждой ветке соответствующий коммит не найден: на ветке по умолчанию main последним остаётся повышение номера версии от 29.09.2026, на ветке dev последним является один merge от 16.10.2026, ветка 0.1.14 — 17.09.2026, 0.1.12 и race — ещё раньше; в потоке событий также нет событий push в окне мониторинга. Исходя из этого, это трактуется как «отсутствие существенных изменений», в текущий выпуск не включается, и природа события не домысливается. Последним официальным релизом этого репозитория остаётся 0.1.15 от 16.10.

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

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

3.1 TileOPs: первый релиз всё ещё в последней ячейке, артефакт по-прежнему dev-версия от 09.10.2025 (#2738)

Отслеживающий тикет первого релиза обновлён 26.10, из восьми пунктов очистки выполнено шесть: именование и параметры манифеста, упаковка и файлы репозитория, Python 3.13, публичный API и точки расширения, соглашения о поведении операторов, генерация проверки входных данных, сайт документации. Остались две ячейки: ночной бенчмарк (исправление двух сломанных базовых линий #2763 уже объединено; подтягивание новых строк выше всех компараторов, изменение накопительного сканирования на однократный запуск #2766 уже объединено, программа D128 для GQADenseWithLSEFwdOp #2774 уже объединена, но строка MHA всё ещё показывает лишь 0.996 от FA3, тюнинг продолжается) и сам релиз (сторонние уведомления, документация по установке через PyPI, группа зависимостей разработчика, CHANGELOG и шаги релиза #2753, а также пробный запуск TestPyPI и первый tag, требующие выполнения мейнтейнером). Проверка PyPI и репозитория: tileops последней остаётся 0.0.1.dev2 от 09.10.2025, всего три версии, в репозитории нет ни одного tag и release. Первый релиз находится в последней ячейке «документация готова, артефакт не загружен», и это продолжается уже много дней.

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

3.2 TileOPs: в плотном прямом проходе D128 разделены key, prefill-rope снизился с 31.04 микросекунды до 17.09 (#2777)

Объединено 26.10 в 22:45, 2 файла +527/-124. Если в плотном вызове D128 плитка из 128 строк оставляет SM простаивающими, то key каждой плитки разбивается на min(число SM, делённое на число плиток, число шагов key самой короткой плитки) частей, и разбиение выполняется только тогда, когда оно позволяет выровнять две ступени key из самой длинной плитки; каждая часть записывает ненормализованный вывод fp32 с максимумом и суммой масштабирования данной строки, а объединяющее ядро взвешивает по наибольшему максимуму, поэтому ни члены сокращения, ни большие logit не теряют точности. На последнем шаге key плитки, если последний необходимый некоторому warpgroup key попадает в младшие 64 key, softmax выполняет маскирование и экспоненцирование только для этих 64, и причинно-диагональный шаг warpgroup 0 обязательно попадает в них. На H200 1500 MHz prefill-rope для f16 и bf16 снизился с 31.04 и 31.01 до 17.09 и 17.09 микросекунды, что даёт 2.324 и 2.328 раза относительно FA3; блочное предзаполнение qwen3-8b снизилось с 31.30 до 17.79. В тестовых узлах чистый прирост 14.

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

3.3 TileOPs: персистентный прямой проход D128 переписан по потокам и записывает LSE (#2774)

Объединено 10-26 18:22, 2 файла +656/-48. Первопричина в том, что GQADenseWithLSEFwdOp ранее работал на непостоянном warp-специализированном ядре, поскольку постоянное ядро отвергало return_lse, и ночной бенчмарк показывал всего 0,73–0,78 от FA3. В этот раз на поточных примитивах TileLang (TMA-отображение, mbarrier, дескрипторы WGMMA, warp shuffle, stmatrix) переписан постоянный плотный форвард D128 с прямым выводом log-sum-exp; одна и та же реализация обслуживает и GQADenseFwdOp, и GQADenseWithLSEFwdOp; когда все key и value превышают половину L2, блоки запросов одного заголовка располагаются смежно, так что параллельные тайлы совместно используют K и V в L2; первый QK каждого тайла запускается в завершающей фазе предыдущего тайла. Публичный интерфейс и манифест не изменены. После слияния строка MHA по-прежнему составляет 0,996 от FA3, трассировочная одноаннотационная тюнинг-работа продолжается, что вместе с #2777 относится к двум независимым линиям оптимизации.

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

3.4 TileOPs: fused-add LayerNorm с построчным разбиением по CTA, дефект точности вынесен отдельной задачей (#2768 / #2773)

#2768 объединено 10-26 15:45, 5 файлов +525/-24. Отправной точкой стал откат ночного бенчмарка (строка декодирования llama-70b выросла с 0,0037 до 0,0044 мс), выяснилось, что причина не в изменениях кода, а в повторном выборе тюнинга: после изменения ключа тюнинга FusedAddLayerNormKernel прошёл повторный тюнинг, тюнер выбирал кандидатов по счётчику L2 спина к спине и выбрал 1024 потока, что на 19% медленнее холодного старта с 256 потоками. Новый FusedAddLayerNormHeldKernel удерживает одну строку (на SM90 — долю одного cluster CTA) в регистрах, при более низкой стоимости на самом занятом SM разбивает строку максимум на 8 CTA, размер блока берётся из числа варпов, делящего объём обращений доли; диспетчеризация выбирается по двум взаимно исключающим условиям — «число строк в очереди» и «регистры свыше 255», тюнинг предоставляет только ширину, которую может заполнить доля тюнингуемой сборки, так что строка декодирования больше не может быть перевыбрана. Декодирование llama-70b — 2,368 микросекунды против 4,800 у flash-attn. Выявленная той же правкой проблема точности вынесена в #2773: при удалении строки от нуля среднее и дисперсия суммы fp32 теряют точность (строка с константой 10000,1 после нормализации даёт около 0,7 вместо смещения), вблизи 1e36 сумма квадратов переполняется, приводя к полному NaN; автор предлагает перейти на Welford или двухпроходный метод с масштабированием.

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

3.5 TileOPs: четыре поэлементных и редукционных операции, строка bf16 для clamp ускорена на 5,9% (#2766 / #2767 / #2769 / #2770)

Четыре объединены в один день, все следуют принципу «сначала найти первопричину, потом менять». clamp (#2767) переведён на сравнение в собственной точности, первопричина в том, что старая реализация расширяла каждый элемент до fp32, для каждой границы выполняла проверку на NaN и выбор — около 5 инструкций на элемент; теперь одна HMNMX2.NAN на пару элементов, квадрат 16384 bf16 снизился с 0,25704 мс до 0,25024, что в 1,0587 раза лучше torch-compile, и на всех 16-битных режимах, NaN, положительной и отрицательной бесконечности и знаковом нуле подтверждено побитовое совпадение с сравнением fp32 в torch. Кумулятивное сканирование (#2766) переводит cumsum и cumprod с dtype=float32 с «сначала расширить в копию, затем сканировать» на один запуск с прямой записью в целевую точность, при рабочей нагрузке hidden-state-f32 пересылка снизилась с 112 MiB до 48 MiB, cumprod — с 47,36 до 14,62 микросекунды, cumsum — с 47,86 до 14,62. Умножение GELU (#2769) переведено на вычисление в fp32 через общий gelu() с единственным округлением, старый путь erf нёс ветку выбора NaN, три 16-битных округления и более длинную цепочку зависимостей; число блоков малой сетки гейтинга снизилось с 56 до 14. Бинарные операции (#2770) поднимают число элементов на поток для двухбайтных операндов с 8 до 16, сохраняя две 16-байтных загрузки в полёте, 34 строки конфигурации по умолчанию соответственно изменились, строка llama-7b для MinimumFwdOp снизилась с 8,42 до 8,00 микросекунды.

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

3.6 TileOPs: проверка выравнивания вызывающих тензоров поднята на уровень оператора, три завершающие правки (#2772 / #2778 / #2780)

#2772 объединён 26.10 в 15:28, 55 файлов +373/-224: исполнитель перед реализацией запуска перекопирует вызывающие тензоры, начало хранилища которых не выровнено по границе 16 байт, в выровненное хранилище, после успеха копирует обратно те несколько записанных и возвращает тензор вызывающей стороны; если скопированный тензор перекрывается с другими аргументами (и хотя бы одна сторона записывается), запуск сразу отклоняется, так как копирование скрыло бы собственные проверки ядра; хранилище без адреса (поддельный тензор или обёртка, отказывающаяся от чтения) не копируется; выравнивание и логика временного хранения на стороне ядра (включая сгруппированный GEMM MoE, состояние причинной свёртки, qkv с fused QK-norm RoPE) полностью удалены. Цена — рост хост-времени выровненного вызова на горячем пути на 1–2 процентных пункта, MoEPostPermuteFwdOp с out — на 4,4 процентных пункта. #2778 встраивает это решение в сгенерированный наблюдатель и позволяет кэшу планов определять копирование, горячий путь только читает это решение, прежний отдельный обход адресов исполнителя удалён. #2780 открыт 26.10: сейчас только четыре ядра проверяют перекрытие out с входом, остальные операторы с out молча конкурируют, и каждая проверка повторно читает data_ptr (MoEPostPermuteFwdOp за один вызов — на 12 раз больше), направление — построчно объявить в списке для каждого оператора, разрешено ли перекрытие, и унифицировать проверку на уровне оператора.

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

3.7 TileOPs: три типа исправлений смещённого хранилища, сгруппированный GEMM SM90 и выходной ABI (#2761 / #2765 / #2762)

#2761 заставляет MMA-реализацию сгруппированного GEMM принимать непрерывные тензоры с началом со смещением и выравнивать векторный доступ, при записи результата в смещённый выход сохранять идентичность и нетронутые строки экспертов, проверять перекрытие перед копированием; прежняя реализация на таких представлениях падала с ошибкой невыровненного адреса CUDA, добавлено 16 регрессий, покрывающих смещения a, b, bias, out для FP16 и BF16 по отдельности, с компактной и маскированной раскладками. #2765 дополняет ту же возможность до сгруппированного GEMM MoE на SM90: когда начало хранилища не на границе 16 байт, копируются a, b и временно сохраняется out, vector_aligned теперь явно фиксирует правило «записываемый тензор нужно временно сохранить и скопировать обратно после запуска», MMA-ядра BF16 и FP16 переходят на него, ядро FP8 наследует от ядра SM90 FP8. #2762 исправляет выходной ABI: когда адаптер TVM FFI объявляет, что выход выделяет вызываемая сторона, это передаётся исполняющему телу TileLang для собственного выделения и возврата, остальные адаптеры сохраняют прежний путь выделения; сообщённая регрессия — старая реализация на новой сборке проваливает все 27 случаев Mamba, кандидатная реализация проходит на публичном пути вызова.

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

3.8 TileOPs на рассмотрении: постраничный префилл запросов FP8 переведён на умножение value в e4m3 (#2775)

Открыт 26.10 в 17:20, всё ещё обновляется 27.10 в 07:00. GQAPagedFP8WSKernel выполняет умножение value для запросов FP8 в e4m3: отвечающий за преобразование warp транспонирует каждый тайл value в операнд e4m3 с K-мажорным порядком (транспонирование ldmatrix, одна байтовая перестановка по словам, stmatrix), упорядочивает ключи внутри каждого 16 так, чтобы счёт каждого потока был ровно его вероятностью, затем упаковкой в регистрах округляет вероятности до e4m3. softmax в форме FP8 как в FA3: смещение на 2 в 8-й степени не даёт малым вероятностям уйти в субнормали e4m3, максимум строки сдвигается с каждым большим счётом, тайлы ключей сканируются с последнего назад. На H200 llama-8b-fp8-prefill снижается с 369,3 до 339,8 микросекунды (относительно FA3 с 1,000 до 1,087), qwen3-fp8-ragged-odd с 130,9 до 123,9. Автор честно отмечает, что допуск точности ослаблен с 2 в минус 4-й степени до 2 в минус 3-й степени, поскольку сами вероятности e4m3 имеют округление порядка половины 2 в минус 4-й степени, и FA3 на подобных входах также выходит за допуск.

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

3.9 TileOPs на рассмотрении: разделённое сканирование префилла GDN для разреженных длинных последовательностей (#2779)

Открыт 26.10. На SM89 для разреженного префилла GDN на длинных последовательностях переиспользуется сетка существующего сканирования состояния и отдельная выходная программа: условия применимости — плотный вход D128 с key-мажорным порядком, не менее 16384 токенов и не более четверти числа SM для CTA сканирования состояния, остальные вызовы сохраняют прежнюю реализацию. Способ декомпозиции записан в идентификатор кэша программ, добавлены постоянные регрессии для хвостового блока и начального состояния. Измерения (CUDA 13, PyTorch 2.13, TileLang 0.1.15): длина 32768 в bf16 — 1,0384 мс против 1,0656 у FLA, 65536 — 2,1178 против 2,1786, а на 16384 FLA всё ещё быстрее (0,4695 против 0,4409). Автор явно указывает, что выигрыш имеет место только на этой группе из четырёх value head, для существующих форм моделей с восемью и шестнадцатью value head этот путь не выбирается, и кандидат также не повсеместно быстрее FLA.

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

3.10 TileFoundry: динамическая разделяемая память размеруется по фактическому представлению ядра (#271)

10-26 создан и находится на рассмотрении, 12 файлов +340/-32. Синтезированная хостовая точка входа ранее передавала каждому TIR-ядру dynamic_smem=0, поэтому ядра, обращающиеся к разделяемой памяти через целочисленные смещения TensorView, работали без динамической разделяемой памяти и при превышении 48 KiB никогда не повышали cudaFuncAttributeMaxDynamicSharedMemorySize. Автор выбрал размер по фактически достигаемому ядром наибольшему числу байт, а не по целевой ёмкости, по двум причинам: CUDA учитывает статическую и динамическую разделяемую память в одном лимите на блок, а данная среда выполнения сама содержит некоторый объём статической разделяемой памяти (ProgramMetaData или разделяемые тензоры размещения), поэтому запрос полной ёмкости просто не даст запуститься; кроме того, объём запроса на блок определяет, сколько блоков поместится на одном SM: на sm_89 при запросе 99 KiB помещается только один блок, даже если используется всего 512 байт. Генерация кода также изменена так, чтобы объявлять smem_base один раз перед первым оператором, поскольку в loop_outputs он ранее оказывался внутри блока первого цикла и был невидим для второго цикла, и nvcc сообщал о неопределённом идентификаторе. Два GPU-теста на H200 и sm_89 проходят; без патча первый даёт недопустимое обращение к памяти, а второй — ошибку компиляции.

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

3.11 TileFoundry: окно передаётся целиком по самому операнду, GEMM на sm_89 снова компилируется (#273)

10-26 создан и находится на рассмотрении, 5 файлов +54/-24. Код, который schedule finalize генерирует для GEMM на sm_89, не компилируется: при понижении весь аккумулятор оборачивается в окно, начинающееся с нуля и покрывающее всю его длину, T.ptr_of(acc[0:0+16, 0:0+8]), а генерация кода CUDA для среза указателя по фрагментированному регистровому тайлу сообщает «требуется обычная компоновка источника»; далее MMA читает регистровые фрагменты как голые срезы вида lhs_frag_dst[0:0+16, 0:0+32], для которых также нет ветки обработки. На самом деле эти окна представляют собой весь операнд целиком. Исправление внесено в понижение: ConvertHIRToTIR один раз определяет, покрывает ли окно свой операнд (начало — статический ноль, размер равен форме операнда), три места построения окон используют это определение совместно, в итоге записывается T.ptr_of(acc), а MMA читает фрагменты напрямую. Из 49 фикстур планирования, имеющих TIR-двойников, у 29 вывод finalize изменился согласно этому; GEMM на sm_89 и представления mesh снова компилируются, два вывода wgmma проходят отказ среза указателя и останавливаются на следующем представлении, причём эта компоновка также отклоняется на main.

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

3.12 Новый дефект TileFoundry: сторона анализа завышает пик разделяемой памяти и не передаёт его в TIR (#272)

10-26 создан, автор обнаружил его при ревью #271; оба пробела ведут к одному и тому же последствию: программа, проходящая анализ, отклоняется при запуске, либо итоговый TIR не может сам сообщить, сколько разделяемой памяти требуется. Первое: анализ памяти сравнивает smem с facts.memory.shared.per_cta (101376 байт на sm_89), но CUDA применяет этот лимит одновременно к статической и динамической разделяемой памяти, а среда выполнения сама содержит 12-байтовый __shared__ ProgramMetaData, и эта занятость записана только в runtime.cuh и не сообщается анализу, поэтому размещение, укладывающееся ровно в пределах 12 байт, проходит анализ и отклоняется при запуске. Второе: анализ записывает верхнюю отметку разделяемой памяти на CTA в RegionMemoryMetadata.peaks, но ConvertHIRToTIR не записывает её в производимый TIR, PrimFunction.metadata не печатается и не разбирается, поэтому после finalize приходится пересчитывать заново (#271 как раз пересчитывал, сканируя все представления разделяемой памяти с целочисленными адресами), а недостаточный запрос в хостовая точка входа, написанной автором, не будет обнаружен до запуска.

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

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

4.1 Языковой сервис: запущена экспериментальная отладка CUDA, подключены cuda-gdb и Nsight (10-26)

tilelang-lsp 10-26 в 17:11 отправил первый коммит отладки (78a5761, заголовок — добавление экспериментальной отладки TileLang CUDA). Содержание — четыре блока: добавлен лаунчер tlsp.debug, выполняющий сопоставление номеров строк Python, компиляцию с -G и -lineinfo, изоляцию кэша и захват сгенерированных исходников; cuda-gdb от Nsight подключён в VS Code с поддержкой точек останова Python, проверок перед запуском и просмотра сгенерированных исходников CUDA; добавлены модульные тесты, конфигурации запуска и опциональные приёмочные GPU-тесты (покрывают отложенные точки останова, пошаговое выполнение на устройстве, инспекцию потоков и числовые значения); документация на китайском и английском и CI обновлены синхронно. Это первый содержательный коммит данного репозитория с августа; предыдущее последнее обновление — CI-тесты от 8-04.

Источник: https://github.com/tile-ai/tilelang-lsp/commit/78a5761404

4.2 Сайт документации: восемь обновлений за один день (10-26)

tilelang.github.io за окно мониторинга отправил восемь пушей (пекинское время 07:32, 10:16, 10:21, 16:51, 17:25, 22:42, на следующий день 00:08, 00:16), заголовки коммитов единообразно «Update docs» — это автоматическая синхронизация сборки документации при слиянии в основной репозиторий. В тот же день был один пуш в сторонний репозиторий локализации (tilelang-docs-l10n) и первый содержательный коммит в узловой репозиторий китайского сообщества RISC-V, что говорит о нормальной работе цепочки синхронизации документации.

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

4.3 Четыре изменения в управлении процессом контрибуции: регламент описания PR, пропуск серверных тестов для чисто документационных изменений, подтверждение заголовка агентом, снижение шума от ревью-бота (#3671 / #3672 / #3673 / #3674)

Четыре изменения слиты в один день, все вокруг «стоимости ревью». #3671 в CONTRIBUTING.md прописывает, что описание PR должно содержать проблему, решение и компромиссы, доказательства должны прикрепляться рядом с утверждением, которое они поддерживают, при исправлении дефекта указывается поведение до и после или намерение и результат регрессионного теста, при изменении производительности приводятся сопоставимые измерения и условия. #3672 позволяет PR с чисто документационными изменениями пропускать сборку и тесты CUDA, ROCm, Metal, Ascend и CuTeDSL: используется проверка существующей lint-задачи, смотрящей на разницу относительно базы слияния, — если изменённые пути полностью состоят из Markdown, reStructuredText или находятся под docs/_static/, шаг пропускается; белый список намеренно сужен, исполняемые примеры, файлы сборки и рабочие процессы в него не входят. #3673 добавляет в навык коммита шаг подтверждения: после готовности итогового diff, проверок и описания PR предлагаемый заголовок в исходном виде показывается контрибьютору для подтверждения или замены. #3674 отключает автоматическое резюме и проверку покрытия docstring у ревью-бота и добавляет общие указания для ревью, требующие привязывать находки к текущему diff и фактическим путям вызова, различать регрессии и уже существующие ограничения, сначала сравнивать с существующей реализацией и лишь затем предлагать исправление.

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

4.4 Новые репозитории сообщества: узловой репозиторий сообщества RISC-V, репозиторий операторов CANN, репозиторий бэкенда Hexagon, репозиторий инструментов отладки (26.10)

За окно мониторинга появилось несколько заслуживающих регистрации внешних инициатив. TongmingLAIC/tilelang-riscv создан 23.10 и отправлен 26.10, позиционируется как узловой репозиторий сообщества TileLang на RISC-V, объединяющий бэкенд компилятора, AI-операторы, инструменты разработки и документацию; пока это пустой каркас, но направление совпадает с запросами отечественного стека RISC-V AI. Atlas-Ascend-Fork/ops-tilelang — репозиторий операторов TileLang от сообщества CANN для NPU Ascend, акцентирующий унифицированное управление реализациями ядер, автоматизированные тест-кейсы, стандартные нагрузки и базовые показатели производительности; 26.10 был пуш. xwhzz/tilelang-hexagon реализует бэкенд для Qualcomm Hexagon, репозиторий 32 MB, 26.10 был пуш. superAngGao/tilelang-debugger — рабочее пространство отладки с привязкой к исходникам на H200. Также обычные обновления RuneFang/tvm_tilelang_cookbook (37 звёзд) и localizethedocs/tilelang-docs-l10n.

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

4.5 Промышленность: репортаж о сокращении разрыва экосистем Ascend и DeepSeek (26.10)

26.10 СМИ пересказали отчёт SemiAnalysis от 22.10, согласно которому тесное взаимодействие DeepSeek и Ascend сокращает разрыв с NVIDIA в аппаратном и программном обеспечении чипов, и прогнозируется, что к 2027 или 2028 году DeepSeek будет в основном опираться на технологический стек Ascend; в репортаже особо отмечается, что этот вывод появился через несколько недель после выпуска DeepSeek программных инструментов для экосистемы Ascend, включающих шесть программных модулей, взаимно однозначно соответствующих платформе NVIDIA. Само сообщение не касается изменений функциональности TileLang, в этом выпуске оно регистрируется лишь как отраслевой фон: инструмент компиляции языков высокого уровня среди этих шести модулей — это tilelang-ascend, и установщик репозитория Ascend за тот день и исправление int4-матричного умножения лежат на одной и той же цепочке.

Источник: https://www.scmp.com/tech/tech-war/article/3371806/jensens-nightmare-huawei-deepseek-narrow-ai-chip-ecosystem-gap-nvidia-report

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

  1. Центр тяжести дефектов корректности смещается от «регрессии производительности» к «не сообщается на этапе компиляции, ошибка только на определённой архитектуре или в определённой области памяти». Все три главные темы дня относятся к этому типу: вывод размещается в неверном потоке (#3649, повреждение только при нестандартном потоке и наличии поставленной в очередь работы в текущем потоке), подсказке кэша TMA не хватает одного аргумента шаблона (#3683, сбой только при пакетной загрузке с eviction_policy), у разрежённых метаданных sm8x нет запасного варианта (#3335, распознаётся только каноническое написание torch, да ещё и применяется ошибочное соглашение о раскладке fp16). Общая черта таких дефектов — наличие аппаратных или параметрических предпосылок вне охвата обычного CI, поэтому их всё чаще выявляет серия фаззинг-тестов — пять из шести закрытых в этой главе за день пришли из этой серии.

  2. Синхронизация и планирование памяти становятся главным полем боя. Появившиеся за день три дефекта ThreadSync в #3680, два исправления в #3682 и #3681, правка прувера в форке tvm #88, регионализованная перестановка kill в основном репозитории #3670, а также пиковое значение разделяемой памяти в TileFoundry #272, не попадающее в TIR, — всё это указывает на одно: решения компилятора о синхронизации и памяти переходят от «эвристик» к «доказуемости». Это также объясняет, почему #3666 важен — он заменяет встроенный TVM на отдельно упакованный TIRx, после чего подобные правки прувера смогут попадать в основную ветку в обычном ритме релизов, а не оставаться в форке.

  3. Адаптация бэкендов вступает в фазу «точного контроля». T.annotate_buffer_addresses для Ascend позволяет адресовать пулы UB, L1, L0A, L0B, L0C и SS побайтно с проверкой перекрытий, выравнивания и ёмкости; регистровый конвейер Hygon, дивергенция warp и двухбанковое разделение scale-буфера — всё это перевод «работает» в «раскладывается согласно структуре оборудования». Эти возможности сосредоточены в последних двух неделях, что говорит о том, что отечественные бэкенды прошли порог работоспособности и начинают решать проблемы потолка производительности.

  4. Релизная инженерия по-прежнему остаётся слабым местом. Первый выпуск TileOPs, когда документация, манифест, упаковка, стабильность API и валидация входных данных уже готовы, несколько дней застрял на пробном прогоне в TestPyPI и первом теге, а артефакт на PyPI остановился на 2025-10-09. Для библиотеки операторов с 196 звёздами и 14 влитыми за день коммитами устанавливаемость определяет скорость внешнего принятия больше, чем ещё один процент оптимизации; единственный незавершённый технический пункт в трекере — строка MHA 0.996 — и руководство не должно позволять ему и дальше блокировать тег.

  5. Внешняя экосистема синхронно расширяется. Центральный репозиторий сообщества RISC-V, репозиторий операторов CANN и репозиторий бэкенда Hexagon появились или обновились в один день, что говорит о том, что TileLang превращается из «проекта организации tile-ai» в экосистему, где каждая платформа строит свой репозиторий. Для отечественных читателей особенно стоит отслеживать направление RISC-V: оно напрямую связано с прежними аппаратными направлениями AI-очков и периферийного инференса, но пока остаётся лишь пустым каркасом, а содержательный прогресс требует появления работоспособных операторов и бэкендов.

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

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

Источник Результат проверки
Организация GitHub tile-ai (29 репозиториев) За окно было обновление в 7 репозиториях: tilelang, TileOPs, TileOPs-nightly, tilelang-hygon, tilelang-metax, tilelang-lsp, tilelang.github.io
Влитые (is:merged) 30 коммитов, время вливания каждого сверено с окном
На рассмотрении (обновления is:open) 49 коммитов, полностью сверены по времени обновления
Новые issue 11, каждая прочитана по тексту
Закрытые issue 12, из них 6 — серия фаззинга
PyPI (tileops / tilelang) tileops последняя 0.0.1.dev2 (2025-10-09); tilelang последняя 0.1.15 (2026-09-30)
Проверка тегов на странице релизов У TileOPs нет тегов и релизов; у tilelang последний v0.1.15 (2026-09-30)
Репозиторий tilelang-ascend За окно нет обновлений (последнее 10-23), но активность PR и issue высокая
Репозиторий tilelang-metax В представлении организации есть метка времени обновления, но ни в одной из шести веток нет коммитов за окно, считается отсутствием существенных изменений
Потребители (TileKernels / FlashQLA) Обновлений за окно нет (10-14)
Множественные запросы Google News на китайском и английском За 24 часа нет точных совпадений по компонентам TileLang; один отраслевой материал см. в 4.5
Hacker News За окно одна фронтенд-запись, не связанная с TileLang, исключена
arXiv Поиск по ключевым словам не выявил новых статей за окно, последняя от 2026-07-24
Поиск по сообществам репозиториев участников и потребителей В шести внешних репозиториях были обновления за окно, см. в 4.4

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

#2775）— https://github.com/tile-ai/TileOPs/pull/2775