Ежедневный отчёт TileLang (2026-10-20)
Окно мониторинга: 2026-10-19 07:00 ~ 2026-10-20 07:00 (пекинское время, 24 часа).
Источники: GitHub (полная проверка пушей в 29 репозиториях организации tile-ai; в окне пуши были в 9 репозиториях: TileOPs, TileOPs-nightly, TileOPs.github.io, tilelang, tilelang.github.io, TileFoundry, tilelang-ascend, tilelang-hygon, tilelang-mlir-ascend; в организации за окно смержено 36 PR 〔TileOPs 19, tilelang 11, TileFoundry 3, TileOPs.github.io 1, tilelang-ascend 1, tilelang-mlir-ascend 1〕, на рассмотрении 52 PR 〔tilelang 24, tilelang-ascend 15, TileOPs 11, tilelang-hygon 2〕; проверка неосновных веток поштучно 〔ascendc_pto, snapshots, stats〕) + проверка страницы релизов и файлов версий + проверка ветки статистики спецификаций + проверка репозиториев потребителей (TileKernels / FlashQLA) + проверка репозиториев сообщества + проверка ночных пайплайнов + множественные запросы Google News на китайском и английском (через прокси) + Hacker News + arXiv.
Индекс выпуска
- Ключевое за сегодня: SM89 за один день два GEMM переведены в рабочий статус, TileOPs отчитался о покрытии спецификаций и бенчмарков 205/205, основной репозиторий включил Windows ARM64 в CI (10-19)
-
- Ход работ по ключевым проектам
- 1.1 Основной репозиторий: Windows ARM64 включён в CI, установка пакетов CUDA 13.4 arm64 и слой совместимости с clang (#2962) (10-19)
- 1.2 Основной репозиторий: T.cumsum на разделяемом буфере дополнен смещением потоков при warp-специализации (#3536) (10-19)
- 1.3 Основной репозиторий: «голые» переменные data в call_extern консервативно помечаются как записываемые (#2618) (10-19)
- 1.4 Основной репозиторий: lower trace разделяет исходный код хоста и устройства, базовая линия переведена на редактируемую (#3496) (10-19)
- 1.5 Основной репозиторий: Autotuner пробрасывает параметры profiler, печать констант ROCm FP4 и два обслуживания CI (#3337 / #3548 / #3565 / #3567) (10-19)
- Ход работ по ключевым проектам
-
- Адаптация под множество бэкендов (Ascend / MetaX / Hygon / Moore Threads)
- 2.1 Ascend: установлена связь между операторами хоста и ядрами устройства под msprof (#3321) (10-19)
- 2.2 Ascend: скалярный sqrt вне домена VF идёт через стандартную математическую библиотеку, PTO Ramp генерируется по типизированным векторам (#3466 / #3535) (10-19)
- 2.3 Ascend: для четырёх поэлементных операторов дополнены китайскоязычные API-документы и примеры использования (#1657) (10-19)
- 2.4 Очередь на рассмотрении Ascend — 15 коммитов: операторы сжатия GLM-5.3 и три обратных оператора (10-19)
- 2.5 Hygon: два коммита на рассмотрении дополняют константу K окна MLS и конвейер регистров (#25 / #18) (10-19)
- 2.6 MetaX / Moore Threads / MLIR-Ascend: с обеих сторон новых пушей нет, линия MLIR исправляет изоляцию TVM (#204) (10-19)
- Адаптация под множество бэкендов (Ascend / MetaX / Hygon / Moore Threads)
-
- Экосистема и стороны-пользователи (TileOPs / TileFoundry / библиотека операторов)
- 3.1 TileOPs: блочно-масштабируемый FP8 групповой GEMM для SM89 и INT8 W8A8 влиты в один день (#2659 / #2655) (10-19)
- 3.2 TileOPs: пять foundry-операторов дополняют пагинацию MLA, запись KV и fused-нормализацию (#2642 / #2635 / #2639 / #2634 / #2651) (10-19)
- 3.3 TileOPs: постраничный FP8 forward для GQA переведён в implemented, 36 примеров сверены с FA3 (#2648) (10-19)
- 3.4 TileOPs: серия точечной оптимизации производительности линейного внимания и GEMM (#2650 / #2645 / #2637 / #2636 / #2641 / #2640 / #2646 / #2652) (10-19)
- 3.5 TileOPs: реализовано выборочное сканирование Mamba-1, в перечисление маршрутизации MoE добавлен предел сходимости (#2638 / #2653) (10-19)
- 3.6 TileOPs: покрытие спецификаций и бенчмарков отчитано как 205/205, в ночной отчёт добавлена колонка операторов (#2656) (10-19)
- 3.7 TileFoundry: поддержка Apple AMX и пример prefill для Qwen запущены на M2 Pro (#244 / #240 / #245) (10-19)
- 3.8 Стороны-пользователи и ночной конвейер: TileKernels / FlashQLA без пушей, снимок продвинут до 10-20 (10-20)
- Экосистема и стороны-пользователи (TileOPs / TileFoundry / библиотека операторов)
-
- Сообщество, руководства и мероприятия
- 4.1 Документация: шесть новых операторов внесены на страницу API, сайт документации основного репозитория обновляется автоматически (#77) (10-19)
- 4.2 Репозиторий сообщества: практический репозиторий W4A16 kernel и репозиторий навыков отладки CUDA попали в поле зрения (10-19)
- 4.3 Линия версий: основной репозиторий по-прежнему на 0.1.15, TileOPs продолжает фиксировать 0.1.12 (10-19)
- 4.4 Медиа и академическая сфера: три попадания новостного поиска не относятся напрямую, на arXiv новых статей нет (10-19)
- Сообщество, руководства и мероприятия
-
- Наблюдения за трендами
- 5.1 SM89 становится точкой по умолчанию для новых операторов
- 5.2 Линейное внимание закрепляется как «одно матричное умножение на chunk + выбор стратегии по волнам SM»
- 5.3 TileFoundry выходит за пределы CUDA: Apple AMX и prefill для Qwen
- 5.4 Наблюдаемость становится основной линией ужесточения компилятора
- 5.5 Библиотека операторов Ascend переходит к операторам, специализированным для диалоговых моделей
- Наблюдения за трендами
- Приложение: материалы и пояснения по проверке
- Полный список источников
Ключевое за сегодня: SM89 за один день два GEMM переведены в рабочий статус, TileOPs отчитался о покрытии спецификаций и бенчмарков 205/205, основной репозиторий включил Windows ARM64 в CI (10-19)
В данном окне активности хостов по-прежнему соответствует ночному времени по пекинскому времени: в организации слито 36 коммитов, из них 19 приходятся исключительно на TileOPs, причём блочно-масштабируемый FP8 сгруппированный GEMM для SM89 (уровня Ada) и INT8 W8A8 GEMM были реализованы в один день; вместе с пятью новыми операторами они за один раз закрыли постраничное внимание MLA, запись в KV-кэш, фьюжн QK-Norm-RoPE и покомпонентно-токенную квантизацию RMSNorm. Ветка статистики спецификаций организации сразу обновилась до «реализовано 205/205, отбенчмаркано 205/205», то есть заявленное самой библиотекой операторов покрытие спецификаций и покрытие бенчмарками одновременно достигли максимума. В основной репозиторий tilelang слито 11 коммитов, наиболее весомый — включение Windows ARM64 в релизный конвейер (12 файлов, скрипт установки CUDA 13.4 arm64 и заголовок совместимости с clang), кроме того два исправления корректности (T.cumsum на разделяемом буфере, аннотация записи call_extern) и одна переработка отладочной инфраструктуры (разделение lower trace на исходники хоста и устройства). На стороне бэкендов Ascend сопоставил операторы хоста и ядра устройства под msprof, два коммита Hygon на рассмотрении дополняют конвейер регистров и адресацию константы K, а TileFoundry впервые запустил пример префилл-фазы Qwen на блоках AMX Apple M2 Pro, обработав 128 токенов.
1. Прогресс ключевых проектов
1.1 Основной репозиторий: Windows ARM64 в CI, установка пакета CUDA 13.4 arm64 и слой совместимости с clang (#2962)
Этот коммит подключает Windows on ARM как официальную цель сборки в релизный конвейер: добавлен скрипт установки PowerShell, который из предрелизного канала NVIDIA загружает локальный установочный пакет CUDA 13.4 для windows_arm64 и размещает его по фиксированному пути; добавлены модуль CMake FindPipCUDAToolkit и заголовок совместимости windows_arm64_clang_compat.h, устраняющие различия заголовочных файлов и встроенных средств инструментария MSVC на ARM64; синхронно скорректированы pyproject и два файла requirements, а также изменены модуль раскладки и вспомогательный модуль MSVC для адаптации к новой платформе. Всего 12 файлов, добавлено 356 строк, удалено 45 строк. Ранее на стороне организации последовательно были добавлены ROCm и Ascend, а этот коммит делает Windows ARM64 третьей поверхностью поставки вне x86, что является прямым улучшением удобства для пользователей, разрабатывающих локальные ядра на ноутбуках с процессорами Snapdragon X.
Источник: tile-ai/tilelang#2962 (19.10)
1.2 Основной репозиторий: T.cumsum на разделяемом буфере дополняется смещением потока при warp-специализации (#3536)
После того как warp-специализация распределяет производителей и потребителей по разным warp, вычисляемый сканирующими примитивами на разделяемом буфере номер потока по-прежнему считается по всему CUDA-блоку, что приводит к смещению принадлежности строк в cumsum. Изменение вводит параметр thread_offset в оба набора шаблонов scan для CUDA и HIP, который передаётся из нижней границы потоков в параметрах нижнего уровня (минимум thread_bounds); шаблон на его основе заново выравнивает lane и индекс строки; фоновая общая реализация и Python-тесты обновлены синхронно, всего 4 файла, добавлена 101 строка, удалено 38 строк. Подобные проблемы проявляются только при фактически включённой warp-специализации и размещении буфера в разделяемой памяти и относятся к граничным условиям, выявляемым после масштабного использования.
Источник: tile-ai/tilelang#3536 (19.10)
1.3 Основной репозиторий: обнажённые data-переменные call_extern консервативно помечаются как записываемые (#2618)
Аннотация параметров только для чтения ранее опиралась на анализ псевдонимов для определения, записывается ли буфер, но при записи через внешний вызов (call_extern) по сырому указателю анализ не видит операции записи, и буфер ошибочно помечается как доступный только для чтения, после чего сворачивается как константа или переиспользуется досрочно. Изменение в этом проходе специально распознаёт call_extern и все его обнажённые data-переменные консервативно помечает как записанные, а также добавляет 69 строк тестов, покрывающих вариант «call_extern пишет в выходной буфер». По направлению это соответствует принципу «лучше меньше оптимизаций, чем ошибочные оптимизации» и согласуется с недавней серией исправлений корректности в основном репозитории.
Источник: tile-ai/tilelang#2618 (19.10)
1.4 Основной репозиторий: lower trace разделяет исходники хоста и устройства, базовая версия становится редактируемой (#3496)
Инструмент отладки lower trace ранее сохранял только один сгенерированный исходник, в котором обёртка хоста и ядро устройства были смешаны, и изменить одно место без возможности отката было сложно. Изменение разделяет обе стороны на независимые артефакты: сторона устройства выводит codegen_device.cpp с его базовой версией .original и последней версией .latest, сторона хоста при выполнении генерации кода хоста дополнительно выводит codegen_host.cpp и файлы базовой версии той же структуры, а report.html через символическую ссылку указывает на последний запуск; документация синхронно переписана под процесс редактирования и перекомпиляции, всего 6 файлов, добавлено 121 строка, удалено 63 строки.
Источник: tile-ai/tilelang#3496 (19.10)
1.5 Основной репозиторий: Autotuner передаёт параметры profiler, печать констант ROCm FP4 и две правки CI (#3337 / #3548 / #3565 / #3567)
#3337 — всего три строки, но исправляет метрику автотюнинга: warmup, rep и timeout в декораторе ранее не передавались тюнеру, поэтому изменение этих трёх параметров в декораторе фактически не действовало; теперь они передаются. #3548 добавляет печать констант с плавающей точкой FP4 (e2m1) на стороне ROCm: поскольку этот тип можно сконструировать только из сырых байтов, константы теперь генерируются на устройстве через вспомогательную функцию преобразованием из float32, что избегает ошибок генерации кода HIP. Ещё две правки — поддерживающие: версия cibuildwheel поднята с 4.2 до 4.3, а для линии ROCm добавлен ещё один ревьюер CodeRabbit.
Источники: tile-ai/tilelang#3337, #3548, #3565, #3567 (19.10)
2. Адаптация под несколько бэкендов (Ascend / MetaX / Hygon / Moore Threads)
2.1 Ascend: установление связи между хост-оператором и устройственным ядром под msprof (#3321)
Профайлер на стороне Ascend ранее сообщал только время выполнения устройственных ядер, и соответствие между хост-операторами и ядрами отсутствовало, поэтому выборки автотюнинга нельзя было отнести к конкретному вызову. Этот коммит добавляет нативный канал профилирования: при сборке проверяется наличие заголовков профилирования CANN (поддерживаются оба варианта установки — msprof и profiling), при наличии они включаются, при отсутствии пропускаются, без введения пользовательского флага компиляции; на стороне выполнения добавлены ascend_profiling.h и хост-модуль profile, которые сопоставляют запуск хост-оператора и выполнение устройственного ядра по порядку вызовов, а путь групповой компиляции использует ту же схему отчётности. Всего 13 файлов, добавлено 676 строк, удалено 13 строк — это структурное дополнение возможностей измерения линии Ascend.
Источники: tile-ai/tilelang#3321 (19.10)
2.2 Ascend: скалярный sqrt вне области VF через стандартную математическую библиотеку, PTO Ramp генерируется как типизированный вектор (#3466 / #3535)
#3466 обрабатывает случай, когда скалярный sqrt оказывается вне области векторизации (VF): нижележащий слой ранее единообразно понижал его до sqrtf, но на скалярном ядре эта реализация недоступна; изменение переключается на AscendC::Std::sqrt, если значение вне области и имеет тип float32, и при необходимости подключает стандартные математические заголовки, добавлено 126 строк вместе с примерами. #3535 исправляет способ генерации выражения Ramp в кодогенерации PTO — теперь типизированный вектор собирается из выражений отдельных lane; примеры охватывают генерацию и выполнение векторизованного сохранения для Ramp с параллельной индексацией и с явным шагом.
Источники: tile-ai/tilelang#3466, #3535 (19.10)
2.3 Ascend: дополнены китайскоязычная документация API и примеры для четырёх поэлементных операторов (#1657)
Репозиторий адаптации Ascend дополнил китайскоязычную документацию API для четырёх поэлементных операторов под T.tile — relu, leaky_relu, sigmoid, silu, — указав назначение, параметры, возвращаемые значения и типичное использование (например, семантику отрицательного наклона в leaky_relu), а также синхронно добавил тестовые примеры и точку входа в документацию: всего 6 файлов, добавлено 451 строка, удалено 19 строк. Это редкое для линии Ascend слияние с документацией в качестве основного содержания, что говорит о том, что набор операторов этого репозитория вступил в стадию, когда требуется публично гарантировать семантику.
Источники: tile-ai/tilelang-ascend#1657 (19.10)
2.4 В очереди на рассмотрении Ascend 15 коммитов: компрессионный оператор GLM-5.3 и три обратных оператора (19.10)
Репозиторий Ascend — место наибольшей концентрации рассматриваемых коммитов в этом окне (15). На стороне операторов заметна явная ориентация на диалоговые модели: компрессионный оператор kpool для GLM-5.3 (#1910), chunk-обратный проход KDA (#1895), обратный проход indexer (#1887), situ_glu и его градиент (#1891 / #1877). На стороне примеров добавлены две позиции cann-bench — lstm и gru (#1892 / #1902). На стороне компилятора две правки относятся к стабильности: #1907 сохраняет аннотации блоков и ND-смещения при промежуточном хранении из UB в L1, #1903 исправляет публикацию кэша параллельных ядер и восстановление после удаления кэша. Ещё пять правок касаются сверки фактов об API побитовых операторов (#1651 / #1655 / #1656 / #1659 / #1893), причём #1893 — повторная отправка после отката на прошлой неделе, остальные четыре — новые сверки.
Источники: tile-ai/tilelang-ascend#1910, #1895, #1887, #1891, #1907, #1903 (19.10)
2.5 Hygon: два PR на рассмотрении дополняют константу окна K для MLS и регистровый конвейер (#25 / #18)
В окне мониторинга ветки по умолчанию репозитория Hygon слияний не было, но оба PR на рассмотрении имеют содержательный прогресс. #25 исправляет проблему потери смещения константного окна при развёртывании цикла в продвижении ресурсов MLS: сначала проверяется непринадлежность оси обновлению, затем принимается решение о применении одноосного обновления или адресации с прямой разностью, при необходимости окно консервативно перестраивается, а в существующие тесты MLS добавляется регрессия. #18 добавляет регистровый конвейер для бэкенда HCU, который может накладываться на существующий конвейер разделяемой памяти, вставляя соответствующие инструкции фиксации и барьеров по ступеням конвейера, а также включает warp_diverge при числе warp больше 4; в описании PR указано, что производительность example_gemm уже близка к rocblas.
Источники: tile-ai/tilelang-hygon#25, #18 (19.10)
2.6 MetaX / Moore Threads / MLIR-Ascend: с обеих сторон новых пушей нет, линия MLIR исправляет изоляцию TVM (#204)
Последний пуш tilelang-metax был 16.10, tilelang-musa — 30.09, в окне мониторинга в обоих коммитов нет, что является нормальным затишьем. В tilelang-mlir-ascend слит один коммит с исправлением изоляции TVM (#204), устраняющий взаимное загрязнение между этим репозиторием и пространством имён и видимостью символов TVM; объём небольшой, но он напрямую влияет на стабильность линковки при сосуществовании нескольких бэкендов.
Источники: tile-ai/tilelang-mlir-ascend#204 (19.10)
3. Экосистема и потребители (TileOPs / TileFoundry / библиотеки операторов)
3.1 TileOPs: блочно-масштабируемый FP8 групповой GEMM для SM89 и INT8 W8A8 слиты в один день (#2659 / #2655)
#2659 расширяет блочно-масштабируемый FP8 групповой GEMM (сценарий MoE) на SM89: четыре явные конфигурации, шесть эталонных случаев с копированием, тесты покрывают неизвестные ключи, недопустимые размеры tile, подсчёт масок и хвостовые блоки slab, тюнинг, псевдонимы вывода, невыровненные входы, сохранение простаивающих экспертов и неизменность входов; оба набора базовых линий TileLang (0.1.12 и разрабатываемая 0.1.15) проходят по 35 пунктов; число случаев выросло с прежних 42 до 77, добавлено 12 новых случаев SM89. #2655 на той же архитектуре включает INT8 W8A8 GEMM, 13 интерфейсных случаев SM89 и прежние 29 эталонных нагрузок INT8 проходят проверку корректности против vLLM, при этом заявлено, что результаты на SM90 не утверждаются. Оба коммита вместе означают, что низкоточные GEMM в TileOPs больше не нацелены исключительно на Hopper.
Источники: tile-ai/TileOPs#2659, #2655 (19.10)
3.2 TileOPs: пять операторов foundry дополняют постраничное внимание MLA, запись KV и слитую нормализацию (#2642 / #2635 / #2639 / #2634 / #2651)
Все эти изменения произведены конвейером спецификаций TileFoundry: для каждого оператора сначала создаются спецификация HIR и эталонная реализация, затем — раскладка, планирование и исполняемый двойник.#2642 добавляет постраничный прямой проход MLA: обход таблицы страниц по запросам, оценка с накоплением в fp32, онлайн-softmax, возврат выхода и lse; шесть строк тестовых примеров совпадают с эталоном в пределах абсолютной погрешности 4e-3; в строке два умножения-накопления по 576 и 512 столбцам определены как ограниченные вычислительной мощностью.#2635 добавляет запись KV-кэша MLA: для каждого токена конкатенация kv_c и k_pe, поворот в fp32 с последующей записью, масштабирование по scale и приведение с ограничением к E4M3; шесть строк тестовых примеров проходят против одноимённого интерфейса vLLM.#2639 объединяет QK-Norm-RoPE и охватывает семь комбинаций строк/точности, поддерживая две раскладки поворота — NeoX и соседние пары.#2634 объединяет поточную квантизацию RMSNorm со сложением остатка; 12 рабочих нагрузок согласованы с одноимённым интерфейсом vLLM; реализация — одна строка на CTA, 128 потоков при 4096 столбцах, 512 потоков для более длинных.#2651 добавляет обёртку оператора INT8 W8A8.
Источники: tile-ai/TileOPs#2642, #2635, #2639, #2634, #2651 (19.10)
3.3 TileOPs: постраничный FP8 прямой проход GQA переведён в implemented, 36 примеров сопоставлены с FA3 (#2648)
Этот коммит позволяет постраничному прямому проходу GQA напрямую принимать Q и KV в FP8 и переводит статус оператора из спецификации в реализованный. Эталонная реализация записывает HIR по измерениям запроса и KV-головы, головы запросов группируются в строки оценки, каждая видимая страница образует ключевой tile, операнды FP8 деквантуются по своим scale и перемножаются; все 36 примеров проходят, включая проверки на NaN и Inf. В качестве объекта сравнения выбрана аналогичная постраничная реализация FA3, развёрнутая в соответствующую форму через descale «на запрос, на KV-голову». Сопутствующие изменения затрагивают 17 файлов, добавлено 2601 строка, включая несколько путей ядер: постраничное декодирование, декодирование с batch 1, декодирование FP8, специализацию warp и предобработку RoPE.
Источники: tile-ai/TileOPs#2648 (19.10)
3.4 TileOPs: серия точечной оптимизации производительности линейного внимания и GEMM (#2650 / #2645 / #2637 / #2636 / #2641 / #2640 / #2646 / #2652)
Восемь коммитов [Perf] продолжают ритм оптимизации прошлой недели, схема весьма единообразна: свести матричное умножение каждого chunk к одному, сдвинуть или объединить точки синхронизации, выбирать стратегию по аппаратным волнам.#2650 заменяет обращение chunk в prefill GDN на решение по группам warp с границей «примерно семь групп warp-решений на SM»; ниже одной волны по-прежнему выгоднее один warp;#2645 заставляет прямой проход состояния DeltaNet выполнять матричное умножение только один раз на chunk;#2637 убирает вставку матричного умножения в цепочку градиента обратного состояния GLA;#2636 переносит сохранение каждого tile пакетного GEMM под первую WGMMA следующего tile для перекрытия записи;#2641 переходит на свёртку из 512 потоков для редукции длинных строк;#2640 обеспечивает чтение строк ядра двумерных индексов max pooling в L2 с evict-last;#2646 обходит полнодиапазонное масштабирование делением bf16, когда коэффициент масштабирования ровно равен 1;#2652 реализует остаток от деления и целочисленное деление с округлением вниз для float32 через приближённую обратную величину.
Источники: tile-ai/TileOPs#2650, #2645, #2637, #2636, #2641, #2640, #2646, #2652 (19.10)
3.5 TileOPs: реализовано выборочное сканирование Mamba-1, перечисление маршрутизации MoE с пределом сходимости (#2638 / #2653)
#2638 Добавлена прямая реализация выборочного сканирования Mamba-1: рабочая нагрузка из девяти строк охватывает префилл 4096 токенов и декодирование одного токена для Mamba-2.8B, декодирование батча 64 для 2.8B, префилл 1.4B, 130M без гейтинга, упакованный префилл с последовательным индексом, групповое состояние и «голое» сканирование; эталонная реализация вместе с раскладкой и генератором расписания также помещены в репозиторий, что означает официальное включение моделей пространства состояний в список операторов. #2653 исправляет перечисление маршрутизации в групповом GEMM для MoE: ранее маршрутизация групп перечислялась полностью, и при превышении лимита комбинаций время тратилось впустую; теперь перечисление прекращается, как только объединение превышает верхний предел.
Источники: tile-ai/TileOPs#2638, #2653 (10-19)
3.6 TileOPs: спецификации и покрытие бенчмарками оба отчитываются 205/205, в ночной отчёт добавлена колонка операторов (#2656)
Ветка статистики спецификаций организации обновилась сразу после последнего слияния, данные бейджа — «реализовано 205/205, покрыто бенчмарками 205/205», оба показателя максимальные, что согласуется с непрерывно добавляемыми в этом выпуске операторами. #2656 добавляет в ночной отчёт две вещи: операторы, не измеренные в текущем периоде, выделяются в отдельную колонку «новые операторы», а чтение исторических данных переведено со старого строкового представления на разбор JSON, причём всё, что не является отображением, отбрасывается; число вариантов использования выросло с 43 до 47. Один фрагмент исторических данных из описания #2656 также стоит зафиксировать: ночной отчёт за 10-18 (окно статистики с 10-05 по 10-17) содержит 109 улучшений, 6 новых операторов, всего 121 строку, из которых W4A8 GEMM занимает 98 строк.
Источники: tile-ai/TileOPs#2656, ветка статистики спецификаций TileOPs (10-19)
3.7 TileFoundry: поддержка Apple AMX и пример префилла Qwen успешно запущены на M2 Pro (#244 / #240 / #245)
#244 включает в TileFoundry четыре операции AMX (загрузка X/Y, внешнее произведение FP16 с накоплением FP32, сохранение Z) и позволяет Copy и Fill размещаться в областях действия ядра Apple и исполнителя AMX; в каталоге примеров представлены HIR всей модели, HIR после планирования, близнец среды выполнения на Python и проекционное ядро на C++, при первом использовании локальное ядро собирается и кэшируется, а веса переиспользуются из заранее подготовленного GGUF Q8_0. В реальных замерах на M2 Pro Qwen3-1.7B и Qwen3.5-4B завершили префилл 128 токенов, выходные logits конечны, номера следующего токена — 1813 и 306 соответственно. Это первый случай в экосистеме TileLang, когда пример вывода целой модели запущен на не-CUDA ускорителе. #240 устраняет четыре пробела в понижении на стороне CUDA (Cast в форме эффекта без порождающего обработчика, запись Fill в голый указатель параметра ядра, координаты среза через арифметику Dim, которые не разрешаются), изменения только в каталоге генерации кода, 45 добавленных строк и 10 удалённых. #245 унифицирует контекст оценки и снимки областей действия на уровне анализа.
Источники: tile-ai/TileFoundry#244, #240, #245 (10-19)
3.8 Принимающие стороны и ночной конвейер: TileKernels / FlashQLA без пушей, снимок продвинулся до 10-20 (10-20)
deepseek-ai/TileKernels по-прежнему находится на слиянии релиза v2.1.0 от 10-14, QwenLM/FlashQLA — на исправляющем коммите от 10-14, в окне у обоих нет пущей, у принимающей стороны сегодня тихо. Ветка snapshots репозитория TileOPs-nightly за окно получила один новый снимок (10-20 05:39), предыдущий был 10-19 05:41, ночной бенчмарк сохраняет ритм раз в день.
Источники: TileOPs-nightly snapshots (10-20)
4. Сообщество, руководства и мероприятия
4.1 Документация: шесть новых операторов внесены на страницу API, сайт документации основного репозитория обновляется автоматически (#77)
Сайт документации TileOPs вместе со слиянием операторов дополнил страницу API шестью новыми реализациями (#77), по каждой изложены соглашения о вызове и текущий статус, в соответствии с правилом «реализовано — значит на странице», установленным на прошлой неделе. Сайт документации основного репозитория tilelang за окно получил одну автоматическую сборку-пуш без ручных изменений содержимого.
Источники: tile-ai/TileOPs.github.io#77 (10-19)
4.2 Репозитории сообщества: репозиторий практики W4A16 kernel и репозиторий навыков отладки CUDA попали в поле зрения (10-19)
В сообществе появились несколько сторонних репозиториев вокруг TileLang: fanyuxiao-ui/w4a16-tilelang создан 10-10, использует TileLang плюс рукописный CUDA для GEMM и GEMV в W4A16-квантованном выводе на A800 (SM80); sablin39/tilelang-cuda-skills создан в мае этого года, уже имеет 145 звёзд, предоставляет наборы навыков для написания на TileLang и отладки с помощью инструментов CUDA, автор также является активным контрибьютором TileFoundry; localizethedocs/tilelang-docs-l10n занимается локализацией официальной документации; chaigua/tilelang-agent создан 10-19. Все четыре — самодеятельные репозитории сообщества, не официальные репозитории организаций.
Источники: fanyuxiao-ui/w4a16-tilelang, sablin39/tilelang-cuda-skills, localizethedocs/tilelang-docs-l10n (10-19)
4.3 Линия версий: основной репозиторий 0.1.15 всё ещё не сменил поколение, TileOPs продолжает фиксировать 0.1.12 (10-19)
Файл версии основного репозитория по-прежнему 0.1.15, что соответствует релизу 09-30, в окне мониторинга нет новых тегов или релизов. В объявлении зависимостей TileOPs по-прежнему жёстко прописано 0.1.12, PR с прошлой недели, пытавшийся поднять до 0.1.15, уже закрыт, новый влитый в этом выпуске тест-кейс SM89 также прогонялся одновременно на 0.1.12 и 0.1.15 dev-версии, что показывает: шлюз обновления по-прежнему предполагает валидацию на двух базовых линиях.
Источники: tile-ai/tilelang VERSION, tile-ai/TileOPs pyproject (10-19)
4.4 Медиа и академическая сфера: три попадания новостного поиска не относятся напрямую, на arXiv новых статей нет (10-19)
Множественные запросы на китайском и английском в Google News (TileLang, tile-ai, TileOPs, TileFoundry, tilelang оператор и т. д.) в окне мониторинга дали лишь три попадания: две отраслевые публикации о DeepSeek и чипах Huawei (The Information, Business Upturn), а также релизный пост Prime Intellect о prime-rl 0.10.0; после построчной проверки в доступном для чтения содержимом не найдено конкретных фактов о TileLang, обработано как не относящееся напрямую, в записи не включено. Поиск tilelang на Hacker News не дал релевантных попаданий, полнотекстовый поиск на arXiv: последняя релевантная статья — всё ещё TileSight от 07-24 (модель профилирования производительности GPU с центричностью на tile), в окне мониторинга новых нет.
Источники: поиск TileLang в Google News, поиск TileLang на arXiv (10-19)
5. Наблюдения за трендами
5.1 SM89 становится точкой приземления по умолчанию для новых операторов
Вчера фокус TileOPs был всё ещё на W4A8 на Hopper и кластерном взаимодействии SM90, в этом окне мониторинга оба новых GEMM стартуют прежде всего на SM89 и в PR явно заявляют, что не претендуют на результаты работы на SM90. Видеокарты уровня Ada (L40S, 4090 и подобные) — основное железо в руках отечественных разработчиков, и приземление низкоточных GEMM сначала на SM89 означает, что библиотека операторов начинает расставлять приоритеты по «доступности», а не только по «пиковой производительности».
5.2 Линейное внимание закрепляется как «одно матричное умножение на chunk + выбор стратегии по волнам SM»
Четыре линии GLA, GDN, KDA, DeltaNet в этом выпуске дали в сумме восемь изменений производительности, приёмы сошлись к двум: сжать матричные умножения в цепочке chunk до одного на блок и определять по числу групп warp, помещающихся в SM, использовать групповое решение или одиночный warp. Такая запись «сначала посчитать волны, потом выбрать ядро» ранее встречалась только в рукописных библиотеках GEMM, теперь она вошла в путь по умолчанию библиотеки операторов.
5.3 TileFoundry выходит за пределы CUDA: Apple AMX и префилл Qwen
TileFoundry всегда позиционировался как «фабрика операторов для агентов», но ранее все примеры были на NVIDIA. Поддержка AMX в этом выпуске плюс прогон префилла целой модели Qwen на M2 Pro показывают, что его трёхуровневая абстракция спецификации, раскладки и планирования не привязана к CUDA, а может опускаться до матричного сопроцессора Apple. Если эта линия продолжится, позиционирование TileFoundry сменится с «фронтенда TileLang» на «кросс-акселераторный компилятор от спецификации к ядру».
5.4 Наблюдаемость становится главной линией закалки компилятора
В этом окне три изменения указывают на одно и то же: сделать результаты компиляции измеримыми и атрибутируемыми. lower trace раздельно сохраняет исходный код хоста и устройства на диск, Autotuner действительно передаёт параметры декоратора в метрики, Ascend сопоставляет хостовые операторы и ядра устройств под msprof. После того как операторный охват почти заполнен (статистика спецификаций уже сообщает 205/205), оставшийся разрыв нужно искать через измерения, а не через добавление операторов.
5.5 Библиотека операторов Ascend переходит на специализированные операторы для диалоговых моделей
В очереди на рассмотрение Ascend появились kpool-сжатие для GLM-5.3, обратный проход KDA и indexer, situ_glu и его градиенты — это не универсальные операторы, а компоненты структур конкретных диалоговых моделей. Вместе с автоматной векторизацией PTO на прошлой неделе центр тяжести работ по линии Ascend сместился с «запуска TileLang на Ascend» на «сначала реализовать на Ascend ключевые операторы отечественных популярных моделей».
Приложение: материалы и пояснения по проверке
- Данное окно является обычным 24-часовым окном, охватывающим период с 10-19 07:00 по 10-20 07:00, и не пересекается с предыдущим выпуском (трёхдневное окно с 10-16 07:00 по 10-19 07:00).
- Список репозиториев организации и статусы push получены полной выгрузкой через GitHub API; из 29 репозиториев 9 имели push в окне; коммиты в основную ветку каждого репозитория проверены построчно через Atom feed, ветки, отличные от ветки по умолчанию (ascendc_pto в tilelang-ascend, snapshots в TileOPs-nightly, stats в TileOPs), проверены отдельно.
- Количество слитых и находящихся на рассмотрении взято из достоверного поиска через GitHub search API (по одному запросу is:merged и is:open с меткой времени начала окна), без использования накопления списков, выгруженных по каждому репозиторию, во избежание повторного счёта.
- Анонимная квота API организации исчерпалась на позднем этапе сбора, поэтому текст PR и статистика изменений считывались из веб-страниц и файлов патчей, что не повлияло на проверку фактов по записям; статус «слито» всегда определяется по результатам поиска is:merged.
- Цифры покрытия спецификаций и бенчмарков взяты из данных бейджей ветки статистики организации; эта ветка обновлена после последнего слияния в окне и относится к самодекларируемым показателям проекта.
- Количество ночных снимков определяется путём пересчёта времени UTC из Atom feed ветки snapshots в пекинское время и проверки попадания в окно.
-
Три совпадения по новостной стороне проверены по отдельности; из них два отраслевых материала не позволяют подтвердить конкретную связь с TileLang по доступному для чтения содержимому и обработаны как непроверенные, не включаясь в записи этого выпуска; по arXiv и Hacker News — ноль новых.
- Отчёт не содержит никакой информации о локальной среде.
Полный список источников
- [1] tile-ai/tilelang PR #2962 Поддержка Windows ARM64 и CI — https://github.com/tile-ai/tilelang/pull/2962
- [2] tile-ai/tilelang PR #3536 T.cumsum смещение потока на разделяемом буфере — https://github.com/tile-ai/tilelang/pull/3536
- [3] tile-ai/tilelang PR #2618 call_extern маркировка записи в переменные с «сырыми» данными — https://github.com/tile-ai/tilelang/pull/2618
- [4] tile-ai/tilelang PR #3496 lower trace разделение исходного кода хоста и устройства — https://github.com/tile-ai/tilelang/pull/3496
- [5] tile-ai/tilelang PR #3337 Autotuner сквозная передача параметров profiler — https://github.com/tile-ai/tilelang/pull/3337
- [6] tile-ai/tilelang PR #3548 ROCm FP4 печать констант с плавающей точкой — https://github.com/tile-ai/tilelang/pull/3548
- [7] tile-ai/tilelang PR #3565 обновление cibuildwheel — https://github.com/tile-ai/tilelang/pull/3565
- [8] tile-ai/tilelang PR #3567 ROCm линия добавление ревьюера — https://github.com/tile-ai/tilelang/pull/3567
- [9] tile-ai/tilelang PR #3321 Ascend msprof связывание хоста и устройства — https://github.com/tile-ai/tilelang/pull/3321
- [10] tile-ai/tilelang PR #3466 Ascend VF скалярный sqrt вне области — https://github.com/tile-ai/tilelang/pull/3466
- [11] tile-ai/tilelang PR #3535 Ascend PTO типизированный вектор Ramp — https://github.com/tile-ai/tilelang/pull/3535
- [12] tile-ai/tilelang-ascend PR #1657 документация на китайском для четырёх поэлементных операторов — https://github.com/tile-ai/tilelang-ascend/pull/1657
- [13] tile-ai/tilelang-ascend PR #1910 GLM-5.3 оператор сжатия kpool — https://github.com/tile-ai/tilelang-ascend/pull/1910
- [14] tile-ai/tilelang-ascend PR #1895 KDA обратный оператор chunk — https://github.com/tile-ai/tilelang-ascend/pull/1895
- [15] tile-ai/tilelang-ascend PR #1887 обратный оператор indexer — https://github.com/tile-ai/tilelang-ascend/pull/1887
- [16] tile-ai/tilelang-ascend PR #1907 UB в L1 аннотация блока сохранения промежуточных данных — https://github.com/tile-ai/tilelang-ascend/pull/1907
- [17] tile-ai/tilelang-ascend PR #1903 исправление выпуска кэша параллельных ядер — https://github.com/tile-ai/tilelang-ascend/pull/1903
- [18] tile-ai/tilelang-hygon PR #25 MLS начало константы K окна — https://github.com/tile-ai/tilelang-hygon/pull/25
- [19] tile-ai/tilelang-hygon PR #18 регистровый конвейер и warp_diverge — https://github.com/tile-ai/tilelang-hygon/pull/18
- [20] tile-ai/tilelang-mlir-ascend PR #204 TVM исправление изоляции — https://github.com/tile-ai/tilelang-mlir-ascend/pull/204
-
[21]
- [22] tile-ai/TileOPs PR #2655 SM89 INT8 W8A8 GEMM — https://github.com/tile-ai/TileOPs/pull/2655
- [23] tile-ai/TileOPs PR #2642 MLA постраничный прямой проход — https://github.com/tile-ai/TileOPs/pull/2642
- [24] tile-ai/TileOPs PR #2635 MLA запись KV-кэша — https://github.com/tile-ai/TileOPs/pull/2635
- [25] tile-ai/TileOPs PR #2639 Фьюзинг QK-Norm-RoPE — https://github.com/tile-ai/TileOPs/pull/2639
- [26] tile-ai/TileOPs PR #2634 Фьюзинг RMSNorm с пошаговым квантованием — https://github.com/tile-ai/TileOPs/pull/2634
- [27] tile-ai/TileOPs PR #2651 Обёртка операторов INT8 W8A8 — https://github.com/tile-ai/TileOPs/pull/2651
- [28] tile-ai/TileOPs PR #2648 Постраничный FP8 прямой проход GQA переведён в реализованное состояние — https://github.com/tile-ai/TileOPs/pull/2648
- [29] tile-ai/TileOPs PR #2650 Решение прямого заполнения GDN по warp-группам — https://github.com/tile-ai/TileOPs/pull/2650
- [30] tile-ai/TileOPs PR #2645 Состояние прямого прохода DeltaNet с одним матричным умножением на chunk — https://github.com/tile-ai/TileOPs/pull/2645
- [31] tile-ai/TileOPs PR #2637 Удаление матричного умножения из градиента обратного состояния GLA — https://github.com/tile-ai/TileOPs/pull/2637
- [32] tile-ai/TileOPs PR #2636 Перекрытие записи тайлов пакетного GEMM — https://github.com/tile-ai/TileOPs/pull/2636
- [33] tile-ai/TileOPs PR #2641 Свёртка длинных построчных редукций на 512 потоках — https://github.com/tile-ai/TileOPs/pull/2641
- [34] tile-ai/TileOPs PR #2640 Стратегия вытеснения L2 для строк индексов max-pooling — https://github.com/tile-ai/TileOPs/pull/2640
- [35] tile-ai/TileOPs PR #2646 Обход масштабирования деления bf16 — https://github.com/tile-ai/TileOPs/pull/2646
- [36] tile-ai/TileOPs PR #2652 Остаток от деления float32 и округление целочисленного деления вниз — https://github.com/tile-ai/TileOPs/pull/2652
- [37] tile-ai/TileOPs PR #2638 Прямой проход выборочного сканирования Mamba-1 — https://github.com/tile-ai/TileOPs/pull/2638
- [38] tile-ai/TileOPs PR #2653 Схождение перечисления групповой маршрутизации MoE — https://github.com/tile-ai/TileOPs/pull/2653
- [39] tile-ai/TileOPs PR #2656 Новая колонка операторов и разбор истории в ночном отчёте — https://github.com/tile-ai/TileOPs/pull/2656
- [40] tile-ai/TileOPs ветка статистики спецификаций (205/205) — https://github.com/tile-ai/TileOPs/tree/stats
- [41] tile-ai/TileFoundry PR #244 Поддержка Apple AMX и пример предзаполнения Qwen — https://github.com/tile-ai/TileFoundry/pull/244
-
[42] tile-ai/TileFoundry PR #240 Четыре пробела в генерации кода CUDA — <https://github.com/tile-ai/T
- [43] tile-ai/TileFoundry PR #245 унификация контекста оценки — https://github.com/tile-ai/TileFoundry/pull/245
- [44] tile-ai/TileOPs-nightly ветка snapshots — https://github.com/tile-ai/TileOPs-nightly/commits/snapshots
- [45] tile-ai/TileOPs.github.io PR #77 шесть новых операторов добавлены на страницу API — https://github.com/tile-ai/TileOPs.github.io/pull/77
- [46] tile-ai/tilelang файл версии 0.1.15 — https://github.com/tile-ai/tilelang/blob/main/VERSION
- [47] tile-ai/TileOPs объявление зависимости tilelang 0.1.12 — https://github.com/tile-ai/TileOPs/blob/main/pyproject.toml
- [48] fanyuxiao-ui/w4a16-tilelang реализация сообщества — https://github.com/fanyuxiao-ui/w4a16-tilelang
- [49] sablin39/tilelang-cuda-skills пакет навыков сообщества — https://github.com/sablin39/tilelang-cuda-skills
- [50] localizethedocs/tilelang-docs-l10n локализация документации — https://github.com/localizethedocs/tilelang-docs-l10n
- [51] Google News поиск по TileLang — https://news.google.com/search?q=TileLang
- [52] arXiv поиск по TileLang — https://arxiv.org/search/?query=TileLang&searchtype=all