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

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

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

  • Ключевое за сегодня: исправления корректности компилятора в основном репозитории по пяти коммитам вошли в тот же день, на стороне Ascend добавлены обратные и рекурсивные операторы пятью коммитами, TileOPs на SM90 перевёл три foundry-оператора внимания на warp-специализированные программы (10-20)
    1. Ход работ по ключевым проектам
      • 1.1 Основной репозиторий: анализ векторизации включён в нижнюю границу ненулевых циклов, исправлены циклы с unit-step и guard смешанной точности (#3579) (10-21)
      • 1.2 Основной репозиторий: обнаружение межпоточного WAW в разделяемой памяти переведено на двухпоточный экземпляр, восстановлен пропущенный барьер (#2821) (10-20)
      • 1.3 Основной репозиторий: неподдерживаемые MMA tile теперь требуют явного заполнения K, чтобы избежать тихого сокращения K и сбоя вывода раскладки (#2818) (10-20)
      • 1.4 Основной репозиторий: fragment reduce поддерживает сохранение целевых измерений, исправлен сбой цели (64,1) (#2815) (10-20)
      • 1.5 Основной репозиторий: предикаты TIR переписываются синхронно с повышением разрядности целых (#3550) (10-20)
      • 1.6 Основной репозиторий: в CuTeDSL scan добавлен параметр смещения участвующих потоков, исправлен сбой построения cumsum (#3576) (10-20)
      • 1.7 CI основного репозитория: задача регрессии производительности закреплена за H100, для машин H20 обязательна блокировка fail-fast (#3569) (10-20)
    1. Адаптация под несколько бэкендов (Ascend / MetaX / Hygon / Moore Threads)
      • 2.1 Основной репозиторий Ascend: локальное копирование L0 поддерживает явный pitch, один GEMM tile может собираться из нескольких копирований (#3511) (10-20)
      • 2.2 CI основного репозитория Ascend: в рамках одной проверки разделены шаги сборки и тестирования, после сбоя остальные наборы продолжают выполняться (#3564) (10-20)
      • 2.3 Репозиторий Ascend: реализован обратный оператор indexer_bwd для дообучения DSA с разреженностью, на двух бэкендах AscendC и PTO (#1887) (10-20)
      • 2.4 Репозиторий Ascend: в репозиторий добавлен обратный оператор SiTU-GLU, сверен с gated-активацией Kimi K3 и перекрёстно проверен на 910B3 (#1877) (10-20)
      • 2.5 Репозиторий Ascend: в cann-bench добавлены LSTM и GRU, 12 / 8 jit kernel охватывают многослойные двунаправленные варианты и проекции (#1892 / #1902) (10-20)
      • 2.6 Репозиторий Ascend: дополнена китайская версия бенчмарка CrossEntropy, сохранена практика парных en/zh (#1801) (10-20)
      • 2.7 Ascend на рассмотрении: на ревью находится среда векторизации PTO (#3512), tvm_ffi планируется как бэкенд исполнения по умолчанию для pto (#3516), выравнивание эвристик DeepGEMM (#3577) (10-20)
      • 2.8 Hygon: исправление начала координат константы K окна MLS влито (#25), две правки — разбиение fp4 и распределённая — на рассмотрении (#26 / #19) (10-20)
      • 2.9 MetaX / Moore Threads / MLIR-Ascend: новых пушей за окно нет, все три линии остаются на существующих версиях (10-20)
    1. Экосистема и стороны внедрения (TileOPs / TileFoundry / библиотеки операторов)
      • 3.1 TileOPs: прямой проход сегментированного GQA переведён на paged warp-специализированные программы, пул FP8 читается через 16-битное зеркало страниц (#2668) (10-20)
      • 3.2 TileOPs: FP8 lightning indexer переведён на два consumer warpgroup, поочерёдно обрабатывающих 64-key tile (#2666) (10-20)
      • 3.3 TileOPs: декодирование MLA идёт по программе смещённых качелей, keys разбиваются на сегменты по числу SM, затем объединяются через log-sum-exp (#2665) (10-20)
      • 3.4 TileOPs: верхняя граница производительности тензорных ядер заменена на огибающую набора операндов, исправлены ложные срабатывания 105%/108% в ночных прогонах (#2661) (10-20)
      • 3.5 TileOPs: бэкенд-тесты и baselines torch-compile могут выполняться на не-CUDA устройствах (#2658) (10-20)
      • 3.6 TileOPs: SM89 для одиночной строки и одиночного столбца переведён на путь векторного GEMM, кандидаты фильтруются по бюджету разделяемой памяти (#2660) (10-20)
      • 3.7 TileOPs: три исправления кэша builder и контрактов тюнинга (#2662 / #2663 / #2664) (10-20)
      • 3.8 TileOPs, 11 на рассмотрении: параллельно идут три линии — NSA top-k для SM89, сегментированный DSA, узкий tile W4A16 (10-20)
      • 3.9 TileFoundry: объявление 0.0.3 внесено в README, рефакторинг тестов переиспользует кэш компиляции (#247 / #246) (10-20)
      • 3.10 Стороны внедрения: у TileKernels и FlashQLA за окно новых пушей нет (10-20)
    1. Сообщество, руководства и мероприятия
      • 4.1 Репозиторий сообщества: создан репозиторий китайского системного руководства tilelang-tutorial, охватывающий операторы, этап компиляции и два основных бэкенда (10-20)
      • 4.2 Репозиторий сообщества: создан репозиторий сообщества CANN ops-tilelang, реализации для Ascend и baselines управляются как операторные активы (10-20)
      • 4.3 Инструментарий: devkit и расширение VS Code активны, в основном репозитории на рассмотрении полный вынос lower trace (#3583) (10-20)
      • 4.4 Репозиторий сообщества: tilelangGym проверяет чувствительность PTO lowering, появился экспериментальный репозиторий бэкенда Hexagon (10-20)
      • 4.5 Документация

и версия: документационный сайт автоматически обновился четыре раза, VERSION основного репозитория по-прежнему 0.1.15 (21.10)

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

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

В окне мониторинга слито 26 коммитов (tilelang 9, TileOPs 9, tilelang-ascend 5, TileFoundry 2, tilelang-hygon 1), новых релизов не было. Три направления продвигаются одновременно: основной репозиторий большую часть слияний за день направил на корректность компилятора (нижняя граница векторизованных циклов, межпоточный WAW в разделяемой памяти, K-атом в MMA tile, сохранение размерности в fragment reduce, переписывание предикатов TIR при повышении разрядности) — это очистка «тихих ошибок», а не расширение функциональности; на стороне Ascend все пять слияний пришлись на обучение и рекуррентные операторы (обратное распространение разреженной донастройки DSA, обратное распространение SiTU-GLU, LSTM, GRU), плюс на рассмотрении находится коммит на 21 тысячу строк с фреймворком векторизации PTO — PTO становится вторым путём исполнения помимо AscendC; TileOPs на SM90 в один день перевёл три foundry-оператора — постраничный прямой проход GQA, FP8 lightning indexer и декодирование MLA — на warp-специализированные программы, а также сменил методику калибровки верхнего предела производительности тензорных ядер на огибающую множества операндов, устранив ложное превышение на 105% в ночных прогонах.

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

1.1 Основной репозиторий: анализ векторизации учитывает ненулевую нижнюю границу цикла (#3579)

Слито 21.10 в 00:27, 10 файлов +818/-198, продолжает исправление guard из #2819, одновременно закрывает две проблемы #2669 и #2616. Исходная реализация сравнивала f(i) с f(floor(i / W) * W), чтобы определить, одинаково ли значение выражения в каждой векторной группе, что неявно предполагало цикл с нуля; при ненулевой нижней границе цикла векторизация с unit-step вычисляла адреса разных lane как одну группу, и доказательства широковещательной загрузки и общего для всей группы условия теряли силу. Новое решение включает минимум цикла в анализ и сохраняет вложенные guard при переносе смешанных по точности преобразований. Это дефект компилятора, способный молча изменить численный результат.

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

1.2 Основной репозиторий: обнаружение межпоточного WAW в разделяемой памяти (#2821)

Слито 20.10 в 23:41, 3 файла +153/-19, исправляет #2698. В разделяемой памяти S[tx] и S[tx + 10] для одного потока — два разных адреса, но записи из разных потоков могут пересекаться; исходная проверка WAW использовала одну и ту же потоковую переменную для обоих обращений, поэтому пропускала нарушение и не вставляла барьер. Исправление использует два разных потоковых экземпляра для WAW в одной итерации и переименовывает соответствующие определения Bind, переиспользуя существующую модель RAW/WAR; WAW с зависимостью по циклу обрабатывается как раньше. Тесты покрывают оба случая — «барьер нужен» и «барьер не нужен» — включая guarded-записи, индексы, производные от Bind, непересекающиеся разделы и перестановки threadIdx.y/z; прошли 95 тестов в наборах синхронизации потоков и issue.

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

1.3 Основной репозиторий: неподдерживаемый MMA tile требует явного заполнения K (#2818)

Слито 20.10 в 16:41, 5 файлов +57/-14, исправляет #2715. Универсальный MMA emitter мог незаметно уменьшить K-атом до построения раскладки операндов, и FP16 при K=8 давал неверный результат; для нерегулярных размеров вроде K=20 и K=28 сборка падала на этапе вывода раскладки разделяемой памяти с невнятным сообщением. Теперь при инициализации emitter проверяет K плитки по фиксированному для архитектуры K-атому и при неделимости выбрасывает ValueError с указанием dtype входных данных, требуемого отношения делимости и рекомендуемого значения заполнения (например, на SM80 для FP16 при K=20 сообщает об атоме 16 и рекомендует дополнить до K=32 с нулевым заполнением лишних элементов).

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

1.4 Основной репозиторий: fragment reduce поддерживает цель с сохраняемой размерностью (#2815)

Слито 20.10 в 16:35, 8 файлов +343/-63, исправляет #2530. T.reduce_sum и аналогичные при свёртке (64, 64) по dim=1 в (64, 1) падали на этапе компиляции (Check failed: vars.size() == InputDim() (2 vs. 1)). При этом документальный контракт reduce_op.py явно допускает обе выходные формы — [X, Y] и [X, 1, Y] — и уже строит соответствующую ожидаемую форму, так что это пробел бэкенда, а не некорректное использование. Исправление обучает оба пути lowering и сохраняет прежнее поведение для целей без сохраняемой размерности.

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

1.5 Основной репозиторий: предикаты TIR переписываются синхронно с повышением разрядности целых (#3550)

Объединено 10-20 09:59, 4 файла +950/-27. После повышения разрядности целочисленного индекса проверка границ, контролирующая доступ к буферу, по-прежнему вычислялась в старой разрядности, что теоретически могло пропустить выход за границы. Добавлен IndexPredicateRewriter, который в ConfigIndexBitwidth и FlattenBuffer записывает информацию о повышении для каждого доступа, а после обработки операторов доступа переписывает соответствующие guard; сопоставление ограничено эквивалентными формами сложения и умножения, сохраняются посторонние условия, выражения значений данных и границы фиксированной ширины. Добавлены тесты пяти категорий: повышение индекса, переписывание guard, область видимости доступа, векторные конечные точки и конфликтующие повышения.

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

1.6 Основной репозиторий: CuTeDSL scan дополнен смещением участвующих потоков (#3576)

Объединено 10-20 12:49, 2 файла +30/-6. Понижение (lowering) scan в разделяемой памяти в #3536 ввело смещение участвующих потоков, но helper на стороне CuTeDSL по-прежнему генерировал вызовы со старыми конструктивными параметрами, и tl.CumSum1D(32, False, 0).run(...) напрямую выдавал ошибку конструктивных параметров; одновременно вышедшие из строя утверждения codegen также были источником двух сбоев CUDA CI в #2815. Исправление добавляет необязательный thread_offset (по умолчанию 0 для совместимости) в helper’ы cumsum/cummax для одного и двух измерений, перед выбором участвующих lane и warp вычитает смещение, а также помечает CUDA-специфичные тесты меткой requires_cuda.

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

1.7 CI основного репозитория: задача регрессии производительности закреплена за H100 (#3569)

Объединено 10-20 16:00, 1 файл +15/-1. Бот регрессии PR сравнивает задержку по каждому ядру с базовой линией, и цифры имеют смысл только на той же модели карты, на которой собиралась базовая линия. Организация внедряет самостоятельно размещённые машины H20 для обычного CI корректности, чей набор меток пересекается с существующим [self-hosted, nvidia]; без различения они будут молча забирать эталонные задачи: плотная производительность FP16 у H20 составляет около 148 TFLOPS против примерно 990 TFLOPS у H100 — разница примерно в семь раз, к тому же путь FP64 сильно урезан, поэтому каждый эталон сообщал бы о несуществующей регрессии. Теперь задача pr-regression обязана иметь метку h100 и перед началом работы проверять, что выделенная карта действительно H100; при ошибке в метке — быстрый отказ вместо публикации вводящих в заблуждение цифр.

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

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

2.1 Основной репозиторий Ascend: локальное копирование L0 с поддержкой явного pitch (#3511)

Объединено 10-20 14:16, 24 файла +1458/-187 — крупнейшее изменение направления Ascend в этом окне мониторинга. Ранее один tile GEMM мог заполняться только одним копированием L1→L0A/L0B, ненулевое целевое начало отклонялось, а каждое копирование выводило компактный pitch из собственной области, поэтому две половины, перенесённые раздельно, не складывались в целый целевой блок; транспонированная загрузка MajorK/MajorMN также была ограничена. Добавлен специфичный для Ascend параметр T.copy(..., l0_pitch=...), задающий L0 pitch по fractal-блокам, при этом вызывающая сторона указывает компактный pitch потребительского GEMM (включая динамический хвостовой блок); при опускании параметра сохраняется прежнее поведение компактной области. В примере обе половины A tile из 64 строк пишут l0_pitch=4. Одновременно обработаны lowering целевого смещения после выравнивания, проверки физических границ и транспонированных групп, а также интервалы версий распределения и буферов.

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

2.2 CI основного репозитория Ascend: разделение сборки и тестирования в рамках одной проверки (#3564)

Объединено 10-20 12:06, 1 файл +61/-9. Результаты сборки и тестирования Ascend собраны в одну проверку PR Test for Ascend с пошаговым отображением этапа сбоя и перечислением выводов каждого этапа в workflow summary. Окружение CANN разворачивается и собирается один раз; примеры AscendC, примеры PTO и юнит-тесты Ascend выполняются как три отдельных шага одного job; при успешной сборке и неотменённом workflow после сбоя одного набора остальные наборы продолжают выполняться, но итог всё равно считается неуспешным. Переменные CANN передаются через GITHUB_ENV.

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

2.3 Репозиторий Ascend: обратный оператор тонкой настройки разреженности DSA indexer_bwd (#1887)

Слито 20.10 в 15:26, 2 файла +1294. Для NPU Ascend добавлен indexer_bwd, обслуживающий разреженную донастройку DSA, одновременно поддерживаются два бэкенда — AscendC и PTO. Оператор вычисляет градиенты функции потерь KL-донастройки lightning-indexer по dIndexQ, dWeights, dIndexK, поддерживает упакованные последовательности переменной длины. Реализация разделена на три уровня: быстрый kernel с персистентным GEMM ориентирован на непрерывный и совместно используемый построчный индексный layout, с конвейеризацией и меньшим числом перемещений данных; универсальный kernel поддерживает произвольный индексный layout, включая накопление градиентов для недопустимых индексов -1; дополнительно предусмотрен cast kernel для преобразования точности.

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

2.4 Репозиторий Ascend: обратный оператор SiTU-GLU (#1877)

Слито 20.10 в 15:22, 5 файлов +854. Семантика согласована с aclnnSituGluGrad из CANN ops-nn, применяется для обратного прохода активации гейта Kimi K3 Stable LatentMoE: t = tanh(gate / beta), s = sigmoid(gate), situ_a = beta * t * s, восходящая ветвь в зависимости от знака linear_beta идёт через tanh или напрямую, на выходе grad_gate и grad_up. Эталон точности — независимая реализация на fp64 PyTorch, выполнена перекрёстная проверка с официальным aclnn на 910B3, максимальная относительная погрешность около 3e-3, что соответствует ожидаемому округлению bf16.

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

2.5 Репозиторий Ascend: cann-bench дополнен LSTM и GRU (#1892 / #1902)

Оба влиты в один день, +2304 и +2267 соответственно. LSTM охватывает четыре гейта i/f/g/o и временную рекурсию cell state, поддерживает многослойное стекирование (numLayers ≤ 8), двунаправленность, layout batchFirst, проекционное понижение размерности (projSize > 0, LSTMP) и опциональный bias, веса передаются как TensorList по принципу «один tensor на слой и направление», выход (y, hn, cn), реализация — 12 kernel’ов @tilelang.jit; основной путь идёт по fused-пути seq, один launch на (слой, направление), внутри ядра по T.serial(S) последовательно по временным шагам выполняются двойной GEMM, broadcast bias, гейтирование и обновление состояния, а не зависящий от рекурсии вклад гейтов со стороны x предвычисляется отдельным xgemm параллельно по всем временным шагам. GRU — 8 kernel’ов, объединённый GEMM трёх гейтов z/r/n на Cube, гейтирование и обновление состояния на Vector, межъядерные данные передаются через GM workspace, и вручную отключён AUTO_SYNC — согласно описанию коммита, автоматический режим порождает около 58 PipeBarrier на каждый временной шаг.

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

2.6 Репозиторий Ascend: китайская версия бенчмарка CrossEntropy (#1801)

Слито 20.10 в 15:26, 2 файла +91. Для документации бенчмарка CrossEntropy, влитого в #1567, добавлена китайская версия по принятому в FA (bench_mark.md / bench_mark_zh.md) и SFA (README.md / README_zh.md) парному обычаю en/zh: в examples/cross_entropy_loss/ добавлен benchmark_summary_zh.md, полностью переведены восемь шагов пути оптимизации, внутреннее сравнение производительности, сравнение с CANN, известные ограничения и покрытие тестами, в верхней части английской версии добавлена ссылка переключения языка. Чисто документальное изменение.

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

2.7 Ascend на рассмотрении: фреймворк векторизации PTO, бэкенд tvm_ffi и выравнивание с DeepGEMM

Три основные линии одновременно проходят проверку. #3512 (50 файлов +21022/-72) позволяет пользователям писать скалярно-параллельные программы T.Parallel (поэлементные, управляющие потоки, редукции) внутри T.SimdVF(lanes=L), которые компилятор векторизует в инструкции PTO VMI, компилирует через PTOAS и исполняет на Ascend 950; полная функциональность реализуется несколькими PR, данный PR — веха для инфраструктуры конвейера и поддержки поэлементных операций. #3516 (18 файлов +1473/-22) позволяет target="pto" с execution_backend="auto" разрешаться в tvm_ffi и становиться значением по умолчанию; на стороне хоста используется MakePackedAPI и сборка через C host, а среда исполнения Ascend запускает скомпилированный PTODSL aibin через aclrtLaunchKernelWithHostArgs. #3577 (29 файлов +1392/-1036) дополняет выбор конфигураций и epilogue для DeepGEMM-Ascend, охватывая fused bias, Alpha, случайное округление BF16 и квантование FP8 для плотных, пакетных и сгруппированных GEMM. Также #3530 (частичное копирование NZ-буфера UB→L1), #3573 / #3580 / #3581 / #3582 — группа исправлений предикатов PTO, анализа направления MAD и знаковости redux.

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

2.8 Hygon: исправление начала координат константы K в окне MLS (#25)

Смёржено 20.10 в 09:53, 2 файла +346. После разворачивания цикла подъём ресурсов MLS теряет смещение константного окна, из-за чего последующие обращения попадают в неверный диапазон. Исправление состоит в том, чтобы перед выбором одноосевого обновления или прямой инкрементальной адресации проверить не обновляемое измерение и при необходимости консервативно переинициализировать окно; также добавлено регрессионное покрытие на существующих тестах MLS. В окне репозитория Hygon на рассмотрении три изменения: #26 обрабатывает warp-разбиение для fp4 MLS (14 файлов +261/-144), #18 добавляет поддержку регистрового конвейера и warp-diverge, #19 — изменения, связанные с распределённостью.

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

2.9 MetaX / Moore Threads / MLIR-Ascend: новых публикаций в окне нет

По всем трём линиям в данном окне коммитов нет. Последний в репозитории MetaX — подготовка версии 0.1.15 от 29.09, репозиторий Moore Threads остановился на подготовке публичного релиза 0.1.15 от 30.09, последний в MLIR-Ascend — исправление изоляции TVM от 19.10 (вне данного окна). Число PR на рассмотрении во всех трёх репозиториях равно нулю, активности новых веток не наблюдается.

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

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

3.1 TileOPs: постраничный GQA forward переведён на paged warp-специализированную программу (#2668)

Смёржено 20.10 в 22:41, 8 файлов +901/-103. Новое ядро GQAPrefillPagedWSKernel (SM90) принимает GQAPrefillPagedWithKVCacheFwdOp: один append launch записывает chunk в страницы (key вращается под fused RoPE, оба квантуются в пуле FP8) и вращает query, после чего warp-специализированная программа GQAPagedFwdOp выполняет внимание над раздутым кэшем. Эта программа обслуживает head dim 256 (один query slot) и groups query шириной более 16 голов на одном KV head (уплощая query (T * H, D) в строки); размерности 64 и 256 пула, не покрытые программой FP8, читаются через 16-битное зеркало страницы с лимитом 4 GiB за раз; при возможном разделении запросов они сортируются по длине кэша, самые длинные обслуживаются первыми, не более 128 запросов.

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

3.2 TileOPs: FP8 lightning indexer переведён на потребление двумя warpgroup (#2666)

Объединено 10-20 22:41, 6 файлов +448/-2. Исходная реализация: один warpgroup умножает 64-key tile, затем выполняется epilogue длиной, равной произведению (каждый head умножается на key scale, FMUL и FADD раздельно); также имеется чистый проход с записью -inf за пределами окна со скоростью 2.5 TB/s — согласно описанию коммита, это занимает 425 микросекунд из 472 микросекунд chunk-prefill-kscale; в вызове декодирования одиночный query занимает один CTA и проходит все 16384 ключа за 289 микросекунд. Новый FP8LightningIndexerWSKernel (SM90) использует TMA warp для потоковой передачи 64-key tile двум consumer warpgroup поочерёдно, на каждый tile один m64n256 FP8 WGMMA, epilogue потребителя свёрнут в csrc helper.

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

3.3 TileOPs: декодирование MLA идёт по программе смещённых качелей (#2665)

Объединено 10-20 22:39, 7 файлов +881/-590. MLADecodeSeesawKernel (SM90, 512 + 64 latent, один KV head) перенимает MLADecodeWithKVCacheFwdOp, следуя идее смещённых качелей MLAPagedFwdOp, по запросу через TMA читает k и k_pe; тело программы перенесено в kernels/attention/mla/seesaw.py, предоставляет два входа — постраничный и плотный, раздельная сборка MLAPagedFwdOp даёт то же хранилище, что и раньше. Ключи одного запроса нарезаются на min(sm_count // blocks, tiles // 2) сегментов и объединяются через log-sum-exp; плотный сегмент сохраняет частичные результаты согласно выходному dtype, 16-битный выход временно сохраняется в собственном key-буфере warpgroup, затем сбрасывается на диск через TMA.

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

3.4 TileOPs: верхняя граница тензорных ядер заменена на оболочку множества операндов (#2661)

Объединено 10-20 22:27, 4 файла +125/-82. Ночной прогон 39299002056 пометил GemmW4A8FwdOp как достигающий 105% и 108% верхней границы tensor_core.fp8; после проверки выяснилось, что сама граница занижена: исходная калибровка брала скорость cuBLAS на случайных операндах, а при лимите мощности 700 Вт значения операндов напрямую определяют тактовую частоту, поэтому kernel с энергопотреблением ниже случайного cuBLAS может работать выше этой скорости. Теперь gemm_throughput.py одновременно замеряет время на полностью нулевых операндах, записывает лучшую скорость каждого множества операндов в calibration_operands и берёт среди них максимум как калибровочное значение — в соответствии с уже существующим правилом оболочки для калибровки HBM; h200.yaml переразмечен соответственно.

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

3.5 TileOPs: бэкенд-тесты и базовые линии torch-compile могут выполняться на не-CUDA устройствах (#2658)

Объединено 10-20 13:17, 10 файлов +241/-48, исправляет три сбоя на стороне TileOPs, выявленных прогоном MTT S5000 в #2595. Во-первых, torch_musa переименовывает PrivateUse1, поэтому reader, зарегистрированные по "privateuseone" в test_backend.py и test_kernel_dispatch.py, не могут его прочитать; используется torch.device("privateuseone").type; во-вторых, torch_musa печатает содержимое при импорте, из-за чего проверка импорта, сравнивающая весь stdout, завершается неудачей; добавлен tests/fresh_interpreter.py, читающий только одну строку с маркером, и семь таких проверок единообразно переключены; в-третьих, путь torch-compile без Triton в среде вне CUDA снова доступен.

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

3.6 TileOPs: SM89 одна строка/один столбец преобразуются в векторный GEMM (#2660)

Объединено 10-20 10:35, 3 файла +87/-11. Существующая реализация GEMV на CUDA-core открыта для SM89: левый операнд — одна строка с транспонированными весами, правый операнд — один столбец без транспонирования, оба случая переиспользуют одно и то же тело kernel, точка пересечения двух строк по-прежнему калибруется по SM90. Одновременно кандидаты для тюнинга фильтруются по бюджету разделяемой памяти выбранного устройства, максимальный принятый staged ring составляет 96 KiB, сгенерированная программа дополнительно объявляет не более 192 байт временного буфера для редукции; если входные данные одной строки/одного столбца непрерывно начинаются с невыровненного смещения элемента, используется существующий helper выравнивания.

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

3.7 TileOPs: три исправления контракта кэша builder и тюнинга (#2662 / #2663 / #2664)

Три коммита слиты одновременно 20.10 в 13:16. #2662 добавляет 32-элементный LRU-кэш в SM89 FP8 grouped GEMM builder, в соответствии с существующим SM90 builder и соглашениями репозитория; повторные вызовы с одинаковыми четырьмя параметрами формы переиспользуют JIT-фабрику; #2663 добавляет такой же 32-элементный LRU в два builder’а плотного GLA subchunk prefill, ключ кэша сохраняет форму, scale, dtype активации и dtype gate, конфигурация потоков launch оставлена параметром внутренней фабрики; #2664 восстанавливает контракт 12 входов обычного tuning supplier для paged GQA — ранее ожидалось 8 входов, а возвращалось 9, терялись три scale query/key/value, а также удалён неподдерживаемый ключ-кандидат num_split. Все три коммита заявляют отсутствие претензий на производительность на стороне устройства, вычислительное тело не изменено.

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

3.8 TileOPs на рассмотрении 11 PR: три линии SM89 развиваются параллельно

Обновлённые в окне мониторинга PR на рассмотрении сосредоточены на SM89: top-k выбор в NSA sparse attention (#2675 расширяет диапазон выбора до 2048 chunk, #2674 режет tile по group48, #2673 дополняет выходные слоты top-k для каждого запроса), paged DSA (#2669 первая поддержка, #2670 переиспользование scale и общих value), W4A16 узкий tile (#2671 более узкий tile для малого M, #2672 включает также M64). Кроме того, #2667 убирает tuning tile GQA, выходящий за пределы shared memory, #2654 сокращает объём работы W4A8 при малом числе token, #2515 добавляет attention sinks в плотный GQA, #2490 планирует вывести из эксплуатации старый forward GQA с append в KV cache.

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

3.9 TileFoundry: анонс 0.0.3 и рефакторинг тестов (#247 / #246)

#247 добавляет в раздел Latest News README три записи за 10/2026: релиз 0.0.3, пример prefill Qwen3.8-27B (#229 / #236), пример prefill Apple AMX (#244), а две записи за 08/2026 переносит в свёрнутый список Earlier, чтобы сохранить пять самых свежих. В сообщении коммита указано, что README — это длинное описание для PyPI (readme = "README.md" в pyproject.toml), поэтому анонс должен быть зафиксирован до проставления тега. #246 — рефакторинг тестов на 51 файл +655/-366: локально и в CI переиспользуются ccache и кэш torch extension (изоляция кэша между PR, статистика по job), два теста переведены на сравнение с эталоном CuTe, скомпилированным на стороне хоста, для проверки раскладки MMA вместо рукописного эталона и тысяч параметризованных случаев, torch переведён на ленивый прокси-импорт с разрешением dtype по требованию.

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

3.10 Адаптеры: TileKernels и FlashQLA без пушей в окне мониторинга

В TileKernels от DeepSeek и FlashQLA от Qwen в данном окне мониторинга коммитов нет; последние соответственно — слияние релиза v2.1.0 от 10-14 и слияние исправления SM120 от 10-14. Число PR на рассмотрении в обоих репозиториях равно нулю.

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

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

4.1 Репозиторий сообщества: создан китайскоязычный систематический учебник tilelang-tutorial

huandu/tilelang-tutorial создан и запушен 20.10, позиционируется как «систематический китайскоязычный учебник по TileLang от нуля до единицы», охватывает написание операторов, механизмы времени компиляции, а также два бэкенда — NVIDIA и Ascend. Ранее китайскоязычные материалы были представлены в основном отдельными блогами и переводами официальной документации; систематизированный учебный репозиторий появился впервые.

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

4.2 Репозиторий сообщества: создан ops-tilelang сообщества CANN

Atlas-Ascend-Fork/ops-tilelang создан 20.10 и позиционируется как репозиторий высокопроизводительных операторов TileLang для NPU Ascend в сообществе CANN, аккумулируя операторные активы по принципу «централизованное управление реализациями Kernel, автоматизированные тестовые примеры, стандартные рабочие нагрузки и базовые показатели производительности»; цель — сделать операторы TileLang на стороне Ascend переиспользуемыми, проверяемыми и способными к устойчивой эволюции. Это контрастирует с прежним состоянием, когда операторы для Ascend были разбросаны по examples/ и cann-bench.

Источник: https://github.com/Atlas-Ascend-Fork/ops-tilelang

4.3 Инструментарий: активность devkit и расширения VS Code, lower trace целиком выносится из основного репозитория

easy-tilelang/tilelang-devkit (создан в июле) и tilelang-devkit-vscode (создан 19.10, пуш 20.10) проявляли активность в этом окне мониторинга; последний позиционируется как средство исследования lowering pass, различий IR и JIT-компиляции CPU C в VS Code. В соответствии с этим рассматриваемый в основном репозитории PR #3583 (13 файлов, +106/-4390) целиком выносит из основного репозитория реализацию lower trace, специализированную конфигурацию окружения и соответствующие тесты, оставляя лишь общий интерфейс инструментирования компилятора, чтобы инструмент трассировки развивался независимо. Это продолжение линии основного репозитория с июля — «вынесение отладочных инструментов вовне».

Источники: https://github.com/easy-tilelang/tilelang-devkit-vscode, https://github.com/tile-ai/tilelang/pull/3583

4.4 Репозитории сообщества: tilelangGym и эксперименты с бэкендом Hexagon

PTO-ISA/tilelangGym (создан 14.10, пуш 20.10) предназначен для проверки чувствительности различных benchmark и микробенчмарков к производительности при PTO lowering, синхронно с продвижением линии PTO для Ascend. SoilSoi1/tilelang-hexagon создан 20.10 и представляет собой эксперимент сообщества с бэкендом Hexagon; в очереди на рассмотрение основного репозитория ему соответствует #2155 [feat] Add Hexagon HMX backend support (автор Calaweh, с мая обсуждались дальнейшие пути в рамках «развязки бэкендов»). Кроме того, sablin39/tilelang-cuda-skills (146 звёзд) в течение окна продолжал обновляться — это набор навыков CUDA для написания и отладки TileLang.

Источники: https://github.com/PTO-ISA/tilelangGym, https://github.com/SoilSoi1/tilelang-hexagon, https://github.com/tile-ai/tilelang/pull/2155

4.5 Документация и версии: четыре автообновления сайта документации, VERSION по-прежнему 0.1.15

tilelang.github.io в течение окна автоматически обновлялся четыре раза (20.10 12:16, 13:01, 16:08 и 21.10 00:49), каждый раз по триггеру слияния в основной репозиторий. VERSION основного репозитория по-прежнему 0.1.15, последняя веха в README остановилась на бэкенде Ascend 950 от 30.09; статистическая ветка TileOPs обновилась с последним слиянием 20.10 22:42, новых tag не обнаружено; TileFoundry 0.0.3 ещё ожидает простановки tag.

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

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

Главной линией этого окна стала корректность компилятора, а не расширение функциональности. Из 9 слияний в основном репозитории 5 — это устранение тихих ошибок: векторизация предполагает цикл с нуля (#3579), пропуск обнаружения WAW в разделяемой памяти (#2821), MMA emitter молча сокращает K-атомы (#2818), сохранение размерности в fragment reduce падает (#2815), предикат не следует за расширением разрядности индекса (#3550). Общее у этих дефектов то, что они не выдают ошибку, а лишь дают неверный результат или пропускают барьер — это категория, которую пользователям труднее всего локализовать в производстве. В тот же день #2661 в TileOPs также исправляет «метрику верхней границы производительности», что показывает: проект ставит «достоверность цифр» на один уровень с «скоростью работы» как инженерную цель.

Линия Ascend переходит от дополнения примеров к тренировочной стороне. Из пяти слияний три — обратные или рекуррентные операторы (indexer_bwd, situ_glu_grad, LSTM/GRU), причём situ_glu_grad напрямую согласуется с гейтовой активацией Kimi K3 Stable LatentMoE и приводит данные перекрёстной проверки с aclnn на 910B3. Вместе с выравниванием #3577 по выбору конфигурации DeepGEMM-Ascend и epilogue охват на стороне Ascend расширяется от прямого прохода инференса к обратному проходу тренировки, рекуррентным состояниям и гейтованию MoE больших моделей.

PTO становится вторым путём выполнения помимо AscendC. Находящийся на рассмотрении #3512 — это фреймворк векторизации на 21 тыс. строк (T.Parallel → PTO VMI, целевая платформа Ascend 950); #3516 предлагает сделать tvm_ffi исполнительным бэкендом по умолчанию для цели pto, а вместе с группой исправлений копирования и предикатов #3530, #3573, #3581, #3582 объём работ по линии PTO уже заметно превышает масштаб точечного патча. Основной репозиторий одновременно через #3511 дополняет возможности локального копирования L0 на стороне AscendC — оба пути развиваются параллельно, а не в отношениях замены.

Инструментарий выплёскивается в отдельные репозитории. За один день появились три новых репозитория (учебник на китайском, репозиторий операторов сообщества CANN, эксперимент с бэкендом Hexagon) и набор расширения VS Code, а основной репозиторий к тому же целиком перенёс lower trace в devkit (#3583, удалено 4390 строк). Скорость расширения на стороне сообщества уже превышает темпы модульного деления основного репозитория, что согласуется с направлением рефакторинга «мультибэкендовых языковых диалектов» июля: основной репозиторий сужается до семантики и IR, а инструменты и эксперименты с бэкендами выносятся наружу.

Явная формализация критериев CI заслуживает внимания. #3569 привязывает регрессию производительности к H100 и быстро падает при ошибочной маркировке карты, #3564 разносит сборку и тестирование Ascend на пошаговые этапы в рамках одной проверки check, #2661 перемаркирует верхний предел тензорных ядер с помощью оболочки множества операндов — все три изменения «делают числа сопоставимыми», напрямую влияя на порог допуска последующих PR.

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

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

Источник Результат проверки
tile-ai/tilelang (main) За окно 8 коммитов, 9 слияний; VERSION по-прежнему 0.1.15; на рассмотрении 167, из них 163 обновлены в окне
tile-ai/TileOPs (main) За окно 9 коммитов, 9 слияний; на рассмотрении 16, 11 обновлены в окне; статистическая ветка обновлена в 22:42
tile-ai/tilelang-ascend (ascendc_pto) За окно 5 коммитов, 5 слияний; на рассмотрении 236, 7 обновлены в окне
tile-ai/tilelang-hygon (main) За окно 1 коммит, 1 слияние; на рассмотрении 3
tile-ai/TileFoundry (main) За окно 2 коммита, 2 слияния; на рассмотрении 1; 0.0.3 ожидает простановки tag
tile-ai/tilelang-metax (dev) / musa / mlir-ascend Во всех трёх репозиториях за окно коммитов нет
tile-ai/TileRT / tilescale / DeepStack Во всех трёх репозиториях за окно коммитов нет (последние соответственно 09-24, 08-03, 10-05)
deepseek-ai/TileKernels / QwenLM/FlashQLA Коммитов нет, последние остановились на 10-14
Документационный сайт tilelang.github.io За окно 4 автоматических обновления
Репозитории сообщества и третьих сторон Новых репозиториев 3 (tilelang-tutorial, ops-tilelang, tilelang-hexagon), активных 3 (devkit, devkit-vscode, tilelangGym)
Google News / Hacker News / arXiv Множественные запросы на китайском и английском не выявили материалов, напрямую связанных с TileLang; последняя релевантная статья на arXiv по-прежнему от 2026-07-24

Дополнительное пояснение: поисковый интерфейс для очереди на рассмотрении выдал updated_at, где 64 записи приходятся примерно на 10-20 11:30; при выборочной проверке трёх из них в соответствующий момент не было ни комментариев, ни ревью, ни событий пуша в ветку — это определено как массовый пересчёт на стороне платформы и не учтено в активности разработки; количество находящихся на рассмотрении определяется по счётчику is:open каждого репозитория.

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

2658 Тестирование бэкенда и базовая линия torch-compile выведены из CUDA — https://github.com/tile-ai/TileOPs/pull/2658