Отчётный период: с 2 ноября 2026 г. (понедельник) по 6 ноября (пятница), всего 5 рабочих дней Источники материалов: 5 выпусков ежедневного отчёта о динамике сообщества TileLang (11-02, 11-03, 11-04, 11-05, 11-06); окна сбора пяти выпусков стыкуются встык и в сумме покрывают 168 часов с 07:00 10-30 до 07:00 11-06 (выпуск понедельника имеет 72-часовое окно и охватывает предыдущие выходные; выпуски со вторника по пятницу — 24 часа), пересечения с предыдущим еженедельным отчётом (10-26 — 10-30) нет; кроме того, проведена дополнительная проверка 168-часового окна по всем репозиториям организации tile-ai, страницам релизов и тегам каждого репозитория, а также повторная проверка по четырём источникам TileLang×RISC-V (подробности в приложении)

1. Главные новости недели

  • Постоянный ядерный модуль FlashMLA для Ascend 950 напрямую добавлен в основную ветку — с использованием подхода, планируемого компилятором, а не портированием в отдельный репозиторий (слито 11-02, #3678): один коммит добавляет 1997 строк, удаляет 17 строк, затрагивает 20 файлов и включает 42 коммита, покрывает два пути — BF16 prefill и компактный FP8 decode, а также поддерживает опциональный дополнительный KV-кэш FP8 и FP4, attention sink, недопустимые индексы и независимую длину для каждого query. Цикл внимания, онлайн-softmax, деквантование и планирование полностью выражены на TileLang; с помощью T.Stage компилятор управляет версиями буферов и синхронизацией; QK и PV совместно используют два различных представления layout одного кольца памяти L0B, размещение в L1 задаётся напрямую через T.annotate_buffer_addresses; в prefill одно выделение аккумулятора делится на четыре непересекающихся выходных среза, и каждый срез выгружается сразу после завершения нормализации. Чтобы такой подход стал возможен, на стороне компилятора одновременно добавлены три возможности: копирование из L0C в UB с сохранением вывода NZ, выбор каждого потока AIV перед собственный конвейером, сохранение исходного порядка для явно планируемых стадий при совпадении физической фазы выпуска. Целевые формы заданы очень узко: число голов 64, Dqk и Dv равны 512, topk кратен 64. Источник: https://github.com/tile-ai/tilelang/pull/3678

  • ThreadSync изменён трижды за неделю: сначала уступил место миграции, затем вернул производительность и доказательства (11-01, 11-04, 11-05): в понедельник #3733 полностью откатил точный подсчёт потоков (#3291), слитый три недели назад; причина указана прямо — тот набор изменений переносил дорогостоящие запросы подсчёта перед проверкой барьера полного блока, доказательство выполнялось даже тогда, когда все потоки достигают барьера, бюджет решателя TVM исчерпывался и возвращался частичный подсчёт, а законные вспомогательные ядра remap и sequence просто не компилировались с ошибкой «Cannot prove a thread rectangle»; три улучшения корректности — разделительная защита, точная граница участвующих потоков и нерисуночные множества потоков — были отозваны вместе с ним, а #3732 оставлен как трекинг-тикет для повторной оптимизации после миграции. В четверг #3761 устранил побочный эффект отката: чтобы исправить повторение номеров именованных барьеров из #3643, одна и та же группа потоков не могла повторно использовать ID между областями видимости; ядро mhc_pre_big_fuse из TileKernels с 96 потоками на блок из-за этого получало дополнительный ID на двухэтапном конвейере cp.async, объявленный в скомпилированном артефакте ресурс барьеров менялся с 4 на 5, а верхний предел занятости барьеров для скрытого измерения 1280 на B300 падал с 16 блоков/SM до 12 блоков/SM; исправление состоит в том, чтобы ThreadSync принимал упорядоченный список областей видимости и перед совместным использованием одного ID между областями видимости сначала доказывал с помощью нового анализатора, что предикат участия всегда истинен; критерий приёмки — возврат объявления ресурсов к 4 и занятости к 16 блокам/SM, при этом исправление #3643 не откатывается. В пятницу находящийся на рассмотрении #3787 продолжает закрывать слепые зоны доказывателя: при 192 потоках форма tx < 160 and tx // 32 % 2 == 0, где «граничный член и член взятия по модулю индекса warp соединяются через and», даёт ложное срабатывание; корневая причина в том, что кодирование FloorDiv и FloorMod при комбинации «верхняя граница плюс взятие по модулю» не позволяет установить связь между двумя делениями; исправление состоит в переписывании threadIdx.x в warp_index * warp_size + lane. Источники: https://github.com/tile-ai/tilelang/pull/3733, https://github.com/tile-ai/tilelang/pull/3761, https://github.com/tile-ai/tilelang/pull/3787

  • Миграция TIRx превратилась из «словесного обоснования» в три твёрдых доказательства за неделю, смена базовой платформы переходит в открытую подготовку (11-03 – 11-06): во вторник #3736 удалил tilelang/language/ast/, tilelang/language/parser/ и tilelang/language/tir/entry.py — остатки раннего форка локального парсера TVMScript; добавлена 1 строка, удалено 3384 строки, 10 файлов; языковой слой переведён на вышестоящий tvm.tirx.script.parser, регистрация переопределений AST и операторов находится в tilelang.language.overrides. В четверг #3752, зависящий от tile-ai/tvm#91, и пятничный #3752 превратили rlimit Z3 из жёстко заданных 100 000 в настраиваемый для каждого ядра параметр; автор дословно говорит: «это косвенно исправило множество проблем, связанных с доказателем Z3». В пятницу #3803 отправил полный черновик миграции (draft): vendored TVM заменяется на независимый Apache TIRx, фронтенд, преобразования IR, генерация кода, адаптация JIT и упаковка полностью переводятся на нативную объектную модель, закреплены 7e49a148 из apache/tvm-tirx и 3e906bc6 из apache/tvm-ffi, дерево зависимостей остаётся чистым вышестоящим, без внедрения патчей. Три линии указывают на одно: приоритет номер один текущего этапа основной ветки — смена базовой платформы, все остальные оптимизации должны уступить ей дорогу. Источники: https://github.com/tile-ai/tilelang/pull/3736, https://github.com/tile-ai/tilelang/pull/3752, https://github.com/tile-ai/tilelang/pull/3803

  • «Тихая генерация неверного кода» стала главным типом дефектов недели, охватив четыре релизные версии (11-03 – 11-05): самый тяжёлый — #3808: после скаляризации, связанной с условиями, в epilogue разделяемой памяти не хватает одного __syncthreads(), и версии 0.1.12, 0.1.13, 0.1.14, 0.1.15 все тихо генерируют неверный CUDA, компиляция не выдаёт ни одной ошибки, автор отчёта приложил полностью воспроизводимый скрипт. Из того же семейства: #3800 (после отключения warp specialization T.Pipelined перезагружает следующую stage до ожидания барьера tcgen05, и MMA читает перезаписанный operand), #3721 (на sm_100a при T.gemm по TMEM, когда threads/32 не делит N нацело, в PTX нет ни одной tcgen05.mma, произведение тихо отбрасывается), #3775 (в production TP2 sglang с бэкендом nvrtc захват CUDA graph повреждается, три побайтово идентичные конфигурации воспроизводят три разные первые точки отказа), #1991 и #1992 (на Ascend 910_9362 при threads=2 AscendVidReduction безусловно уменьшает вдвое и внедрённый tmp workspace, sin/cos над tile косвенной адресации считаются неверно для 3978/4096 элементов, то есть 97,1%). Всё это семейство не выдаёт ошибок и проходит обычные тесты, его можно поймать только целенаправленным воспроизведением. Источники: https://github.com/tile-ai/tilelang/issues/3808, https://github.com/tile-ai/tilelang/issues/3800, https://github.com/tile-ai/tilelang-ascend/issues/1991

  • TileFoundry переводит синхронизацию из «доказуемого анализом» в «эмиссия должна соответствовать доказательству» (02.11–06.11): #283 оформляет вопрос «как следует упорядочивать обращения к разделяемой памяти» в слой доказуемого анализа — проходит по каждой функции в порядке эмиссии lowering, планирует commit, wait и барьеры в пределах области видимости перед конфликтующими обращениями; из 83 фикстур планирования запланировано 52, отклонено 31 (24 используют TMA-копирование, 6 — wgmma, семантика завершения обоих ещё не объявлена); поскольку по умолчанию ни один вызывающий не запрашивает это, байты 75 фикстур, доходящих до finalize, полностью совпадают с main. Слитый в четверг #285 — его исполнительная сторона: CUDA lowering эмитирует согласно плану синхронизации и построчно проверяет в обратном порядке, при несоответствии выбрасывает LoweringError; на одной sm_89 из 7 собираемых программ 6 полностью совпадают с эталоном, synccheck и memcheck полностью чисты, тогда как те же 7 программ, пониженные способом main, дают 4 ошибочных результата, непосредственный симптом — максимальное отклонение E4M3 GEMM на sm_89 составляет 659. Слитый в пятницу #286 дополняет асинхронную семантику завершения WGMMA, слитый в среду #284 встраивает семантику обучения FSDP и HSDP в вычислитель HIR и выравнивает на четырёхкарточной sm90 с реальным 2×2 PyTorch HSDP. Автор сохраняет явное ограничение для #286: эмиттер матричных инструкций всё ещё не подключён, этот коммит не доказывает корректность исполнения и не даёт производительности. Источники: https://github.com/tile-ai/TileFoundry/pull/283, https://github.com/tile-ai/TileFoundry/pull/285, https://github.com/tile-ai/TileFoundry/pull/286

  • TileOPs за пять окон слил 15, 7, 1, 7, 9 коммитов, точка конкуренции полностью сместилась с «количества операторов» на «селектор и стоимость подготовки» (вся неделя): из 15 коммитов понедельничного окна 13 приходятся на один день 01.11, почти каждый сопровождается таблицей измерений на H200 и заявлением «при отсутствии измеренной строки в manifest — откат»; 7 коммитов в четверг и 9 в пятницу возвращают ритм к самой плотной послепраздничной динамике. Три твёрдые цифры обобщают направление этой недели: плотный прямой проход GQA выбирает программу decode или prefill по одному измеренному правилу, сожаление выбора относительно самого быстрого кандидата снижается с 1.818 до 1.017 (#2808); decode, специализированный по warp, перекрывает softmax с QK и PV, геометрическое среднее задержки затронутых 478 вызовов сканирования составляет 0.952 от main (#2817); чтение широковещательного операнда на месте снижает время строк широковещания для пяти классов операторов lerp, where, clamp, masked_fill до 0.13–0.44 от основной линии (#2832). В тот же период — переписывание тарификации roofline (три физически невозможные строки пропускной способности исправлены с 7.48, 7.50, 5.30 TB/s на 0.62, 0.62, 2.65 TB/s, #2830) и отдельная подгонка трёхступенчатого тайлового фактора (сожаление выбора снижено с 0.89% до 0.06%, #2836). Источники: https://github.com/tile-ai/TileOPs/pull/2808, https://github.com/tile-ai/TileOPs/pull/2832, https://github.com/tile-ai/TileOPs/pull/2830

  • Отечественные бэкенды переходят от «собственных веток» к «публичному ревью», Hygon впервые открывает официальный PR (05.11): главная ветка репозитория Hygon остановилась на #32 от 02.11 (совместимая обработка gfx948 по gfx946); на этой неделе движение сначала накопилось в ветке feat/hcu-scaled-cast — четыре коммита (преобразование FP8/BF8/FP4 с масштабированием E8M0, фаза чтения LDS с 64 банками, доказательство не более двух путей для смешанного LDS wrap, смешанная проверка по физической ширине чтения), в пятницу добавились ещё три исправления фазы LDS, а в 09:08 и 19:52 открыты два PR — #33 (scaled cast) и #34 (XOR shuffle для reduction с использованием DPP и lane swap) — это первый среди отечественных бэкендов, прошедший весь путь «от собственной ветки до публичного ревью». mlir-ascend в тот же период перенёс адаптер инструментальной цепочки BishengIR в contrib (#209) и тремя небольшими PR #203, #211, #212 перенёс tl.sub/mul/div/max/min, tl.shl, tl.shr по одному в диалект TLIR. Moore Threads 05.11 сделала один документационный коммит README, завершив полное молчание с 30.09. Источники: https://github.com/tile-ai/tilelang-hygon/pull/33, https://github.com/tile-ai/tilelang-hygon/pull/34, https://github.com/tile-ai/tilelang-mlir-ascend/pull/209

  • Аппаратная поддержка расширяется одновременно в обе стороны: новые SKU Blackwell и отечественные ускорители (вся неделя): с одной стороны — новые карты. Заявленный для B300 sm_103a спотыкается на обоих путях по одному разу — tcgen05.mma.kind::i8 существует только в sm_100a и sm_110a, и ядро после lowering получает от ptxas ошибку Feature '.kind::i8' not supported (#3770); вспомогательные функции TMA в CuTeDSL ведут белый список архитектур, и B300 отбрасывается ещё до входа в DSL (#3772; после добавления sm_103a и sm_107a для gr100 на B300 проходят 8 и пропускается 1, на gr100 проходят 13 и пропускается 1). На SM120 GEMM с FP8 и накоплением в fp32 всё время использовал атомы mma.sync от SM89; после перехода на block-scaled MMA с единичным масштабированием ue8m0 на RTX 5090 куб 8192 вырос с 463 до 727 TFLOPS (на рассмотрении #3790). С другой стороны — отечественные карты: Hygon выполняет микроархитектурную настройку фазы LDS с 64 bank, Ascend меняет фиксированное соотношение ядер Cube и Vector с 1:2 на конфигурируемое (#3783) и ужесточает нижнюю границу периода AutoSchedule, снижая время решения Z3 на 10.04% (#3776). Источники: https://github.com/tile-ai/tilelang/pull/3772, https://github.com/tile-ai/tilelang/pull/3790, https://github.com/tile-ai/tilelang/pull/3783

2. Прогресс основных проектов (основной репозиторий tilelang)

Версии и релизы

  • Последний тег и релиз основного репозитория по-прежнему v0.1.15 (выпущен 30.09); за всю неделю не было ни новых тегов, ни новых записей на странице релизов, и к пятнице этой недели прошло уже 37 дней.
  • Обе новые версии в организации остановились на прошлой неделе: TileOPs — v0.1.0 (27.10), TileFoundry — 0.0.4rc1 (25.10); на этой неделе ни одна не стала стабильной.
  • Теги линий адаптации за всю неделю не менялись: tilelang-ascend v0.1.2.000 (09.09), tilelang-metax v0.1.15 (16.10), tilelang-musa v0.1.15+musa.1 (30.09), tilelang-mlir-ascend остановился на серии v0.1.2.020.
  • Объём и ритм слияний: окно понедельника составило 72 часа (30.10 07:00 — 02.11 07:00), слито 35 коммитов; четыре 24-часовых окна со вторника по пятницу дали соответственно 18, 8, 12 и 18 коммитов, всего по пяти окнам — 91 коммит. Обновления на рассмотрении — 82, 44, 48, 42 и 57 соответственно, всего 273; новых issue — 11, 8, 7, 4 и 5, всего 35. Распределение коммитов в основных ветках: tilelang — 7, 5, 5, 3, 5; TileOPs — 15, 7, 1, 7, 9; TileFoundry — 3, 1, 1, 1, 1; tilelang-ascend — 5, 4, 0, 1, 2; tilelang-mlir-ascend — 1, 0, 1, 2, 0; основная ветка tilelang-hygon — 2, 1, 0, 0, 0 (плюс коммиты в функциональных ветках); tilelang-lsp и расширение VS Code — по 1; tilelang-musa — 1.
  • По ритму неделя имела форму «расчистка в начале недели, провал в середине, подъём к концу»: в среду было всего 8 слияний — минимум недели, в пятницу вернулись к 18; очередь на рассмотрении снижалась с 82 в начале недели до 42, а в пятницу поднялась до 57, что говорит о том, что накопившиеся в начале недели ревью были в значительной степени обработаны за эту неделю.

Компилятор, языковой уровень и пути исполнения

  • Ядро FlashMLA для Ascend 950 и сопутствующие возможности компилятора (#3678, смёржен 11-02, 20 файлов +1997/-17): см. первый пункт в разделе главных новостей. Примечательно, что этот коммит добавляет в основную ветку сразу три низкоуровневые возможности, необходимые для «явного планирования», а не патчит их в ветке Ascend — копирование из L0C в UB с сохранением вывода NZ, выбор потока для каждого AIV до конвейера, сохранение порядка исходного кода при совпадении фаз физического выпуска.
  • Два последовательных исправления регрессии векторной ширины, вызванной развязкой cast (#3734 смёржен 11-02, #3759 смёржен 11-05): #3734 исправляет ошибку Cannot convert type float32x16 to CUDA type при прямом проходе ViT на SM100 — планировщик векторизации выбирал 16 lane для 256-битного хранения BF16 и откладывал ограничения cast, а после того как #3579 позволил развязке обходить проверку при чтении памяти, финальная легализация по-прежнему использовала отложенную стратегию, и остаточное выражение FP32 раздувалось до 512-битного вектора; исправление состоит в том, чтобы сделать defer_cast_constraints C++-параметром планировщика, а не свойством IR, — тогда FP32→BF16 в проверке опускается до 8 lane, а успешно развязанные копии BF16 по-прежнему сохраняют 16. #3759 устраняет хвост этой проблемы: int64 cast в условии ранее сжимал ширину вектора данных с 16 lane до 4 lane, даже если условие полностью согласовано на 16 lane; исправление состоит в повторном использовании анализа согласованности условий для освобождения от ограничения ширины cast — условия if, tirx.if_then_else и Select со скалярным предикатом обращаются отдельно, а cast в подвыражениях по-прежнему ограничен.
  • Трёхэтапная эволюция ThreadSync (#3733, #3761, #3787): см. второй пункт в разделе главных новостей. #3733 — явно обозначенная временная мера, в тексте написано «уступить место миграции TIRX», сопутствующее предложение по оптимизации форка TVM tile-ai/tvm#89 закрыто. В 401-строчном изменении #3761 наиболее ценным является эксперимент по атрибуции: автор, изменив только число объявленных в скомпилированном артефакте ресурсов барьера и не трогая ни одной исполняемой инструкции, воспроизвёл то же замедление, превратив «регрессию производительности» из ощущения в одну идентифицируемую переменную.
  • Вычитание на уровне языка и фронтенда: #3736 удаляет 3384 строки унаследованного форка парсера TVMScript (см. третий пункт в разделе главных новостей); риск сохранения описан конкретно — если кто-то напрямую импортирует tilelang.language.parser.operation, он зарегистрирует конфликтующие операторы на tirx.PrimExpr и tirx.IterVar; #3760 вводит ограниченный внутрипроцессный кэш преобразований для вложенных T.macro, сохраняющий только результаты преобразований, не зависящие от контекста, при этом экземпляр макроса привязывается к текущим globals и значениям замыкания, что решает проблему «ядро уже попало в кэш, но макрос по-прежнему разбирается и компилируется повторно»; #3802 добавляет T.annotate_host_mapped_buffers, записывая индексы аннотированных параметров в свойство tl.host_mapped_params PrimFunc, во время выполнения состояние регистрации проверяется через cuMemHostGetFlags, а device-видимый адрес разрешается через cuMemHostGetDevicePointer, обычная страничная память по-прежнему отвергается.
  • Отладка и наблюдаемость: #3703 меняет директиву #line с «привязки только к первой строке каждого оператора TIR» на уровневый выпуск (0 — выключено, 1 — уровень оператора, 2 и выше — перепривязка построчно), реализация находится в общем потоке вывода C-семейства, используемом CUDA, HIP, CPU C, C host, Metal, Ascend, добавлена TILELANG_EMIT_LINE_DIRECTIVES как значение по умолчанию на стороне Python; автор отмечает, что тесты CUDA и HIP были пропущены на его машине, а фактическое поведение cuda-gdb напрямую не проверялось. В тот же день последовал коммит в репозитории LSP, выравнивающий отладку CUDA по новому сопоставлению строк; в четверг ещё один коммит в репозитории расширения VS Code обновил документацию и переименовал артефакт в vsix.
  • Метрики автотюнинга: #3546 меняет замыкание подачи set_autotune_inputs с «параметров, никогда не передаваемых в точке вызова» на get_current_device(); ранее на машинах без CUDA запасная ветка переносила каждый захваченный тензор на CPU, и каждый раунд тюнинга завершался ошибкой из-за несовпадения устройства лаунчера; на двухкарточном Ascend950 три регрессии автотюнинга Ascend до исправления падали полностью, после исправления проходят все. Из того же семейства #37

84 В CI AscendC перед и после пробы копирования между устройствами добавлена синхронизация устройства, чтобы асинхронные ошибки попадали в область действия try ... except RuntimeError — ранее они всплывали на шаг позже и выглядели как не связанный с этим torch.arange(), сообщавший о несоответствии размера memcpy в 4 байта.

  • Ресурсы решателя Z3 из жёстко заданных стали настраиваемыми (#3752, влито 11-06, зависит от tile-ai/tvm#91): rlimit по умолчанию по-прежнему равен 100 000, без настройки поведение не меняется. Это те же две стороны одного направления, что и ужесточение нижней границы периода AutoSchedule на стороне Ascend в четверг (#3776, время решения Z3 снижено на 10,04%): одно меняет бюджет поиска, другое — начальную точку поиска.

Корректность и исправление дефектов

  • ThreadSync и барьеры: помимо полного отката #3733 (см. главное) и кросс-доменного повторного использования #3761 (см. главное), находящийся на рассмотрении #3765 явно принимает «частичный барьер с защитой границ по индексу отдельного потока», когда размерность потоков кратна целому warp; #3767 заменяет внутреннее утверждение StructuralEqual при доступе к фрагменту с несколькими индексами на ValueError с двумя выражениями индексов; #3786 превращает выход целевой области T.copy за границы в ошибку времени компиляции; #3788 исправляет быстрый путь ProveFragmentContains — он возвращал истину напрямую при полном копировании большого фрагмента, пропуская проверку принадлежности потока, из-за чего случай «малый фрагмент владеет потоками с 0 по 127, большой фрагмент владеет только с 0 по 95» ошибочно признавался содержащим.
  • Семейство молчаливо неверного кода: помимо #3808, #3800, #3721, #3775, #1991/#1992 (см. четвёртый пункт главного), на этой неделе также #3774 (перестановка в разделяемой памяти 16×16 выполняет только сдвиг строк, теряя смещение столбцов, 223 из 256 элементов неверны, тогда как аналогично написанная 32×32 корректна), #3773 (в равномерном цикле while три копирования в разделяемую память плюс GEMM: последовательный контроль 100 запусков — 0 ошибок, версия с while — 21 ошибка из 100 и 10 различных выходов), #3763 (вложенное целочисленное деление с округлением вниз (x // -2) // 4 при int16, 64 элементах и 32 потоках выдаёт Let cannot bind the same var to two different values). Эти три вместе с #3775 образуют концентрированную поверхность сообщений об ошибках для комбинации «потребительская карта + WSL + установка через pip».
  • Первопричина анализа зависимостей и вставки синхронизации: #3749 — отчёт о дефекте с наибольшей технической глубиной на этой неделе: две последовательные SIMD-загрузки потеряли зависимость от более раннего копирования GM в UB, первопричина в том, что логическая область, восстановленная MemoryAccessDetector::LogicalPointerRegion, упускает часть физического диапазона загрузки: в минимальном воспроизведении work.shape равно (2, 128), сгенерированный дескриптор чтения покрывает физический [64, 192), а восстановленная область — [0, 128), поэтому анализ зависимостей считает входной DMA не связанным с загрузкой, AutoSchedule переносит копирование входа после загрузки, а InsertSync также не выдаёт set/wait MTE2 к V. В том же окне в основном репозитории добавлено ещё 8 issue, указывающих на генерацию кода и вставку синхронизации двух бэкендов AscendC и PTO (с #3740 по #3749), из которых 7 поданы одним и тем же репортёром, все с воспроизводимыми prim_func и построчно сопоставленным сгенерированным кодом.
  • Сломанное граничное условие сразу даёт внутреннюю ошибку: #3738 исправляет «T.copy из глобальной области меньше разделяемого tile после padding (например, K=100 копируется в BK=112) ошибочно определяется как цельное копирование TMA и жёстко падает»; #3766 исправляет подсчёт байтов TMA для cta_group::2 — байты пары CTA попадают на барьер чётного CTA, тогда как tma_copy с cluster-барьером ранее ожидал cluster_size × tile_bytes только для rank 0, что случайно верно для (2,1,1) и (1,2,1), но неверно для (4,1,1) и (2,2,1).
  • Границы операторов и примеров: #3709 заставляет скалярный GEMM на CPU очищать только целевую область, #3712 заставляет LegalizeNegativeIndex пропускать недостижимые ветви; #3714 добавляет округление к ближайшему чётному для преобразования fp32 в tf32 (cvt.rn.tf32.f32, несатурирующее требует SM90 и выше, сатурирующее — SM100 и выше); #3702 заставляет несколько примеров внимания выводить tile из shared_memory_per_block_optin, на RTX PRO 6000 (sm_120) 9 элементов CUDA из «2 прошли, 7 провалились» становятся 9 полностью прошедшими, на H100 есть опциональный предел 227 KiB, по-прежнему используется исходный tile.
  • Два семантических отклонения на стороне Ascend: #3762 сообщает, что T.simd.vcvt(src, "uint8", sat=False) для вектора int32 молча игнорирует sat=False и всегда сатурирует, документированная семантика заворачивания (-1 становится 255, -128 становится 128, 300 становится 44) на практике оказывается усечением (-1 становится 0, -128 становится 0, 300 становится 255), что влияет на пути преобразования f32 в int8 вроде завершающей стадии int8-квантования; #3789 сообщает, что бэкенд PTO для целочисленного T.atomic_add выдаёт знаковые pto.si32/si64, тогда как атомарный интерфейс PTO принимает только беззнаковые носители, и в тот же день появились исправления в двух направлениях #3792 и #3794.

Среда выполнения, CI и инструментальная цепочка

  • Оба типа проблем со стабильностью CI проработаны: #3784 исправляет кейс автоматической подстройки provisioning в AscendC CI (см. выше), #2833 поднимает RUNNER_VERSION в TileOPs с 2.336.0 до 2.338.2 (после того как GitHub отключил старые runner, все gpu-smoke встали в очередь, а entrypoint.sh с --disableupdate не может самообновиться), #2827 переносит набор off-CUDA в TileOPs в ночные задания (раньше эти два шага в preflight занимали 13–18 минут на ubuntu-latest, после переноса в preflight остались только проверки сбора, отбрасывания тестовых отчётов и обязательной проверки падения cuda_only, около 3 минут).
  • Параллельное слияние: отсутствие post-merge контроля подтверждено на практике: #2824 при слиянии в 21:06 изменил строку roofline группового GEMM так, что теперь обязательно объявлять sizes и выводить из неё G; #2830 слит в 21:08 и по-прежнему содержит только G без sizes; два коммита с интервалом около 2,5 минут, у каждого CI зелёный, но после слияния GPU Smoke в main сразу выдал RowError: row misses ['sizes'], в тот же день исправлено через #2837. Это самая примечательная процессная запись недели: прохождение CI одного коммита не покрывает «семантическое пересечение двух коммитов».
  • Документация и локализация: tilelang.github.io в четырёх окнах мониторинга на этой неделе имел 3, 5, 3 и 3 автоматические сборки Update docs, содержимое соответствует изменениям интерфейсов — иерархическое отображение #line, список областей действия ThreadSync, аннотации host-mapped, округление TF32 и конфигурация Z3; в TileOPs.github.io был один пуш, в TileOPs-nightly только пуш ветки, в основной ветке новых коммитов нет.
  • Детали инструментария на стороне Ascend: #1956 обеспечивает совместимость дампа исключений AI Core с CANN 9.2.0 (добавлены обязательные поля acldumpTensorInfo, пропускаются пустые адреса и тензоры нулевого размера, регистрационный интерфейс заменён с устаревшего aclrtSetExceptionInfoCallback на aclrtExceptionInfoCallbackRegister, порог версии снижен с 9.3 до 9.2); #1932 выравнивает навыки производительной настройки A2/A3 со структурой рабочего процесса Ascend950 и добавляет два навыка: распознавание типа чипа и поиск информации о чипе.

Наблюдение за очередью ревью

  • Самый крупный по объёму по-прежнему полная миграция репозитория Ascend (#1975): на момент обновления 11-04 00:34 достигнуто 607317 добавленных строк, 97595 удалённых строк, 2887 файлов; TileLang-Ascend переносится в основную многобэкендовую архитектуру, существующие компоненты Ascend перемещаются под ascend/a2a3, разделение с основным путём A5 обеспечивается через arch="dav-2201". Автор весьма конкретно описал в описании список задач: различия файлов в трёх каталогах с исходными файлами ещё предстоит свести, статические проверки и CI форматирования предстоит восстановить, документацию и связанное с Agent предстоит перенести.
  • Черновик миграции TIRx (#3803) — самый значимый пункт на ревью за всю неделю: он определяет, в каком темпе последующие исправления на стороне TVM будут попадать в основную ветку, а также нужно ли будет менять связующий слой Ascend, Hygon и mlir-ascend.
  • Очередь ревью TileOPs сосредоточена на селекторах и раскладках: #2836 отдельно подбирает коэффициент (1.36) для трёхфазных тайлов coop2_swap, #2838 исправляет «трактовку K как непрерывного измерения, из-за чего отклонялся легальный вызов TN-раскладки при K=1», #2839 заставляет FP8 групповой GEMM загружать масштабы A через TMA и заменяет у нефьюзированного плотного группирования при K ≤ 1536 тайлы ping-pong 64×256, #2840 добавляет prefill sinks для SM89 и включает наличие sink в идентичность экземпляра компиляции.
  • Динамика числа на ревью: 82, 44, 48, 42, 57. Высокий уровень в начале недели связан с накоплением за прошлую неделю, к середине недели снижение до 42, затем в пятницу рост, синхронный с ростом объёма слияний к выходным, что является нормальным ритмом «закрытие счетов — открытие новых».

3. Многобэкендовая адаптация (Ascend / Hygon / Moore Threads / MetaX)

Ascend: из «цели адаптации» в «первоочередную цель»

  • Сторона первичной реализации: ядро FlashMLA (#3678) нацелено на Ascend 950 и написано прямо в основной ветке, используя низкоуровневые возможности вроде T.Stage и T.annotate_buffer_addresses, предназначенные для явного планирования, а не создавая отдельный репозиторий для портирования; целевая форма — 64 головы, Dqk и Dv равны 512, topk кратен 64. Сопутствующая возможность пакетной пересылки дополнена 11-03: #3689 позволяет записать область [B, M, N] одним T.copy, а при lowering генерируются пакетные дескрипторы ND2NZ, сохраняются страйд批次 и строк источника, транспонированный пакет идёт через DN2NZ, исходные пакеты, выходящие за границу GM, пропускаются; проверка выполнена на Ascend950DT с CANN 9.2, прошли 50 примерных тестов AscendC и 1733 юнит-теста昇腾, из них 28 устройственных用例 BF16/FP32 покрывают хвостовые блоки пакетов и M/N, выход всего пакета за границы, ненепрерывный страйд источника, случай, когда назначение хранения больше источника, и прямоугольное транспонирование.
  • Корректность синхронизации: #1981 исправляет аннотацию синхронизации в copy_pipe_to_l1 — межъядерная коммуникация V→C через PTO на A2/A3 идёт с промежуточной передачей через GM, три этапа пересылки относятся к MTE3 AIV, MTE2 AIC и MTE1 AIC, а в operation_config.h она помечена как PIPE_MTE3, поэтому автоматическая синхронизация ждёт MTE3 собственного ядра AIC, но эта очередь пуста, событие устанавливается немедленно, и чтение L1 в MTE1 может быть пропущено до завершения записи в MTE2; два доказательства воспроизведения: после замены PIPE_MTE3 на PIPE_MTE2 в сгенерированном исходном коде 10/10 полностью восстанавливают точность, локально на последнем главном стволе 10/10 запусков с ошибкой, за один запуск максимум 16384/16384 несовпадающих элементов. Исправление — разделить имена, а не просто заменить: copy_pipe_to_l1 помечается как PIPE_MTE2, добавляется copy_pipe_to_l1_bind с меткой PIPE_MTE1 (прямой путь A5), чтобы не сломать A5, исправив A2/A3.
  • Сведение к основной ветке: RFC #1971 даёт маршрут замены реализации синхронизации — портировать ядро синхронизации основной ветки и отдельный сервис доказательств Z3, подключить через адаптер старого TIR к существующему конвейеру компиляции, сохранить TVM, планирование, планирование памяти, codegen AscendC/PTO и ABI времени выполнения, после приёмки переключить реализацию по умолчанию и удалить старый алгоритм; сопутствующий #1965 переносит z3-версию insert sync из основной ветки в репозиторий Ascend; #1975 — миграция скелета уровня репозитория. Три коммита — три уровня одного пути. В тот же период #3752 позволяет настраивать rlimit Z3 по kernel, #3735 предлагает доказывать непересекаемость модульных индексов в thread sync — это как раз обработка с двух концов проблемы бюджета доказательств, оставшейся от откатов на прошлой неделе: «настройка ресурсов решателя» и «уменьшение форм, которые нужно доказывать».
  • Линия накладных расходов хоста: #1948 заменяет безусловно создаваемый при каждом вызове в CythonKernelWrapper.forward() TVM Analyzer() на ленивое создание с повторным использованием по требованию, на 910B3 vec add для одного блока снижается с 442–467 микросекунд до 105.1 микросекунды, для 32 блоков — с 441–493 до 108.0, сжатие kpool — с 447–450 до 111.3, вызов голого ctypes около 12.5 микросекунды практически не меняется; продолжающий #1967 кэширует в cython wrapper результаты разбора символьных форм, продолжая снижать стоимость на символьном пути, в том же семействе находится рассматриваемый #3757 (ключ по объекту Var, а не по печатному имени, чтобы избежать совпадения разных переменных с одинаковым именем как одного символа).
  • Поставка операторов: #1827 добавляет прямой проход WY Fast, #1828 переносит и дополняет сегментированный обратный проход wy_fast_bwd_split (оба слиты 11-02 около 02:33 с разницей менее минуты), доделывая эту цепочку операторов линейного внимания на Ascend до обучаемого состояния; #1866 переносит прямой проход Chunk O (chunk-гейтированное линейное внимание), используя режим разработчика с четырьмя конфигурациями pass, одним ядром @tilelang.jit и grid-stride циклом с фиксированными 20 ядрами, достигая внутрикристального слияния, нулевого GM workspace и нулевых хостовых операторов ATen; #1879 добавляет пример прямого прохода MHA (раскладка BSHD, трёхдиспетчерский конвейер MMA в режиме Expert, wrapper не содержит никакого кода T.*), все 20用例 cann-bench проходят по точности, среднее ускорение относительно базовой линии torch-npu — 0.97×, совокупный балл 73.13 (локальный показатель 910B2) — 0.97× означает, что чистого выигрыша пока нет, но скелет уже дан; #1966 — рассматриваемый пример GLM-5.3 kpool fp8_mqa_logits с двумя kernel и разделением труда AIV/Cube, 32 головы, размерность 128, размер страницы 64.
  • Инженерные стандарты: #1954 превращает ошибку сегментации на этапе компиляции для T.tile.bitwise_rshift на int64/uint64 в ValueError с фактическим dtype и списком поддерживаемых (разрешены только int16/uint16/int32/uint32), явно фиксируя «не добавлять 64-битные

«Поддержка сдвига вправо»; #3783 позволяет явно конфигурировать или выводить по устройству соотношение ядер Cube и Vector (vector_core_num / cube_core_num, при неудаче — откат к 2), и проносит выбранное число групп через домен sid, распределение задач, синхронизацию и генерацию кода; #3776 добавляет для интервала запуска AutoSchedule три необходимых условия — границу одиночной задачи, границу ресурсной группы и границу самозависимости, инициализирует бинарный поиск максимумом из трёх, все примеры на Ascend 950PR с недельной сборкой CANN 9.3.0 компилируются, результат планирования не меняется, суммарное время этапа решения Z3 снижается на 10.04%.

  • Состояние регрессии: в понедельник в окне ежедневные тесты — 2970 пунктов, все пройдены, 11-01 был один сбой workflow (#1976); в среду и четверг ежедневные тесты зафиксированы как 3030 пунктов, все пройдены, проходимость 100% (#1980, #1988). Объём тестов за неделю вырос с 2970 до 3030, что говорит о продолжающемся расширении покрытия, но проблемы #1991/#1992 и #3762, обнаруженные только по ручным сообщениям, показывают, что «проходимость» нельзя считать заменой показателя корректности.

Hygon: из собственной ветки в публичное ревью

  • Основная ветка: #31 слит 10-30, улучшена динамическая граница MLS и обработка stride в MatrixStore, исправлено опускание clear_accum в MMAC и blockscaled GEMM, на пути постоянного WASP использованы M-разделённые потребители и независимые барьеры A/B с сохранением чередования slot и phase; #32 слит 11-02 в 10:46, gfx948 трактуется как совместимый с gfx946, переиспользуются существующие MLS atoms, MMAC, barrier и опции компиляции, цель — прежде чем подключать gather load, заставить существующие builtins работать на новой модели. Это типичный двухшаговый подход: сначала повесить новый номер архитектуры на существующий путь и получить измеримую базовую линию, затем добавить то, что действительно требует новых инструкций.
  • Функциональная ветка: feat/hcu-scaled-cast за восемь часов 11-04 получила четыре коммита — преобразование FP8/BF8/FP4 с масштабированием E8M0 (скалярный и упакованный lowering, добавлены 223 строки HCU scaled_cast.h и lower_scaled_cast.cc, основной файл преобразований изменён на 650 строк), вывод swizzle для plain-copy по геометрии bank-word и явное обратное преобразование для узких b16 tile (12 файлов +1536/-112, добавлено 5 тестовых файлов), доказательство «не более двух путей» для гибридного LDS wrap, переход проверки гибридного wrap на физическую ширину чтения; 11-05 добавлены ещё три коммита (адаптация swizzle по умолчанию к 64 bank, восстановление планирования swizzle на потребителе MN-срезов, layout-осведомлённое чтение K-раздела).
  • Состояние ревью: 11-05 в 09:08 открыт #33 (scaled cast, то есть основное тело ветки), в 19:52 открыт #34 (оптимизация XOR shuffle для reduction через DPP и lane swap, плюс исправление «вывод layout сначала распространяет известные ограничения, затем распределяет свободные кандидаты»). Переход от «на функциональной ветке накоплено четыре коммита, PR не открыт» к двум официальным PR — первый случай прохождения этого пути среди отечественных бэкендов; тематика сосредоточена на конфликтах банков разделяемой памяти и выравнивании фаз, что показывает: бэкенд перешёл от «компилируется» к этапу «настройка под микроархитектуру».
  • Есть также находящийся на рассмотрении #19 (заголовок только Distributed, описание пустое, обновлён 11-03 в 11:06), детали реализации в этом отчёте не интерпретируются.

Moore Threads (MUSA) и MetaX: один восстановил поддержку документации, у другого только синхронизация с upstream

  • Moore Threads: 11-05 в 19:46 в основную ветку попал коммит docs: update README installation and operator development references, это первый коммит в данном репозитории с 09-30; тогда же обновлены метки времени 6 оставшихся PR (обновление зависимостей, подмодуль TileOPs, документ о пороге доказательств интеграции MUSA), но содержание — старые предложения с апреля по август. В оценке следует быть осторожными: коммит документации говорит о том, что репозиторий кем-то поддерживается, но недостаточен для доказательства возобновления разработки бэкенда, настоящий сигнал — появление последующих коммитов операторов или lowering.
  • MetaX: ежедневный отчёт во всех пяти окнах фиксировал «нет push», при перепроверке в 168-часовом окне подтверждено, что в основной ветке на этой неделе действительно были коммиты, но при поштучном просмотре все являются синхронизацией upstream tilelang (#3760, #3752, #3759, #3802, #3714, #3784, #3772, #3534, #3761, #3689, #3546, #3736, #3758, #3748, #3678, #3734, #3702, #3703 и др.), авторы — все мейнтейнеры upstream, собственных коммитов со стороны MetaX не обнаружено. Поэтому точная формулировка: основная ветка синхронизируется с upstream, собственная разработка по-прежнему в период тишины, уже третье окно подряд. Ветка dev последний раз останавливалась на 10-16.

mlir-ascend и TLIR: послойное воплощение семантики в диалекты

  • Структурная организация: #206 смёржен 31.10, удалены ссылки на внутренний кэш-сервис в CI — очистка для воспроизводимости; #209 смёржен 04.11 в 16:35, код адаптации инструментальной цепочки BishengIR перенесён в contrib, «клеевой слой, который требует частых изменений по мере эволюции вышестоящей инструментальной цепочки», отделён от ядра lowering, и при последующих обновлениях BishengIR область изменений сужается до contrib.
  • Рабочий процесс Agent: #208 смёржен в ветку harness (не в основную — при ссылке не следует писать «смёржен в основную»), добавлено 2715 строк, 66 файлов; для harness операторов добавлен первый набор рабочего процесса Codex: при инициализации проекта используются итеративные вопросы, подключение существующего репозитория операторов, восстановление разработки из персистентного контекста, настройка или анализ; три набора навыков совместно используют редактируемые Markdown-записи в .workspace/, при передаче задач различаются «текущий кандидат» и «проверенный лучший кандидат», сохраняются неудачные и откатанные эксперименты. Автор явно указывает, что в этот раз реальных измерений на NPU не было; адаптация выполнения на NPU и распространение через npm отнесены к последующей работе.
  • Диалектные операторы: #203 добавляет в TLIR tl.sub/mul/div/max/min и дополняет lowering из TL в HIVM, #211 добавляет tl.shl, #212 добавляет tl.shr; три коммита продвигались последовательно 05.11 в 01:42, 19:44, 20:58. Эта линия идёт медленно, но направление единое: семантика TileLang слой за слоем переносится в диалект MLIR.

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

TileOPs: третья неделя после 0.1.0 — переход от настройки к инженерии селекторов

  • Ритм слияния: в понедельник в окне 15 коммитов (13 из них — 11-01), во вторник 7, в среду 1, в четверг 7, в пятницу 9, за неделю всего 39. Во вторничной партии появились два типа непроизводительных изменений — тестовая инфраструктура (цель tileops_test в #2816, 1281 строка, 59 файлов, благодаря чему дымовые тесты tests/ops больше не обязаны выполняться на CUDA) и CI-контракт (#2818), — что говорит о том, что библиотека операторов начинает структурную подготовку к мультибэкендности и поддерживаемости; четверг и пятница вернулись к чистой производительности.
  • Внимание: #2808 свёл выбор программы для GQADenseFwdOp и GQADenseWithLSEFwdOp в одно правило AttentionCall.dense_decode_program; на входе — длины query и key, batch, группировка голов, head dim, RoPE и число SM; от 2 до 64 позиций query идут через warp-специализированный decode; сожаление выбора относительно самого быстрого доступного варианта программы (геометрическое среднее и максимум) снизилось с 1.818 и 25.2 до 1.017 и 1.81, а также исправлены две реальные ошибки в основной ветке: «score scale, умноженный на log2(e), исчезает в FP32, что приводит к NaN» и «вызов D256 с несколькими позициями на SM90 не запускается из-за нехватки разделяемой памяти». #2817 перекрывает softmax warp-специализированного decode с QK и PV (вероятности остаются в регистрах, для K и V используются отдельные барьеры, разделяемый трёхслотовый кольцевой буфер); на H200 для (16, 8, 4096, 64, 8, 128) float16 — 0.074610 мс против 0.077490 у fa3 (в 1.039 раза, 3.654 TB/s), bfloat16 — 0.074260 против 0.077220 (в 1.040 раза, 3.671 TB/s). #2807 расширяет paged MHA decode на каждый head dim, при котором «разделяемой памяти одного блока хватает»; #2797 удаляет из GQAPagedFwdOp старый paged decode kernel.
  • MoE и GEMM: #2815 заставляет MoE-вход с tight-раскладкой собираться по токенам, одна и та же hidden-строка переиспользуется между несколькими её маршрутами; на (4096, 256, 8, 7168, 2048) bfloat16 — 7.4417 мс против 7.6982 у flashinfer-cutlass (в 1.034 раза, 387.844 TFLOPS), float16 — 7.5040 против 7.9869 (в 1.064 раза, 384.624 TFLOPS). #2805 переводит решётку частичного однократного W4A8 prefill, заполняющую не менее трёх четвертей SM, с stream-K на «один CTA — один tile»; llama-8b-up-b33 сокращается с 23.52 мкс до 17.98 (vLLM 20.42), в 44 сканированиях N/K для M от 33 до 128 вне списка изменение формы планирования быстрее на 5%–36%, медиана 18%, ни один вариант не стал медленнее. #2834 создаёт для NT GEMM с малым M плитку нативной ширины, способную заполнить одну волну; decode-down (128×7168×2048) в основной ветке использует лишь 112 из 132 SM, после изменения опережает самый быстрый вариант сравнения на 3%, gate-up-b32 быстрее основной ветки на 1.4%.
  • Поэлементные операции и свёртки: #2814 переводит bfloat16 Sigmoid и 16-битный SiLUAndMul на обратную величину с масштабом четверти, обходя деление с проверкой диапазона; #2813 обрабатывает 8-байтовые операнды при длинном запуске как «один поток — один вектор» (сначала с помощью roofline-анализа TileFoundry подтверждено, что все такие побитовые операторы ограничены доступом к памяти); #2832 заставляет пять классов операторов — lerp, where, clamp, masked_fill — больше не разворачивать broadcast-тензоры со stride 0 в форму выхода заранее, время выполнения broadcast-строки снижается до 0.13–0.44 от основной ветки; #2829 сжимает унарные поэлементные операции «не более двух дешёвых инструкций в теле, суммарный объём чтения и записи не более трёх четвертей L2» в одну волну из 256 потоков и попутно исправляет старую проблему ReLU, возвращающего 0 на вход NaN; #2828 для построчного блочного ArgMax/ArgMin при условии «одна волна резидентно» даёт каждому потоку восемь 16-байтовых векторов, bf16 ArgMin для (2048, 4096) сокращается с 8.704 мкс до 7.393 (FlagGems 8.543). #2835 заставляет перестановку осей и форму выхода у reduction-kernel вычисляться один раз на экземпляр; полное время горячего вызова для шести конфигураций FP32/BF16 VarMeanFwdOp снижается на 4.33%–6.17%, сгенерированные CUDA- и host-исходники полностью совпадают — экономится стоимость подготовки вызова, а не время GPU.
  • Дефекты класса точности исправлены напрямую: #2811 добавляет к каждой fp32-величине, построенной из строк, инвариант «ограничена только разбросом строки, не смещением», реализуется взятием некоторого элемента внутри строки за опорный и последующим вычитанием (для warp-доминирующих ядер — первый элемент warp, для фрагментированных ядер — первый элемент строки), устраняя переполнение вблизи 1e36; #2821 инициализирует слияние дисперсий от первой реальной частичной величины, а не от пустого аккумулятора, дисперсия константного входа больше не даёт NaN (tests/ops/test_reduce.py вырос с 155 до 158). Эти

Подобные «пустому аккумулятору + большому смещению» числовые ловушки проявляются только на формах, где сохраняемые оси сворачиваются одновременно до и после, обычное свёртывание по всей оси их не ловит.

  • Метрическая база становится продуктовой характеристикой: #2830 переписал roofline по принципу «сколько этот вызов реально прочитал» — для плотного, varlen и постраничного прямого прохода GQA K/V считаются только для «строк ключей, которые может увидеть какой-либо query», а таблица RoPE — для фактически достигнутой позиции усечения ключей; GroupedGEMM через новый grouped_gemm_roofline оценивает группу для строки по batch_sizes. После исправления три физически невозможные строки превратились из 7.48, 7.50, 5.30 TB/s в 0.62, 0.62, 2.65 TB/s; #2822 изменил сброс L2 в bench_kernel с «zero_() для int8-буфера» на «чтение приватного int32-буфера размером вдвое больше заявленного L2», поскольку старый flush оставлял около 120 MB грязных строк, обратная запись которых попадала в следующий вызов таймера (Sum для 2048×11008 f32 снизился с 27.81 до 21.38 микросекунд).
  • Селектор и модель затрат: #2825 изменил _select_config для W4A16 на взятие argmin _config_cost по _legal_configs, удалил три исключающих правила — приоритет tile, окно равенства по самому широкому K и joint_switch_margin — и добавил член stage_fill; #2826 вписал контракт калибровки плотного селектора в docstring (три шага: сканирование, подгонка, проверка, включая мягкий отжиг и holdout-проверку); #2836 отдельно подогнал коэффициент (1.36) для трёхступенчатого tile в coop2_swap, regret снизился с 0.89% до 0.06% геометрического среднего, максимум — с 5.04% до 0.99%.
  • Форма релиза и критерии тестирования: tileops_test из #2816 снял smoke-тесты с CUDA, автор при этом явно очертил границу — сравнение с общим эталоном доказывает только бэкенд-независимый контракт и переносимость, но не числовую корректность реального бэкенда, эта линия по-прежнему висит на #2611; #2827 перенёс набор off-CUDA в ночной прогон; #2833 обновил runner.

TileFoundry: синхронный анализ, семантика тренировки и управление именованием

  • Синхронизация: помимо аналитического слоя #283 и исполнительного конца lowering #285 (см. пятый пункт главных новостей), #286 был влит 11-06 13:40: он объявляет условия завершения для каждого эффекта операнда, подвешивает fence, commit и wait WGMMA к HIR-узлам, которым они нужны, lowering потребляет эти шаги в точке выпуска инструкции, а если последующая область не вмещает коллективное ожидание, оно выполняется в области выпуска; одновременно исправлен алиас регистра Reshard, чтобы элемент оставался в одном и том же регистре на каждом потоке, держащем копию. Автор оставил два предупреждения о рисках: эмиттер матричных инструкций всё ещё не подключён, этот коммит не доказывает корректность исполнения и не даёт производительности; семантика завершения TMA рассматривается отдельно. 11-06 20:12 открыт #288, посвящённый учёту «уже завершённого частичного свёртывания» в анализе.
  • Семантика тренировки: #284 был влит 11-05 16:36; прежний вычислитель трактовал reshard как глобальный тензор и упускал разницу вкладов разных участников; этот коммит вычисляет локальные значения по участникам, координатам GPU, вызовам и циклам, gather/scatter и свёртыванию Partial(sum) вдоль оси, делая доказательство алиасов Reshard чувствительным к принадлежности. Приёмка: 775 пунктов основного набора, 778 пунктов бэкенд-набора (13 пропущено на двухкарточном sm90), десять конфигураций HIR прошли на CPU, две стратегии резидентности на четырёхкарточном sm90 отработали два шага обновления с настоящим 2×2 PyTorch HSDP — локальные градиенты, параметры и моменты полностью совпали.
  • Планирование и релиз: #280 и #282 собрали решения о том, «что выпускать на каждый вызов», «каким mesh выпускать» и «как нарезать окно», в одно место — analysis/emission.py; приёмка была побайтовой сверкой — из 83 планировочных фикстур 75 finalize-выходов полностью совпали с основной веткой, остальные 8 точно так же падают на основной ветке; сайтовый репозиторий 10-30 выпустил статью о планировании версии 0.3 «Scheduling, Designed for Agents», перенеся «планирование, ориентированное на агентов» из кода и release notes в публичное повествование.
  • Управление именованием: #287 — issue документационного характера, указывающий на «пять имён одного понятия» в кодовой базе: член mesh называется placement в 133 местах, program в 110, instance в 92, participant в 44, а слово unit одновременно обозначает пустой тип, аппаратный вычислительный блок и члена mesh; кроме того, сигнатуры local_layout в Python и CUDA-рантайме не согласованы. Такие issue обычно не ведут к немедленным изменениям кода, но это неизбежный шаг для компиляторного проекта на пути от «способен писать» к «способен долго читаться людьми».

TileSight: модель производительности от статьи к репозиторию

  • tile-ai/TileSight (24 звезды) — открытая реализация статьи arXiv от 24.07 «TileSight: A First-Principles Tile-Centric Analytical GPU Performance Model from Cores to Clusters»: исходя из формы тайлов, перемещения данных, программного конвейера и профиля аппаратуры, оценивает на CPU стоимость оператора, трафик кэша и время выполнения. 05.11 появился первый PR репозитория (#1): H20 — это Hopper/GH100, план с requires_arch=("wgmma",) ранее выбрасывал UnsupportedError, поскольку профиль архитектуры не объявлял поддержку WGMMA; этот коммит устанавливает support_wgmma=True и явно сохраняет support_utcmma=False, в результате такие планы принимаются, а требования tcgen05 и Tensor Memory по-прежнему отклоняются. Проверка — все 847 тестов на CPU пройдены, автор явно заявляет, что это «не проверка производительности на GPU».
  • Значение этой линии в том, что она образует треугольник с roofline-тарификацией TileOPs и аналитическим слоем TileFoundry: один даёт аналитическую модель, другой — построчную тарификацию, третий — доказуемо корректный план синхронизации. Все три на этой неделе имели содержательные действия.

Принимающие стороны: TileKernels и FlashQLA

  • Оба репозитория принимающих сторон всю неделю не имели коммитов: 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 при скрытой размерности 1280 потеряло 25% бюджета занятости барьеров. Это показывает, что ядра принимающих сторон уже стали регрессионными зондами на стороне компилятора: в те недели, когда TileKernels не пушит код, его ядра всё равно пассивно «тестируются» изменениями основной ветки. Для тех, кто следит за экосистемой, это обратный сигнал, который стоит иметь в виду.

Ветки сообщества и сторонние репозитории

  • Формируется исследовательская цепочка RISC-V: исследователь сообщества zishu0131-png 31.10 за один день оформил три репозитория — tilelang-riscv как общую точку входа (описывает исследовательскую цель, две цепочки компиляции и комбинации версий), numba2tilelang статически разбирает Numba CUDA Python и преобразует в исходники TileLang, tilelang2rvv преобразует поддерживаемые реальные tvm.tirx.PrimFunc в RVV intrinsic C, затем через Clang компилирует в RV64GCV ELF и запускает на QEMU virt в режиме одного харта. Автор чётко очерчивает границы: это не официально зарегистрированный бэкенд TileLang; адресация, зависящая от threadIdx, межхартовая синхронизация и неподдерживаемый IR отвергаются, блоки выполняются последовательно; исследовательская среда — TileLang 0.1.15, Clang 23.1.3, QEMU 11.1.0. Обход четырёх источников в 168-часовом окне на этой неделе не выявил новых изменений.
  • PTO-ISA/tilelangGym выделен в отдельный репозиторий: этот набор тестов чувствительности и регрессии выделен из examples/ascend/t_parallel_st_suite репозитория PTO-ISA/tilelang с целью позволить набору тестов независимо фиксировать версии TileLang и PTOAS; текущие три основных пути — Ascend SoftLib SimtVF (target=ascend, cython, основные регрессионные ворота), VectorizeParallelToPTO и vload/vstore для VL, а также ручные эталонные реализации SIMD DSL на @pto.jit в benches; каждый каталог тест-кейса содержит config.toml и VERDICT.md.
  • Небольшие сторонние репозитории: sablin39/tilelang-cuda-skills (149 звёзд, набор навыков для написания на TileLang и отладки с помощью инструментов CUDA), xwhzz/tilelang-hexagon (2 звезды, перенос TileLang на DSP Qualcomm Hexagon), Index4NekoMeteor/Tilelang-AttnRes (1 звезда, реализация на TileLang структуры attention residual из Kimi K3), Tianchen-ai/tilelang-weft-bench (0 звёзд, унифицированный движок бенчмарков операторов на стороне CPU), flytiger-eco/tilelang-wrapper-for-sail (0 звёзд, описания пока нет). Первые три направления интересны — одно занимается опытом разработки на стороне инструментальной цепочки, два других — портированием на новую целевую архитектуру и передовые структуры моделей; последние два очень ранние, пока заносим в список наблюдения, не стоит писать о них как о техническом прогрессе.

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

  • Поверхность сообщений о сбоях сосредоточена на форме установки: в среду из 7 новых issue четыре поступили от одного и того же внешнего репортёра, использующего pip-пакет, работающего в WSL и использующего потребительскую карту sm_89 (RTX 4060 Laptop), все с воспроизводимыми скриптами и полным выводом; технически наиболее содержательный #3775 — это комбинация в производственной среде «pip-пакет + две H200 NVLink + множественные контексты CUDA + бэкенд nvrtc + захват CUDA graph», и репортёр полностью описал контрольные условия трёх запусков — три байтово идентичные конфигурации воспроизвели три различные первые точки отказа, а различие между /proc/environ и чистым контролем состояло лишь в одной переменной окружения TILELANG_EXECUTION_BACKEND=nvrtc. Это не глубокие дефекты компилятора, но они определяют опыт нового пользователя при первой попытке.
  • На стороне Ascend возник замкнутый цикл «репортёр он же рецензент»: авторы #1991/#1992 были рецензентами #1897, а исправление в #1897 покрывало только семантику со стороны count (HasModifiedBuffer пропускает слоты workspace), сторона workspace осталась необработанной — другая грань той же логики компоновки была указана лишь две недели спустя. Это показывает, что инварианты компоновки бэкенда Ascend ещё не оформились в проверяемый письменный контракт.
  • Цепочка документации и локализации работает нормально: tilelang.github.io за всю неделю получил несколько автоматических синхронизаций Update docs, TileOPs.github.io — один пуш; репозиторий LSP выровнял отладку CUDA по новому отображению строк (71e762e0), репозиторий расширения VS Code обновил документацию и переименовал артефакт в vsix, вслед за «выравниванием отладки CUDA по отображению строк TileLang» от 11-02 — построчная привязка, вошедшая в основную линию прошлого выпуска, изменила источник номеров строк генерируемого кода, и конфигурацию отладки на стороне расширения пришлось менять следом.
  • Учебные материалы и интеллектуальные активы: на этой неделе новых учебных репозиториев не появилось, усилия сообщества сосредоточены на поддержке существующих репозиториев и двух небольших репозиториев по направлению RISC-V, упомянутых выше; выделение PTO-ISA/tilelangGym в отдельные репозитории — первый случай, когда тестовые активы получили отдельное версионирование.
  • На новостной стороне за всю неделю ноль содержательных попаданий: в пяти окнах сбора по запросам Google News на китайском и английском (охватывающих точные термины, GPU kernel, китайские операторы и компиляторы, Ascend и MetaX/Hygon/Moore Threads) самая свежая релевантная публикация по-прежнему — SemiAnalysis «The Huawei-DeepSeek Symbiosis» от 10-21, остальные сосредоточены вокруг волны 09-30 о «DeepSeek с открытым исходным кодом версии TileLang для Ascend, DeepGEMM, DeepEP», все раньше текущего периода; поиск по tilelang на Hacker News за 24 часа дал 0 записей; последние попадания на arXiv — «Beyond Compression Ratio: High-Throughput KV Compression for Disaggregated LLM Serving» от 11-03 (работа, цитирующая TileLang, вне окна) и сам TileSight от 07-24. Основное содержание отчёта поэтому полностью исходит со стороны кода, и это уже устойчивая форма на протяжении нескольких выпусков подряд.
  • Расхождение формулировок на промышленной стороне стоит отметить: в окне среды Wccftech сообщил, что DeepSeek будет обучать V5 с нуля на Atlas 950 SuperPoD, построенном на Ascend 950DT, тогда как The Information в тот же период сообщил о «начале предобучения на 950DT не ранее конца 2026 года», а ранние публикации Bloomberg и Financial Times утверждали, что заказ на 160 тысяч чипов 950DT предназначен в основном для инференса, а после неудач обучения на Ascend произошёл откат к Nvidia. Три версии не согласуются в вопросе «будет ли и когда использоваться для обучения» и ни одна не подтверждена официально Huawei или DeepSeek; настоящий отчёт фиксирует это лишь как фон и не делает на этом основании выводов о темпах реализации вложений на стороне Ascend.

6. Наблюдение трендов

Во-первых, миграция TIRx уже фактически стала первым приоритетом основной линии и начинает проявляться в форме «вычитания». За неделю три свидетельства: в понедельник ради миграции откатили целый набор улучшений корректности ThreadSync (#3733, одновременно закрыв предложение об оптимизации на стороне форка TVM #89), во вторник удалили 3384 строки унаследованного форка парсера (#3736), в пятницу отправили полный черновик миграции (#3803, с пином двух апстрим-коммитов и явным отказом от инъекции патчей). Между ними также прошло превращение ресурсов Z3 из жёстко заданных в настраиваемые (#3752 и tvm #91). Общее у этих действий — снятие сложности с критического пути: лучше временно вернуться к консервативной реализации, чем позволять внутренней структуре вендорированной зависимости продолжать расти в языковом слое. Цена — наличие в основной ветке в окне миграции ряда «заведомо неоптимальных» путей, выгода — после завершения смены базы все нижележащие бэкенды смогут адаптироваться к стабильной границе IR.

Во-вторых, риск корректности сменил форму: от «не компилируется» к «компилируется, работает без ошибок, но даёт неверный результат», причём скрытно через версии. На этой неделе как минимум семь пунктов относятся к этой группе: #3808 — отсутствие одного __syncthreads() на четырёх версиях от 0.1.12 до 0.1.15; #3800 — на пути tensor memory в Blackwell MMA читает перезаписанный operand; #3721 — на sm_100a молча отбрасывается всё произведение GEMM; #3775 — повреждение захвата CUDA graph в production; #3749 — анализ зависимостей не видит физический диапазон загрузки; #1991/#1992 — при threads=2 на Ascend sin/cos ошибочно вычисляются для 97,1% элементов; #3774 и #3773 — неверные результаты на потребительских картах. Эта группа наносит пользователям наибольший ущерб, поскольку превращает предпосылку «компилятор прав» в гипотезу, требующую проверки. Направление реагирования на этой неделе — сделать корректность опровержимым доказательством: #283/#285 в TileFoundry — сначала доказательство, затем генерация, после генерации — обратная проверка; #3786 и #3787 — замена молчаливой генерации неверного кода на ошибку времени компиляции или явный отказ, — а не дальнейшее наращивание количества модульных тестов.

В-третьих, роль Ascend сместилась с «цели адаптации» на «цель первого выпуска», инженерия продвигается одновременно по двум направлениям. Одно — архитектура: такое ключевое ядро, как FlashMLA на собственном пути декодирования DeepSeek, напрямую нацелено на Ascend 950 и написано в основной ветке, причём используются низкоуровневые возможности вроде T.Stage и annotate_buffer_addresses, предназначенные для явного планирования; вместе с тремя сопутствующими изменениями на стороне компилятора путь Ascend переходит от «может работать» к «может работать согласно замыслу планирования». Другое — каркас: #1975 переносит весь TileLang-Ascend в многобэкендовый каркас основной ветки — 2887 файлов, 600 тысяч строк добавлений, — с разделением бэкендов через arch="dav-2201"; #208 в mlir-ascend записывает три набора рабочих процессов Codex — инициализацию, подключение и восстановление — вместе с передачей задач в .workspace/. Рассматривая «репозиторий миграции» и «репозиторий миграции с помощью агентов» на одной неделе, можно увидеть, что узкое место на стороне Ascend сместилось с «нехватки операторов» на «нехватку согласованного инженерного каркаса и переиспользуемого процесса миграции».

В-четвёртых, конкурентная точка библиотеки операторов полностью сместилась к качеству селектора и стоимости подготовки; оставшееся пространство разрыва в производительности — не в том, «как вычислять». Среди девяти слияний TileOPs в пятницу действительно меняют GPU kernel только два: широковещательное чтение на месте и нативные широкие тайлы; кэш компоновки reduction экономит подготовку вызова (с 4,33% до 6,17%, сгенерированный исходный код не меняется), а трёхэтапный коэффициент тайлов меняет одну константу (regret снижается с 0,89% до 0,06%). Вместе с переписыванием оценки roofline и переводом сброса L2 в режим только для чтения просматривается чёткий маршрут: сначала перечислить допустимые конфигурации для каждой формы, затем упорядочить их единой моделью стоимости, а процесс калибровки записать как воспроизводимый контракт (#2826 уже вписал три шага — «сканирование, подгонка, проверка» — в строку документации). Если этот путь пройден правильно, производительность больше не зависит от ручной настройки отдельных ядер; если неправильно — он выродится в «некоторые формы особенно быстры, остальные — как повезёт».

В-пятых, отечественные бэкенды переходят от «разового слияния большого бэкенда» к «продвижению в публичный репозиторий минимальных рецензируемых единиц». Hygon прошёл от собственной ветки до двух официальных PR — #33 и #34; Moore Threads завершил двухмесячное молчание одним README; mlir-ascend тремя небольшими PR переносит tl.sub/mul, tl.shl, tl.shr по одному в диалект TLIR; Ascend же поочерёдно дополняет инженерные детали — коэффициенты ядра, дампы исключений, средства производительности. Общее у трёх линий — отказ от ожидания одного большого слияния: для upstream это хорошо, рецензенты могут рано высказаться о дизайне раскладки и lowering, а не смотреть на тридцать тысяч строк разового слияния; для downstream тоже есть цена — основная ветка в такие окна находится в состоянии «известно, что кое-что не работает», что требует более строгой дисциплины регрессии перед релизом.

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

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

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

Источники Результаты проверки
GitHub организация tile-ai по репозиториям (29 репозиториев) В пяти окнах сбора по 10, 9, 9, 8, 10 репозиториев с пушами; при повторной проверке в 168-часовом окне — 14 репозиториев с пушами
Поиск по фактам слияния (is:pr is:merged) В пяти окнах 35, 18, 8, 12, 18 PR, итого 91; первое окно — 72 часа (с 10-30 07:00 до 11-02 07:00)
Поиск по фактам на рассмотрении (is:pr is:open updated) В пяти окнах 82, 44, 48, 42, 57 PR, итого 273
Проверка новых issue В пяти окнах 11, 8, 7, 4, 5 записей, итого 35
Проверка на уровне коммитов основной ветки tilelang 7, 5, 5, 3, 5; TileOPs 15, 7, 1, 7, 9; TileFoundry 3, 1, 1, 1, 1; tilelang-ascend 5, 4, 0, 1, 2; tilelang-mlir-ascend 1, 0, 1, 2, 0; tilelang-hygon основная ветка 2, 1, 0, 0, 0; tilelang-lsp и расширение VS Code по 1; tilelang-musa 1
Страницы релизов и теги tilelang последний v0.1.15 (09-30), за всю неделю новых версий нет; TileOPs последний v0.1.0 (10-27); TileFoundry последний 0.0.4rc1 (10-25); tilelang-ascend v0.1.2.000 (09-09); tilelang-metax v0.1.15 (10-16)
Принимающие стороны TileKernels остановился на 10-14; FlashQLA остановился на 10-30, последний релиз по-прежнему v0.1.3 от 09-30
Обход четырёх источников TileLang×RISC-V За 168-часовое окно без изменений
Сторонние и экосистемные репозитории Три репозитория исследовательской цепочки, tilelangGym, cuda-skills, hexagon, AttnRes, weft-bench, wrapper-for-sail, TileSight — во всех есть пуши или новые создания
Google News (четыре–шесть групп на китайском и английском) Во всех пяти окнах нет существенных совпадений внутри окна; последняя релевантная публикация — статья SemiAnalysis от 10-21
Hacker News Поиск по tilelang за 24 часа дал 0 записей, найдены не относящиеся к теме Show HN
arXiv Последнее совпадение — статья о сжатии KV от 11-03 (ссылается на TileLang, вне окна) и сам TileSight от 07-24

Основные цитирования

Пояснения

  • Факты, номера версий и даты в настоящем отчёте взяты из оригинальных текстов пяти ежедневных отчётов за период и их первоисточников, без предположительных дополнений; пять окон сбора данных идут встык (с 10-30 07:00 по 11-06 07:00, всего 168 часов) и не пересекаются с предыдущим еженедельным отчётом, поэтому числа смёрженных, находящихся на рассмотрении и новых issue за пять окон суммируются напрямую.
  • Числа смёрженных, находящихся на рассмотрении и issue взяты из фактического вывода поиска GitHub; число коммитов в основной ветке — результат порепозиторной проверки atom, эти две метрики имеют разную основу (один мёрж может соответствовать нескольким коммитам, коммиты не в ветке по умолчанию не учитываются в числе смёрженных), в приложении они перечислены по двум категориям раздельно и не суммируются.
  • «15 коммитов» TileOPs приходятся на 72-часовое окно понедельника, из них 13 приходятся на 11-01, поэтому при цитировании не следует писать «15 за один понедельник».
  • Заголовки и текст #1991 и #1992 совпадают, они зафиксированы как «две записи с одинаковой темой» и не объединены в одну; семь коммитов функциональной ветки Hygon и документационные коммиты Moore Threads описаны как динамика на уровне ветки и не означают, что они уже смёржены в основную линию.
  • MetaX во всех пяти окнах отмечена как «без пушей»; проверка за 168 часов выявила коммиты в основной ветке, но поштучная сверка показала, что все они — синхронизация с вышестоящим tilelang (авторы — мейнтейнеры вышестоящего проекта), поэтому формулировка — «в основной ветке идёт синхронизация с вышестоящим, собственная разработка всё ещё в период затишья». Изменения WGMMA в TileSight и TileFoundry самими авторами ограничены метаданными возможностей или планом синхронизации, отчёт следует их ограничениям.