ROCm ядра для PyTorch: создание и публикация кастомных ядер на AMD GPU | AiManual
AiManual Logo Ai / Manual.
01 Фев 2026 Гайд

Пишем и делимся высокопроизводительными ROCm-ядрами для PyTorch: полный гайд на примере GEMM

Пошаговое руководство по разработке и публикации высокопроизводительных ROCm-ядер для PyTorch на примере GEMM kernel. Инструменты, оптимизации, ошибки.

Зачем вообще писать свои ядра? (Спойлер: производительность)

Представьте ситуацию: у вас есть свежая Radeon AI Pro R9700, вы запускаете обучение модели, а оно идет медленнее, чем на трехлетней RTX 4090. Знакомо? Это не потому, что AMD хуже. Это потому, что PyTorch по умолчанию использует универсальные ядра, которые работают везде, но не оптимально нигде.

Вот вам цифры на 01.02.2026: стандартный torch.matmul на матрицах 8192x8192 дает около 45 TFLOPS на MI300X. Кастомный GEMM kernel, написанный под конкретную архитектуру - до 90 TFLOPS. Разница в два раза. Это не магия, это просто знание железа.

Важно: я говорю не про CUDA. ROCm - это отдельная вселенная со своими особенностями. Если вы привыкли к CUDA, забудьте половину привычек. Здесь все работает иначе, иногда удобнее, иногда - раздражающе.

Проблема: почему стандартные ядра такие медленные?

PyTorch использует библиотеки типа MIOpen или rocBLAS для базовых операций. Они хороши для общего случая, но не учитывают:

  • Конкретный размер ваших тензоров
  • Паттерны доступа к памяти в вашей модели
  • Возможности новых инструкций (например, matrix core на CDNA 3)
  • Оптимизации под конкретный datatype (FP8, BF16 с разными rounding modes)

В статье про обучение Llama 3.2 я уже показывал, как кастомные ядра ускоряют обучение в разы. Но там речь шла о готовых решениях. Сегодня - о том, как сделать свои.

Решение: kernel-builder от Hugging Face

До недавнего времени написание ROCm-ядер было адом. Нужно было знать ассемблер, разбираться в регистрах, писать километры boilerplate кода. В 2024 году Hugging Face выпустили kernel-builder - инструмент, который меняет правила игры.

💡
Kernel-builder - это Python-библиотека, которая генерирует высокооптимизированный HIP-код (ROCm аналог CUDA) из высокоуровневого описания. Вы описываете, что хотите сделать, а она генерирует код, который работает быстрее рукописного ассемблера в 80% случаев.

Почему это важно? Потому что теперь вы можете:

  1. Писать ядра на Python, а не на HIP/C++
  2. Автоматически тестировать их на разных размерах тензоров
  3. Сравнивать производительность с эталонными реализациями
  4. Делиться ядрами через kernels-community репозиторий

Пошаговый план: от идеи до публикации

1 Установка и настройка окружения

Первое правило ROCm разработки: никогда не используйте системные пакеты. Всегда ставьте через Docker или conda. На 01.02.2026 актуальная версия ROCm - 6.3.2.

# Создаем отдельное окружение
conda create -n rocm-kernels python=3.11 -y
conda activate rocm-kernels

# Устанавливаем ROCm совместимый PyTorch
pip install torch torchvision torchaudio --index-url https://download.pytorch.org/whl/rocm6.3

# Kernel-builder и зависимости
pip install kernel-builder==0.9.4  # последняя версия на 01.02.2026
pip install transformers datasets

Предупреждение: если у вас гибридная система с NVIDIA и AMD (как в этой статье), убедитесь, что ROCm видит только AMD карты. Иначе будет конфликт драйверов.

2 Пишем свой первый GEMM kernel

GEMM (General Matrix Multiply) - это основа всего. Линейные слои, внимание, свертки - все сводится к матричному умножению. Начнем с простого случая: умножение двух float16 матриц.

Как НЕ надо делать:

# Наивная реализация - медленно
import torch

def naive_gemm(A, B):
    return torch.matmul(A, B)

Почему это плохо? Потому что torch.matmul вызывает общую библиотеку, которая не знает про:

  • Выравнивание памяти
  • Локальность данных
  • Векторизацию инструкций
  • Использование shared memory

А вот как надо:

import kernel_builder as kb
import torch
import numpy as np

# Определяем размеры tile'ов для оптимизации
# Для CDNA 3 архитектуры (MI300X, R9700) оптимальные размеры:
TILE_M = 128
TILE_N = 128
TILE_K = 32

