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

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

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

  • Главное за сегодня: из 35 принятых коммитов основной репозиторий сосредоточился на жизненном цикле анализатора; выход за границы записи в селекторе top-k DeepSeek V3.2 локализован как CUDA_ERROR_ILLEGAL_ADDRESS; раскладка NZ для Ascend обеспечила частичное копирование UB→L1 (18 файлов); сообщество CANN открыло репозиторий операторов ops-tilelang
    1. Ход работ по ключевым проектам
      • 1.1 Основной репозиторий: клонирование анализатора перед векторизационным переписыванием (#3612) (22.10)
      • 1.2 Основной репозиторий: после разрезания параллельного цикла привязка старого анализатора становится недействительной (#3226) (22.10)
      • 1.3 Основной репозиторий: выход за границы буфера кандидатов в селекторе top-k deepseek_v32 (#3280) (22.10)
      • 1.4 Основной репозиторий: дополнены 16-битные обратные тригонометрические и бинарные математические понижения CUDA (#3172) (22.10)
      • 1.5 Основной репозиторий: для разреженных метаданных sm8x формы, не являющиеся биективными, теперь отклоняются (#3125) (22.10)
      • 1.6 Основной репозиторий: eager-фронтенд сохраняет позиции в исходном коде Python (#3558) (22.10)
      • 1.7 На рассмотрении: проверка возможностей pass, беззнаковый mulhi и предикат assume (22.10)
    1. Адаптация под множество бэкендов (Ascend / Hygon / MetaX / Moore Threads)
      • 2.1 Ascend: буферы NZ поддерживают частичное копирование UB→L1 (#3530) (22.10)
      • 2.2 Ascend PTO: экспоненциальный рост предикатов направления MAD приводит к остановке компиляции (#3581) (22.10)
      • 2.3 Ascend PTO: локальные переменные векторов SIMT распределяются по lane (#3523) (22.10)
      • 2.4 Репозиторий Ascend: cann-bench дополнен Dilation2D, в тот же день выполнен откат depthwise_conv_2d (#1916 / #1919) (22.10)
      • 2.5 Репозиторий Ascend, на рассмотрении: прямой и обратный проход sparse_mla (#1929 / #1928) (22.10)
      • 2.6 Hygon: в объединённой ветке удалён enable_register_pipeline (22.10)
      • 2.7 MetaX / Moore Threads / MLIR-Ascend: в окне мониторинга отправок не было (22.10)
      • 2.8 Промышленность: SCMP перепечатал отчёт SemiAnalysis, TileLang — единственный пригодный к использованию из шести репозиториев (22.10)
    1. Экосистема и стороны внедрения (TileOPs / TileFoundry / библиотеки операторов)
      • 3.1 TileOPs: в FP8-ядрах с постраничной специализацией по warp совмещение имён разделяемой памяти приводит к NaN (#2708) (22.10)
      • 3.2 TileOPs: dense prefill переведён на постоянный обход, D64 с тремя потребителями (#2706) (22.10)
      • 3.3 TileOPs: SM90 dense GQA decode использует специализацию по warp (#2694) (22.10)
      • 3.4 TileOPs: причинное скользящее окно обслуживает всю группу query head двумя key tile одного KV head (#2696) (22.10)
      • 3.5 TileOPs: при автотюнинге постраничного decode ранее использовалась общая block table (#2703) (22.10)
      • 3.6 TileOPs: переполнение FMA в masked softmax и нечётная половинная ширина в varlen RoPE (#2699 / #2700) (22.10)
      • 3.7 TileOPs: многоуровневое временное хранение scale в FP8 grouped GEMM и пинг-понг warp group для gate/up (#2704 / #2693) (22.10)
      • 3.8 TileOPs: широковещательная передача предикатов и булевы операнды переведены на 16-байтовое векторное чтение (#2702 / #2697) (22.10)
      • 3.9 TileOPs: три пути SM89 продвигаются в один день (#2705 / #2691 / #2690) (22.10)
      • 3.10 TileOPs: однократный выпуск свёртки редукции и проверка выбора пустого DSA (#2689 / #2698) (22.10)
      • 3.11 TileOPs: эталонные вычисления переведены на float32 (#2692) (22.10)
      • 3.12 TileFoundry: установлена архитектура sm_89 и устройство RTX 6000 Ada (#259) (22.10)
      • 3.13 TileFoundry: расчёт цены gmem переведён на footprint и рабочий mesh с привязкой к размерностям вызываемой стороны (#258) (22.10)
      • 3.14 TileFoundry: при отказе устройства запускать kernel должно выдаваться исключение (#256) (22.10)
      • 3.15 Дистрибутив: Sunrise выпустил PR релиза 0.1.15+sunrise.1.2.0 (#9, на рассмотрении) (22.10)
      • 3.16 Стороны внедрения: TileKernels и FlashQLA в окне мониторинга отправок не было (22.10)
    1. Сообщество, руководства и мероприятия
      • 4.1 Сообщество CANN открыло репозиторий операторов ops-tilelang (22.10)
      • 4.2 tilelang/agent: оптимизация операторов оформлена как рабочий процесс SKILL для agent (22.10)
      • 4.3 Новые репозитории сообщества: обёртка Sail, agent удалённого автотюнинга и репозиторий руководств (22.10)
      • 4.4 Документация и версии: сайт документации обновлён трижды, VERSION по-прежнему 0.1.

15 (10-22)

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

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

В этом окне мониторинга организация tile-ai слила 35 коммитов (TileOPs 20, tilelang 9, TileFoundry 3, tilelang-ascend 2, сайт документации TileFoundry 1), ещё 71 PR были обновлены в течение окна. Девять коммитов основного репозитория сместили акцент с «проверки форм и геометрии» на прошлой неделе к жизненному циклу анализатора: клонирование analyzer перед перезаписью векторизации (#3612, 2 строки в одном файле), устаревание старых привязок после разбиения параллельных циклов (#3226, 6 файлов), сохранение позиций исходного кода Python во фронтенде eager (#3558) — ни один из этих трёх не приведёт к неверным вычислениям kernel, но заставит этап компиляции читать устаревшее состояние после перезаписи. На численной стороне наиболее заслуживает внимания #3280: селектор top-k в deepseek_v32 временно размещает radix-кандидатов в фиксированном разделяемом буфере [2, 4096], но не ограничивает размер порогового бакета, выбранного на первом раунде; FP32-оценки с одинаковыми младшими битами мантиссы или полностью совпадающие входы приводят к записи за границы буфера и падению с CUDA_ERROR_ILLEGAL_ADDRESS. На стороне Ascend #3530 — самый крупный коммит дня (18 файлов, +993/-256): сохраняет сжатое заполнение NZ вплоть до нормализации хранения и явно задаёт строчный шаг группы C0 для части копирований UB→L1; #3581 устраняет на PTO экспоненциальный рост предикатных выражений, который мог задерживать компиляцию более чем на десять минут. На стороне экосистемы две новые линии: сообщество CANN 20-10 создало репозиторий операторов ops-tilelang, 22-10 был впервые отправлен фикс; дистрибутив Sunrise выпустил PR релиза 0.1.15+sunrise.1.2.0, включив в дистрибуцию wheel бэкенда TANG.

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

1.1 Основной репозиторий: клонирование analyzer перед перезаписью векторизации (#3612)

Слито 22-10 в 23:43, изменён только один участок исходного кода векторизации, +2/-1. Векторизация циклов при перезаписи IR повторно использует переданный арифметический анализатор, а анализатор кэширует константы и ограничения диапазонов до перезаписи; как только векторизация меняет границы циклов или индексные выражения, последующие pass получают старые выводы. Исправление — копировать анализатор перед входом в перезапись векторизации, перезапись действует только на копию. Этот коммит не сопровождается регрессионным тестом; его ценность в том, что он закрывает с другого входа проблему, родственную #3226, — в тот же день тот коммит разбирался с устаревшими привязками после разбиения параллельных циклов; рассматривая оба вместе, можно сказать, что жизненный цикл контекста анализа очищается как отдельный класс дефектов.

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

1.2 Основной репозиторий: устаревание привязки analyzer после разбиения параллельных циклов (#3226)

6 файлов, +55/-19, затрагивает три pass — loop_partition, loop_vectorize, lower_tile_op — и их заголовочные файлы. После разбиения параллельного цикла переменная итерации тела цикла перепривязывается, но pass по-прежнему ищет в анализаторе по переменной до разбиения и получает диапазон и информацию о делимости старой переменной, из-за чего упрощение и определение выхода за границы искажаются. Исправление изолирует контекст анализа по pass и упрощает путь инициализации снимка анализа циклов, а также добавляет набор минимальных регрессионных случаев анализа параллельных циклов. Среди авторов коммита указан соавтор Codex.

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

1.3 Основной репозиторий: выход за границы буфера кандидатов селектора top-k в deepseek_v32 (#3280)

2 файла, +91/-10, исправляет #1351. tl_topk_impl упаковывает radix-кандидатов в фиксированный разделяемый буфер [2, 4096], но ничто не ограничивает пороговый бакет, выбранный на первом раунде radix, размером не более 4096 элементов — при одинаковых младших битах мантиссы FP32-оценок или полностью совпадающих входах атомарное добавление выходит за конец буфера и падает с CUDA_ERROR_ILLEGAL_ADDRESS. Исправление — в слое-обёртке tl_topk выделяется глобальный workspace (batch, 2, seq_len), чтобы оба radix-буфера ограничивались seq_len; публичная сигнатура tl_topk() не меняется. Приложен 79-строчный воспроизводящий пример, построенный по входу из #1351.

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

1.4 Основной репозиторий: дополнение понижения для CUDA 16-битных обратных тригонометрических и бинарных математических функций (#3172)

3 файла, +463/-8. Для asin, acos, asinh, acosh, atanh, log1p, atan2, hypot, copysign, ldexp, nextafter отсутствуют специализированные правила понижения для CUDA; общие правила обрабатывают только float32/float64, поэтому вызовы с float16 и bfloat16 на этапе codegen никем не перехватываются. Исправление регистрирует их в CUDAMath и добавляет перегрузки half_t / bfloat16_t: десять операторов вычисляются как float32 с обратным преобразованием, nextafter переведён на целочисленное пошаговое сравнение по 16-битному кодированию, чтобы сохранить субнормальные числа bfloat16 при fast math. Имена функций float32/float64 остаются неизменными, на стороне common.h добавлено 182 строки шаблонной реализации.

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

1.5 Основной репозиторий: для sm8x небиективные формы разреженных метаданных теперь отклоняются (#3125)

Один файл tilelang/layout/gemm_sp.py +6/-0, исправляет #2606. Interleave в make_cutlass_metadata_layout_sm8x является биективным только тогда, когда M кратно размеру группы, а K чётно; при невыполнении условий ранее молча генерировались некорректные метаданные, и результат разреженного GEMM ошибался без предупреждения. Теперь оба условия явно отклоняются. Ранее добавленный вместе с PR тест формы разреженных метаданных по требованию ревью был вынесен и не попал в основную ветку.

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

1.6 Основной репозиторий: eager-фронтенд сохраняет позиции в исходном коде Python (#3558)

4 файла. После того как eager-фронтенд понижает операторы Python for / with до генерируемых вызовов, сообщения об ошибках и отладочные отображения указывают на строки сгенерированного кода. Исправление отображает сгенерированные вызовы циклов обратно на операторы for, обёртки контекста — на операторы with, а также добавляет позиции распаковки итерируемых аргументов; как и в #3125, ревью потребовало вынести прилагавшийся регрессионный тест позиций в исходном коде.

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

1.7 На рассмотрении: проверка возможностей pass, беззнаковый mulhi и предикаты assume (22.10)

Очередь на рассмотрение основного репозитория обновилась за окно мониторинга на 50 единиц. Стоит отслеживать несколько: [Transform] Locate VerifyBufferInit reads outside tile ops (#3620), [Codegen] Add unsigned mulhi intrinsic (#3617), [Transform] Keep assumption checks with local dependencies in scope (#3611) и [Transform] Preserve tl.assume predicates that simplify to true (#3610), [CPU] Support clz in C and C host code generation (#3609); на стороне Ascend [BugFix][Ascend] Reject odd FP4 view geometry in GM->L1 copies (#3608) и [Ascend] Support exact placement of versioned on-chip buffers (#3607) обновлены в тот же день. Кроме того, #3613 явно помечен как «не ревьюить» — черновик нативной редукции argmin/argmax.

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

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

2.1 Ascend: поддержка частичного копирования UB→L1 для NZ-буферов (#3530)

Самое крупное изменение за день: 18 файлов +993/-256. Заполнение сжатого NZ-лейаута ранее сглаживалось на этапе логического лейаута, и при частичном копировании UB→L1 нельзя было выразить собственные шаги строк групп C0 источника и назначения. Исправление сохраняет компактное NZ-заполнение до этапа нормализации хранилища и затем точно материализует его один раз по ёмкости; при понижении области копирования явно задаются шаги строк групп C0 источника/назначения блоками по 32 байта, и шаги строк проводятся через шаблон scatter, DMA codegen Ascend/PTO и футпринт доступа; также добавлены проверки геометрии, выравнивания и выхода за границы, сохранён владелец физического представления, вызывающая сторона переведена на логическую форму распределения с явными шагами строк. В src/ascend/op/nd2nz.h добавлено 83 строки.

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

2.2 Ascend PTO: экспоненциальный рост предикатов направления MAD приводит к остановке компиляции (#3581)

2 файла +48/-5. PTO-ядро с повторяющимися изменяемыми охранными условиями может более десяти минут занимать анализ, даже если в нём нет матричного умножения — это остановка на этапе компиляции до выполнения на NPU. Корневая причина в том, что PTOMmadDirectionAnalyzer консервативно сохраняет состояние входа, когда охранное условие читает изменяемую память; если обе ветви сохраняют направление MAD, Merge получает одинаковый предикат S, а реализация всегда строит и упрощает S || S; арифметический упроститель в воспроизводящем примере не устраняет этот дубликат, и последовательные ветви удваивают выражение предиката кратно. Исправление пропускает перестроение при совпадении предикатов с обеих сторон и добавляет 20 строк направленных тестов.

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

2.3 Ascend PTO: локальные переменные SIMT-векторов распределяются по lane (#3523)

4 файла, +273/-29. Локальные векторные переменные SIMT распределялись как скалярные буферы PTODSL без учёта числа lane, но при этом использовался доступ по индексам Python; AllocatedBufferValue не поддерживает поэлементное присваивание, поэтому запись по индексу цикла завершалась ошибкой ещё раньше, на этапе переписывания AST PTODSL. Исправление выделяет для каждого векторного слота достаточное число скалярных элементов и преобразует индекс слота в скалярное смещение, повторно используя существующий путь последовательных load/store, сохраняя семантику локальных векторных переменных AscendC, не изменяя развёртывание циклов и не добавляя новых pass понижения. Добавлены две группы тестов кодогенерации — float32x2 и float32x4, всего 209 строк.

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

2.4 Репозиторий Ascend: в cann-bench добавлен Dilation2D, в тот же день откат depthwise_conv_2d (#1916 / #1919)

#1916 добавляет прямой проход Dilation2D (задача уровня level3 в cann-bench): NHWC, fp16/fp32, произвольные kernelSize/stride/rate и SAME/VALID; вклад смещений за границей берётся как -inf + filter и побитово совпадает с golden, распространение NaN на входе и +inf в блоках фильтра превращают вклад за границей в NaN; реализовано как чистый ядерный Vector-оператор плюс Python-обвязка, с развёртыванием с учётом ограничения выравнивания строчного шага 32B механизма копирования Ascend, сетка масштабируется адаптивно относительно геометрии без жёстко заданного числа ядер. В тот же день #1919 выполняет откат всех изменений, ранее принятых 21 октября для depthwise_conv_2d (#1915); в сообщении коммита причина не указана, и общий репозиторий на примерах Претерпевает состояние «один входит, один выходит».

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

2.5 Репозиторий Ascend на рассмотрении: прямой и обратный проход sparse_mla (#1929 / #1928)

#1929 добавляет sparse_mla_fwd (прямой проход DSA разреженного MLA: varlen packed случайный gather по строкам, онлайн softmax, слияние двух GEMM QK/PV), размещён в examples/dsa_sparse_finetune/, два ядра (специализированный fast и универсальный portable) вместе с host-маршрутизацией; допустимая область fast-пути жёстко зафиксирована: H=128 и G=1, dim/tail 512/64, topk кратно 256, NI от 1 до 6; вне этой области используется универсальный путь. #1928 добавляет sparse_mla_bwd (цепочка GEMM по dQ, векторная цепочка P/dP, рассеивание dKV по случайным строкам), размещён в examples_experiment/dsa/ и внесён в список тестов оператора; прямой результат k1 уже включён в настоящий оператор. Всего в репозитории Ascend на рассмотрении 11 патчей.

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

2.6 Hygon: в сливаемой ветке удалён enable_register_pipeline (22.10)

В gcy_merge_main 22.10 в 11:11 отправлен коммит очистки на 6 файлов, +10/-24, полностью удаляющий переключатель enable_register_pipeline, что затрагивает два pass — внедрение программной конвейеризации регистров HCU и планирование конвейера, а также pipeline_utils.h и расширение loop на уровне языка; примеры и скрипты производительности同步 изменяют вызовы. Удаление переключателя обычно означает, что регистровый конвейерный путь уже действует по умолчанию и поддержка двух веток более нерентабельна; однако ветка main остаётся ожидать реализации в preserve constant K origins in hoisted MLS windows (#25) от 20 октября, эта партия изменений всё ещё ожидает реализации в ветке слияния.

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

2.7 MetaX / Moore Threads / MLIR-Ascend: нет обновлений в окне наблюдения (22.10)

Последняя отправка в tilelang-metax приостановилась 16.10, в tilelang-musa — 30.09, в tilelang-mlir-ascend — 19 octobre; в течение окна наблюдения ни в одной из трёх линий адаптации не было коммитов. Все изменения внутренних частей в этом окне пришлись на линии Ascend и Hygon.

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

2.8 Со стороны индустрии: SCMP перепечатал отчёт SemiAnalysis, TileLang — единственный доступный среди шести репозиториев (22.10)

Южно-Китайская утренняя почта (SCMP) 22 октября опубликовала статью «Jensen’s nightmare», пересказывающую опубликованный в четверг отчёт SemiAnalysis о Huawei-DeepSeek: последний программный релиз DeepSeek вместе с дополнительными экспериментами с Ascend 950PR свидетельствуют о сокращении разрыва с США, и делается вывод, что к 2027 или 2028 году DeepSeek будет «в основном» полагаться на «уже достаточно зрелый» технологический стек Huawei. Оригинал этого отчёта — тот же, что был включён в предыдущий выпуск, суждение «из шести репозиториев, получивших поддержку Ascend, пригоден только TileLang» не изменилось, также упоминается, что бэкенд Ascend 950 был включён в основной репозиторий TileLang, а не остался на отдельном уровне адаптации. Это вторичное распространение того же источника, оно не является новым техническим достижением, но показывает, что данное суждение проникает в мейнстримную медийную повестку.

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

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

3.1 TileOPs: NaN из-за алиасинга общей памяти в специализированном по warp kernel FP8 с постраничной обработкой (#2708)

2 файла +80/-30. GQAPagedFP8WSKernel возвращает NaN в комбинации «семь однотокенных запросов стоят перед запросом на 300 токенов, 4099 строк FP8 cache, 48 строк страниц». Корневая причина — планировщик общей памяти считает буферы ролей warp активными в разные промежутки времени, из-за чего размещает split handover по адресу key- или value-кольца, который всё ещё заполняется TMA, а running maxima handover’а также оказываются сжатыми до бита флага кольца. Исправление явно разделяет эти буферы в общей памяти и добавляет 24 строки регрессионных тестов (в терминах qwen3-fp8-ragged).

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

3.2 TileOPs: dense prefill переведён на постоянный обход, D64 с тремя потребителями (#2706)

6 файлов +590/-5, добавлен GQADensePersistentWSKernel (внешнее имя gqa_dense_persistent, приоритетнее gqa_dense): когда число query tile превышает число SM, один CTA занимает один SM и поочерёдно проходит свои tile, загружая query следующего tile во время epilogue текущего. На D64 каждый tile — это 192 строки query, распределённые на три warpgroup-потребителя по 64 строки, повторное чтение K/V сокращается на треть; на D128 сохраняются 128 строк.

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

3.3 TileOPs: SM90 dense GQA decode переведён на специализацию по warp (#2694)

Добавлен GQADecodeWSKernel: warp-производитель TMA заполняет двухслотовое K/V-кольцо, один warpgroup WGMMA поддерживает онлайн-softmax группы head; число split определяется заполнением всех слотов резидентных CTA за один раз, при этом гарантируется, что каждый split содержит не менее четырёх KV tile; при единственном split нормализация выполняется внутри CTA и вывод пишется напрямую, без прохода через combine kernel; KV head меняется между CTA быстрее всего.

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

3.4 TileOPs: причинное скользящее окно обслуживает всю группу query head двумя key tile одного KV head (#2696)

Добавлен GQADenseSlidingWindowBandKernel: GQADenseFwdOp ранее для каждого query head повторно загружал одни и те же два key/value tile, и один CTA делал только одну (query tile, head); новый kernel для одной пары (query tile, KV head) загружает два key/value tile band’а только один раз, последовательно выполняя редукцию для этой группы query head, одновременно позволяя query tile следующего head быть в полёте. Ориентирован на сценарии, где число строк скользящего окна не превышает 128 key.

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

3.5 TileOPs: автоподстройка постраничного decode ранее использовала общую таблицу блоков (#2703)

4 файла +24/-13. Подача autotune для постраничного decode заставляет каждый запрос получать одну и ту же block table, что эквивалентно тестированию страниц только одного запроса, причём его страницы постоянно находятся в L2, а ранжирование кандидатов отражает фактический трафик, оплачиваемый обслуживаемым пакетом. В реальных замерах на H200 с частотой 1500 MHz тюнер MHA decode на opt-1.3b-b16 выбрал block_N 128 (холодный 80.0 us) и отказался от block_N 64 (72.5 us), опустившись до 0.994 от flashinfer. Исправление заставляет тюнинг читать разные страницы cache для каждого запроса.

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

3.6 TileOPs: переполнение FMA в masked softmax и нечётная полуширина varlen RoPE (#2699 / #2700)

#2699 (4 файла +54/-6) исправляет NaN-строки в dense GQA forward со скользящим окном: при причинном окне тайл тюнинга 128x64 заставляет каждый query tile начинаться с окна первой строки, и последние несколько строк подряд сталкиваются с двумя полностью замаскированными key tile; после первого блока оба running maxima находятся на зажиме -1e38, fast math сливает prev * scale - curr * scale в FMA, ошибка округления (около 2 в 99-й степени) переполняется в NaN, исправление состоит в вычитании running maxima до масштабирующего пересчёта. #2700 исправляет границы цикла varlen query RoPE при нечётной ширине полувращения (rotary_dim 10, 14 и т. д.): когда параллелизм rows умножить на chunks не делится на число потоков ветви вращения, TileLang теряет верхнюю границу цикла, последующие query-строки пути warp-специализации вычисляются неверно, а общий путь (dim 256) просто зависает.

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

3.7 TileOPs: поуровневое промежуточное хранение scale в FP8 grouped GEMM и пинг-понг warp group gate/up (#2704 / #2693)

#2704 добавляет ключ конфигурации stage_scales: производитель через cp.async копирует A/B scale каждого уровня в разделяемую память и достигает барьера этого уровня при фиксации; stage_scales=0 по-прежнему генерирует старый kernel. Причина: masked FP8 grouped GEMM в математическом warpgroup читает A scale каждого шага K из глобальной памяти, углубление цикла сжимает L1 и останавливает эти load, из-за чего ds-v3-e8-up-masked достигает лишь 0.979 от DeepGEMM. #2693 добавляет в GemmTemplate пинг-понг форму: два математических warpgroup каждый выполняют полный tile из 64 строк, по очереди берут tile CTA, epilogue одной группы попадает под mainloop другой, упорядочивание через ротационный барьер.

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

3.8 TileOPs: вектроное чтение 16 байт для предикатного вещания и булевых операндов (#2702 / #2697)

#2702 заставляет предикат построчного вещания (Eq/Ne/Gt/Lt/Ge/Le и LogicalAnd/Or) читать stride-1 операнды каждого целого блока из 1024 столбцов 16-байтовыми векторами через load16_evict_last, для частичных блоков и остальных staged body сохраняется копирование fragment; linear-bias-float32 (1024x4096 против 1x4096) для шести сравнений поднялся с 0.99 до 1.06, логические операторы до 1.04. #2697 меняет в бинарных kernel сравнения и логики полноразмерные операнды на 16-байтовые векторные чтения evict-first, строка float32 ранее была лишь на уровне 0.99–1.00 от torch.compile.

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

3.9 TileOPs: три пути SM89 продвигаются в один день (#2705 / #2691 / #2690)

#2705 снижает число уровней конвейера по умолчанию для SSD chunk scan на SM89 до двух, уменьшая потребность в буферах scan kernel, для шести форм FP16/BF16 замеры берут медиану трёх раундов по 200 сэмплов CUPTI в каждом. #2691 при выборе существующим планировщиком нескольких разделов разбивает интервал paged DSA key на SM89, переиспользуя имеющееся слияние частичного вывода FP32 и log-sum-exp; при одном разделе сохраняются вывод BF16 и один kernel, а цель компиляции включается в идентификатор кэша слияния, для каждой цели создаётся отдельный JIT builder. #2690 на SM89 для выбора NSA top-k при группах query head 32/64 и head dim 64/128 переходит на tile из 32 столбцов, сохраняя tile из 64 строк и ту же реализацию ранжирования по оценкам, медиана GPU busy для восьми用例 FP16/BF16 снижается на 2.1%–15.0%.

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

3.10 TileOPs: однократный запуск свёртки сокращения и проверка выбора пустого DSA (#2689 / #2698)

#2689 добавляет ReduceFoldKernel, который сворачивает непрерывный суффикс сегмента lead для каждого сохраняемого индекса за один запуск (ранее сокращение префикса и суффикса по осям требовало двух запусков), и передаёт ему все вызовы ReduceEdgeKernel и VectorNormEdgeKernel, где строки могут удерживаться в полёте блоками из 512 потоков; предыстория — Sum отстаёт от torch-compile на 3d-multidim-float16 (0.977) и mlp-intermediate-float32 (0.990). #2698 исправляет проверку постраничного DSA: когда строка query не имеет ни одного валидного выбранного key, эталонная реализация возвращает нулевой вывод и LSE минус бесконечность, а вычитание эталонного LSE даёт границу ошибки NaN и признаёт корректный результат ошибочным; исправление задаёт пустым строкам точную границу нулевого LSE и добавляет контрольный пример, где при «все строки пусты» валидатор обязан отвергать ненулевой вывод и конечный LSE.

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

3.11 TileOPs: эталонные вычисления бенчмарка переводятся на float32 (#2692)

ref_program ранее служил и эталоном корректности, и объектом хронометража базовой линии torch; многие эталоны вызывали torch напрямую с 16-битным dtype входа, и результат зависел от того, как torch на стороне устройства округляет на каждом шаге. CUDA скрывал это различие, а CPU torch и torch_musa — нет: GroupNorm и InstanceNorm падали на MTT S5000, тогда как тот же kernel проходил проверку против эталона float64. Исправление заставляет ref_program вычислять на float32 или шире, а базовую линию torch хронометрировать по dtype входа. Это редкий случай, когда проблема выявлена бэкендом отечественного GPU, но исправлять её пришлось на уровне эталонной реализации.

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

3.12 TileFoundry: добавлены архитектура sm_89 и устройство RTX 6000 Ada (#259)

11 файлов +392/-22. Добавлены два описания целей nvidia_sm89.toml и nvidia_rtx_6000_ada.toml (114 и 124 строки соответственно), в docs/spec/target.md синхронно добавлено 42 строки. Этот коммит расширяет моделирование устройств TileFoundry с Hopper до Ada и представляет собой другую сторону того же явления, что и выявленное в тот же день в #256: тесты, скомпилированные под цель sm_90, на однокарточной системе sm_89 дают только неверные значения; попутно добавлено объявление e4m3 warp mma для sm_89 (#261, на рассмотрении).

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

3.13 TileFoundry: учёт gmem переведён на занятый объём и рабочий mesh, а также привязку измерений вызываемой стороны (#258)

Самый масштабный по объёму изменений коммит дня: 74 файла +1128/-553, исправляет #239, #252, #253. Ранее, когда из 132 CTA исследовались только 40, учёт вёлся как «одно перемещение на каждую позицию каждого окна», и under-fill был совершенно не виден; измерения, объявляемые только вызываемой стороной, не привязывались; смысл выходной цели analyze и построчных строк памяти по функциям был неясен. Исправление сосредоточивает привязку измерений check и analyze в instantiate, и теперь измерения, объявленные вызываемой стороной или подмодулем, могут быть привязаны. Коммит инициирован участником сообщества и переписан мейнтейнером.

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

3.14 TileFoundry: при отказе устройства запустить kernel необходимо выбрасывать ошибку (#256)

7 файлов +69/-11. Ранее лаунчер после выдачи kernel<<<grid, block, smem, nullptr>>>(...) никогда не читал состояние запуска; когда устройство отказывало в запуске, вызов возвращался как обычно, вывод сохранял прежнее содержимое, и ошибка не доходила до вызывающей стороны; на однокарточной системе sm_89 тесты, скомпилированные под связанную цель sm_90, проявлялись лишь неверными значениями, и только compute-sanitizer показывал, что cudaLaunchKernel возвращает cudaErrorNoKernelImageForDevice. Исправление заставляет CUDA-лаунчер выбрасывать ошибку после чтения cudaGetLastError(), модуль на стороне CPU синхронно изменён, спецификация и тесты обновлены вместе.

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

3.15 Дистрибутив: Sunrise выпускает PR релиза 0.1.15+sunrise.1.2.0 (#9, на рассмотрении)

10-22 16:48 создана ветка candidate/sunrise-0.1.15-1.2.0-20261019 и открыт PR, который обновляет дистрибутив Sunrise до 0.1.15+sunrise.1.2.0 и включает TileLang 0.1.15 вместе с сопутствующими проверенными изменениями TileOPs, TileKernels, TVM и tvm-ffi; импорт исходников и слияние, а также адаптация дистрибутива выполнены отдельными коммитами, в финале адаптирована конфигурация Sunrise wheel для бэкенда TANG и исправлена регрессия slice GEMM (оба флага аккумулятора проверяются на публичном шаблоне, сохранены проверки stride, origin и offset). Самопроверка: 55 контрактных тестов на локальной машине пройдены, 54 теста codegen конвейера TANG пройдены, два lint job и проверка границ доверия пройдены; однако после сборки пакета S2 wheel установочная дымовая проверка импорта падает на tangGetDeviceCount. Три коммита в ветке — это адаптация дистрибутива 0.1.15, сужение CI до бэкенда TANG и тест флагов аккумулятора slice GEMM.

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

3.16 Принимающие стороны: TileKernels и FlashQLA без пушей в окне (10-22)

Последний пуш deepseek-ai/TileKernels остановился на 10-14, QwenLM/FlashQLA также остановился на 10-14, у обеих принимающих сторон в этом окне нет коммитов.

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

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

4.1 Сообщество CANN открыло репозиторий операторов ops-tilelang (10-22)

Atlas-Ascend-Fork/ops-tilelang создан 10-20 и позиционируется как репозиторий высокопроизводительных операторов TileLang сообщества CANN для Ascend NPU: централизованно управляет реализациями ядер, автоматизированными тест-кейсами, стандартными рабочими нагрузками и базовыми показателями производительности, цель — сформировать переиспользуемые, проверяемые и устойчиво развивающиеся активы операторов. В репозитории уже есть четыре уровня src, tests, scripts, docs и файлы спецификаций открытого атомарного сообщества, такие как OAT.xml, SECURITY.md; 10-22 запушен коммит, исправляющий “переполнение после комбинирования легального верхнего seed с хешем тест-кейса, из-за чего падал CPU-тест”; в более ранних коммитах есть “замена ссылок Ascend A5 на Ascend 950” и “сессии с неудачным или прерванным benchmark всё ещё перезаписывают базовые показатели производительности”. Репозиторий операторов уровня сообщества с единственным DSL в качестве единственного языка реализации ранее не встречался, за этим направлением стоит следить еженедельно.

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

4.2 tilelang/agent: оптимизация операторов как SKILL-воркфлоу агента (10-22)

tilelang/agent (создан 08-28, 3 звезды) запушен 10-22, сам README — это исполняемый шаблон промпта: внедряет доменные знания через @tilelang-operator-dev/SKILL.md, на входе только имя оператора, эталонная реализация GPU/CPU, shape и ограничения “tilelang npu поставка, китайская поставка”; в комментариях указано, что агент должен выполняться в среде A5 и после загрузки проходить замкнутый цикл из 10 шагов, Skill содержит дерево решений по выбору, правило 8-CSV, стратегии от L0 до L4 и 11 ловушек. Если смотреть вместе на репозиторий операторов из 4.1, инструкционно-уровневое планирование из 3.15 и этот репозиторий, на стороне Ascend складывается разделение труда “человек пишет правила, агент тюнит, в репозитории оседают базовые показатели”.

Источник: https://github.com/tilelang/agent

4.3 Новые репозитории сообщества: обёртка для Sail, агент удалённого автотюнинга и репозитории руководств (10-22)

flytiger-eco/tilelang-wrapper-for-sail создан 10-21, запушен 10-22 (Python, без описания); cc-c122/TileLang-Remote-Autotuning-Agent (Python, запушен 10-22) и easy-tilelang/tilelang-devkit, easy-tilelang/tilelang-tutorials запушены в тот же день. Их много, звёзд у отдельных репозиториев мало — это естественный рост уровня руководств и инструментов, технического направления, по которому можно судить, пока нет.

Источник: https://github.com/flytiger-eco/tilelang-wrapper-for-sail

4.4 Документация и версии: сайт документации обновлён трижды, VERSION всё ещё 0.1.15 (10-22)

tilelang.github.io трижды Update docs, TileFoundry.github.io один коммит render TileFoundry v0.0.3, TileOPs.github.io один коммит. VERSION основного репозитория остановился на 0.1.15, последний релиз по-прежнему v0.1.15 от 09-30; у репозитория 8742 звезды, 915 форков, 389 открытых issue.

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

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

1. Типы дефектов в основном репозитории сместились на другой уровень. На прошлой неделе в основном репозитории устраняли «проверку форм и геометрии» — аннотации округления cast молча отбрасывались, CUDA target без arch молча откатывался к sm_50, формы разреженных метаданных не проверялись. В этом окне четыре из девяти коммитов касаются состояния времени компиляции, такого как анализатор и позиции в исходном коде: анализатор переиспользуется после перезаписи, переменные перепривязываются после разбиения, eager-фронтенд теряет позиции Python. Общее свойство таких дефектов — «сгенерированный kernel может оставаться корректным, но то, что компилятор говорит вам, не совпадает с тем, что он фактически делает», а TileLang сейчас подключает agent для автоматического написания kernel (4.2), и согласованность состояния времени компиляции напрямую определит, сможет ли замкнутый цикл итераций agent сойтись.

2. Прогресс на стороне Ascend сместился с количества операторов на низкоуровневые пути передачи данных. #3530 — один коммит на 18 файлов, реализует выражение межстрочного расстояния для сжатого layouts NZ в части копирований UB→L1; #3581 и #3523 соответственно исправляют взрыв предикатов времени компиляции и семантику размещения локальных переменных SIMT-векторов. Ни один из трёх не добавляет операторов — все они делают существующие пути компилируемыми и корректно исполняемыми на большем числе форм. Вместе с тем, что сообщество CANN заводит отдельный репозиторий операторов ops-tilelang (4.1), можно прочитать разделение труда: операторные активы несут репозитории сообщества, основной репозиторий сосредоточен на семантике бэкенда и codegen.

3. TileOPs — 20 коммитов за два дня подряд, центр тяжести работы — специализация SM90 и доработка SM89. В окне добавлены три специализированных warp-kernel для SM90 (dense decode, причинно-скользящее окно band, перекрытие FP8 varlen), одновременно продвинуты три маломасштабных пути SM89: SSD chunk scan, paged DSA, NSA top-k. Примечательно, что два исправления NaN (#2708, #2699) оба появились на путях warp-специализации: определение времени жизни в планировщике разделяемой памяти и слияние FMA при fast math — две наиболее подверженные ошибкам области в таких вручную оркестрируемых kernel, и именно это TileLang обещает взять под контроль за пользователя.

4. Обратный индикатор — уровень бенчмарков. #2692 переводит эталонную реализацию на вычисления float32, причина — сбой GroupNorm и InstanceNorm на MTT S5000: проблема в том, что эталонная реализация следует за dtype входа, а не в самом kernel. В тот же день TileFoundry #256 превращает «устройство отказалось запускаться» из молчаливого в выбрасывающее ошибку. Оба коммита в одном направлении: превратить «выглядит прошедшим» в «действительно прошло»; для команд, зависящих от этих инструментов при адаптации чипов, практическая ценность таких исправлений выше нескольких цифр производительности.

5. Линии дистрибуции и индустрии ускоряются. Дистрибутив Sunrise упаковывает 0.1.15 и wheel бэкенда TANG в PR релиза, что показывает: у TileLang уже появилось более одной формы нижестоящей дистрибуции; на индустриальной стороне SCMP перепечатал оценку SemiAnalysis, вынеся «из шести репозиториев пригоден только TileLang» в мейнстримные медиа. По техническому ритму это по-прежнему спокойный день (в основном репозитории нет релизов, нет breaking change), но организационных действий на уровне экосистемы заметно больше, чем на уровне кода.

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

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

Источник Результат проверки
Репозитории организации tile-ai 29 репозиториев, 10 с пушами в окне; влито 35 коммитов, 71 обновление на рассмотрении
Основной репозиторий tilelang 9 влитых коммитов, 1 однострочное исправление, 1 изменение Ascend на 18 файлов; VERSION 0.1.15, новых релизов нет
TileOPs 20 влитых коммитов, все — kernel и тесты Attention / MoE / Elementwise / Reduction
TileFoundry 3 влитых коммита, ещё 4 на рассмотрении; пояснение о релизе 0.0.3 от 10-21 уже включено в прошлый выпуск
tilelang-ascend 2 влитых коммита (1 новый пример, 1 откат примера), 11 на рассмотрении
tilelang-hygon в main нет коммитов; в слитой ветке gcy_merge_main один коммит очистки переключателей
tilelang-sunrise создана ветка candidate и открыт PR релиза 0.1.15+sunrise.1.2.0, не влит
tilelang-metax / musa / mlir-ascend пуши в окне отсутствуют, остановились соответственно на 10-16 / 09-30 / 10-19
TileRT / tilescale / DeepStack пуши в окне отсутствуют
Принимающие стороны TileKernels / FlashQLA обе остановились на 10-14, пушей нет
Google News (несколько наборов на китайском и английском, через прокси) 3 попадания, 2 — перепечатки одного и того же отчёта SemiAnalysis, 1 — о финансировании DeepSeek, не по теме, исключено
Hacker News обсуждений TileLang нет
arXiv поиск по TileLang в заголовке или ключевых словах аннотации, 0 в окне
Новые репозитории сообщества 5 новых или активных репозиториев, даты создания и время пушей сверены
Сайты членов или производителей к данной теме неприменимо

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