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

Источники: порепозиторная проверка организации GitHub tile-ai (29 репозиториев, 8 с пушами в окне) + поиск подтверждённых данных по слитым и находящимся на рассмотрении изменениям (12 слитых, 42 обновления на рассмотрении) + 4 новых issue + поштучная проверка коммитов в основную ветку (tilelang 3, TileOPs 7, TileFoundry 1, tilelang-ascend 1, mlir-ascend 2) + проверка неосновных веток (4 коммита в ветке feat/hcu-scaled-cast репозитория tilelang-hygon) + проверка страниц релизов и тегов (5 репозиториев, новых релизов в окне нет) + два репозитория adopters + четыре группы запросов Google News на китайском и английском (через прокси) + Hacker News + arXiv.

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

  • Главное за день: изменение архитектуры фреймворка PTO для Ascend, слитое 10-14 (#3512), вступило в фазу концентрированной расчистки — за день всплыли пять независимых друг от друга патчей; TileOPs слил за день 7 производительных изменений — самый плотный день после праздников; все три слитых изменения основного репозитория связаны с «метрическими критериями», сброс L2 переведён с записи нулей на только чтение
    1. Основной прогресс проекта
      • 1.1 Основная линия: архитектурный шлюз TMA в CuTeDSL принимает sm_103a и sm_107a, шесть тестовых случаев восстановлены на B300 и gr100 (#3772) (11-04)
      • 1.2 Основная линия: сброс L2 переведён с записи нулей на редукцию только чтением, из do_bench удалён параметр fast_flush (#3534) (11-04)
      • 1.3 Основная линия: тестовые случаи автоподбора в CI AscendC явно синхронизируют два устройства, ошибки асинхронного копирования больше не маскируются под сбой arange (#3784) (11-04)
      • 1.4 На рассмотрении: доказательство согласованности warp в ThreadSync дополнено составным предикатом, собранным из индексов warp (#3787) (11-04)
      • 1.5 На рассмотрении: выход целевой области за границы в T.copy теперь становится ошибкой времени компиляции, доказательство включения полностью копируемого fragment дополнено проверкой принадлежности потоку (#3786 / #3788) (11-04)
      • 1.6 На рассмотрении: FP8 GEMM на SM120 переведён на block-scaled MMA, на RTX 5090 куб 8192 вырос с 463 до 727 TFLOPS (#3790) (11-04)
      • 1.7 На рассмотрении: проверки критериев Profiler и автоподбора отправлены группой, шесть за один день (#3797 / #3798 / #3799 / #3716 / #3717 / #3718) (11-04)
    1. Адаптация под множество бэкендов
      • 2.1 Ascend PTO: одно архитектурное изменение породило четыре класса регрессий, пять независимых друг от друга патчей отправлены в один день (#3792 / #3793 / #3794 / #3795 / #3796) (11-04)
      • 2.2 Ascend: соотношение ядер Cube и Vector стало настраиваемым, больше не фиксировано 1:2 (#3783) (11-04)
      • 2.3 Ascend: нижняя граница по циклам в AutoSchedule подтянута, время решения Z3 снизилось на 10.04% (#3776) (11-04)
      • 2.4 Ascend: дамп исключений AI Core совместим с CANN 9.2.0, порог версии снижен с 9.3 (#1956) (11-04)
      • 2.5 Hygon: четыре коммита в ветке feat/hcu-scaled-cast, преобразование масштабирования E8M0 и планирование чтения LDS по 64 банкам (11-04)
      • 2.6 mlir-ascend: адаптер инструментальной цепочки BishengIR перенесён в contrib (#209) (11-04)
      • 2.7 MetaX и Moore Threads: пуши в окне отсутствуют (11-04)
    1. Экосистема и adopters
      • 3.1 TileOPs: 7 слияний за день — упаковка GQA varlen, восемь векторов на поток в ArgMax/ArgMin, стратегия кэша по умолчанию (11-04)
      • 3.2 TileOPs на рассмотрении: унарный поэлементный запуск в одну волну, непрерывная раскладка FP8 сгруппированного GEMM, опциональный bias в GEMM (#2829 / #2824 / #2823) (11-04)
      • 3.3 TileFoundry: семантика обучения FSDP и HSDP вошла в вычислитель HIR (#284) (11-04)
      • 3.4 TileFoundry на рассмотрении: план синхронизации wgmma (#286),另 открыт issue об унификации именования (#287) (11-04)
      • 3.5 Adopters и третьи стороны: в TileKernels и FlashQLA пушей нет, три сторонних репозитория активны в течение 24 часов (11-04)
    1. Сообщество, руководства и мероприятия
      • 4.1 Добавлено 4 новых issue: доказательство для warp размером 64, знаковые атомарные операции PTO, квантизатор NVFP4 в CuTeDSL (11-04)
      • 4.2 Новостная сторона: четыре группы запросов, HN и arXiv — попаданий в окне нет (11-04)
    1. Наблюдения за трендами
  • Приложение. Материалы и пояснения к проверке

Главное за день

Основная линия этого дня — не «что ещё слили», а «вчерашний фундамент начали принимать». Слияние #3512 от 10-14 векторизовало T.Parallel в конвейер VMI PTO — это структурное изменение бэкенда Ascend: целочисленные типы перешли от беззнаковой семантики i* к явно знаковым si*, lowering SIMD перешёл под управление нового трёхступенчатого процесса PTO — проверка, легализация, векторизация, — а определение явных параллельных SIMD-циклов было переписано заново. Спустя три недели это изменение на реальных формах операторов было по одному выявлено в четырёх классах регрессий: целочисленные атомарные операции падают при компиляции из-за непринятия pto.si32 (сообщение #3789, исправление #3792), T.clear в T.SimdVF уходит на путь широковещания, поддерживающий только числа с плавающей точкой (#3795), граничный факт скалярного cast слишком рано связывается при легализации памяти VMI (#3793), явные параллельные SIMD-циклы отвергаются как скалярный доступ к lane (#3796), плюс ещё один коммит адаптирует целочисленные атомарные операции к беззнаковому носителю (#3794). Все пять патчей основаны на основной ветке и не зависят друг от друга; автор в описаниях повторяет одну и ту же фразу «This PR is based directly on main and does not depend on the other PTO regression fixes» — это намеренно: каждый из пяти коммитов должен сливаться независимо, чтобы избежать накопления гигантского пакета отката, как бывало раньше. Верификация тоже больше не «просто запустилось», а аппаратные измерения с конкретным числом тестовых случаев: 47, 18, 16, 37 пунктов, четыре независимых оператора (SwiGLU, Top-k, Cache, обратный проход Engram) в качестве репрезентативных случаев.

Второй пункт — самокоррекция методики измерения, причём два репозитория сделали одно и то же в один день. Основной репозиторий #3534 изменил сброс L2 в do_bench с «записи нулей в буфер» на «однократное чтение буфера»: запись нулей оставляет в L2 целые грязные строки, и во время следующего раунда замера эти строки нужно записать обратно — то есть стоимость сброса кэша раскладывается на измеряемое ядро. #3534 читает блоками через cache.view(-1, 4096).sum(dim=1) и заодно удаляет публичный параметр fast_flush — тот параметр лишь выбирал, создавать буфер через int32 или int8. Бэкенд msprof на стороне Ascend изменён радикальнее: раньше цикл SimtVF записывал нули в UB, затем следовало сохранение WTS_FV, теперь выполняется чтение из глобальной памяти в UB по NORMAL_FV, а тип зафиксирован как int32. #2822 в TileOPs — это близнецовое изменение той же идеи, дополнительно записывающее метод сброса и сравнивающее с историческими показаниями только при совпадении метода. Три изменения вместе — сигнал: производительные цифры начинают требовать «сопоставимости между версиями», а не просто «в этот раз быстрее, чем в прошлый».

Третий пункт — запоздалая оплата аппаратного фронтира. B300 сообщает целевую архитектуру sm_103a, а вспомогательные функции TMA в CuTeDSL отсекают архитектуру белым списком, из-за чего шесть случаев CuTeDSL на B300 выбрасывают ValueError ещё до входа в DSL; #3772 добавляет sm_103a вместе с сообщаемым gr100 (cc 10.7) sm_107a в белый список, автор воспроизвёл на узле B300: 8 пройдено, 1 пропущено, на gr100 — 13 пройдено, 1 пропущено. В том же окне #3790 идёт другим путём: GEMM с FP8 и накоплением fp32 на SM120 всё время использовал атомарные операции mma.sync от SM89, на RTX 5090 составляя лишь 68% от cuBLASLt; после перехода на block-scaled MMA с единичным масштабированием ue8m0 куб 8192 вырос с 463 до 727 TFLOPS, а выход побитово совпадает со старыми атомарными операциями. Природа двух коммитов различна — один «даёт новой карте работать», другой «даёт новой карте работать на полную», — но оба указывают на одно: граница аппаратной поддержки расширяется одновременно от Hopper к Blackwell и к отечественным ускорителям.

Четвёртый пункт посвящён ритму коммитов TileOPs. Из 7 слияний за день 5 — производительность, 1 — CI, 1 — документация: GQA varlen на SM90 упаковывает запросы по KV-группам, снижая число чтений key/value для коротких запросов до одного раза на KV-голову; ArgMax/ArgMin при условии размещения за один проход даёт каждому потоку восемь 16-байтовых векторов, bf16 ArgMin (2048, 4096) снижается с 8.704 микросекунды до 7.393, обгоняя 8.543 у FlagGems; унарные поэлементные и нормализационные ядра возвращают операндам, читаемым «только один раз», стратегию кэширования по умолчанию, сохраняя evict-last/evict-first только там, где чтение действительно повторяется. Общее у этих изменений не «быстрее», а «знание, почему быстрее» — каждый коммит несёт таблицу сравнения и отмечает строку, где он всё ещё отстаёт от сравниваемого (Softmax (32768, 4096) откатывается с 121.792 до 131.648 микросекунды).

1. Основной прогресс проекта

1.1 Основная ветка: архитектурные ворота TMA в CuTeDSL принимают sm_103a и sm_107a, шесть случаев на B300 и gr100 восстановлены (#3772)

Слито 11-04 12:46 (пекинское время, далее то же), это самое раннее слияние в текущем выпуске. Целевая архитектура в отчёте B300 — sm_103a, а вспомогательная функция TMA в CuTeDSL ведёт белый список архитектур и при встрече со значением, которого нет в таблице, выбрасывает ошибку ещё до входа в DSL; в сообщении об ошибке допустимые значения перечислены как sm_90, sm_90a, sm_100a, sm_110, sm_120, sm_120a, sm_121, sm_121a. На основной ветке fa69f7fa, с nvidia-cutlass-dsl 4.7.0 и CUDA 13.3, на B300 из-за этого падают шесть тестовых примеров: три на пути tma_load в test_tilelang_jit_cutedsl.py (по одному для GEMM JIT, многопоточного ядра и динамических форм), а также три варианта tmem_copy_roundtrip на пути tma_store в другом тестовом файле CuTeDSL для tcgen05. Сам DSL принимает sm_103a — его операции tcgen05 принимают всё семейство sm_100f — так что проблема чисто в этом шлюзе. После добавления этого пункта оба файла на B300 (cc 10.3, CUDA 13.3, cutlass-dsl 4.8.1) дают 8 прохождений и 1 пропуск. В тот же коммит добавлен и sm_107a, поскольку плата bring-up для gr100 (cc 10.7) сообщает это значение и упирается в тот же шлюз; после добавления обоих пунктов те же файлы на gr100 дают 13 прохождений и 1 пропуск, включая файл gather/scatter, внесённый в #3685.

Стоит отметить очерченную автором границу верификации: аппаратуры sm_103 под рукой не было, изменения проверены на машине B300, предоставленной другим репортёром (0z5a, впервые сообщившим и выполнившим bisect в #3685), а CI не покрывает sm_103a. Подобные случаи «код слит, проверка на чужой машине» в период догоняющей разработки архитектур будут учащаться, и читателю нужно знать, какие возможности подкреплены реальными измерениями, а какие лишь добавлены в белый список.

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

1.2 Основная ветка: сброс L2 заменён с записи нулей на read-only редукцию, у do_bench удалён параметр fast_flush (#3534)

Слито 11-04 11:27, изменено 7 файлов, добавлено 36 строк, удалено 49 строк, из них три места — собственно реализация, четыре — подогнанные тесты. Раньше было cache.zero_(): L2 вытеснялся записью нулей в буфер размером cache_size MB, проблема в том, что после записи нулей буфер полон грязных строк, и при выполнении тестируемого ядра в следующем раунде эти строки нужно записывать обратно в память, а полоса записи обратно засчитывается во время тестируемого ядра — чем больше буфер, тем заметнее отклонение. Теперь стало cache.view(-1, 4096).sum(dim=1, dtype=cache.dtype): чтение блоками по 4096 столбцов, строки читаются в L2 и порождают результат редукции, проблемы грязных строк нет; разбиение на блоки нужно, чтобы избежать глобальной синхронизации при редукции целого блока; сохранение dtype нужно, чтобы не запускать преобразование всего буфера в int64. Это изменение одновременно внесено в цикл estimate в bench.py, в три пути event, cupti и cudagraph в torch_bench.py, а также в специализированное ядро сброса для Ascend в msprof.py — последнее раньше с помощью SimtVF записывало нули в UB и затем сохраняло их обратно в глобальную память через WTS_FV, что является типичным сбросом записью, а теперь читает из глобальной памяти в UB через NORMAL_FV, и тип элементов тоже зафиксирован как int32 вместо выбора между int8/int32.

Изменение публичного интерфейса требует внимания: параметр fast_flush у do_bench удалён. Раньше он лишь переключал тип элементов буфера между int32 и int8 (int32 быстрее создаёт буфер), теперь буфер фиксирован как int32, и параметр не имеет смысла. Для скриптов, напрямую вызывающих do_bench(..., fast_flush=False), это ломающее изменение; тесты внутри репозитория уже синхронно изменены, а внешним пользователям придётся напоминать через release note.

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

1.3 Основная ветка: в CI AscendC тестовый пример autotuning supply явно синхронизирует два устройства, ошибка асинхронного копирования больше не маскируется под сбой arange (#3784)

Слито 11-04 17:01. Симптом: в задании CI AscendC в тестовом примере set-inputs для autotuning параметризованный элемент [other_npu] при следующем вызове torch.arange() сообщает srclen=4, dstlen=4 — несовпадение размера memcpy в 4 байта, а это ровно форма пробы меж device-копирования. Настоящая причина в том, что сбой меж device-копирования сообщается асинхронно, и ошибка всплывает на шаг позже места возникновения, поэтому точка сообщения об ошибке выглядит как совершенно посторонний arange. Исправление простое: добавить по одной синхронизации устройства (текущего устройства и npu:other) до и после пробы, чтобы асинхронная ошибка вернулась в область действия try ... except RuntimeError. Проверка — запуск этого файла тестового примера на одной машине с двумя картами, все четыре пункта проходят, и при реальном асинхронном сбое меж device-копирования параметризованный элемент [other_npu] пропускается, а не валит весь тестовый пример.

Этот коммит сам по себе небольшой, но он является продолжением предыдущего #3546 — тот коммит перевёл подачу входных данных для автотюнинга на разрешение по текущему устройству, и при приёмке упоминалась регрессия на двух картах Ascend; здесь же дополняется стабильность того же сценария в CI. Одна и та же функциональность исправляется в двух последовательных окнах мониторинга, что указывает на то, что автотюнинг на стороне Ascend по-прежнему находится в состоянии «работает, но границы нестабильны».

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

1.4 На рассмотрении: в доказательство warp-согласованности ThreadSync добавлен составной предикат, собранный из индексов warp (#3787)

Этот коммит продолжает ту же область, оставленную предыдущими #3761 и #3765: ThreadSync должен доказать, что «потоки-участники барьера согласованы в пределах всего warp», иначе lowering отклоняется. Триггерная форма — при 192 потоках составной предикат tx < 160 and tx // 32 % 2 == 0 — warp 0, 2, 4 участвуют целиком, warp 1, 3, 5 не участвуют целиком, ни один warp не разделён, что в принципе должно быть доказуемо, но основная ветка выбрасывает внутреннюю ошибку «потоки-участники не warp-согласованы». Автор составил таблицу сравнения, точно очертив область дефекта: одиночные предикаты tx // 32 % 2 == 0, tx % 64 < 32, tx < 160 в основной ветке проходят lowering, и только форма, где «граничное условие и условие по модулю индекса warp соединены через and», ломается; а tx < 168 (который действительно разделяет 5-й warp) корректно отклоняется в обеих ветках. Корневая причина — в кодировании доказывателя: он переименовывает номер потока в два гипотетических потока и заставляет арифметический решатель доказать, что из linear1 / warp_size == linear2 / warp_size следует равенство предикатов с обеих сторон, но кодирование FloorDiv и FloorMod в комбинации «верхняя граница + взятие по модулю» не позволяет установить связь между двумя делениями. Исправление состоит в переписывании threadIdx.x как warp_index * warp_size + lane: два гипотетических потока совместно используют индекс warp и имеют каждый свой свободный lane, что в точности соответствует отношению «один и тот же warp»; тогда предикаты, читающие номер потока только через индекс warp, схлопываются в одно и то же выражение с обеих сторон, согласованность доказуема, а предикаты, действительно зависящие от lane, по-прежнему не равны, и tx < 168, перечисленный в #3786, по-прежнему отклоняется — доказательство не ослаблено, просто добавлен пропущенный класс.

Тот же автор в тот же день отправил #3786: когда целевая область T.copy выходит за пределы буфера, вместо генерации обращения за границы выдаётся ошибка времени компиляции. Оба коммита вместе отражают один и тот же подход — превратить «компилятор молча генерирует неверный код» в «компилятор говорит нет».

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

1.5 На рассмотрении: в доказательство включения полностью реплицированного fragment добавлена проверка принадлежности потоков (#3788)

ProveFragmentContains при определении «покрывает ли большой fragment малый» имеет быстрый путь: если большой fragment полностью реплицирован, сразу возвращается истина. Этот путь возвращает результат до проверки принадлежности потоков, из-за чего возникает такое ошибочное суждение: малый fragment владеет потоками с 0 по 127, большой fragment полностью реплицирован, но владеет только потоками с 0 по 95, а 96–127 вообще не покрыты, однако функция утверждает, что включение выполняется; при включённом check_forward_index два полностью реплицированных fragment с разными физическими слотами также принимаются. Следствие — вывод раскладки принимает несовместимые отображения за совместимые, это относится к типу «ошибка молча пропускается»; сообщивший обнаружил это попутно при проверке #2399. Исправление: сначала посмотреть на физический индекс, затем пересчитать отображение потоков малого fragment в систему координат большого fragment (с учётом ограниченного смещения диапазона потоков), с помощью арифметического анализатора доказать попадание в диапазон репликации и диапазон потоков, а при невозможности доказать вернуть ложь. Проверка выполнена на A100 на новой сборке основной ветки upstream, через мост, скомпилированный на время тестирования, напрямую вызывается реальная реализация на C++, покрываются отклонение при 128 против 96 потоков, принятие 96 против 128 и 96 против 96, физические слоты, несмежную принадлежность, ограниченные диапазоны, символьные индексы и частичную репликацию.

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

1.6 На рассмотрении: FP8 GEMM для SM120 переведён на block-scaled MMA, на RTX 5090 8192 в кубе с 463 до 727 TFLOPS (#3790)

Это самый значимый материал с точки зрения производительности в текущем выпуске. Проблема в выборе инструкций: на SM120 T.gemm с FP8 и fp32-накоплением на RTX 5090 на формах с 8192 в кубе даёт лишь около 465 TFLOPS, тогда как cuBLASLt — 681; настройка tile и числа стадий конвейера не помогает, поскольку узкое место — сама инструкция: TileLang выдаёт mma.sync ... f32.e4m3.e4m3.f32 эпохи SM89, а такой атом на этой карте работает лишь на половинной скорости. Автор сделал набор проб, измеряющих только регистры: e4m3 в f32 через mma.sync даёт 512 TFLOPS, через kind::mxf8f6f4.block_scale (единичное масштабирование) — 1021, через mma.sync с f16-накоплением — 1022; вся разница в инструкциях. Поскольку при масштабе ue8m0, равном 0x7f (то есть 2 в степени 0), block-scaled и обычный MMA дают одинаковый результат, можно при включённом TL_SM120_BLOCK_SCALED_MMA_ENABLED и цели sm_120a/f, sm_121a/f перевести четыре пути диспетчеризации FP8 с fp32-накоплением на block-scaled атомы, а остальные цели, f16-накопление и явный sm_120 без суффикса a оставить как есть. Побитовая идентичность проверена на уровне инструкций: четыре комбинации e4m3 и e5m2, случайное содержимое регистров, включая кодировки NaN и Inf.

Цифры: при e4m3 GEMM, выводе bf16, tile 128×128×128, 3 стадиях конвейера, 256 потоках — 4096 в кубе с 408 до 673 (cuBLASLt 530), 8192 в кубе с 463 до 727 (681), 4096×14336×4096 с 449 до 696 (498), 256×8192×8192 с 348 до 606 (457). Единственное замедление — форма M=16 с tile в 128 строк: с 31,5 микросекунды до 33,5 микросекунды; автор поясняет, что эта форма ограничена пропускной способностью и с tile в 16 или 64 строки будет быстрее. Новый тест напрямую проверяет PTX: sm_120a должен получать block-scaled атомы, sm_89 — обычные.

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

1.7 На рассмотрении: группа коммитов по сверке критериев Profiler и автотюнинга, шесть за один день (#3797 / #3798 / #3799 / #3716 / #3717 / #3718)

Все шесть от одного автора (Afloat16), тема — согласование «критериев валидатора и проверяемого объекта». assert_allclose теперь отклоняет случай, когда ядро возвращает None там, где должен быть тензор, — прежняя реализация пропускала сравнение, если одна из сторон None, так что ядро, возвращающее пустоту, тоже могло «пройти» (#3797); assert_consistent раньше напрямую делал zip по возвращаемым значениям, на скалярных тензорах выбрасывал TypeError, при отсутствии или лишнем выходе молча игнорировал, а при разных формах могла выручить обрезка строк или broadcast; теперь сначала нормализует в список выходов, затем сверяет количество и формы, включая пустые тензоры (#3798); в бенчмарках CUPTI suppress_stdout_stderr раньше вызывал fileno() на Python-потоках, а у ноутбука и StringIO нет файловых дескрипторов — падение; теперь отдельно перенаправляются нативные дескрипторы 1/2 и Python-потоки, а ExitStack гарантирует восстановление и на путях исключений (#3799). Ещё три — на стороне автотюнинга: при использовании словаря цели как ключа кэша бралось строковое представление, и эквивалентные словари с разным порядком вставки давали разные ключи — теперь сериализация в отсортированный JSON (#3716); детект Tensor Core сначала использует явную вычислительную способность, затем явную или текущую целевую архитектуру и лишь потом локальное устройство, иначе кросс-компиляция без GPU ложно сообщает о неподдержке (#3718); в Profiler замена динамических форм теперь охватывает составные размерности вроде m + 1 и при отсутствующей переменной выдаёт ValueError вместо TypeError (#3717).

Техническая ценность этой группы не в самих изменениях, а в том, какие формы они обнажают: валидатор долго был слабым звеном по принципу «прошло — значит верно», а такие инфраструктурные вещи, как ключи кэша автотюнинга и разбор целей, при малейшем расхождении проявляются в виде «неповторяемых результатов», что очень трудно атрибутировать.

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

2. Мультибэкендовая адаптация

2.1 Ascend PTO: одно архитектурное изменение порождает четыре класса регрессий, пять независимых патчей отправлены в один день (#3792 / #3793 / #3794 / #3795 / #3796)

Сначала о предпосылках. Слияние #3512 ([Ascend][PTO] Vectorize T.Parallel to PTO VMI), вошедшее 10–14, перевело векторизацию параллельных циклов со старого канала SIMD на конвейер VMI в PTO: целочисленные типы сменились с беззнаковых i* на явные знаковые si*, Lowering проходит три этапа — Verify, Legalize, Vectorize, — а определение явных SIMD-циклов реализовано заново. Каждый из этих пяти патчей соответствует одной проблеме, выявленной на реальных операторах:

  • #3792 Ошибка компиляции целочисленных атомиков: codegen_pto для целых чисел генерирует pto.si32 / pto.si64, тогда как атомарный интерфейс PTO принимает только беззнаковые i32 / i64, поэтому T.atomic_add, T.atomic_max, T.atomic_min на целочисленных буферах просто не компилируются (соответствует issue #3789). Исправление — обёртка совместимости вокруг tl.atomic_*: тип элемента указателя и операнды переписываются на беззнаковые той же разрядности, целочисленным min/max передаётся правильная знаковость, а возвращаемое значение приводится обратно к исходному типу. Набор тестов расширен с 15 до 18 позиций, покрывая знаковые и беззнаковые 32/64-битные сложение, вычитание и взятие экстремума, отрицательные операнды, возврат предыдущего значения и счётчики разрядностью более 32 бит.
  • #3794 Адаптация знакового носителя: тот же корень, что у #3792, но другой путь — семантика tl.atomic_* не меняется, вместо этого на границе генерируемых вызовов PTODSL переписываются представление указателя и операнды, адресное пространство сохраняется, атомики с плавающей точкой идут прежним путём. Показательный тест — счётчик clamp в SwiGLU (bf16, hidden 2048, clamp 0.5, раскладка psum); все 18 атомарных тестов проходят.
  • #3795 Потеря SIMD-заполнения: int32 T.clear внутри T.SimdVF попадал на путь широковещания, поддерживающий только плавающую точку, потому что новый конвейер PTO больше не запускает старый проход SIMD, а тот проход отвечал одновременно за заполнение и параллельные циклы. Исправление — выделен отдельный вход, делающий только заполнение, с переиспользованием прежнего LowerFill, размещённый после нормализации версий буфера и до параллельной проверки PTO; по-прежнему генерируется прежняя последовательность vdup/vsts и сохраняются проверки непрерывных областей, выравнивания, целых векторов и срезов многобуферных версий. Проходят 16 тестов, показательный — Top-k (512 token, 1 эксперт, top-1).
  • #3793 Граничные факты скалярного cast привязывались слишком рано: случаи вида u = uint32(i - 1), где i лежит в [0, 63] — скаляр, не связанный с адресом VMI, — бросали ошибку на легализации памяти; поспешный Bind к тому же оставлял заниженный интервал, и ошибка всплывала позже на анализе cast в IntSet. Исправление — сначала проверять представимость целочисленного cast, затем устанавливать факты, давая движку доказательств безопасное скалярное определение, а IntSet — консервативный числовой интервал, чтобы избежать прямого вычисления концов cast. Проходят 47 тестов, показательный — Cache (block_size=64, n=65536).
  • #3796 Явный SIMD-параллельный цикл трактовался как скалярная lane: инициализация в обратном проходе Engram использовала вложенность Parallel, где явное SIMD-сохранение продвигается целыми векторными блоками, а новый вывод раскладки lane счёл это скалярным доступом и отверг (M=2, N=32, L=128). Исправление — классификация явного SIMD разделяется между выводом раскладки и обнаружением параллельных единиц: такая вложенность сохраняет общий путь раскладки, существовавший до #3512, а автоматические поэлементные единицы по-прежнему идут новым путём lane PTO. Проходят 37 тестов, показательный — обратный проход Engram (512 token, hc 4, hidden 2048). Автор также указывает ограничение: смешанные случаи явного и автоматического SIMD с FP32/128 lane покрыты только на уровне генерируемого исходного кода — текущий компилятор устройств PTOAS/Bisheng на этой комбинации превышает бюджет стека VF в 6 KiB, результатов запуска на устройстве нет.

Рассматривая эти пять изменений вместе, можно прочитать инженерный ритм на стороне Ascend: структурные изменения вроде #3512 сначала вливаются, а затем на реальных операторах шаг за шагом выявляются границы; каждый патч сопровождается аппаратными измерениями и счётчиком тестов и намеренно сохраняет независимость от остальных. Подход «сначала слить каркас, потом разгрести регрессии» уже применялся в прошлый цикл миграции архитектуры (#1975), а теперь стал обычным процессом — ценой является то, что основная ветка в такие окна находится в состоянии «известно, что несколько мест не работают», и вероятность того, что внешние пользователи на это наткнутся, немала.

Источники: https://github.com/tile-ai/tilelang/pull/3792, https://github.com/tile-ai/tilelang/pull/3793, https://github.com/tile-ai/tilelang/pull/3794, https://github.com/tile-ai/tilelang/pull/3795, https://github.com/tile-ai/tilelang/pull/3796

2.2 Ascend: соотношение ядер Cube и Vector становится настраиваемым и больше не фиксировано на 1:2 (#3783)

AI Core в Ascend делится на два типа — Cube и Vector; AscendC поддерживает группировки mix(1,2), mix(1,1), mix(1,0), mix(0,1) и т. д., тогда как TileLang ранее всегда запускался в пропорции 1:2, что фактически фиксировало верхнюю границу одной аппаратной степени свободы и затрудняло адаптацию к чипам с иным соотношением ядер в будущем. Данное изменение превращает соотношение ядер в явный параметр или параметр, выводимый по устройству: T.MixedKernel(..., sids=1) или sids=2 работают в соответствии с запрошенным размером группировки и проходят проверку возможностей; при пропуске или передаче None используется vector_core_num / cube_core_num текущего устройства, а при недоступности информации об устройстве происходит откат к 2; T.Kernel сохраняет прежний интерфейс и определение типа ядра, при распознавании смешанного ядра использует подготовленное число векторных ядер, путь одиночного ядра не меняется, семантика группировки явного T.Vector(1/2) сохраняется. Выбранное число групп передаётся сквозь цепочку в домен sid, распределение задач, синхронизацию и генерацию кода, избегая повторных запросов к устройству и повторной нормализации; конфликтующие объявления и запросы, превышающие возможности устройства, отклоняются напрямую.

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

2.3 Ascend: нижняя граница периода AutoSchedule ужесточена, время решения Z3 снижено на 10,04% (#3776)

AutoSchedule использует двоичный поиск для нахождения интервала инициализации конвейера (II); прежняя нижняя граница бралась как максимум II всех задач, что было слишком слабым и требовало лишних итераций двоичного поиска. Данное изменение добавляет в ресурсную модель три необходимых условия: граница одной задачи (сохраняется P >= max_i II_i), граница ресурсного клика (для клик попарно конфликтующих ресурсов сумма ранжируется по остаткам, получается P >= sum_i II_i), граница самозависимости (S_u - S_u >= wait - d*P даёт P >= ceil(wait/d), расстояния, связанные с версиями, усекаются по числу ступеней конвейера). Двоичный поиск инициализируется максимумом из трёх, остальная логика не меняется. Тестовая среда — Ascend 950PR, еженедельная сборка CANN 9.3.0, кэш отключён, один прогрев и три замера поочерёдно: examples/ascend полностью компилируется, результаты планирования не меняются, суммарное время этапа решения Z3 снижается на 10,04%. Подобные изменения «не меняют результат, а только начальную точку поиска» обладают минимальным риском, но дают глобальный выигрыш — время решения на этапе компиляции платит каждый пользователь.

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

2.4 Ascend: дамп исключений AI Core совместим с CANN 9.2.0, порог версии снижен с 9.3 (#1956)

Этот коммит влит в ветку ascendc_pto репозитория tilelang-ascend (11-04 10:18) и является единственным слиянием в репозитории Ascend за текущий период. Функция дампа исключений AI Core изначально заявляла зависимость от CANN 9.3.0 и выше; автор провёл проверку работоспособности на 9.2.0 и исправил три места: дополнены обязательные поля acldumpTensorInfo (addrType=ACL_DUMP_ADDR_RAW, placement, shapeNum, originShapeNum), поскольку 9.2.0 строго проверяет addrType, а прежняя реализация получала ACL_DUMP_ADDR_PTR через обнуление calloc, что отклонялось; пропускаются пустые адреса и тензоры нулевого размера, так как 9.2.0 не принимает tensorAddr равным NULL; интерфейс регистрации заменён с устаревшего aclrtSetExceptionInfoCallback на aclrtExceptionInfoCallbackRegister. На стороне инструментов parse_exception_dump() совместим с обоими форматами вывода 9.2.0 — .npy и старым .bin; требования к версии в китайском и английском руководствах также изменены на CANN 9.2.0 и выше. Для пользователей, всё ещё работающих на 9.2, это переход от «функция недоступна» к «доступна».

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

2.5 Hygon: четыре коммита в ветке feat/hcu-scaled-cast, масштабирующее преобразование E8M0 и планирование чтения LDS с 64 банками (11-04)

Основная ветка репозитория Hygon остановилась на 11-02; вся активность в данном окне мониторинга сосредоточена в функциональной ветке, четыре коммита уложились в восемь часов (с 16:26 до 20:16 по пекинскому времени), авторы из Hygon (huangteng@hygon.cn и zhouweiwang@hygon.cn):

  • feat(hcu): add scaled_cast and unify GEMM LDS read planning: добавлено преобразование FP8/BF8/FP4 с масштабированием E8M0, скалярный и упакованный варианты lowering, добавлены src/tl_templates/hcu/scaled_cast.h (223 строки) и lower_scaled_cast.cc; одновременно метаданные потребителя GEMM распределены между выбором чтения матриц и стратегией LDS, поддержана одноуровневая нарезка LDS, основной файл преобразований изменён на 650 строк.
  • fix(hcu): align GEMM LDS strategies with 64-bank read phases: swizzle для plain-copy выводится из геометрии bank-word, для узкого b16-тайла, общего для матриц и двух собственных размеров упаковки, предоставляется явное обратное преобразование; асинхронный потребитель со смешанными или ограниченными фазами в приоритетном порядке использует совместимое общее хранилище; при выборе собственного wrap полная фаза чтения b128 учитывается полностью. 12 файлов, добавлено 1536 строк, удалено 112 строк, добавлено 5 тестовых файлов.
  • fix(hcu): prefer mixed LDS wraps proven at most two-way: для смешанного LDS wrap добавлено доказательство «не более двух путей», приоритет отдаётся только при выполнении доказательства.
  • fix(hcu): use physical read width for mixed wrap checks: проверка смешанного wrap переведена на физическую ширину чтения, убрана прежняя запись с выводом 8 байт на lane по kPack.

Тема — банковые конфликты разделяемой памяти и выравнивание фаз — это ключевой путь производительности GEMM на Hygon DCU, а также признак того, что бэкенд перешёл от «способности компилировать» к этапу «тюнинга под микроархитектуру». Все четыре коммита находятся в функциональной ветке, PR ещё не открыт; по времени возможна синхронизация с основной линией TileLang перед отправкой.

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

2.6 mlir-ascend: адаптеры инструментальной цепочки BishengIR перенесены в contrib (#209)

Смержено в основную ветку (11-04 16:35), код адаптации инструментальной цепочки BishengIR перенесён с прежнего места в каталог contrib. Это чисто структурная упорядоченность: «клеевой слой, который приходится часто менять по мере развития вышестоящей инструментальной цепочки», отделён от ядра lowering, и при последующих обновлениях BishengIR зона изменений ограничивается contrib. В той же ветке есть основной коммит refactor(contrib): move BishengIR toolchain adapters into contrib, merge-коммит — второй. Такого рода упорядочивание само по себе не имеет функциональной ценности, но это начало восприятия «адаптационного слоя» как независимого жизненного цикла — версии инструментальной цепочки Ascend обновляются быстро, и заменяемость клеевого кода важнее элегантности.

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

2.7 MetaX и Moore Threads: в окне мониторинга нет активности (11-04)

Последняя активность tile-ai/tilelang-metax остановилась на 10-26 15:08, tile-ai/tilelang-musa — на 09-30 17:50; в обоих репозиториях в этом окне нет коммитов и нет обновлений PR на рассмотрении. Из четырёх отечественных бэкендов Ascend и Hygon — две активные линии этого окна, MetaX и Moore Threads находятся в тихом периоде. В прошлом отчёте у MetaX также не было активности, это уже второе окно подряд; у Moore Threads движения нет с 09-30, стоит обратить внимание в последующих окнах, не перешли ли они к внутренней разработке.

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

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

3.1 TileOPs: 7 слияний за один день, упаковка GQA varlen, восемь векторов на поток в ArgMax/ArgMin, стратегия кэширования по умолчанию (11-04)

Все 7 слиты в окне мониторинга, по темпу это самый плотный день квартала; по порядку:

#2819 добавил для SM90 в GQAVarlenFwdOp упакованную warp-специализированную программу: когда группа KV-голов одного запроса помещается в один tile по позициям, запросы заполняют tile по KV-группам, и тогда каждая KV-голова читает key/value только один раз; применять ли упаковку, решает модель стоимости — только если «объём, который гарантированно удаётся сэкономить по форме» превышает измеренные 9% накладных расходов на одну голову; широкие и узкие запросы поочерёдно получают работу в одном запуске, без чтения с хоста и без синхронизации; когда подсчитанное на устройстве число неделимых элементов меньше числа SM, упакованный запрос одного tile разбивает сканирование key на несколько сегментов и объединяет их через рабочую область fp32. Также исправлены три дефекта основной линии: строки без видимых key возвращали NaN, когда масштабирование добавки sink превышало примерно 2.4; путь отката читал таблицу RoPE до первого key запроса; а негативный контроль оставлял кадр стека вызывающего в цикле ссылок, из-за чего CUDA graph уничтожался досрочно между тестами. Тестовые узлы выросли с 173 до 197.

#2828 Для построчного разбиения ArgMax/ArgMin при условии «одной волны пребывания» каждый поток получает восемь 16-байтовых векторов, а блок из 64 потоков добавлен в кандидаты настройки; решение о пребывании определяется числом потоков SM, верхним пределом постоянно резидентных CTA и регистрами. Для bf16 ArgMin размера (2048, 4096) время снижено с 8,704 микросекунды до 7,393 (эталон FlagGems 8,543, ускорение в 1,155 раза), в сценарии подряд идущих вызовов — с 7,808 до 6,000; для f16 ArgMax — с 6,816 до 6,400 (эталон 9,344, ускорение в 1,46 раза). В 12 строках бенчмарка 4 строки ускорились, 0 замедлились.

#2820 Возвращает операнды операций «каждый элемент читается один раз» к стратегии кэширования по умолчанию: поэлементные, rope, norm, softmax и ядра сокращения больше не используют повсеместно evict-last/evict-first, а сохраняют их только там, где действительно есть повторное чтение (широковещательные строки бинарных ядер, строки rope cos/sin, общие для нескольких пакетов и голов, два чтения в двухпроходном RMSNorm, первые два прохода в fused-сложении). Одновременно для RMSNorm добавлены две реализации — парное ядро и поэтапное ядро, CTA warp-строчного Softmax уменьшен ровно до числа warp, заполняющих SM, 2D neox RoPE переведён на 128 потоков, а верхний предел резидентных CTA для каждой архитектуры сведён в RESIDENT_CTAS. Цена — Softmax (32768, 4096) откатился с 121,792 до 131,648 микросекунды (эталон quack 126,080); автор честно отметил это в таблице.

#2822 — изменение только для чтения при сбросе L2, имеет тот же источник, что и #3534 в основном репозитории, 168 тестов пройдено, дополнительно фиксируется метод сброса, и сравнение с историческими показаниями ведётся только при совпадении метода. #2825 использует единую переобученную модель стоимости для ранжирования допустимых конфигураций W4A16; CTA split-K хранит только свою порцию метаданных квантования, а учёт разделяемой памяти включает запас на запуск. #2826 — сопутствующая документация, где ясно описан контракт калибровки для плотного и tiny-M режимов (способ измерения, ограничения подгонки, отложенная проверка), без изменений времени выполнения. #2827 переносит набор off-CUDA в ночные задания: изначально в preflight эти два шага на ubuntu-latest занимали от 13 до 18 минут и давали результат на десять минут позже остальных проверок, причём дымовой прогон tests/ops использует эталонные реализации операторов и проверяет независимость от устройства и диспетчеризацию, а не результаты ядер; после переноса в preflight остаются только проверка сбора, отбрасывание отчётов тестов и обязательная проверка cuda_only на неудачу (около 3 минут), а 1856 плюс 2789 тестовых узлов переходят в ночной прогон, при этом релизный контроль по-прежнему требует успеха ночного прогона.

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

3.2 TileOPs на рассмотрении: одноволновой запуск унарных поэлементных операций, непрерывная раскладка FP8 grouped GEMM, опциональный bias у GEMM (#2829 / #2824 / #2823)

#2829 Заставляет запуск унарных поэлементных операций, «в теле которых всего одна-две одноинструкционные операции», работать в одну волну: число читаемых и записываемых байт не превышает трёх четвертей L2, результат и операнд имеют одинаковую ширину, в теле не более двух операций, каждая из которых сводится к одной инструкции; только при выполнении всех трёх условий блок сжимается до одной волны из 256 потоков; условие на тело определяется подсчётом узлов на переменных TIR, а тела с cast, делением, математическими вызовами, охраной или не вычислимые на голой переменной всегда сохраняют запуск по умолчанию. Сейчас этому соответствуют relu, abs, neg, bitwise_not и logical_not над bool. Попутно исправлена настоящая проблема: ReLU для входа NaN возвращал 0, теперь возвращает NaN (для 16-битных используется битовый шаблон torch, для float32 — канонический NaN). В бенчмарке три строки ReLU и одна строка LogicalNotFwdOp ускорились, остальные конфигурации не изменились; автор также отметил, что строка decode со сбросом на GPU 6 лишь сравнялась.

#2824 Добавляет SM90 FP8 grouped GEMM непрерывную раскладку (две упаковки — компактная и выровненная, два варианта метаданных экспертов — хвост сегмента или построчный), SM89 по-прежнему поддерживает только маскированную раскладку, а также дополняет измеренной нагрузкой со смещением по размеру эксперта и сравнением с непрерывной раскладкой DeepGEMM. #2823 Добавляет в GEMMFwdOp опциональный bias[N], охватывая три пути — плотный, GEMV и split-K, с приложенным сравнением с cuBLASLt BIAS. Оба изменения соответствуют пробелам, перечисленным в #2130, что показывает: TileOPs продвигается по списку «поэлементного выравнивания с экосистемой torch», а не добавляет операторы случайным образом.

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

3.3 TileFoundry: семантика обучения FSDP и HSDP в вычислителе HIR (#284)

Совпадает с 11-04 16:36. Прежний вычислитель обрабатывал reshard как глобальный тензор, теряя различия вкладов разных участников, поэтому семантика обучения (особенно HSDP со смешанным шардированием и репликацией) не могла корректно вычисляться на уровне HIR. Этот коммит реализует обучающее подмножество #255: вычисление локальных значений по участникам, координаты GPU, вызовы и циклы, gather/scatter, а также свёртку Partial(sum) вдоль оси — всё идёт через общие правила проекции; доказательство синонимов Reshard изменено на чувствительное к принадлежности, подключено к структурированному жизненному циклу из #170, кортежные аргументы валидируются по тензорным листьям. В комплекте добавлено исследование двухблочного GQA, вдохновлённое modded-nanogpt: явный прямой и обратный проход, momentum SGD, независимый эталон на PyTorch, покрытие FSDP, HSDP 2×2 и 4×4, две стратегии резидентности, отчёт о группах коллективной коммуникации, полезной нагрузке и занимаемой памяти, а также встроенный переиспользуемый SVG; каждая таблица в README выводится из кода.

Проверка: переработанный основной набор — 775 тестов, набор бэкендов — 778 тестов (13 пропущено, две карты sm90), 10 конфигураций HIR проходят на CPU; обе стратегии резидентности на четырёх sm90 согласуются с реальным 2×2 PyTorch HSDP (два шага обновления, сравнение локальных градиентов, параметров и моментума).

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

3.4 TileFoundry на рассмотрении: план синхронизации wgmma (#286), а также отдельный issue о единообразии именования (#287)

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

#287 — issue документационного характера, указывающий, что в кодовой базе «у одного понятия пять имён» (член mesh называется placement в 133 местах, program в 110, instance в 92, participant в 44, плюс написания с unit вроде wave_units), при этом само слово unit одновременно обозначает пустой тип, аппаратный вычислительный блок и члена mesh; сигнатуры local_layout в Python и CUDA-рантайме всё ещё не совпадают. Такие issue обычно не влекут немедленных изменений кода, но это обязательный шаг для компиляторного проекта на пути от «работает» к «читается людьми долго».

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

3.5 Адоптеры и третьи стороны: TileKernels и FlashQLA без пушей, три сторонних репозитория активны за 24 часа (11-04)

Два репозитория-адоптера в этом окне не имеют коммитов: deepseek-ai/TileKernels остановился на 10-14 05:49, QwenLM/FlashQLA — на 10-30 03:02 (последний релиз по-прежнему v0.1.3 от 09-30). Следует отметить, что отсутствие пушей у адоптера не означает отсутствия влияния — именно исправление #3761 из прошлого выпуска касалось регрессии занятости барьеров ядра mhc_pre_big_fuse из TileKernels на B300, что показывает: изменения компилятора воздействуют на их формы в виде «пассивных регрессий».

На стороне третьих лиц есть три активных элемента, обнаруженных сканированием репозиториев, все созданы или запушены за 24 часа: sablin39/tilelang-cuda-skills (149 звёзд, набор навыков для написания на TileLang и отладки с помощью инструментов CUDA), xwhzz/tilelang-hexagon (2 звезды, перенос TileLang на DSP Hexagon от Qualcomm), Index4NekoMeteor/Tilelang-AttnRes (1 звезда, реализация структуры attention residual Kimi K3 на TileLang). Все три малы, но направления интересны: один — опыт разработки на стороне инструментария, два — портирование на новые целевые архитектуры и структуры передовых моделей; это относится к тому типу работы, где сообщество самостоятельно расширяет границы TileLang.

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

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

4.1 Добавлено 4 новых issue: доказательство размера warp 64, знаковые атомарные операции PTO, квантизатор NVFP4 для CuTeDSL (11-04)

Основной репозиторий, три пункта. #3791 — «[BUG][ice-on-valid-code]»: ThreadSync может доказать warp-согласованный предикат только при размере warp 32, и тот же IR отклоняется, когда thread_warp_size равен 64 — этот пункт тесно связан с находящимся на рассмотрении в этом выпуске #3787, что указывает на жёстко закодированное в прувере предположение о размере warp, тогда как AMD и часть отечественных ускорителей используют 64-широкий warp. #3789 сообщает, что бэкенд PTO для целочисленного T.atomic_add генерирует знаковые pto.si32/si64, что приводит к сбою валидации; в тот же день появились два варианта исправления — #3792 и #3794. #3781 сообщает, что NVFP4 block-scaled квантизатор CuTeDSL не компилируется: переинтерпретация упакованного FP4 и отсутствующие ceil/floor. Вместе с #287 в TileFoundry (унификация именования) в этом окне всего 4 новых issue, и все они относятся к типу «детали внутренней реализации протекают в пользовательскую плоскость», ни одна не является жалобой на производительность.

Источники: https://github.com/tile-ai/tilelang/issues/3791, https://github.com/tile-ai/tilelang/issues/3789, https://github.com/tile-ai/tilelang/issues/3781

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

В четырёх группах запросов Google News (английские TileLang GPU kernel, "TileLang" OR "tile-lang", китайские TileLang 算子, TileLang 编译器) последнее попадание остановилось на отметке 352 часа назад (репортаж SemiAnalysis о Huawei и DeepSeek от 10-21), остальные сосредоточены в волне публикаций от 09-30 «DeepSeek открывает исходный код версии TileLang для Ascend, DeepGEMM, DeepEP» (856 часов назад) — все раньше текущего окна; поиск на Hacker News по tilelang за последние 24 часа дал 0 результатов; последнее попадание arXiv по all:tilelang — всё ещё TileSight от 07-24, новых статей в окне нет. Из 9 кандидатов, обнаруженных при поиске репозиториев на GitHub, 7 — сам репозиторий, 1 — не связанный с темой Show HN вроде Swarm World, 1 — посторонний проект типа 2048. Можно подтвердить, что новостная сторона в этом окне нулевая, и основная часть отчёта может происходить только из коммитов и PR — это уже устойчивая форма на протяжении нескольких выпусков подряд.

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

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

Во-первых, модель расчистки PTO для Ascend уже устоялась, и её выгоды и издержки стоит зафиксировать. #3512 был влит 10-14, и в течение трёх недель реальные операторы постоянно выводили его за границы возможностей, а сегодняшние пять дополнений — лишь последняя партия. Такой ритм «сначала заложить архитектуру, потом по операторам расчищать регрессии» уже наблюдался при полной миграции репозитория в #1975: плюс в том, что архитектура итерируется быстро и каждое отдельное изменение контролируемо, минус в том, что основная ветка в течение довольно длительного окна находится в состоянии «известно, что есть несколько неработающих мест». Пять патчей намеренно оставлены независимыми друг от друга, что говорит о сознательном снижении гранулярности отката со стороны мейнтейнеров — это зрелее, чем собирать всё в один большой пакет отката, но означает и потребность в более строгой дисциплине предрелизной регрессии, иначе «независимо сливаемо» превратится в «независимо пропускаемо».

Во-вторых, измерение производительности начинают воспринимать как продуктовую характеристику. Основной репозиторий #3534 и TileOPs #2822 в один и тот же день сменили очистку L2 с записи нулей на только чтение, по одной и той же причине — «запись грязных строк загрязняет следующий раунд измерения времени»; TileOPs пошёл дальше: фиксирует метод очистки и сравнивает исторические показания только при совпадении метода. Это не мелкая правка: если эталонное число несопоставимо между разными версиями, то все выводы «насколько быстрее» придётся пересчитать. Включение самого метода измерения в управление версиями — предпосылка доверия к производительной инженерии.

В-третьих, точка конкуренции операторных библиотек смещается с «количества операторов» на «качество селектора». Влитые сегодня в TileOPs модель затрат с единственной переподгонкой для W4A16 и документ контракта калибровки для плотного GEMM, вместе с находящимися на рассмотрении непрерывным layout для FP8 grouped GEMM и bias для GEMM, образуют чёткую линию: сначала перечислить допустимые конфигурации для каждой формы, затем отсортировать их единой моделью затрат, а процесс калибровки оформить как воспроизводимый контракт. Если эта линия выдержана, производительность операторной библиотеки перестаёт зависеть от ручной настройки отдельных ядер; если нет — она выродится в «некоторые формы особенно быстры, остальные полагаются на удачу».

В-четвёртых, границы поддержки архитектур расширяются с двух сторон. С одной стороны — новые карты: sm_103a для B300 и sm_107a для gr100 попадают в белый список, FP8 для SM120 заменяется на block-scaled MMA; с другой — отечественные ускорители: Hygon проводит микроархитектурную оптимизацию фазы LDS с 64 банками, Ascend закрывает инженерные детали вроде соотношения ядер и дампа исключений. Структура трудозатрат у этих двух направлений расширения совершенно разная: для новых карт главное — «заставить путь работать», для отечественных карт — скорее «заставить тот же путь выдавать эффективность, близкую к аппаратной», что требует гораздо более высокой плотности реальных измерений, и именно поэтому в тех четырёх коммитах Hygon такая большая доля тестовых файлов.

Последний аспект риска: в рамках одного окна 42 обновления PR находятся на рассмотрении, из них 12 с метками Ascend/PTO, 7 с метками CUDA/SM/CuTeDSL, 4 сосредоточены в Profiler — эта структура показывает, что давление изменений в репозитории смещается от «добавления функций» к «исправлению границ». Для пользователей хорошая новость в том, что дефекты указываются по одному; плохая новость в том, что большинство этих дефектов находится в таких «обычно невидимых» местах, как валидаторы, ключи кэша, пруверы, и когда они проявляются, это часто уже влияет на достоверность результатов.

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

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

Источник Результат проверки
Организация tile-ai (29 репозиториев) В окне выполнены push в 8 репозиториев
Основной репозиторий tilelang В основную ветку 3 коммита; открыто 3 новых issue; 38 обновлений на рассмотрении (Ascend/PTO 12, CUDA и CuTeDSL 7, Profiler 4)
TileOPs В основную ветку 7 коммитов; 3 обновления на рассмотрении; новых issue нет
TileFoundry В основную ветку 1 коммит; 1 обновление на рассмотрении; открыт 1 новый issue
tilelang-ascend В ветку по умолчанию ascendc_pto 1 коммит; список на рассмотрении не получен из-за лимита анонимного интерфейса, проверка выполнена по основной ветке и атомарным источникам коммитов
tilelang-mlir-ascend В основной ветке 2 коммита (включая слияние #209)
tilelang-hygon В основной ветке новых коммитов нет (остановка на 11-02); в ветке feat/hcu-scaled-cast 4 коммита
tilelang-metax / tilelang-musa В окне push отсутствуют
TileOPs-nightly Только push веток, в основной ветке новых коммитов нет
Страницы релизов (tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-hygon) В окне новых версий нет, последняя версия tilelang по-прежнему v0.1.15 от 09-30
FlashQLA / TileKernels В окне push отсутствуют
Google News (четыре группы запросов на китайском и английском) / Hacker News / arXiv В окне существенных совпадений нет, последнее совпадение — 352 часа назад

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

*[24]** tilelang-ascend: дамп исключений совместим с CANN 9.2.0 (#1956) — https://github.com/tile-ai/tilelang-ascend/pull/1956

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

Все включённые, находящиеся на рассмотрении, issue и количество коммитов в этом выпуске взяты из фактического вывода поиска GitHub и атомарных источников без оценок; список PR на рассмотрении в репозиториях Ascend и Hygon не удалось получить из-за исчерпания квоты анонимного интерфейса, соответствующие выводы основаны только на коммитах в основную ветку и атомарных источниках уровня веток, что отмечено в таблице проверки источников. Четыре коммита tilelang-hygon находятся в функциональной ветке без открытого PR и относятся к динамике уровня ветки, а не означают, что они уже включены в основную линию. Новостные каналы Google News, Hacker News и arXiv в этом окне мониторинга не дали совпадений, поэтому основная часть отчёта полностью получена со стороны кода.