AIKraft Agents
ВойтиТарифы и покупка

NVIDIA/cutlass

CUDA-шаблоны и Python DSL для высокопроизводительных линейных алгебраических вычислений

Инструменты⭐ 10 696C++последний релиз: v4.8.1
Открыть на GitHub ↗Сайт проекта ↗

Что это за инструмент

CUTLASS — это библиотека CUDA C++ и Python DSL для создания высокопроизводительных ядер линейной алгебры, оптимизированных под архитектуру NVIDIA GPU.

Зачем нужен

Инструмент решает задачу эффективной реализации матричных умножений (GEMM) и связанных вычислений за счет модульной декомпозиции и тонкой настройки алгоритмов. Он полезен для разработчиков, которым нужно писать кастомные ядра с максимальной производительностью, используя Tensor Cores, без написания низкоуровневого кода с нуля.

Что можно реализовать

  • Разработка кастомных ядер для операций линейной алгебры с поддержкой смешанной точности (FP64, FP32, TF32, FP16, BF16, FP8, FP4).
  • Создание оптимизированных реализаций FlashAttention и других алгоритмов внимания для архитектур Ampere, Hopper и Blackwell.
  • Быстрое прототипирование GPU-ядер с помощью Python DSL (CuTe DSL) для ускорения разработки и компиляции.
  • Интеграция оптимизированных вычислений в DL-фреймворки (PyTorch, cuDNN-Frontend) без написания связующего кода.

Ключевые возможности

  • Поддержка широкого спектра типов данных, включая новые форматы (NVFP4, MXFP4, MXFP6, MXFP8, e2m1).
  • Введение Python DSL (CuTe DSL) для написания высокопроизводительных ядер с более простым синтаксисом и быстрой компиляцией.
  • Глубокая оптимизация под Tensor Cores современных архитектур NVIDIA (Volta, Turing, Ampere, Ada, Hopper, Blackwell).
  • Наличие фреймворка планирования задач (Task Scheduling) и улучшенной диагностики компилятора для отладки ядер.

👤 Кому подойдёт: разработчики системного ПО, исследователи в области машинного обучения, performance-инженеры

Релизы

v4.8.1patchbreaking
🕐 15 дн назад · релиз на GitHub ↗

Исправления в CuTe DSL, защита от ошибок памяти на Blackwell/Rubin и предупреждения о багах компилятора CTK, влияющих на генерацию кода.

  • Fixed: Исправлена ошибка в cute.extern для функций без аргументов и устранено некорректное изменение sys.path.
  • Fixed: tensormap_replace теперь безопасно обрабатывает уменьшение дескрипторов TMA, предотвращая ошибки доступа на Blackwell/Rubin.
  • Fixed: Добавлены четкие ошибки для CopyUniversalOp с указанными атрибутами памяти.
  • Breaking: Избегать использования упакованных инструкций вычитания FP8 (MMA_K=128) из-за багов CTK, приводящих к неверным результатам.
  • Breaking: Некоторые упакованные операции смешанного типа (mul, sub, add) на Rubin (sm107a/sm107f) генерируют неверный код при определенных условиях.
v4.8.0major
🕐 1 мес назад · релиз на GitHub ↗

Добавлена предварительная поддержка архитектуры Rubin, расширены возможности FP8/FP4 и TMA, улучшен профилировщик IKET.

  • Added: Добавлена начальная поддержка архитектуры Rubin для ускорения плотных GEMM с новыми инструкциями Tensor Core.
  • Added: Введены расширения CuTe DSL (cute_ext) с автоматическими CTA-V картами, атомарными операциями и планированием TMEM.
  • Added: Реализована поддержка более высокой пропускной способности FP8 (MMA_K=64) и FP4 (MMA_K=128), включая 2:4 разреженность.
  • Added: Добавлен профилировщик IKET для ядер Rubin (sm107) с возможностью выборочного дампа данных по кластерам.
  • Added: CuTe DSL теперь поддерживает x86_64 Windows, а сборка AoT — целевую платформу QNX 8.0.
v4.8.0devmajor
🕐 2 мес назад · релиз на GitHub ↗

