AIKraft/AI/Поиск…⌘KВойти

NVIDIA/cutlass

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

Инструменты⭐ 10 461C++последний релиз: v4.8.0dev
Открыть на 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.0devmajor
🕐 25 дн назад · релиз на 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
🕐 1 мес назад · релиз на 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
🕐 1 мес назад · релиз на 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
🕐 1 мес назад · релиз на 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 раза без него.