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

Источники: порепозиторная проверка организации GitHub tile-ai (29 репозиториев, пуши в 10 за окно) + поиск истинных значений по слитым и находящимся на рассмотрении (18 слито за окно, 57 обновлений на рассмотрении) + 5 новых issue + поштучная проверка коммитов основной ветки (tilelang 5, TileOPs 9, TileFoundry 1, tilelang-ascend 2, tilelang-musa 1) + проверка неосновных веток (ветка feat/hcu-scaled-cast в tilelang-hygon, 3 коммита) + проверка страниц релизов (tilelang / TileOPs / TileFoundry, новых версий за окно нет) + два репозитория внедряющих сторон + проверка сайта документации и ночного репозитория + четыре группы запросов Google News на китайском и английском (через прокси) + Hacker News + arXiv. Специальный обход четырёх источников по RISC-V за это окно без изменений.

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

  • Главное за сегодня: TileOPs за день слил девять коммитов, одновременно продвинув три вещи — «чтение широковещательного операнда на месте, запуск в одну волну, выбор тайла»; помимо пяти коммитов в основном репозитории, обнаружен дефект CUDA с тихой ошибкой, существующий с 0.1.12; Ascend исправил segmentation fault, а также было указано, что при threads=2 sin/cos даёт 97,1% неверно вычисленных элементов
    1. Основные результаты по проекту
      • 1.1 Основная линия: поддержка преобразования TF32 с round-to-nearest-even, раздельная диагностика SM90 и SM100 (#3714) (11-05)
      • 1.2 Основная линия: host-mapped буферы можно явно аннотировать, CPU-тензоры больше не отклоняются проверкой устройства (#3802) (11-05)
      • 1.3 Основная линия: cast в скалярных условиях больше не ограничивает ширину вектора данных (#3759) (11-05)
      • 1.4 Основная линия: rlimit для Z3 настраивается по kernel, по умолчанию по-прежнему сто тысяч (#3752, зависит от tile-ai/tvm #91) (11-05)
      • 1.5 Основная линия: вложенный T.macro переиспользует шаблон преобразования, накладные расходы на повторную компиляцию во фронтенде сокращаются (#3760) (11-05)
      • 1.6 Новый issue: после условной скаляризации отсутствует __syncthreads, с 0.1.12 по 0.1.15 тихо генерируется неверный CUDA (#3808) (11-05)
      • 1.7 Новый issue: без warp specialization T.Pipelined перезагружает stage перед ожиданием tcgen05 (#3800) (11-05)
      • 1.8 На рассмотрении: полный черновик миграции на нативный Apache TIRx (#3803); общий transpose переведён на ldmatrix/stmatrix (#3809) (11-05)
    1. Адаптация под множество бэкендов
      • 2.1 Ascend: 64-битный dtype в bitwise_rshift заменён на явную ошибку, больше не segmentation fault (#1954) (11-05)
      • 2.2 Ascend: навыки производительности A2/A3 приведены в соответствие с рабочим процессом Ascend950 (#1932) (11-05)
      • 2.3 Ascend: при threads=2 tmp workspace безусловно уменьшается вдвое, sin/cos на тайлах с косвенной адресацией даёт 97,1% неверно вычисленных элементов (#1991 / #1992) (11-05)
      • 2.4 Hygon: три исправления фазы LDS в ветке, и впервые открыты два PR (#33 / #34) (11-05)
      • 2.5 Moore Threads: обновление документации README, завершение тишины с 09-30 (11-05)
      • 2.6 MetaX и mlir-ascend: у MetaX третий цикл подряд без пушей; в mlir-ascend три оператора TLIR на рассмотрении (11-05)
    1. Экосистема и внедряющие стороны
      • 3.1 TileOPs: девять слитых коммитов, чтение широковещательного операнда на месте, однократное заполнение GEMM, кэш раскладки для reduction (11-05)
      • 3.2 TileOPs: переписаны правила оценки roofline, три невозможные строки пропускной способности вернулись в разумный диапазон (#2830) (11-05)
      • 3.3 TileOPs: два слитых коммита с интервалом в 2,5 минуты привели к падению GPU Smoke в main, исправлено в тот же день (#2837) (11-05)
      • 3.4 TileOPs на рассмотрении: трёхэтапный фактор тайлинга, одноэлементное сжатие TN, TMA-загрузка масштабирования для FP8 grouped GEMM (#2836 / #2838 / #2839) (11-05)
      • 3.5 TileFoundry: слит план синхронизации WGMMA (#286), семантика завершения частичной редукции на рассмотрении (#288) (11-05)
      • 3.6 TileSight: исправлено объявление возможностей WGMMA для H20, первый PR в репозитории (#1) (11-05)
      • 3.7 Внедряющие стороны и третьи стороны: TileKernels, FlashQLA без пушей; появились два новых сторонних репозитория (11-05)
    1. Сообщество, обучающие материалы и мероприятия
      • 4.1 Сайт документации обновлён трижды; в ночном репозитории нет коммитов в основную ветку; ежедневные тесты Ascend — все 3030 пунктов пройдены (11-05)
      • 4.2 Новостная сторона: четыре группы запросов, Hacker News и arXiv — попаданий за окно нет (11-05)
    1. Наблюдения за трендами
  • Приложение, материалы и пояснения к проверке

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

Самый плотный объём выхода за этот день пришёлся на библиотеку операторов, а не на компилятор. TileOPs за период с 21:06 до 22:46, чуть более чем за час, последовательно влил шесть коммитов, плюс два коммита в 14:01 и 21:31 — девять за день, все с тегами perf, fix, ci. Одновременно продвигаются три вещи: первое — считывание транслированных операндов на месте, пять категорий операторов lerp, where, clamp, masked_fill больше не разворачивают транслированный тензор со stride 0 в форму выхода предварительно; измеренное время выполнения транслированной строки снизилось до 0,13–0,44 от значения основной ветки (#2832); второе — одноволновый запуск, унарные поэлементные операции при условии «в теле не более двух дешёвых инструкций, общий объём чтения-записи не превышает трёх четвертей L2» сжимаются в одну волну из 256 потоков, и попутно исправлена старая проблема ReLU, возвращающего 0 для NaN (#2829); третье — селектор тайлов, для NT GEMM с малым M создан тайл нативной ширины, ширина которого способна заполнить одну волну, decode-down из отстающего превратился в опережающий самый быстрый сравниваемый вариант на 3% (#2834). Общая черта трёх вещей — не «быстрее», а то, что каждый коммит сопровождается воспроизводимой A/B-таблицей и отмечает ту строку, в которой он всё ещё отстаёт.

Второе — вопрос достоверности, причём охватывающий два репозитория. В основном репозитории открыт новый issue #3808 с сообщением: с 0.1.12 по 0.1.15 после скаляризации, связанной с условиями, в epilogue разделяемой памяти пропущен один __syncthreads(), компилятор не сообщает об ошибке, генерируемый CUDA неверен; автор отчёта приложил полный скрипт воспроизведения и четыре номера версий — это относится к «тихому неверному коду», а не к деградации производительности. В том же окне #3800 сообщает о другом тихом неверном коде: при выключенной warp specialization T.Pipelined перезагружает следующую стадию до ожидания барьера в T.tcgen05_gemm, в результате operand, читаемый MMA, уже перезаписан. Эти два сообщения вместе со влитым в основной репозиторий в тот день #3759 как раз образуют контрастную пару — #3759 исправляет «int64 cast в скалярном условии сжимает ширину вектора данных с 16 lane до 4 lane», на одной и той же границе скаляризации/векторизации: с одной стороны исправляют ширину, с другой — пропускают синхронизацию.

Третье — на стороне Ascend. В tilelang-ascend влит #1954, превращающий ошибку сегментации на этапе компиляции для T.tile.bitwise_rshift на int64/uint64 в явный ValueError, разрешены только int16/uint16/int32/uint32; это не добавление поддержки 64-битного правого сдвига, а превращение «компилятор падает» в «компилятор отказывает». В тот же день в этом репозитории открыты два issue по той же теме (#1991, #1992): сообщается, что на 910_9362 / CANN 9.0.0 при threads=2 (C:V=1:2) AscendVidReduction безусловно уменьшает вдвое все распределения UB внутри ядра, включая внедрённый tmp workspace, в результате чего sin/cos при применении к skip-set buffer с косвенной адресацией вычисляют неверно 3978/4096 (97,1%) элементов, а sigmoid/round нарушают контракт AscendC calCount <= tmp, и запись за границы попадает в scratch позади layout. Рядом — отчёт о ежедневном тестировании #1988 за тот же день: 3030 пунктов пройдено, проходимость 100%. Эти два набора чисел вместе говорят об одном — существующее тестовое покрытие не достигает комбинации «threads=2 плюс косвенная адресация», и проходимость не может служить заменителем показателя корректности.

Четвёртое — движения в экосистеме. Репозиторий Hygon из состояния «на функциональной ветке накоплено четыре коммита, PR не открыт» перешёл к официальному ревью: ветка feat/hcu-scaled-cast в тот день получила ещё три исправления фазы LDS (адаптация swizzle по умолчанию под 64 bank, восстановление планирования swizzle на потребительской стороне MN-среза, чтение с учётом layout для K-раздела), и открыты два PR — #33 (scaled cast) и #34 (lane swap и распространение ограничений в reduction). Moore Threads прервала полное молчание с 09-30, в основную ветку попал один коммит с обновлением README по установке и разработке операторов. Вместе с появлением первого PR в репозитории TileSight (исправление заявления о возможностях WGMMA для H20) и черновика миграции TIRx в основном репозитории #3803 можно видеть, что на организационном уровне одновременно разворачиваются три линии: отечественные бэкенды переходят с собственных веток в публичное ревью, моделирование производительности движется от статьи к репозиторию, а основа компилятора готовится заменить vendored TVM.

1. Ход работ по ключевым проектам

1.1 Основная ветка: поддержка преобразования TF32 с round-to-nearest-even, раздельная диагностика SM90 и SM100 (#3714)

Слито в 12:09. Ранее T.cast в tfloat32 имел только путь округления TF32 по умолчанию с округлением к ближайшему; для ties-to-even приходилось самостоятельно импортировать CUDA helper. Этот коммит добавляет нативный lowering для round="rn": несатурирующая форма выдаёт cvt.rn.tf32.f32, сатурирующая форма выдаёт cvt.rn.satfinite.tf32.f32, при этом обходится конструктор округления TF32 по умолчанию и напрямую сохраняется преобразованный битовый шаблон. Покрыты оба пути — скалярный и векторный; доступ к lane для TF32 идёт через упакованный носитель. Диагностика архитектуры разделена по формам: несатурирующая форма требует SM90 и выше, сатурирующая форма требует SM100 и выше; при невыполнении выдаётся ошибка на этапе компиляции; поведение cast без явного указания rounding не меняется. Докладчик — заказчик, закрывший #3713; это относится к категории сходимости интерфейсов «пользователю больше не нужно приносить свой helper».

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

1.2 Основная линия: host-mapped буферы можно явно аннотировать, CPU-тензоры больше не отклоняются проверкой устройства (#3802)

Создано и слито в тот же день (коммит в 12:08, слияние в 16:02). CUDA kernel может напрямую обращаться к host-памяти, отображённой в адресное пространство GPU, но исходная проверка устройства в TileLang предполагала, что параметры-тензоры находятся на целевом устройстве, поэтому такие CPU-тензоры отклонялись до запуска. Новый T.annotate_host_mapped_buffers(*buffers) записывает индексы аннотированных параметров в атрибут tl.host_mapped_params PrimFunc, поддерживает несколько буферов, объединяет повторные объявления и может переживать сериализацию IR и восстановление из кэша kernel. Оба launcher’а — Cython и TVM-FFI — принимают CPU-тензоры для аннотированных параметров, остальные проверки dtype, формы, ранга и шага остаются прежними; проверки параметров на стороне устройства по-прежнему соответствуют текущему устройству исполнения, а для host-тензоров номер CUDA-устройства не предоставляется. В runtime добавлен нативный helper, который проверяет статус регистрации через cuMemHostGetFlags и разрешает видимый устройству адрес через cuMemHostGetDevicePointer; обычная странируемая память по-прежнему отклоняется.

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

1.3 Основная линия: cast в скалярных условиях больше не ограничивает ширину вектора данных (#3759)

Слито в 19:27, в метках присутствует ascend. Причина: #3734 после DecoupleTypeCast корректно ужесточил ограничение ширины для data cast, но заодно учёл cast в условиях: фрагмент копирования BF16, защищённый T.assume(n % 16 == 0), из-за наличия T.Cast("int64", ...) в условии был сжат с 16 lane до 4 lane на b5cb02633, даже если условие согласовано на 16 lane. Исправление состоит в переиспользовании уже имеющегося в planner анализа согласованности условий и расширении исключения по ширине cast с индексов буферов на условия: условия операторов if, tirx.if_then_else и Select со скалярными предикатами обращаются отдельно, ветвления по-прежнему ограничены; на целях с векторными предикатами уровня lane ограничение на cast в условиях сохраняется. Группировка rbits для stochastic rounding также сохраняется при пропуске ограничения ширины. Этот коммит и два новых issue в окне мониторинга попадают в одну и ту же область, что показывает: стык скаляризации и векторизации — зона повышенной частоты текущих дефектов.

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

1.4 Основная линия: rlimit Z3 настраивается для каждого kernel, по умолчанию по-прежнему сто тысяч (#3752, зависит от tile-ai/tvm #91)

Слито в 23:05, предварительно tvm #91 в fork (слито в 23:03). Изменение превращает rlimit решателя Z3 из жёстко заданных 100 000 в настраиваемое для каждого kernel значение: без настройки сохраняется прежнее значение по умолчанию. Автор формулирует прямо — «это косвенно исправляет множество проблем, связанных с прувером Z3». Смысл таких изменений в том, чтобы превратить «превышение бюджета доказательства» из нерегулируемого поведения окружения в настраиваемый параметр компиляции; это два подхода в одном направлении с вчерашним ужесточением нижней границы периода AutoSchedule на стороне Ascend (снижение времени решения Z3 на 10,04%): один меняет начальную точку поиска, другой — бюджет поиска.

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

1.5 Основная линия: вложенные T.macro переиспользуют шаблон преобразования, накладные расходы повторной компиляции фронтенда сокращаются (#3760)

Совпадает с 23:38. Проблема в том, что при каждом запуске внешней функции вложенное определение T.macro заново выполняет разбор исходного кода, переписывание AST и компиляцию Python, даже если сам kernel уже попал в кэш. Решение — ввести ограниченный внутрипроцессный кэш, хранящий только не зависящие от контекста результаты преобразования (несвязанный код и переписанный исходный код), а каждый экземпляр макроса затем привязывает текущие globals и значения замыкания; преобразования, зависящие от контекста, идут по исходному пути, обёрнутые функции обходят кэш. Точки вызова макросов и кэш kernel не меняются. Регрессия покрывает новые замыкания, globals, захваченные буферы, аннотации, контекстно-зависимый откат и жизненный цикл замыканий; проверка фронтенда на стороне CPU пройдена, выполнение на GPU локально не проверено — автор чётко обозначил эту границу.

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

1.6 Новый issue: после условной скаляризации отсутствует __syncthreads, версии с 0.1.12 по 0.1.15 молча генерируют ошибочный CUDA (#3808)

Открыт в 21:14, это самый серьёзный отчёт о дефекте в данном окне мониторинга. Воспроизведение — kernel с epilogue на разделяемой памяти: версии 0.1.12, 0.1.13, 0.1.14 и 0.1.15 после условно-зависимой скаляризации теряют барьер __syncthreads(), компилятор не выдаёт никаких ошибок, и сгенерированный CUDA сразу считает неверно. Автор отчёта приложил полностью запускаемый скрипт воспроизведения на Python и окружение (Python 3.14.9, CUDA 2.15.0+cu134). Разрушительность таких проблем в том, что они охватывают четыре релиза без каких-либо предупреждений — с ними может столкнуться любой kernel, полагающийся на разделяемую память для передачи данных между стадиями, а пользователь видит лишь «неверный результат».

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

1.7 Новый issue: без warp specialization T.Pipelined перезагружает stage перед ожиданием tcgen05 (#3800)

Открыт в 08:44. Условие срабатывания — после отключения warp specialization (tl.disable_warp_specialized=True) запустить цикл с T.Pipelined и T.tcgen05_gemm: stage перезаписывается копированием следующей итерации ещё до mbarrier_wait_parity, и MMA читает уже переписанный operand. Автор отчёта привёл исходный код пониженного kernel и указал на отсутствие ограничения порядка между записью As и чтением tcgen05mma_ss. Это та же категория, что и #3808 — «сгенерированный код выглядит корректным, но фактически неверен», — с той разницей, что здесь проблема на пути tensor memory в Blackwell.

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

1.8 На рассмотрении: полный черновик миграции на нативный Apache TIRx (#3803); общий transpose переводится на ldmatrix/stmatrix (#3809)

#3803 — самый весомый пункт на рассмотрении в этом выпуске: заменить вендоринговый компилятор TVM на независимый Apache TIRx, перенеся фронтенд, преобразования IR, генерацию кода, адаптацию JIT и упаковку полностью на нативную объектную модель, с фиксацией двух коммитов вышестоящих репозиториев — apache/tvm-tirx на 7e49a148 и apache/tvm-ffi на 3e906bc6; дерево зависимостей остаётся чистым вышестоящим, без внедрения патчей. Автор пометил это как draft, что само по себе является публичной подготовкой «смены основания».

#3809 (коммит в 01:27) решает другую конкретную задачу: T.transpose между двумя разделяемыми буферами всегда шёл через общий цикл SIMT, 8-битные и 16-битные тайлы не использовали пропускную способность разделяемой памяти, и в examples/deepseek_v32 FP8-разреженный MLA приходилось обходить ручным сохранением в 16x16 регистровый буфер; к тому же после того, как T.Pipelined добавил к исходной области ось stage, несовпадение ранга приводило к прямому сбою MakeIndices. Решение — на целях, поддерживающих stmatrix (SM90 и выше, исключая CuTeDSL), сначала пробовать warp-matrix lowering: для 16-битных тайлов каждый блок 8x8 использует один ldmatrix и один stmatrix.trans, для 8-битных тайлов по 16x16-байтовым блокам два ldmatrix.trans и __byte_perm для выбора и сборки четырёх строк, остальные области откатываются на исходный путь SIMT.

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

2. Адаптация множества бэкендов

2.1 Ascend: для 64-битного dtype bitwise_rshift теперь явная ошибка вместо segfault (#1954)

Объединено в 16:36 (ветка ascendc_pto). При передаче операндов int64/uint64 через публичный API T.tile.bitwise_rshift на A2/A3 компилятор на этапе OptimizeForTarget завершается с segfault. Исправление заключается в проверке dtype dst и src0 перед построением intrinsic: допускаются только int16/uint16/int32/uint32, для остальных выбрасывается ValueError с фактическим dtype и списком поддерживаемых; для величины сдвига scalarValue dtype выражения не ограничивается; Buffer и BufferRegion единообразно проходят через _underlying_buffer(). Добавлены диагностические тесты прямого API и регрессионный тест на падение в дочернем процессе. Автор явно указал границы: в этот раз поддержка 64-битного правого сдвига не добавляется, и путь прямого построения низкоуровневого intrinsic не изменяется. Связано с #1720.

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

2.2 Ascend: выравнивание навыков производительности A2/A3 с рабочим процессом Ascend950 (#1932)

Смержено в 15:36, чисто рефакторинг навыков и процессов: навыки настройки производительности A2/A3 выровнены и унифицированы со структурой рабочего процесса A5 (Ascend950), добавлены два навыка — определение типа чипа и получение информации о чипе. Само по себе не имеет функциональной ценности, но означает, что «руководство по эксплуатации» настройки производительности начинает ветвиться по поколениям чипов — один и тот же набор skill должен уметь определять, на каком поколении он запущен, иначе рекомендации по тюнингу будут ошибочно применяться между поколениями.

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

2.3 Ascend: при threads=2 tmp workspace безусловно уменьшается вдвое, sin/cos на tile с косвенной адресацией вычисляются неверно для 97,1% элементов (#1991 / #1992)

В 17:03 и 17:25 подряд открыты две issue по одной теме, заголовок и текст совпадают (по содержанию можно рассматривать как одну, отправленную дважды). Среда отчёта — 910_9362 / CANN 9.0.0, threads=2 (C:V=1:2). Корневая причина в AscendVidReduction: он безусловно уменьшает вдвое все выделения UB внутри kernel, включая tmpspace, инжектируемый InjectTmpBuffer, тогда как форма skip-set buffer (косвенная адресация с загрузкой UB по индексу GM, например T.copy(KV[b_i, idx_ub[bi_i], g_i, :D], kv_ub)) сохраняется целой и изначально не участвует в уменьшении вдвое для vid. Последствия двух уровней: sin/cos вычисляются напрямую неверно (на upstream HEAD ошибочны 3978/4096 элементов, 97,1%); sigmoid/round нарушают контракт AscendC calCount <= tmp — численно случайно корректны, но запись за границы попадает в scratch позади раскладки, и при изменении раскладки это тихое загрязнение данных. Автор отчёта указывает, что это обнаружено при review #1897, а исправление #1897 покрывало только семантику стороны count (пропуск слота workspace через HasModifiedBuffer), сторона workspace по-прежнему не обработана.

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

2.4 Hygon: три коммита исправления фаз LDS в ветке и первые два PR (#33 / #34)

Основная ветка по-прежнему на 11-02, все три коммита за день находятся в ветке feat/hcu-scaled-cast, автор — со стороны Hygon (домен коммитов hygon.cn):

  • 10:06 fix(hcu): adapt default LDS swizzle to 64-bank phases
  • 14:04 fix(hcu): restore LDS swizzle planning for MN slice consumers
  • 19:49 fix(hcu): support K-partitioned layout-aware LDS reads

Важнее изменение статуса review: в 09:08 открыт PR #33 (scaled cast, то есть основное содержимое ветки), в 19:52 открыт PR #34 (оптимизация XOR shuffle для reduction с помощью DPP и lane swap, плюс исправление «при выводе раскладки сначала распространить известные ограничения, затем распределить свободные кандидаты»). В прошлом отчёте было написано «четыре коммита в функциональной ветке, PR ещё не открыты»; в этом окне эти два PR продвинули Hygon от собственной ветки к публичному review — первым среди отечественных бэкендов, прошедшим этот путь.

Источники: https://github.com/tile-ai/tilelang-hygon/pull/33, https://github.com/tile-ai/tilelang-hygon/pull/34, https://github.com/tile-ai/tilelang-hygon/commit/07451160e5aa758419201adb598725cbef70e3a4

2.5 Moore Threads: обновление документации README, завершение молчания с 30.09 (05.11)

19:46 в основную ветку попал коммит docs: update README installation and operator development references — первый коммит в этом репозитории с 30.09. Одновременно обновились временные метки шести оставшихся PR (обновления зависимостей среди #1–#7, подмодуль TileOPs, документ с критериями доказательств интеграции MUSA), но их содержимое — это старые предложения с апреля по август, изменений кода в этом окне мониторинга нет. Оценка должна быть консервативной: коммит документации говорит о том, что репозиторий кем-то поддерживается, но недостаточен для доказательства возобновления разработки бэкенда; настоящий сигнал — появление последующих коммитов с операторами или lowering.

Источники: https://github.com/tile-ai/tilelang-musa/commit/4cac45b859108a86ac9a0b1cbfdeb33da7711697

2.6 MetaX и mlir-ascend: у MetaX третий подряд период без пушей; в mlir-ascend на рассмотрении три оператора TLIR (05.11)

Последний пуш tile-ai/tilelang-metax по-прежнему 26.10 — третье подряд окно мониторинга молчания. В основной ветке tile-ai/tilelang-mlir-ascend новых коммитов нет, но продвигаются три PR с операторами на диалекте TLIR: #203 добавляет в TLIR tl.sub/mul/div/max/min и дополняет lowering из TL в HIVM, обновлён в 01:42; #211 добавляет tl.shl, #212 добавляет tl.shr, оба коммита отправлены в 19:44 и 20:58 соответственно. Суть этой линии — послойно перенести семантику TileLang в диалект MLIR; темп медленный, но направление единое.

Источники: https://github.com/tile-ai/tilelang-metax、https://github.com/tile-ai/tilelang-mlir-ascend/pull/203、https://github.com/tile-ai/tilelang-mlir-ascend/pull/211、https://github.com/tile-ai/tilelang-mlir-ascend/pull/212

3. Экосистема и стороны внедрения

3.1 TileOPs: девять слияний — чтение операндов broadcast на месте, заполнение GEMM одним волновым фронтом, кэш раскладки reduction (05.11)

Девять слияний сосредоточены в двух интервалах: 14:01 и с 21:06 до 22:46; рассмотрим основную линию по каждому:

  • #2832 Чтение широковещательных операндов на месте: пять категорий операторов LerpTensorFwdOp, WhereFwdOp, ClampTensorFwdOp, MaskedFillScalarFwdOp, MaskedFillTensorFwdOp больше не разворачивают широковещательный тензор со stride 0 в форму выхода, а читают его напрямую по stride 0 в одном kernel. Один общий builder обслуживает все пять, формы операндов объединяются в один цикл вложенности, векторы до 16 байт остаются в одной строке внутреннего уровня, каждый оператор пишет только своё поэлементное тело. Время выполнения широковещательных строк составляет от 0,13 до 0,44 от основной линии, на уровне torch-compile или лучше; остальные поэлементные строки побайтно совпадают с основной линией.
  • #2829 Унарные поэлементные операции в одну волну: когда в теле не более двух операций, каждая из которых сводится к одной инструкции, суммарный объём чтения и записи не превышает трёх четвертей L2, а результат и операнды имеют одинаковую ширину, операция сворачивается в одну волну из 256 потоков. Попутно исправлена проблема ReLU, возвращающего 0 для NaN (теперь возвращается NaN: для 16 бит используется битовый шаблон torch, для float32 — канонический NaN). На H200 все 327 GPU-регрессий пройдены, 13 новых тест-кейсов прошли Compute Sanitizer memcheck.
  • #2834 Нативные по ширине тайлы для NT GEMM с малым M: ds-v3-decode-down (128x7168x2048, NT) на основной линии использует только 112 из 132 SM, тайл шириной 112 в 0.1.12 разбивается на семь n16 WGMMA. После перехода на отправку WGMMA собственной ширины тайла каждые 16 K decode-down опережает самый быстрый сравниваемый вариант на 3%, gate-up-b32 быстрее основной линии на 1,4% (всё ещё отстаёт от cuBLAS на 0,5–1%). Новый тайл fill включается только для форм с 96 < M <= 128 и N, делящимся на 8, структурный фактор 0,58 подбирается отдельно.
  • #2835 Кэш раскладки reduction: kernel редукции при каждом вызове заново строит перестановку осей и форму выхода, тогда как ReduceCall фиксирован. Теперь это вычисляется один раз на экземпляр kernel, с общим чистым shape-хелпером с существующим адаптером тензоров, без глобального кэша. Полное время горячего вызова для шести конфигураций VarMeanFwdOp FP32/BF16 снизилось на 4,33%–6,17%, сгенерированный CUDA-код и исходники хоста полностью совпадают — экономится стоимость подготовки вызова, а не время GPU.
  • #2833 Исправление CI: после того как GitHub объявил runner 2.336.0 устаревшим, все gpu-smoke встали в очередь, а entrypoint.sh с --disableupdate не может самообновиться, RUNNER_VERSION поднят с 2.336.0 до 2.338.2.

Источники: https://github.com/tile-ai/TileOPs/pull/2832, https://github.com/tile-ai/TileOPs/pull/2829, https://github.com/tile-ai/TileOPs/pull/2834, https://github.com/tile-ai/TileOPs/pull/2835, https://github.com/tile-ai/TileOPs/pull/2833

3.2 TileOPs: переписаны правила оценки roofline, три невозможные по пропускной способности строки вернулись в разумный диапазон (#2830)

Ночная задача 41873301494 выявила три физически невозможные строки roofline: два конфигурации gpt-oss-serve-window у GQAVarlenFwdOp показывали 7,48 и 7,50 TB/s, qwen3-30b-ep8-empty у GroupedGEMMFwdOp — 5,30 TB/s. Этот коммит переписывает правила оценки по принципу «сколько фактически прочитано за этот вызов»: для плотного, varlen и постраничного forward GQA K/V учитываются только для «строк key, видимых хотя бы одним query», q — для «строк query, видящих key», таблица RoPE — для фактически достигнутых key позиций усечения, масштабы FP8 — для запросов, видящих key; GroupedGEMM через новый grouped_gemm_roofline оценивает группу строки по batch_sizes; сжатие NSA, разреженный NSA, DSA decode, постраничный MHA и FP8 lightning indexer учитываются только в реально читаемых позициях, DSA также учитывает is_causal=False. После исправления три строки стали 0,62, 0,62 и 2,65 TB/s. Дополнительный выигрыш — устранено деление на ноль для пустого хранилища K/V и пустого пула страниц. Тестовых узлов добавлено 16.

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

3.3 TileOPs: два слияния с интервалом 2,5 минуты привели к сбою GPU Smoke в main, исправлено в тот же день (#2837)

В 22:46 слитый коммит сам по себе является записью о дисциплине. #2824 при слиянии в 21:06 изменил строку roofline для grouped GEMM так, что теперь требуется объявлять sizes и выводить из неё G; #2830 слит в 21:08, и в нём по-прежнему только G без sizes. Между двумя коммитами около 2,5 минут, CI у каждого зелёный, но после слияния в main GPU Smoke сразу выдал RowError: row misses ['sizes']. Исправление — заменить G: 2 у _grouped на sizes: [1, 0], байты в утверждениях (24, 32) не меняются. Причина таких инцидентов не в том, что кто-то ошибся, а в отсутствии post-merge контроля при параллельных слияниях: прохождение CI отдельного коммита не покрывает случай «семантического пересечения двух коммитов», помочь может только немедленный полный прогон после слияния.

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

3.4 TileOPs на рассмотрении: трёхстадийный коэффициент тайлов, одноэлементное сокращение TN, TMA-загрузка масштабов для FP8 grouped GEMM (#2836 / #2838 / #2839)

#2836 отдельно подбирает коэффициент (1,36) для трёхстадийных тайлов coop2_swap, поскольку по FLOP он не быстрее четырёхстадийного, но по байтам и числу K-шагов завышается примерно на 10%, в результате вытесняется тайл шириной 160 и ночной результат падает с 1,005–1,011 до 0,985–0,995 от QuACK. Коэффициент подбирается только на тех 42 off-manifest формах (плато 1,344–1,392), на сканировании 44 форм regret селектора снижается с 0,89% до 0,06% по геометрическому среднему, максимум — с 5,04% до 0,99%. #2838 исправляет то, что диспетчер считал K непрерывным измерением и потому отклонял допустимый вызов TN-раскладки с K=1 (A=[1,64], B=[1,64], вычисление AᵀB вполне можно было выполнить ядром cp.async), заменяя это откатом к проверке, что физические непрерывные измерения обоих операндов имеют ширину строки не менее 4 байт; добавлено 24 тестовых узла. #2839 (коммит в 00:56) заставляет FP8 grouped GEMM загружать масштабы A через TMA (четыре K-шага на блок, stage_scales 2, калибровка на трёх строках), а для нефузированных компактных grouped BF16/FP16 GEMM с K <= 1536 переключает на ping-pong тайл 64x256; FP8 qwen3-30b-ep8-skew-up поднимается с 0,973 до 1,015 от DeepGEMM, BF16 skew-down — с 0,962 до 1,099 от QuACK. #2840 (коммит в 06:51) добавляет prefill sinks для SM89 и включает наличие sink в идентичность экземпляра компиляции.

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

3.5 TileFoundry: план синхронизации WGMMA слит (#286), семантика завершения частичного сокращения на рассмотрении (#288)

Слитый в 13:40 #286 — это пункт, бывший на рассмотрении в прошлом выпуске; он касается асинхронности WGMMA: после выдачи инструкции она продолжает читать разделяемую память и регистры, и последующие обращения должны быть упорядочены после её эффектов. Подход — объявить условия завершения для эффекта каждого операнда, привязать fence, commit и wait к тем узлам HIR, которым они нужны, а lowering потребляет эти шаги в точке выдачи инструкции; если в последующей области видимости не помещается коллективное ожидание, оно выполняется в области выдачи. Одновременно исправлен алиасинг регистрового Reshard, чтобы каждый элемент оставался в одном и том же регистре у каждого потока, держащего копию, и зависимость по регистрам отслеживалась по этой идентичности хранения; на SM90 CUDA-модуль, выдающий групповые шаги WGMMA, компилируется отдельно с sm_90a. Автор сохранил два предупреждения: эмиттер матричных инструкций всё ещё не подключён, поэтому этот коммит не доказывает корректность исполнения и не даёт производительности; семантика завершения TMA рассматривается отдельно. В 20:12 открыт #288, касающийся учёта «завершённых частичных сокращений» в анализе.

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

3.6 TileSight: исправление объявления возможностей WGMMA для H20, первый PR в репозитории (#1)

14:37 коммит, 17:26 обновление — это первый PR в tile-ai/TileSight (24 звезды). Репозиторий представляет собой открытую реализацию статьи arXiv от 07-24 «TileSight: A First-Principles Tile-Centric Analytical GPU Performance Model from Cores to Clusters»: исходя из формы тайлов, перемещения данных, программного конвейера и аппаратного профиля, на CPU оцениваются стоимость оператора, трафик кэша и время выполнения, а CUDA/HIP-пробы отдельно измеряют целевой GPU. Изменение небольшое, но направление ясное: H20 — это Hopper/GH100, и планы с requires_arch=("wgmma",) ранее выбрасывали UnsupportedError, поскольку профиль архитектуры не объявлял поддержку WGMMA; теперь устанавливается support_wgmma=True и явно сохраняется support_utcmma=False, благодаря чему такие планы принимаются, а требования tcgen05 и Tensor Memory по-прежнему отклоняются. Автор также отмечает, что ёмкость следует обрабатывать отдельно по SKU (NVIDIA одновременно указывает H20 на 96 GB и 141 GB), в этот раз изменены только метаданные возможностей, а значения ёмкости, пропускной способности, частоты, throughput и утилизации остаются неизменными. Верификация — все 847 тестов на CPU пройдены, с явным заявлением «это не проверка производительности на GPU».

Источники: https://github.com/tile-ai/TileSight/pull/1, https://arxiv.org/abs/2607.22432

3.7 Принимающие стороны и третьи стороны: TileKernels, FlashQLA без пушей; появились два новых сторонних репозитория (11-05)

deepseek-ai/TileKernels по-прежнему на 10-14, QwenLM/FlashQLA на 10-30, в данном окне коммитов не было, последний релиз FlashQLA — всё ещё v0.1.3 от 09-30. На стороне третьих сторон появились два новых небольших репозитория: Tianchen-ai/tilelang-weft-bench («TileLang-Weft full CPU operator benchmarks using the shared Weft Bench engine», 0 звёзд, создание единого движка для бенчмарков операторов на стороне CPU), flytiger-eco/tilelang-wrapper-for-sail (0 звёзд, без описания). Оба очень ранние, пока заносятся в список наблюдения, не стоит подавать их как технический прогресс.

Источники: https://github.com/deepseek-ai/TileKernels, https://github.com/QwenLM/FlashQLA, https://github.com/Tianchen-ai/tilelang-weft-bench, https://github.com/flytiger-eco/tilelang-wrapper-for-sail

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

4.1 Три обновления документационного сайта; в ночном репозитории нет коммитов в основную ветку; ежедневные тесты Ascend — все 3030 пройдены (11-05)

tile-ai/tilelang.github.io в 16:07, 19:33 и 23:43 сделал по одному пушу Update docs, соответствующих изменениям интерфейсов основного репозитория за день (host-mapped аннотации, TF32 rounding, конфигурация Z3). TileOPs-nightly — только пуши веток, в основной ветке новых коммитов нет; TileOPs.github.io в 09:02 был один пуш. Ежедневный тестовый issue #1988 в репозитории Ascend фиксирует, что 11-05 в 08:16 раунд из 3030 пунктов полностью пройден, проходимость 100%, что контрастирует с открытыми в тот же день #1991/#1992 (см. 2.3).

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

4.2 Новостная сторона: четыре группы запросов, Hacker News и arXiv — совпадений в окне нет (11-05)

Четыре группы запросов в Google News (английские "TileLang" OR "tile-lang", TileLang GPU kernel, китайские TileLang 算子, TileLang 编译器) — самая свежая находка относится к материалу SemiAnalysis «The Huawei-DeepSeek Symbiosis» 376-часовой давности, остальные сосредоточены в волне публикаций от 09-30 о «DeepSeek с открытым исходным кодом версии TileLang, DeepGEMM и DeepEP для Ascend» (около 880 часов назад), и все они раньше текущего окна. Поиск по tilelang на Hacker News за 24 часа дал 0 результатов; три найденных Show HN — это терминальный просмотрщик данных, движок рабочих процессов и файловый менеджер для macOS, к теме не относятся. Самый свежий результат all:tilelang на arXiv — от 11-03, «Beyond Compression Ratio: High-Throughput KV Compression for Disaggregated LLM Serving» (вне окна, работа ссылается на TileLang), до него — сам TileSight от 07-24. Основное содержание отчёта поэтому по-прежнему целиком идёт со стороны кода.

Источники: https://news.google.com/, https://news.ycombinator.com/, https://arxiv.org/

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

Во-первых, точка конкуренции в библиотеках операторов полностью сместилась с «количества операторов» на «стоимость селектора и подготовки». Из девяти коммитов TileOPs сегодня реально меняют GPU kernel только #2832 (широковещательное чтение на месте) и #2834 (собственные широкие тайлы); #2835 экономит на подготовке вызова (с 4,33% до 6,17%, сгенерированный исходный код не меняется), #2836 меняет один коэффициент подгонки, regret снижается с 0,89% до 0,06%. Это показывает, что оставшийся запас производительности кроется в основном не в том, «как считать», а в том, «какой тайл выбрать и какую подготовительную работу сократить». Дойдя до конца, этот путь приведёт к модели стоимости, описанной калибровочным контрактом, а не к набору вручную настроенных параметров — сегодняшний #2836 с «отдельной подгонкой и изменением одной константы» как раз зачаток такой формы.

Во-вторых, основной тип дефектов на стороне компилятора сейчас — «молчаливая генерация неверного кода», а не ошибка компиляции или недостаточная производительность. #3808 на протяжении четырёх версий с 0.1.12 по 0.1.15 пропускает один __syncthreads(), #3800 на пути Blackwell заставляет MMA читать перезаписанный operand, #1991/#1992 при threads=2 на Ascend заставляют sin/cos вычислять неверно 97,1% элементов — все три не выдают ошибок, проходят обычные тесты и требуют, чтобы человек написал целевое воспроизведение, чтобы их поймать. Такие дефекты наносят пользователям наибольший ущерб, потому что превращают предпосылку «правильно ли компилятор» в гипотезу, требующую проверки. Соответствующее инженерное действие должно заключаться в расширении дифференциального тестирования по комбинациям форм (с поэлементным сравнением с эталонной реализацией в контрольной группе), а не в дальнейшем наращивании количества юнит-тестов.

В-третьих, ритм на стороне Ascend сместился с «зачистки регрессий после архитектурных изменений» на «зачистку границ корректности». Четыре типа регрессий, вызванных структурными изменениями вроде #3512, были только что закрыты в прошлом окне; два слияния текущего окна (превращение segfault в сообщение об ошибке, разделение performance skills по поколениям) и две новые issue (вдвое уменьшенный tmp workspace, нарушение calCount <= tmp) относятся к сходимости границ. Примечательно, что автор #1991 — ревьюер #1897, а исправление #1897 покрывало только сторону count — другая грань той же логики раскладки была указана лишь две недели спустя, что говорит о том, что инварианты раскладки бэкенда Ascend ещё не оформились в проверяемый письменный контракт.

В-четвёртых, «обобществление» отечественных бэкендов ускоряется. Hygon прошёл путь от собственной ветки до двух официальных PR, Moore Threads возобновила поддержку документации, mlir-ascend тремя небольшими PR по одному переносит арифметические операторы в диалект TLIR. Общее у всех трёх линий — они больше не ждут «одноразового слияния большого бэкенда», а выдвигают в публичный репозиторий минимальные рецензируемые единицы. Для вышестоящего проекта это хорошо: ревьюеры могут рано высказаться о дизайне раскладки и lowering, а не смотреть на него при одноразовом слиянии тридцати тысяч строк.

Последний аспект риска: если миграция TIRx в основном репозитории #3803 будет продвигаться, vendored-структура 3rdparty/tvm будет заменена на два закреплённых апстрим-коммита, и тогда все бэкенды, зависящие от внутренних API TVM (клеевые слои Ascend, Hygon, mlir-ascend), должны будут последовать за изменением. В черновике явно написано «не внедрять патчи зависимостей» — это и обещание, и ограничение: сегодняшний перенос адаптера BishengIR в contrib в mlir-ascend и разделение performance skills по поколениям чипов в Ascend можно рассматривать как заменяемые поверхности, зарезервированные под такого рода смену основания.

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

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

Источники Результаты проверки
Организация tile-ai (29 репозиториев) За окно мониторинга push в 10 репозиториев
Основной репозиторий tilelang 5 слияний в основную ветку; 2 новых issue (#3808, #3800); обновления на рассмотрении сосредоточены в четырёх кластерах: скаляризация, CuTeDSL, Profiler, Ascend
TileOPs 9 слияний в основную ветку; 4 новых или обновлённых на рассмотрении; новых issue нет
TileFoundry 1 слияние в основную ветку (#286); 1 на рассмотрении (#288)
TileSight Первый PR (#1), коммитов в основную ветку нет
tilelang-ascend 2 слияния в ветку по умолчанию ascendc_pto (#1932, #1954); 3 новых issue (#1991, #1992 на одну тему, #1988 — ежедневный отчёт о тестировании)
tilelang-mlir-ascend Новых коммитов в основную ветку нет; продвижение трёх PR с операторами TLIR: #203, #211, #212
tilelang-hygon Новых коммитов в основную ветку нет (последний — 11-02); 3 коммита в ветке feat/hcu-scaled-cast; новые PR #33, #34
tilelang-musa 1 коммит документации в основную ветку (первый с 09-30); обновлены временные метки 6 оставшихся PR, содержимое не изменилось
tilelang-metax За окно мониторинга push отсутствует (третье окно подряд)
TileOPs-nightly / TileOPs.github.io / tilelang.github.io В ночном репозитории коммитов в основную ветку нет; 1 push в сайт документации TileOPs; 3 Update docs в сайт документации tilelang
Страницы релизов (tilelang / TileOPs / TileFoundry) Новых версий за окно мониторинга нет; последняя версия tilelang по-прежнему v0.1.15 от 09-30, TileOPs — v0.1.0 от 10-27, TileFoundry — v0.0.4rc1 от 10-25
TileKernels / FlashQLA За окно мониторинга push отсутствует
Обзор четырёх источников TileLang×RISC-V Без изменений (ссылки и статусы PR в четырёх upstream совпадают со снимком предыдущего выпуска)
Google News (четыре группы на китайском и английском) / Hacker News / arXiv Существенных попаданий за окно мониторинга нет, последнее попадание — 376 часов назад

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

Пояснения к проверке

Все включённые, находящиеся на рассмотрении, issue и коммиты в этом выпуске получены из фактического вывода поиска GitHub и атомарных источников: включённые и находящиеся на рассмотрении — через search API (18 включённых, 57 обновлённых находящихся на рассмотрении, 5 новых issue), заголовки коммитов — через commits atom каждой ветки (12 веток, 21 коммит попал в окно), релизы — через releases atom. Анонимный API уровня организации на позднем этапе сбора достиг почасового лимита; список организаций за пределами основной ветки сверялся перекрёстно с веб-страницами и atom-источниками, оценочные значения не использовались. Заголовки и тексты двух issue #1991 и #1992 совпадают, они записаны как «две записи с одинаковой темой» и не объединены в одну. Три коммита Hygon и один коммит Moore Threads являются невключёнными коммитами, описаны как динамика на уровне ветки и не означают, что они уже включены в основную линию. Изменения TileSight и TileFoundry, связанные с WGMMA, самими авторами ограничены метаданными возможностей или планом синхронизации; заявлений о корректности выполнения или производительности не делалось, отчёт следует этим ограничениям. Три новостных канала Google News, Hacker News и arXiv в данном окне не дали совпадений.