Добавлена предварительная поддержка архитектуры Rubin (SM107), расширены возможности CuTe DSL и улучшена профилировка IKET.

  • Added: Добавлена начальная поддержка архитектуры Rubin (SM107) для ускорения плотных GEMM, включая новые форматы FP8/FP4 и увеличенный размер TMEM.
  • Added: Расширены возможности CuTe DSL: автоматический вывод CTA-V карт, атомарные TMA-операции, планирование буферов TMEM и оптимизация производительности.
  • Added: Введена опциональная предварительная версия конвейера компиляции cute_ext, позволяющая смешивать API в коде @cute.jit и @cute.kernel.
  • Added: Улучшен инструмент профилирования IKET: добавлена поддержка Rubin, возможность профилирования отдельных кластеров и детализация планировщика задач.
  • Added: Добавлены новые примеры для архитектур Rubin, Blackwell и Ampere, включая GEMM с повторным использованием B-коллектора и blockscaled GEMM.
  • Fixed: Исправлена регрессия 4.6.0, из-за которой cute.autovec_copy генерировал поэлементные инструкции вместо векторизованных для динамических стридов.
v4.7.0major
🕐 2 мес назад · релиз на GitHub ↗

CUTLASS v4.7.0: новый Primitives API для CuTe, оптимизация FP8 на Blackwell, поддержка CUDA 12.6 и улучшенная диагностика компиляции.

  • Added: Введён экспериментальный Primitives API для низкоуровневого программирования Tensor Core через SIMT и фреймворк планирования задач с анализом расписаний.
  • Added: Добавлена реализация бэкапда-ядра FMHA для FP8 с улучшением производительности на ~25% по сравнению с одноядерным решением на Blackwell SM103.
  • Added: Расширены возможности кастомных эпилогов: поддержка скалярных редукций и настройка стратегии перемещения данных для каждой операнды.
  • Added: Улучшена диагностика компилятора: теперь выводятся номера строк для переполнения регистров, локальной памяти и ошибок, а также обнаруживаются риски синхронизации.
  • Added: Добавлена поддержка CUDA 12.6, структурных привязок в NVRTC и конвертации форматов в NumericArrayConverter, включая FP4 (e2m1).
  • Fixed: Исправлены эвристика StreamK, неоднозначность в алгоритмах кортежей CuTe и баг драйвера создания TMA, связанный с маппингом памяти.
4.7.0major
🕐 2 мес назад · релиз на GitHub ↗

CUTLASS 4.7.0: новый Primitives API для CuTe, планировщик задач, поддержка FP4 и оптимизация FMHA на Blackwell.

  • Added: Введен экспериментальный Primitives API для низкоуровневого программирования Tensor Core через SIMT.
  • Added: Добавлен фреймворк планирования задач для статического анализа расписаний warp-specialized ядер.
  • Added: Улучшена диагностика компилятора: отчеты о переполнении регистров и ошибках теперь включают номера строк.
  • Added: Реализован 2-ядерный бэкуард для FMHA на FP8 с приростом производительности ~25% на Blackwell.
  • Added: Добавлена поддержка преобразования в FP4 (e2m1) и оптимизация декодирования LUT.
  • Fixed: Исправлены баги TMA, эвристики streamk и неоднозначности в алгоритмах кортежей CuTe.
v4.6.2patch
🕐 2 мес назад · релиз на GitHub ↗

Исправлены ошибки CuTe DSL (TMA, fp32->f8, ptxas), снижен JIT-overhead и ускорен импорт библиотеки.

  • Fixed: Возвращено поведение TMA bulk copy elect_one как в версии 4.5.x для устранения регрессии.
  • Fixed: Исправлена ошибка векторизованного преобразования fp32->f8 и проблема компиляции kernel fp8 grouped_gemm_dglu.
  • Fixed: Устранена проблема с настройкой уровня оптимизации ptxas и добавлена полная кастомизация имени ядра через set_name_prefix.
  • Added: Снижен накладной расход JIT-компиляции примерно на 50 мс на вызов cute.compile.
  • Added: Ускорен импорт cutlass.cute в 3.8 раза при наличии Torch и в 1.25 раза без него.