# Создаем спецификацию ядра
gemm_spec = kb.KernelSpec(
    name="custom_gemm_fp16",
    inputs=[
        kb.TensorSpec(name="A", dtype=kb.float16, shape=["M", "K"]),
        kb.TensorSpec(name="B", dtype=kb.float16, shape=["K", "N"]),
    ],
    outputs=[
        kb.TensorSpec(name="C", dtype=kb.float16, shape=["M", "N"]),
    ],
    # Используем tile-based подход для лучшей локализации
    tile_sizes={
        "M": TILE_M,
        "N": TILE_N,
        "K": TILE_K,
    },
    # Включаем использование matrix core если доступно
    features=["matrix_core" if kb.has_matrix_core() else "vector_units"],
)

# Генерируем код ядра
kernel_code = kb.generate_kernel(gemm_spec)

# Компилируем в PyTorch extension
gemm_kernel = kb.compile_kernel(kernel_code, "custom_gemm")

3 Оптимизации: от 2x до 10x ускорения

Сгенерированное ядро уже будет быстрее стандартного, но не максимально. Вот что нужно добавить:

Оптимизация Эффект Сложность
Prefetching данных +20-30% Низкая
Double buffering +15-25% Средняя
Использование LDS +40-60% Высокая
Векторизация MFMA +100-200% Очень высокая

Добавим prefetching в наше ядро:

# Улучшенная спецификация с prefetching
gemm_optimized = kb.KernelSpec(
    name="optimized_gemm_fp16",
    inputs=[
        kb.TensorSpec(name="A", dtype=kb.float16, shape=["M", "K"]),
        kb.TensorSpec(name="B", dtype=kb.float16, shape=["K", "N"]),
    ],
    outputs=[
        kb.TensorSpec(name="C", dtype=kb.float16, shape=["M", "N"]),
    ],
    tile_sizes={
        "M": 128,
        "N": 128,
        "K": 32,
    },
    features=[
        "matrix_core",
        "prefetch",  # Автоматический prefetching
        "double_buffer",  # Double buffering для скрытия latency
    ],
    # Настройки оптимизации под конкретную архитектуру
    arch_specific={
        "cdna3": {
            "vector_width": 64,
            "lds_size": 65536,
            "prefetch_distance": 2,
        },
        "rdna3": {
            "vector_width": 32,
            "lds_size": 32768,
            "prefetch_distance": 1,
        }
    }
)

4 Тестирование и бенчмаркинг

Написали ядро? Отлично. Теперь проверьте, что оно:

  1. Дает правильный результат
  2. Быстрее стандартной реализации
  3. Не ломается на edge cases
import torch
import time

# Тест корректности
def test_correctness():
    M, N, K = 1024, 1024, 1024
    
    # Создаем тестовые данные
    A = torch.randn(M, K, dtype=torch.float16).cuda()
    B = torch.randn(K, N, dtype=torch.float16).cuda()
    
    # Референсный результат
    C_ref = torch.matmul(A, B)
    
    # Наше ядро
    C_custom = gemm_kernel(A, B)
    
    # Проверяем точность
    error = torch.max(torch.abs(C_ref - C_custom))
    assert error < 1e-3, f"Ошибка слишком большая: {error}"
    print(f"✓ Корректность проверена, ошибка: {error:.2e}")

# Бенчмарк производительности
def benchmark():
    sizes = [512, 1024, 2048, 4098, 8192]
    
    for size in sizes:
        A = torch.randn(size, size, dtype=torch.float16).cuda()
        B = torch.randn(size, size, dtype=torch.float16).cuda()
        
        # Warmup
        for _ in range(10):
            _ = torch.matmul(A, B)
            _ = gemm_kernel(A, B)
        
        # Измеряем torch.matmul
        torch.cuda.synchronize()
        start = time.time()
        for _ in range(100):
            C_ref = torch.matmul(A, B)
        torch.cuda.synchronize()
        torch_time = time.time() - start
        
        # Измеряем наше ядро
        torch.cuda.synchronize()
        start = time.time()
        for _ in range(100):
            C_custom = gemm_kernel(A, B)
        torch.cuda.synchronize()
        custom_time = time.time() - start
        
        speedup = torch_time / custom_time
        print(f"Size {size}x{size}: torch={torch_time:.3f}s, custom={custom_time:.3f}s, speedup={speedup:.2f}x")

if __name__ == "__main__":
    test_correctness()
    benchmark()

5 Публикация в kernels-community

Вы создали ядро, которое дает 3x ускорение на матрицах 8192x8192. Что дальше? Делитесь! Hugging Face создали репозиторий kernels-community специально для этого.

Шаги публикации:

# Клонируем репозиторий
git clone https://github.com/huggingface/kernels-community.git
cd kernels-community

# Создаем структуру для нашего ядра
mkdir -p kernels/gemm/custom_fp16
cd kernels/gemm/custom_fp16

