Benchmark SAXPY на cuda-core v1.0.0: A100 vs RTX 4090 vs RTX 3090 — реальные цифры

После выхода cuda-core v1.0.0 мы решили проверить практическую сторону: насколько эффективно новый Python API использует GPU-память? Классический тест — SAXPY (Single-precision A·X Plus Y), операция вида out[i] = alpha * x[i] + y[i].

SAXPY читает два массива и пишет один — суммарно 3 операции с памятью на элемент. При известном размере массива легко посчитать эффективную полосу пропускания (GB/s) и сравнить с теоретическим пиком GPU.

Методология

Что меряем: эффективная пропускная способность памяти (Memory Bandwidth)

Формула: BW (GB/s) = 3 × N × 4 байта / время (секунды) / 1e9

Где:
- 3 — три обращения к памяти (read x, read y, write out)
- 4 байта — размер float32
- 5 прогонов, берём среднее и минимум (пиковое значение)

Прогрев (warmup): один запуск перед измерением — исключает эффекты холодного старта JIT и кэша L2.

Синхронизация: stream.sync() после каждого запуска — точное измерение именно времени выполнения ядра.

from cuda.core import (
    Device, DeviceMemoryResource, DeviceMemoryResourceOptions,
    Program, LaunchConfig, launch
)
from cuda.bindings import driver as cuda_driver
import numpy as np, time

SAXPY_SRC = r"""
__global__ void saxpy(float alpha, float* x, float* y, float* out, size_t n) {
    size_t tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) out[tid] = alpha * x[tid] + y[tid];
}
"""

def bench(n, n_repeat=5):
    d = Device(0); d.set_current()
    stream = d.create_stream()
    mr = DeviceMemoryResource(d.device_id, DeviceMemoryResourceOptions())

    prog = Program(SAXPY_SRC, code_type="c++")
    mod = prog.compile("cubin", ("saxpy",))  # v1.0.0: name_expressions обязателен
    kernel = mod.get_kernel("saxpy")

    nbytes = n * 4  # float32
    x_buf = mr.allocate(nbytes, stream=stream)
    y_buf = mr.allocate(nbytes, stream=stream)
    out_buf = mr.allocate(nbytes, stream=stream)

    # Инициализация нулями
    cuda_driver.cuMemsetD32Async(x_buf.handle, 0, n, stream.handle)
    cuda_driver.cuMemsetD32Async(y_buf.handle, 0, n, stream.handle)
    stream.sync()

    block = 256
    grid = (n + block - 1) // block
    cfg = LaunchConfig(grid=grid, block=block)
    alpha = np.float32(2.0)

    # Warmup
    launch(stream, cfg, kernel, alpha, x_buf.handle, y_buf.handle, out_buf.handle, np.uint64(n))
    stream.sync()

    times_ms = []
    for _ in range(n_repeat):
        t0 = time.perf_counter()
        launch(stream, cfg, kernel, alpha, x_buf.handle, y_buf.handle, out_buf.handle, np.uint64(n))
        stream.sync()
        times_ms.append((time.perf_counter() - t0) * 1000)

    x_buf.close(); y_buf.close(); out_buf.close()

    bytes_rw = 3 * nbytes  # read x, read y, write out
    avg_ms = sum(times_ms) / len(times_ms)
    bw_gbs = bytes_rw / 1e9 / (avg_ms / 1000)
    return avg_ms, bw_gbs

Полный код бенчмарка: github.com/IntelionCloud/research-examples/cuda-core/

Стенды

cuda-core версия: 1.0.0 на всех стендах.

Результаты

A100 80GB PCIe — детальные данные

Компиляция ядра: 50.9 мс (первый запуск; из кэша — ~0.05 мс)

*1M — попадает в L2-кэш A100 (40 MB). L2-эффект на 5M уже частичный.

Итог A100 80GB PCIe: 1584 GB/s при 100M элементах — 82% от теоретического пика 1935 GB/s.

RTX 4090 — детальные данные

Компиляция ядра: 15.5 мс (первый запуск; из кэша — ~0.05 мс)

*1M и 5M попадают в L2-кэш GPU (72 MB у 4090) — показатели нереалистично высоки.

Итог RTX 4090: 929 GB/s при 100M элементах — 92% от теоретического пика 1008 GB/s.

RTX 3090 — детальные данные

Компиляция ядра: 26.4 мс

Итог RTX 3090: 848 GB/s при 100M элементах — 91% от теоретического пика 936 GB/s.

Сравнение GPU

A100 опережает RTX 4090 по абсолютной пропускной способности в 1.7× (1584 vs 929 GB/s) за счёт HBM2e. RTX 4090 и 3090 показывают более высокую эффективность (92% vs 82%) — у GDDR6X ниже латентность строчных операций, что хорошо для таких же шаблонов как SAXPY.

Для задач с терабайтами данных (обучение больших моделей, инференс с batch > 100) разрыв A100 vs RTX 4090 только растёт — HBM2e рассчитан на устойчивую высокую пропускную способность без троттлинга.

Эффект L2-кэша

На малых размерах измеренная полоса аномально высока — данные помещаются в L2-кэш GPU:
- A100 80GB L2: 40 MB
- RTX 4090 L2: 72 MB
- RTX 3090 L2: 6 MB

Примерные границы L2-эффекта (три массива × float32):
- 1M элементов = 12 MB → у всех трёх GPU в L2
- 5M элементов = 60 MB → у 4090 в L2, у A100 и 3090 — нет

Поэтому для надёжных цифр реальной пропускной способности VRAM нужны массивы 10M+ элементов (40 MB+).

Время компиляции

Одна из ключевых фич cuda-core v1.0.0 — InMemoryProgramCache. Как это работает на практике:

A100 с sm_80 компилирует дольше всех — JIT генерирует больше инструкций для HBM. Но при повторных запусках разница исчезает: из кэша все получают одинаковые ~0.05 мс.

Почему у A100 эффективность ниже, чем у RTX 4090?

Важное уточнение: все три GPU в этом тесте использовали 100% своих SM. При 100M элементах и block=256 получается 390 625 блоков — с избытком хватает для загрузки 108 SM (A100), 128 SM (RTX 4090) и 82 SM (RTX 3090). SM-утилизация тут ни при чём.

Разные проценты отражают, насколько хорошо SAXPY-паттерн (простой последовательный стриминг) насыщает архитектуру памяти каждого GPU:

• GDDR6X (RTX 4090, RTX 3090) хорошо работает с последовательным доступом — 91–92%
• HBM2e (A100) спроектирован под массивно-параллельный доступ: тысячи одновременных разрозненных запросов, матричные операции. Один линейный SAXPY не насыщает все HBM-стеки одновременно — 82%

Это не значит, что A100 «менее эффективна» в целом. На transformer attention, крупных matmul и scatter/gather операциях картина будет другой. SAXPY — синтетический тест пиковой BW, а не задача, под которую A100 оптимизирован.

Выводы

1. A100 даёт в 1.7× больше абсолютной пропускной способности (1584 vs 929 GB/s). Для задач с большими батчами и длинными контекстами — выбор A100.

2. Проценты эффективности отражают совместимость паттерна доступа с архитектурой памяти, а не SM-утилизацию. Все три GPU в тесте полностью загружены.

3. InMemoryProgramCache критичен для продакшн-кода — 300–1000× ускорение при повторных компиляциях одного ядра.

4. Для надёжных цифр реальной BW нужны массивы > L2: 10M+ элементов (40 MB+).

Хотите запустить свои GPU-задачи? Арендуйте RTX 4090 или RTX 3090 в Intelion Cloud — оплата по минутам, без обязательств.