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

Источники: порепозиторная проверка организации GitHub tile-ai (tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-mlir-ascend / tilelang-hygon / tilelang-metax / tilelang-musa / TileRT / tilescale / DeepStack / сайт документации) + поиск подтверждённых данных по смёрженным и находящимся на рассмотрении изменениям (24 смёрженных, 48 обновлений на рассмотрении) + полная проверка новых issue (6) + проверка diff на уровне коммитов и списков файлов + проверка страницы релизов и тегов + проверка сторонних репозиториев бэкендов + шесть групп запросов Google News на китайском и английском (через прокси) + Hacker News + arXiv.

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

  • Ключевое за сегодня: бэкенд NPUIR для Ascend 910 A2/A3 влит в основную ветку mlir-ascend, все 504 самотеста пройдены; селектор TileOPs переведён на калибровку по занятости, regret манифеста снижен с 5.79% до 0.48%; за 24 часа влито 24 коммита, 48 обновлений на ревью, открыто 6 новых issue
    1. Основные результаты проекта
      • 1.1 Основной репозиторий: индексный коэффициент скалярного чтения больше не подмешивается в число lane, устранено падение компиляции при ширине строки выше 32768 (#3453) (28.10)
      • 1.2 Основной репозиторий: пакетный all-reduce разбирает диапазон потоков по скалярной ветке, исправлена ошибка, при которой при batch>1 все 2048 выходов были неверными (#3341) (28.10)
      • 1.3 Основной репозиторий: добавлена проверка формы атомарных областей SIMT, области за границами отклоняются на этапе компиляции (#3465) (28.10)
      • 1.4 Основной репозиторий: хвостовой блок K разреженного MMA больше не пропускается молча (#3367) (28.10)
      • 1.5 Основной репозиторий: целевой dtype atomic_addx4 сужен до float* (#3369) (28.10)
      • 1.6 Основной репозиторий: суффикс handle не-restrict параметров больше не обрезается (#3366) (28.10)
      • 1.7 Основной репозиторий: дескриптор TMA фиксирует только уже проверенные имена (#3484) (28.10)
      • 1.8 Основной репозиторий: унифицирована проверка выходного индекса profiler, для скалярной и списочной форм действуют одни правила границ (#3486) (28.10)
      • 1.9 Основной репозиторий: coalesced_width в T.Parallel единообразно приводится к IntImm (#3374) (28.10)
      • 1.10 Основной репозиторий: в тестовой инфраструктуре чувствительные к cache кейсы переведены на without_kernel_cache (#3695 / #3696) (28.10)
      • 1.11 Добавлены три дефекта: 13 примеров падают на SM120, экспоненциальная компиляция IfStmtBinding, without_kernel_cache не форсирует перекомпиляцию (#3701 / #3700 / #3699) (28.10)
    1. Адаптация под несколько бэкендов (Ascend / Hygon / ROCm)
      • 2.1 Ascend: бэкенд NPUIR для 910 A2/A3 влит в основную ветку mlir-ascend, все 504 самотеста пройдены, из 72 примеров только 1 средовой сбой (#207) (28.10)
      • 2.2 Основной репозиторий Ascend: идентичность сборки Bisheng включена в ключ бинарного кеша (#3691) (28.10)
      • 2.3 Основной репозиторий Ascend: в рукописном конвейере сохранена фаза регистра конфигурации (#3624) (28.10)
      • 2.4 Репозиторий Ascend на ревью: корневая причина и исправление неверных результатов W4A4 int4×int4 GEMM (#1731) (28.10)
      • 2.5 Репозиторий Ascend на ревью: миграция планировщика a2a3, 100 файлов, 18453 строки (#1951) (28.10)
      • 2.6 Репозиторий Ascend на ревью: Carver читает бюджет ресурсов A3 по выбранному SDK (#1955) (28.10)
      • 2.7 Репозиторий Ascend на ревью: зависимость AI Core Exception Dump снижена до CANN 9.2.0 (#1956) (28.10)
      • 2.8 Репозиторий Ascend на ревью: документация переведена на T.Tensor, интерфейс DSL T.Buffer объявлен устаревшим (#1957) (28.10)
      • 2.9 Hygon: ветка возможности LDS swizzle для обычного копирования HCU, 18 файлов, включая 5 новых тестов (#30) (28.10)
      • 2.10 ROCm: влита поддержка ROCm 10 в pip-версии SDK, MI355X больше не определяется ошибочно как gfx900 (#3549) (28.10)
    1. Экосистема и потребители (TileOPs / TileFoundry)
      • 3.1 TileOPs: селектор переведён на калибровку по занятости, regret манифеста снижен с 5.79% до 0.48% (#2784) (28.10)
      • 3.2 TileOPs: для не-NT раскладки задействован транспонированный постоянный ядро, 2112×7168×4096 TN снижен с 195.3 до 175.3 микросекунд (#2789) (28.10)
      • 3.3 TileOPs: для длинных строк cumsum применено двухэтапное seed-сканирование, (8,65536) снижен с 12.75 до 5.18 микросекунд (#2790) (28.10)
      • 3.4 TileOPs: редукция extrema переведена на PTX max.NaN, распространение NaN не изменилось, а регрессия исчезла (#2792) (28.10)
      • 3.5 TileOPs: эффективный строчный tile параллельного cumsum и узкий предзаполняющий блок GDN, два коммита (#2788 / #2786) (28.10)
      • 3.6 Новая тема TileOPs: при последовательных вызовах крупный GEMM по строкам отстаёт от самой быстрой базовой линии (#2791) (28.10)
      • 3.7 TileFoundry: привязка стадий ring и процесс компиляции унифицированы, 70 файлов (#277) (28.10)
      • 3.8 TileOPs на ревью: релизный образ объединяет TileOPs 0.1.0 и TileFoundry 0.0.3 (#2793) (28.10)
    1. Сообщество, руководства и мероприятия
      • 4.1 Два автоматических обновления сайта документации (28.10)
      • 4.2 Пакетная проверка фактов документации и исправлений установки в репозитории Ascend на ревью, семь пунктов продвигаются одним пакетом (28.10)
      • 4.3 В репозитории Ascend на ревью новые примеры операторов: GroupNorm и fla

фрагмент Два оператора (#1607 / #1953 / #1952) (28.10)

  • 4.4 Сторонние: автоматический агент настройки сообщества обновил китайское описание (#74), в основной ветке репозитория бэкенда Hexagon в окне мониторинга нет новых коммитов (28.10)
  • 4.5 Новостная сторона: шесть групп запросов на китайском и английском в Google News, Hacker News и arXiv — без существенных совпадений в окне мониторинга (28.10)
    1. Наблюдения за трендами
  • Приложение: материалы и пояснения по проверке

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

Центр тяжести сегодняшнего окна — на стороне Ascend. tilelang-mlir-ascend целиком влил в основную ветку бэкенд NPUIR для Ascend 910 (A2/A3): один коммит на 55 файлов и 21490 добавленных строк. На стороне Python — пакет tilelang/ascend910 (backend / adapter / codegen / compiler / execution_backend / jit / pipeline / runtime / target, а также уровни диалектов kernel, copy, gemm, npuir в language), на стороне C++ — два набора эмиттеров разработчика и эксперта в src/ascend910/target, уровень op и три прохода LowerNpuirBlock, NpuLoopVectorize, RewriteFDivToMul, в рантайм добавлены npu_launcher.h и npu_utils, установка через install_npuir.sh, псевдонимы target принимают ascend910 / ascend-910 / a2 / a3 / npuir. Приёмка, заявленная автором, — два набора тестов, поставляемых с репозиторием: pytest testing/npuir — все 504 пункта пройдены, из 72 примеров в examples прошёл 71, 0 таймаутов и 0 ошибок; единственный упавший deepseek_v32/fp8_lighting_indexer.py на той же машине точно так же падает и на старом бэкенде — это вопрос окружения и допусков. Второй коммит построчно выравнивает различия в поведении со старым бэкендом, с акцентом на легализацию bf16 (HIVM не принимает операнды bf16, вычисление переводится в fp32 с явным преобразованием) и на опускание копирования для динамического extent.

В TileOPs сегодня влито семь коммитов, обе основные линии — на стороне производительности. Селектор больше не подгоняется вплотную к manifest, а перекалибруется по времени занятости (device_busy_ms) на 160 плотных shape, 60 shape с m 9–32 и 264 shape с m≤8, с покоординатным спуском по softmin ожидаемого логарифмического сожаления; сквозное сожаление по времени занятости на отложенном manifest снизилось с 5.79% до 0.48%, а на отложенном наборе m≤8 — с 14.18% до 1.79%. На стороне GEMM добавлен coop2_swap: не-NT раскладка заменяется на транспонированное постоянное ядро, два warpgroup-потребителя держат по 128 строк, хвостовая запись возвращается через stmatrix.trans и транспонирование TMA, поддерживается смешивание широких и узких tile (например, 2112 = 13×160 + 32 разбивается на tile шириной 160 плюс узкий tile 128), 2112×7168×4096 TN — с 195.3 микросекунды до 175.3 микросекунды, на другой H200 на той же машине — с 174.3 до 158.5 микросекунды, обе быстрее QuACK с 187.9 и 167.0 микросекунды.

На стороне редукций в тот же день сначала сломали, потом починили. Утром #2787 реализовал распространение NaN в редукции extrema на ключах int32, #2788 изменил строковый tile параллельного cumsum — оба коммита приземлились около 07 часов, после чего #2790 и #2792 по отдельности закрыли два регресса: cumsum переведён на сегментирование по seed и два запуска ((8,65536) float32 — с 12.75 микросекунды до 5.18 микросекунды, заодно исправлен сбой запуска в основной ветке при tile выше 65535), extrema переведён на PTX max.NaN / min.NaN и отдельное ядро для строк с нецелыми 16-байтовыми векторами (amax (2048,4096) bf16 — с 8.352 микросекунды обратно до 6.976 микросекунды, быстрее FlagGems). Такой ритм «внедрили и починили в один день» показывает, что ночные бенчмарки TileOPs уже перехватывают регрессы в тот же день.

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

1.1 Основной репозиторий: индексные коэффициенты скалярного чтения больше не примешиваются к числу lane, падение компиляции при ширине строки выше 32768 закрыто (#3453)

Смержено 28.10 в 11:06, один файл +8/-5. Ядро, которое читает целую строку векторизованным способом и одновременно читает отдельные элементы этой строки по индексу, при ширине строки по модулю 65536 не меньше 32768 падает при компиляции с InternalError: Can’t fetch the lanes of a scalable vector. Корневая причина в VectorTypeAccessChecker::OnArrayAccess: он записывает скалярное чтение как access_dtype.with_lanes(me->coeff), а me->coeff — это модульный коэффициент индекса чтения, для x[row, 0] это ширина строки; число lane у DataType занимает всего 16 бит и читается обратно как int16, поэтому 32768–65535 воспринимаются как масштабируемый вектор, lanes() в get_preferred_dtype сразу выбрасывает ошибку, а более широкие кратные тихо переполняются с возвратом. Исправление — хранить коэффициент отдельно в int64_t, использовать его для взятия модуля при проверке совместимости, а в dtype скалярного чтения оставить только тип элемента. Исходный сценарий срабатывания этого дефекта — GroupNorm из TileOPs: обе ширины групп Stable Diffusion 40960 и 163840 его задевают. Проверка двумя сборками wheel для сравнения: в основной ветке компиляция падает на обеих ширинах 32768 и 40960, этот коммит собирается на всех тестовых ширинах.

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

1.2 Основной репозиторий: пакетный all-reduce разбирает диапазон потоков по скалярной ветке, исправлена ошибка, при которой при batch>1 все 2048 выходов неверны (#3341)

Смержено 28.10 в 05:46, 2 файла +170/-8. Скалярная ветка all-reduce перед вызовом запуска разбирает диапазон участвующих потоков (при reducing_threads больше 32 и поддержке целевой платформой именованных барьеров идёт через ResolveAllReduceThreadRange), а пакетная ветка берёт границы блока как смещение и ширину барьера. Когда потоки приведения — это локальный диапазон [32, 96), не начинающийся с 0, пакетный вызов индексирует с потока 0 и ставит барьер по целому блоку, неинициализированные регистры неучаствующих потоков сворачиваются в приведение: при batch=1 корректно, при batch=2 на том же ядре все 2048 выходных элементов неверны, максимальная абсолютная ошибка 1083, и без какой-либо диагностики. Исправление — пакетная ветка использует тот же разбор, что и скалярная, передавая разобранные смещение и диапазон в MakeBatchAllReduce, stride рабочей области по-прежнему равен размеру блока. Добавленный тест запускает batch=1 и batch=2 на одном ядре и сравнивает с эталоном torch.

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

1.3 Основной репозиторий: добавлена проверка формы SIMT-атомарных областей, области с выходом за границы отклоняются на этапе компиляции (#3465)

Смержено 28.10 в 08:14. Операции SIMT atomic ранее не проверяли форму переданной области, и написание с несовпадающей формой доходило до codegen, давая неверный результат или падение. Этот коммит дополняет проверку формы области на обоих путях CUDA и ROCm, отклоняя некорректные области на этапе компиляции.

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

1.4 Основной репозиторий: хвостовой блок K в разреженном MMA больше не пропускается молча (#3367)

Смержено 28.10 в 16:01. Когда в разреженном MMA по направлению K оставался необработанный хвостовой блок, ранее это не вызывало ошибки, а использовались параметры переноса и вычисления для целого блока, давая тихое числовое отклонение. Этот коммит явно отклоняет такие хвостовые блоки K, перенося проблему на этап компиляции.

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

1.5 Основной репозиторий: целевой dtype atomic_addx4 сужен до float* (#3369)

Смержено 28.10 в 16:08, 3 файла +27/-9. Общие реализации CUDA atomic_addx4 и atomic_addx4_ret ранее принимали любой целевой указатель, и цели int32 или int64 вычислялись по семантике float с ошибкой. Этот коммит ограничивает две общие перегрузки целевым указателем float*, существующие специализации fp16 и bf16 не трогаются; API на стороне Python остаётся типонейтральным, целочисленные цели бэкенда CPU не затронуты. На стороне CUDA добавлен регрессионный пример отказа компиляции для целочисленной цели, в C-бэкенде CPU добавлены примеры корректности накопления для двух целочисленных ширин.

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

1.6 Основной репозиторий: суффикс handle у не-restrict параметров больше не отсекается (#3366)

Смержено 28.10 в 13:19. При дедупликации параметров сравнение шло по имени с удалённым суффиксом _handle, из-за чего два разных параметра A и A_handle в одном ядре считались одним и кодогенерация выдавала только один экземпляр. Исправление — дедупликация по точному имени плюс идентичности переменной, а также добавлен регрессионный пример на стороне CUDA с совместным присутствием A и A_handle.

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

1.7 Основной репозиторий: дескриптор TMA записывает только уже проверенные имена (#3484)

Совпадает с 10-28 17:12. Регистрация дескрипторов TMA ранее выполнялась до проверки; имена дескрипторов, не прошедших проверку, оставались в таблице, и последующие легальные дескрипторы с тем же именем ошибочно определялись как дубликаты. Данный коммит переносит регистрацию после проверки.

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

1.8 Основной репозиторий: унифицирована проверка индексов вывода profiler, скалярная и списочная формы подчиняются одному набору граничных правил (#3486)

Совпадает с 10-28 19:24, 2 файла +15/-9, исправляет #2884. Ранее обработка result_idx на стороне profiler и adapter различалась: для скалярной формы проверка границ выполнялась, для списочной — нет, отрицательные индексы также не нормализовались. Данный коммит заставляет BaseKernelAdapter и Profiler использовать один набор правил: принимается None, целое число или список целых чисел, проверка выполняется по диапазону от -len(params) до len(params)-1, легальные отрицательные индексы нормализуются, выход за границы и нецелые элементы вызывают ValueError. Добавлено 28 регрессионных тестов на стороне CPU, покрывающих оба пути — скалярный и списочный.

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

1.9 Основной репозиторий: coalesced_width в T.Parallel единообразно преобразуется в IntImm (#3374)

Совпадает с 10-28 16:27. Когда coalesced_width передавался как Python-целое, некоторые пути напрямую передавали нативный целый в pass, и проверка типов падала в последующих преобразованиях. Данный коммит на уровне языка единообразно преобразует его в IntImm, устраняя сбои из-за различий в записи.

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

1.10 Основной репозиторий: тестовая инфраструктура переводит чувствительные к кэшу случаи на without_kernel_cache (#3695 / #3696)

Совпадает с 10-28 01:59 и 09:19. В продолжение зафиксированного вчера issue #3694 (попадание в кэш ядра ломает тесты, предполагающие холодную компиляцию), два коммита переводят случаи, проверяющие артефакты компиляции, исходный код и диагностическую информацию, на without_kernel_cache() или явный .compile(…), чтобы утверждения видели результаты текущей сборки, а не старый kernel из внутрипроцессного кэша.

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

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

1.11 Три новых дефекта: 13 примеров падают на SM120, экспоненциальное время компиляции IfStmtBinding, without_kernel_cache не форсирует перекомпиляцию (#3701 / #3700 / #3699)

Три новых issue указывают на тесты и время компиляции. Во-первых, #3701: при запуске полного набора examples на RTX PRO 6000 Blackwell (sm_120, лимит динамической разделяемой памяти 99 KiB) оба бэкенда — CUDA и CuTeDSL — показывают 13 падений, 199 прохождений, 27 пропусков, причём набор падений дословно совпадает: 7 случаев — это примеры с жёстко закодированными размерами tile под Hopper на пути без –tune (например, attention_sink требует 163840 байт, декодирование deepseek_mla требует 231424 байта), 4 случая — отсутствие ядра sm_12x в эталонной реализации flash-attn 2.5.8, ещё два случая — невозможность доказать диапазон потоков в понижении ReduceOp и в ThreadSync соответственно. Во-вторых, #3700 сообщает, что при повторной записи x = x + x в цикле с добавлением guard i < N время tl.LayoutInference растёт с 0.087 секунды до 118.5 секунды, а при отключении инлайнинга replayable bind возвращается к 0.100 секунды. В-третьих, #3699 указывает, что without_kernel_cache() лишь переключает CacheState (который читается только дисковым кэшем), тогда как обёртка @tilelang.jit перед JITImpl.compile() также проверяет внутрипроцессные _call_form_cache и _kernel_cache, поэтому обычный вызов внутри этого блока не гарантирует перекомпиляцию.

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

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

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

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

2.1 Ascend: бэкенд NPUIR для 910 A2/A3 влит в основную ветку mlir-ascend, все 504 самотеста пройдены, из 72 примеров только 1 сбой средового характера (#207)

Объединено 10-28 08:04, 55 файлов +21490/-1, автор CeleNewYear, на основе a35f8ddf от v0.1.15, два коммита отвечают соответственно за реализацию бэкенда и согласование поведения. Задняя часть бэкенда включает backend, adapter, codegen, compiler, execution_backend, jit, pipeline, runtime, target под tilelang/ascend910, а также диалектную часть kernel, copy, gemm, npuir под language; на стороне C++ — два набора эмиттеров для разработчика и эксперта в src/ascend910/target, слой операторов src/ascend910/op и LowerNpuirBlock, NpuLoopVectorize, RewriteFDivToMul в src/ascend910/transform; во время выполнения требуются npu_launcher.h, npu_utils.cpp, npu_utils.py для генерации совместимой с torch пусковой обёртки и оценки каждого блока scratch по *_infer_workspace_shape_function; на стороне сборки подключается CMakeLists и добавляется опция –bishengir-path в install_npuir.sh; псевдоним target принимает ascend910, ascend-910, a2, a3, npuir, переключатель режима TILELANG_ASCEND_MODE принимает expert или developer, по умолчанию expert. Данные приёмки: pytest testing/npuir 504 пройдено, 0 провалов; examples/run_all.py –sequential 72 примера, 71 пройден, 1 провал, 0 таймаутов, 0 ошибок, единственный провал deepseek_v32/fp8_lighting_indexer.py на той же машине со старым бэкендом также проваливается. Различия, согласованные вторым коммитом, включают легализацию bf16 (HIVM отклоняет операнды bf16, вычисления переводятся в fp32 с явным преобразованием) и опускание копирования для областей с динамическим extent.

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

2.2 Основной репозиторий Ascend: идентичность сборки Bisheng в ключе кэша бинарников (#3691)

Объединено 10-28 10:26. Кэш бинарников на стороне Ascend ранее строился по ключу только из ядра и target, поэтому после смены версии инструментальной цепочки Bisheng всё равно попадал в старые бинарники. Этот коммит включает идентичность сборки в ключ кэша, и при изменении инструментальной цепочки кэш инвалидируется и происходит перекомпиляция.

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

2.3 Основной репозиторий Ascend: ручной конвейер сохраняет фазу конфигурационного регистра (#3624)

Объединено 10-28 09:57. Ручной конвейер на этапе перезаписи терял информацию о фазе конфигурационного регистра, из-за чего одна и та же конфигурация не срабатывала на правильном такте. Этот коммит сохраняет фазу при перезаписи и восстанавливает семантику синхронизации, согласованную с автоматическим конвейером.

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

2.4 Репозиторий Ascend на рассмотрении: корневая причина и исправление неверных результатов W4A4 int4×int4 GEMM (#1731)

На рассмотрении, 7 файлов +634/-24, исправляет #438. Корневая причина в том, что весь стек шаблонов C++ выводит фрактальную геометрию из sizeof(AscendC::int4b_t) == 1, тогда как int4b_t — это тип с полубайтовой упаковкой: правильное число элементов в 32-байтовом блоке C0 равно 64, а не 32, поэтому все stride копирования, число повторов и подсчёт байтов имеют двукратное отклонение. Три слоя накладываются: параметры LoadData из L1 в L0A/L0B выводятся из этой константы, данные складываются в неверное место; транспонирование zN в nZ для операнда B молча отключается специализированной инструкцией s4 load_cbuf_to_cb_s4, необходимо перейти на LoadDataWithTranspose к load_cbuf_to_cb_transpose_s4, а семантика его параметров не документирована в заголовках CANN, поэтому автор может реверсить её только на устройстве с помощью пробного ядра (правильные параметры repeatTimes = N/64, srcStride = roundUp16(K)/64, dstGap = 3, причём версия s4 работает блоками по 2048 байт, а не фракталами по 512 байт, как написано в документации); длина в байтах от UB к GM считается как 1 байт на элемент, каждая строка переносится вдвое больше и смешивается с соседними строками. Сама инструкция mad_s4 корректна, симптом — 99,6% несовпадений, вывод -501/627 вместо правильных 641/1785.

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

2.5 Репозиторий Ascend на рассмотрении: миграция планировщика a2a3, 100 файлов, 18453 строки (#1951)

На рассмотрении, 100 файлов +18453/-666, автор erhsh. Миграция планировщика A2/A3 на новую архитектуру — самое крупное изменение на рассмотрении в окне мониторинга. Таблица в описании PR не заполнена, конкретные границы необходимо сверять по списку файлов.

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

2.6 Репозиторий Ascend на рассмотрении: Carver считывает бюджет ресурсов A3 согласно выбранному SDK (#1955)

На рассмотрении, 3 файла +203/-27. На Ascend910_9392 Carver откатывается к Ascend910A и размещает FP32-аккумулятор 256×256 в бюджете L0C размером 256 KiB, тогда как выбранный SDK предоставляет только 128 KiB. Исправление заключается в распознавании двух конкретных SKU — 9392 и 9382 — и чтении INI соответствующей платформы с использованием существующего выбора ASCEND_HOME (дополнительно оставлена опция явного файла для кросс-компиляции), а также в обрезке ёмкости SDK до текущей арены бэкенда с сохранением исходной ёмкости, идентичности конфигурации, поддержки BF16, числа ядер и атрибутов L2. Шесть целевых тестовых случаев охватывают определение устройства, target и явного чипа, выход за пределы и допустимый L0 tile, а также независимую конфигурацию 9382 с отключённым BF16; автор приводит в качестве регрессии: основной набор нативного CI Ascend — 1352 пройдено, 3 пропущено; набор примеров — 153 пройдено.

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

2.7 Репозиторий Ascend на рассмотрении: зависимость AI Core Exception Dump понижена до CANN 9.2.0 (#1956)

На рассмотрении, 5 файлов +49/-32. Ранее эта функция заявляла зависимость от CANN 9.3.0; автор проверил работоспособность на 9.2.0 и исправил два места совместимости: в acldumpTensorInfo добавлены четыре обязательных поля — addrType=ACL_DUMP_ADDR_RAW, placement, shapeNum, originShapeNum (в 9.2.0 acldumpSaveExceptionInfo строго проверяет addrType, а прежняя реализация заполняла его нулём через calloc до ACL_DUMP_ADDR_PTR и получала отказ), а также пропускаются тензоры с адресом NULL или нулевым размером; интерфейс регистрации заменён с устаревшего aclrtSetExceptionInfoCallback на aclrtExceptionInfoCallbackRegister. Сторона разбора поддерживает оба вида артефактов msaicerr.py (9.2.0 выдаёт .npy, старые версии — .bin); требование к версии в документации синхронно изменено на CANN не ниже 9.2.0.

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

2.8 Репозиторий Ascend на рассмотрении: документация переведена на T.Tensor, интерфейс DSL T.Buffer объявлен устаревшим (#1957)

На рассмотрении, один файл +124/-122. Определение параметров kernel в разделе 3.1 руководства по программированию приведено в соответствие с вышестоящим API: текст и примеры единообразно используют T.Tensor, добавлено пояснение, что устаревание T.Buffer(…) и T.Buffer[…] касается только интерфейса TileLang DSL, а тип и понятие Buffer в нижележащем TVM не устарели. Основанием является декоратор @deprecated в tilelang/language/proxy.py ветки ascendc_pto.

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

2.9 Hygon: функциональная ветка LDS swizzle для обычного копирования в HCU, 18 файлов, включая 5 новых тестов (#30)

На рассмотрении, 18 файлов +2000/-84, ветка Feat/hcu plain copy lds swizzle отправлена 10-28 в 13:59. Изменения сосредоточены в новых gemm_lds_strategy_utils (+148/-0 и +29/-0), annotate_mls_gemm_dep.cc (+332/-45) и ds_read_format.cc (+61/-22); в copy и merge_shared_memory_allocations есть сопутствующие небольшие правки; на стороне тестов сразу добавлены test_hcu_plain_copy_lds_swizzle (396 строк), _formula (237), _formula_b8 (220), fp8 (200) и _b8_write_phases (80). Направление — фазы записи LDS и стратегия swizzle на пути обычного копирования, с двумя классами разрядности данных: b8 и fp8.

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

2.10 ROCm: принята поддержка ROCm 10 в pip-версии SDK, MI355X больше не определяется ошибочно как gfx900 (#3549)

Смержено 10-28 07:38, 2 файла +106/-19. Образы rocm/pytorch для ROCm 10.x поставляют ROCm в виде Python wheel-пакетов (rocm-sdk и _rocm_sdk_devel), без дерева каталогов /opt/rocm и без apt-пакетов rocm-. Функция get_rocm_arch() ранее жёстко задавала префикс /opt/rocm и при отсутствии пути откатывалась к gfx900, из-за чего на MI355X (gfx950) статическая проверка rocwmma сразу падала; установка ROCM_PATH также не помогала, поскольку tvm_callback_rocm_get_arch в TVM вызывает эту функцию с префиксом по умолчанию. Исправление — разрешение в порядке: префикс вызывающей стороны, find_rocm_path(), gcnArchName из torch, запасной вариант gfx900; каждая проверка rocminfo получает таймаут 10 секунд; дополнительно добавлен docker/Dockerfile.rocm10 на базе rocm/pytorch:rocm10.1.0_ubuntu24.04_py3.12, убрана apt-установка rocm-, ROCM_PATH указывает на rocm-sdk path –root, сохранена сборка torch из базового образа и добавлена самопроверка import. Автор на MI355X измерил, что определение архитектуры сменилось с gfx900 на gfx950, набор testing/python/amd собирается под правильную архитектуру и проходит.

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

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

3.1 TileOPs: селектор перекалиброван по занятому времени, regret манифеста снижен с 5.79% до 0.48% (#2784)

Смержено 10-28 10:51, 4 файла +259/-203. Оценка плотного селектора изменена на произведение tensor_core_tflops на структурный коэффициент; ранее структурные различия отражались только в одном max(), и shape с ограничением по памяти получали одинаковые оценки для basic, simple и coop2; член launch_us удалён (на занятом времени аппроксимируется нулём). Для NT-полос m≤8 ручные пороги заменены потоковой моделью: скорость на CTA берётся как min(ns·rows·bk·2/latency, cta_gbps), волны считаются по резидентности с ограничением SMEM, добавляется нижняя граница DRAM n·k·2/dram_tbps и член редукции split-K. Путь персистентного ядра для TN/TT задействован (SMEM-тайлы A/B следуют раскладке хранения, transpose_A/B), ширины 176/192, узкие тайлы, ping-pong и split-K по-прежнему ограничены NT. Калибровка выполнена на занятом времени: три сетки из 160/60/264 shape, 7 чередующихся раундов CUPTI с медианой, ожидаемый логарифмический regret softmin минимизируется покоординатным спуском, отжиг tau 0.1→0.05→0.02. Сквозные результаты: отложенный manifest с 5.79% до 0.48%, калибровочная сетка с 9.28% до 2.65%, m 9–32 с 3.13% до 2.09%, m≤8 с 11.20% до 1.17%, отложенный набор m≤8 с 14.18% до 1.79%.

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

3.2 TileOPs: не-NT раскладки переведены на транспонированное персистентное ядро, 2112×7168×4096 TN с 195.3 до 175.3 микросекунд (#2789)

Смержено 10-28 21:06, 4 файла +620/-15, решает #2785. Добавлен coop2_swap: для TN/TT/NN вычисляется Cᵀ = op(B)ᵀ·op(A)ᵀ, тайлами 256×block_n через персистентное ядро, две warpgroup-потребителя держат по 128 строк, хвостовая запись идёт через stmatrix.trans и транспонированную запись TMA. Смешанные ширины тайлов (narrow_n и row_mix) позволяют не тратить M на padding, например 2112 = 13×160 + 32 разбивается на широкие 160 плюс узкие 128, широкие идут первыми, узкие завершают, точка разбиения использует метрику _busiest_cta узких тайлов coop2. Селектор вне NT перечисляет block_n 128 и 160 для coop2_swap, структурный коэффициент 1.2 и перекалиброван по занятому времени. Занятое время на H200: 2112×7168×4096 TN с 195.3 до 175.3 микросекунд (QuACK 187.9), 24576×1536×4096 TN с 421.9 до 416.4 микросекунд; на другом H200 — соответственно с 174.3 до 158.5 (QuACK 167.0) и в диапазоне 384–402 микросекунд. Правило групп строк изменено с фиксированного 4 на 16, на 24 калибровочных shape TN/TT/NN отличается от group 4 в 0.989–1.022 раза; ещё на 9 shape TN смешанная раскладка быстрее равномерного coop2_swap на 3.6%–11%.

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

3.3 TileOPs: cumsum с длинными строками переведён на двухэтапное seed-сканирование, (8,65536) снижен с 12,75 до 5,18 микросекунды (#2790)

Смержен 28.10 в 20:45, 4 файла +255/-246. Для cumsum с небольшим числом длинных строк применено сегментное сканирование с seed, число запусков снижено с трёх до двух, что устраняет регрессию, внесённую #2788, и исправляет сбой раскладки. Правило: если длина строки не менее 64 KiB и строк немного, используется сегментное сканирование, cumprod не изменён; дополнительно исправлен сбой запуска в основной ветке при числе строк более 65535 tile (например, (1, 2**24) float32). H200 @1500 MHz в занятом режиме: (8,65536) float32 с 12,75 до 5,18 микросекунды, (16,65536) с 19,14 до 7,01 микросекунды, (1,65536) bf16 с 9,03 до 3,45 микросекунды, (3,8193) fp16 с 7,52 до 3,29 микросекунды. Среди 384 shape-ов сканирования ни один не стал медленнее, чем до #2788; строки manifest для CumsumFwdOp и CumprodFwdOp быстрее самой быстрой контрольной реализации; добавлена нагрузка few-long-seq (8×65536 float32) 5,20 микросекунды против 6,30 микросекунды у FlagGems.

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

3.4 TileOPs: приведение extrema переведено на PTX max.NaN, распространение NaN не изменилось, регрессия исчезла (#2792)

Смержен 28.10 в 23:31, 4 файла +346/-119. amax/amin переведены на PTX max.NaN и min.NaN для fp32, больше не используется схема с int32-ключами из #2787, семантика распространения NaN не изменилась, а регрессия, возникшая ранее в тот же день, устранена; строки, не являющиеся целыми 16-байтовыми векторами, обрабатываются отдельным строковым ядром extrema, поскольку T.reduce_max на fp32 теряет NaN. H200 @1500 MHz: amax (2048,4096) bf16 до #2787 составлял 7,008 микросекунды, в основной ветке 8,352, в этом коммите 6,976; fp16 с 7,008 до 7,775 до 6,944; amax (64,32768) bf16 с 3,552 до 4,544 до 3,520; amin той же формы с 3,552 до 4,224 до 3,520.

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

3.5 TileOPs: эффективные строковые tile параллельного cumsum и узкие блоки предзаполнения GDN, два коммита (#2788 / #2786)

#2788 смержен 28.10 в 07:13, параллельный cumsum теперь обрабатывает только эффективные строковые tile, это самый ранний коммит в окне мониторинга, позднее #2790 заменил проблемный путь двухэтапным сканированием. #2786 смержен в 07:09, 2 файла +30/-2, на SM89 план предзаполнения GDN с 32 столбцами переведён на существующее сканирование состояния с fused-выходом, целый блок и хвостовой блок обрабатываются раздельно, сохранены три интерфейса: подготовки, вывода и финального состояния; условия применения — плотный вход, batch 1, 16 value head, ширина состояния 128, key major, длина последовательности больше 16384 и не более 65536. При H=4 медианное время выполнения на устройстве снизилось на 13,65%–18,31% относительно базовой линии TileOPs, размах активности устройства снизился на 13,48%–17,07%; в четырёх чередующихся сравнениях с FLA при H=16, длине 65536, BF16 различие размаха составило от -2,57% до +1,77%, автор отмечает, что эта строка находится в пределах колебаний измерений.

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

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

3.6 Новый issue TileOPs: при последовательных вызовах строка большого GEMM отстаёт от самой быстрой базовой линии (#2791)

Открыт заново, продолжает #2133 и #2789. bench_gemm считывает один вызов за раз, между вызовами выполняется сброс L2 и синхронизация, вызов стартует ниже предела мощности, и затраты энергии маскируются; после перехода на последовательные непрерывные вызовы, как при обучении, шесть из десяти крупных строк manifest отстают от самой быстрой базовой линии. Типичная строка: prefill-o-proj 4096×7168×16384 NT данная реализация 1412,7 микросекунды, QuACK 1371,1, cuBLAS 1360,1, DeepGEMM 1359,2, отношение 0,962. Методика: одна карта H200, разблокированные частоты, прогрев 1 секунда, затем замер 1,5 секунды, четыре чередующихся раунда с взятием медианы.

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

3.7 TileFoundry: привязка этапов ring и унификация процесса компиляции, 70 файлов (#277)

Объединено 10-28 16:55, 70 файлов +1567/-1298, заменяет #276. Мотивация двойная: несвязанные представления ring и выбор стадии протаскивают представления без SSA-имён в операнды CUDA; три точки входа компиляции собирают каждый свой список проходов. Решение — сохранять стадию ring как набор связанных TensorView, связывать один раз на каждую область потребления, сделать так, чтобы default_pipeline использовался совместно в finalize, compile и scaffold, проход ядра хостовой точки входа разрешается через Target.resolve и реализацию @bind предоставляет CUDA; Module.entry становится хостовой точкой старта, пиковое значение разделяемой памяти из анализа передаётся в запуск хоста; все CPU-хосты в цепочке вызовов участвуют в линковке, корневой хост идёт первым. Изменения контракта: link_program лишается entries, загрузчик связывает первый слинкованный хост, build требует наличия корневого CPU-хоста. Остаточные риски перечислены в PR: #278 отслеживает нормализацию раскладки Broadcast, #255 отслеживает SymbolRef и ссылки на функции.

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

3.8 TileOPs на рассмотрении: релизный образ объединяет TileOPs 0.1.0 и TileFoundry 0.0.3 (#2793)

На рассмотрении, 9 файлов +120/-44. В .github/runner/Dockerfile добавлена цель release, поверх dev-образа устанавливаются PyPI-релизы tileops и tilefoundry, имя образа формируется из TILEOPS_VERSION и TILEFOUNDRY_VERSION, публикуется как ghcr.io/tile-ai/tileops-release:tileops0.1.0-tilefoundry0.0.3-cuda13.2-torch2.13-tilelang0.1.12; сопутствующий scripts/verify_release_image.sh после вытягивания образа запускает на GPU проверки стека и базовой линии, упаковывает тесты и указанные файлы –bench, сценарии берутся по git tag, соответствующему версии tileops внутри образа (в wheel не входят тесты и бенчмарки). Одновременно удалена базовая линия sgl-kernel: ни один образ больше его не устанавливает, bench_moe_permute_align всегда его пропускает.

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

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

4.1 Два автоматических обновления документационного сайта (10-28)

tilelang.github.io 10-28 в 08:20 и 13:31 получил по одному коммиту Update docs от github-actions, содержимое — синхронизация артефактов сборки документации с основной веткой, самостоятельных изменений содержимого нет.

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

4.2 Пакетная проверка фактов документации и исправления установки в репозитории Ascend на рассмотрении, семь пунктов продвигаются одной партией (10-28)

Конвейер проверки фактов документации tilelang-ascend обновил в окне мониторинга семь пунктов: проверка пяти API T.tile.bitwise_lshift, bitwise_or, bitwise_and, fill, clear, clamp работает на PTO и сверен с документацией (#1620), а также усиление docstring и тестов axpy (#1600). Кроме того, #1943 меняет установку SHMEM так, чтобы останавливаться при ошибке сборки, и вместе с проверкой анализатора снимков #1942 относится к коммитам робота.

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

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

4.3 Новые примеры операторов в репозитории Ascend на рассмотрении: два оператора GroupNorm и fla_chunk (#1607 / #1953 / #1952)

Три PR категории примеров обновлены в окне мониторинга: #1607 добавляет 20 сценариев cann-bench для GroupNorm, #1953 добавляет оператор fla_chunk_o и тесты, #1952 добавляет оператор fla_chunk_intra_token_parallel и тесты. Это также исходный сценарий дефекта краха скалярного чтения из 1.1 (GroupNorm читает первый элемент каждой строки).

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

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

4.4 Сторонние: агент автотюнинга сообщества обновил китайское описание (#74), в основной ветке репозитория бэкенда Hexagon в окне мониторинга новых коммитов нет (10-28)

Репозиторий сообщества cc-c122/TileLang-Remote-Autotuning-Agent 10-28 в 09:21 принял #74, обновив китайский README согласно текущему рабочему процессу V2; ранее, 10-24, в репозитории также произошло слияние, касающееся отмены базового состояния и атрибуции mcprofiler. Другой сторонний бэкенд-репозиторий xwhzz/tilelang-hexagon в организационном представлении показывает отметку времени пуша 10-29 06:56, но последний коммит основной ветки остановился на 08-24 (обновление зависимостей CI); в пределах окна в основной ветке не было коммитов кода, вчерашние 15 относятся к активности в ветке функций, и этот пункт не учитывается как новое содержимое.

Источник: https://github.com/cc-c122/TileLang-Remote-Autotuning-Agent/pull/74

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

4.5 Новостная сторона: шесть групп запросов Google News на китайском и английском, Hacker News и arXiv — существенных попаданий в пределах окна нет (10-28)

Шесть групп запросов включают точное слово TileLang плюс два дня, китайские слова об операторах плюс семь дней, английский GPU kernel DSL плюс три дня, название организации плюс семь дней, китайские слова производителей (Ascend / MetaX / Hygon) плюс семь дней, а также слова о библиотеках операторов плюс два дня. Среди возвращённых результатов связанная с TileLang и попадающая в окно — только одна перепечатка (сравнение матричной производительности Ascend 950PR, комментарий вторичного рынка без исходных данных, не включается), остальное — общие материалы о финансировании и технологиях DeepSeek, семантика заголовков которых не содержит связанного с TileLang содержимого. Английские запросы все вернули пусто. Запись Hacker News за день — выпуск нового языка под названием Vu, не связанного с TileLang. В arXiv нет новых в пределах окна статей о tilelang или tiled DSL. Разрыв в силе сигналов между новостной стороной и стороной GitHub в этот день продолжает увеличиваться, основной корпус ежедневного отчёта по-прежнему опирается на репозитории.

Источник: https://news.google.com/

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

Способ продвижения на стороне Ascend за эти два дня стал ясен: основной репозиторий ежечасно принимает мелкие исправления (ключи кеша, фазы регистров), специализированный репозиторий mlir-ascend одним слиянием на двадцать тысяч строк переносит путь NPUIR для 910 из ветки функций в основную, а tilelang-ascend в установленном ритме двух бэкендов PTO и AscendC занимается дополнением документации и операторов. Три места имеют разный ритм, но указывают на одно: путь компилятора для 910 переходит от «работает» к «приемлемо для приёмки», и 504 самотеста и 71/72 примера — это критерии приёмки, данные на этот раз. Примечательно, что в очереди на рассмотрение миграция планировщика a2a3 занимает 100 файлов, что говорит о том, что объём изменений следующего этапа находится в слое планирования, а не в слое операторов.

Интерес на стороне производительности заключается в самом критерии измерения. Эти два изменения TileOPs переносят выводы селектора и GEMM на время занятости (device_busy_ms) и в том же PR приводят разницу между одиночным вызовом и вызовами подряд: при одиночном чтении данная реализация опережает QuACK, а в сценариях обучения с вызовами подряд и ограничением мощности из десяти крупных строк шесть отстают. Это более важный сигнал, чем «на сколько микросекунд быстрее та или иная строка», потому что обучение на практике и есть вызовы подряд. #2791 выделяет это в отдельный проект, что показывает: команда уже рассматривает вопрос критерия как продуктовый вопрос, а не замазывает его как отклонение базового состояния.

В инженерной гигиене сегодня стоит отметить три пункта: во-первых, самоповреждение и самопочинка в одной партии (#2787/#2788 приняты утром, #2790/#2792 отозваны днём), что говорит о том, что ночная страховка базового состояния работает, но также и о том, что семантические границы этого кода редукции остаются хрупкими; во-вторых, семантическая дыра without_kernel_cache() (#3699) была выявлена в процессе миграции — два уровня кеша внутри процесса не учитывают CacheState, и тест полагает, что выполняет холодную компиляцию, а на самом деле получает старый kernel; такие дефекты лишают смысла утверждение «все наши тесты проходят»; в-третьих, отказ 13 примеров на SM120 (#3701) выявил проблему жёстко закодированных размеров Hopper в примерах; по мере роста установленной базы Blackwell такой долг «пример — не продукт» будет учитываться снова и снова.

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

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

Источники Результаты проверки
По репозиториям организации GitHub tile-ai (tilelang / TileOPs / TileFoundry / tilelang-ascend / tilelang-mlir-ascend / tilelang-hygon / tilelang-metax / tilelang-musa / TileRT / tilescale / DeepStack / сайт документации) В окне мониторинга есть пуши: tilelang, TileOPs, TileFoundry, tilelang-mlir-ascend, tilelang-hygon, tilelang.github.io; в tilelang-metax, tilelang-musa, TileRT, tilescale, DeepStack, tilelang-puzzles нет коммитов в основную ветку в окне мониторинга
Проверка фактов слияния (is:pr is:merged) 24: tilelang 15, TileOPs 7, TileFoundry 1, tilelang-mlir-ascend 1
Проверка фактов на рассмотрении (is:pr is:open updated) 48, включая пакет проверки документации репозитория Ascend и миграцию a2a3
Проверка новых issue 6, из них 1 — ежедневный тестовый бот репозитория Ascend
Проверка на уровне коммитов в основную ветку tilelang 15, TileOPs 7, TileFoundry 1, сайт документации 2; число слияний совпадает с числом коммитов в основную ветку
Страница релизов и tag Последняя версия tilelang по-прежнему v0.1.15 (релиз 09-30), новых версий в окне мониторинга нет; последняя версия TileOPs — v0.1.0
Репозитории сторонних бэкендов Последний коммит в основную ветку репозитория бэкенда Hexagon датируется 08-24, новых коммитов в окне мониторинга нет; слияние агента автотюнинга обновило китайское описание; в шести ветках репозитория MetaX нет коммитов в окне мониторинга
Google News (шесть групп на китайском и английском) Упоминается только одна вторичная перепечатка, оригинальных данных нет, не включается
Hacker News В этот день нет материалов, связанных с TileLang
arXiv Нет новых статей о tilelang или tiled DSL в окне мониторинга

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