# Сохраняем код ядра и метаданные
cat > kernel.py << 'EOF'
# Наш код ядра
import kernel_builder as kb

spec = kb.KernelSpec(
    name="custom_gemm_fp16",
    # ... наша спецификация
)

kernel_code = kb.generate_kernel(spec)
EOF

# Создаем README с бенчмарками
cat > README.md << 'EOF'
# Custom GEMM FP16 Kernel

## Performance

| Size | torch.matmul | custom | Speedup |
|------|--------------|--------|---------|
| 1024 | 1.23 ms      | 0.41 ms | 3.00x   |
| 2048 | 9.87 ms      | 3.12 ms | 3.16x   |
| 4096 | 78.5 ms      | 24.8 ms | 3.16x   |
| 8192 | 625 ms       | 198 ms  | 3.16x   |

## Usage
python
from kernels.gemm.custom_fp16 import gemm_kernel

result = gemm_kernel(A, B)

EOF

# Добавляем тесты
mkdir tests
cat > tests/test_gemm.py << 'EOF'
# Тесты для нашего ядра
import torch
import sys
sys.path.append('..')
from kernel import gemm_kernel

# ... тесты
EOF

Ошибки, которые сломают вам день

Я потратил десятки часов на отладку ROCm-ядер. Вот список самых болезненных ошибок:

1. Неправильное выравнивание памяти
ROCm требует, чтобы данные в глобальной памяти были выровнены по 128 байт. Если ваш тензор не выровнен - производительность упадет в 10 раз. Всегда используйте .contiguous() перед передачей в ядро.

2. Игнорирование occupancy
Каждая AMD GPU имеет ограниченное количество wavefronts (аналог warps в CUDA). Если ваше ядро использует слишком много регистров или LDS, occupancy упадет, и GPU будет недогружен. Используйте rocminfo чтобы проверить.

3. Неучет bank conflicts в LDS
Local Data Store (LDS) имеет 32 банка. Если две wavefronts пытаются читать из одного банка - происходит конфликт. В kernel-builder есть автоматическая оптимизация, но проверяйте.

4. Проблемы с синхронизацией
ROCm не имеет аналога __syncthreads() из CUDA. Вместо этого используется amd_wave_barrier() или workgroupBarrier(). Неправильное использование приведет к race conditions.

ROCm vs CUDA: что быстрее?

Вечный спор. На 01.02.2026 ситуация такая:

Операция NVIDIA H100 AMD MI300X Примечания
FP16 GEMM 8192x8192 145 ms 198 ms Стандартные библиотеки
FP16 GEMM 8192x8192 92 ms 112 ms Оптимизированные кастомные ядра
FP8 GEMM 16384x16384 205 ms 228 ms С поддержкой новых инструкций
Attention с flash-attn 18 ms 22 ms Batch size 32, seq_len 4096

Выводы? AMD догоняет, но для максимальной производительности нужны кастомные ядра. Хорошая новость: с kernel-builder писать их стало в разы проще.

Когда это действительно нужно?

Не каждый проект требует кастомных ядер. Вот случаи, когда они окупаются:

  • Тренировка больших моделей - как в случае с Llama 3.2, где каждый процент производительности - это тысячи долларов
  • Инференс в реальном времени - когда задержка критична
  • Специфичные операции - нет в стандартных библиотеках
  • Эксперименты с новыми форматами - типа FP8, которые еще плохо поддерживаются

А вот когда не стоит заморачиваться:

  • Прототипирование - используйте стандартные функции
  • Маленькие модели - разница будет незаметна
  • Если у вас нет времени на отладку - кастомные ядра всегда требуют больше времени

Что дальше? Будущее ROCm экосистемы

К 2026 году ROCm достигла зрелости, но работа продолжается. Вот что нас ждет:

  1. Автоматическая оптимизация - kernel-builder учится подбирать параметры под конкретные размеры тензоров
  2. Кросс-платформенные ядра - один код для CUDA и ROCm (мечта, но уже близко)
  3. Поддержка новых инструкций - AMD анонсировала новые matrix core инструкции для CDNA 4
  4. Интеграция с компиляторами - MLIR-based компиляция прямо из PyTorch

Мой прогноз: к концу 2026 года написание кастомных ядер станет рутинной задачей для любого ML инженера, работающего с AMD GPU. Инструменты упростятся настолько, что разница между CUDA и ROCm практически исчезнет.

А пока - начинайте с простого GEMM ядра, тестируйте на своих данных, делитесь результатами в kernels-community. Каждое новое ядро делает экосистему сильнее. И помните: производительность - это не только железо, но и умение его использовать.

P.S. Если столкнетесь с проблемами при запуске двух Radeon AI Pro R9700 в одной системе, проверьте эту статью - там есть решения самых частых проблем.