ИИ и технологии

Практическая инженерия GPU: шейдеры Metal, профилирование и развертывание на периферийных устройствах

2026-11-15

Иллюстративный вид с орбиты на огни городов северо-запада Америки, а не вычислительная система ИИ.

Практическая инженерия GPU охватывает кроссплатформенное программирование шейдеров, диагностику по аппаратным счетчикам и развертывание на периферийных устройствах. Ниже подробно рассматриваются все три рабочих процесса.

1. Пользовательский шейдер плиточного матричного умножения для Apple Metal (MSL)

На Apple Silicon (M1–M4 / A-series) вычислительные ядра пишутся на Metal Shading Language (MSL) и используют Threadgroup Memory (аналог CUDA Shared Memory у Apple), чтобы избежать медленного чтения глобальной памяти.

Ниже приведен шейдер плиточного матричного умножения (C = A × B) на MSL с плитками групп потоков 16 × 16.

#include <metal_stdlib>
using namespace metal;

constant uint TILE_SIZE = 16;

struct MatrixDims {
    uint M; // Rows of A and C
    uint N; // Cols of B and C
    uint K; // Cols of A and Rows of B
};

kernel void matmul_tiled(
    device const float* A        [[buffer(0)]],
    device const float* B        [[buffer(1)]],
    device float* C              [[buffer(2)]],
    constant MatrixDims& dims    [[buffer(3)]],
    uint2 gid                    [[thread_position_in_grid]],
    uint2 tid                    [[thread_position_in_threadgroup]]
) {
    // Fast on-chip Threadgroup (SRAM) tiles
    threadgroup float tileA[TILE_SIZE][TILE_SIZE];
    threadgroup float tileB[TILE_SIZE][TILE_SIZE];

    float acc = 0.0f;
    uint numTiles = (dims.K + TILE_SIZE - 1) / TILE_SIZE;

    for (uint t = 0; t < numTiles; ++t) {
        uint tiledCol = t * TILE_SIZE + tid.x;
        uint tiledRow = t * TILE_SIZE + tid.y;

        // Collaboratively load 16x16 tiles into Threadgroup memory
        tileA[tid.y][tid.x] = (gid.y < dims.M && tiledCol < dims.K)
                              ? A[gid.y * dims.K + tiledCol] : 0.0f;
        tileB[tid.y][tid.x] = (tiledRow < dims.K && gid.x < dims.N)
                              ? B[tiledRow * dims.N + gid.x] : 0.0f;

        // Synchronize all threads in the threadgroup before reading tiles
        threadgroup_barrier(mem_flags::mem_threadgroup);

        // Compute partial dot product entirely inside registers/SRAM
        for (uint k = 0; k < TILE_SIZE; ++k) {
            acc += tileA[tid.y][k] * tileB[k][tid.x];
        }

        // Wait before overwriting tiles in the next loop iteration
        threadgroup_barrier(mem_flags::mem_threadgroup);
    }

    if (gid.y < dims.M && gid.x < dims.N) {
        C[gid.y * dims.N + gid.x] = acc;
    }
}

Ключевое соответствие оборудования: [[thread_position_in_grid]] напрямую соответствует blockIdx * blockDim + threadIdx в CUDA, а threadgroup_barrier(mem_flags::mem_threadgroup) соответствует __syncthreads() в CUDA. Для производственных нагрузок FP16 Apple Silicon также предоставляет встроенные аппаратные функции simdgroup_float8x8 для прямого доступа к блокам матричного умножения Apple.

2. Пошаговый разбор оптимизации Roofline с ncu

Предположим, вы написали простое ядро CUDA или Triton для построчного LayerNorm на тензоре FP16 размером (4096, 4096), и оно работает заметно медленнее встроенной реализации PyTorch.

Шаг 1: сбор аппаратных счетчиков

Запустите NVIDIA Nsight Compute (ncu), указав только имя своего ядра, чтобы избежать накладных расходов профилирования кода подготовки:

ncu --kernel-name regex:layernorm_naive --metrics \
sm__throughput.avg.pct_of_peak_sustained_elapsed,\
dram__throughput.avg.pct_of_peak_sustained_elapsed,\
l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum,\
smsp__warp_issue_stalled_long_scoreboard_per_warp_active.pct \
python run_layernorm.py

Шаг 2: диагностика узкого места на Roofline

Изучите выходные метрики:

Производительность вычислений (SM): 11.4% от пикового значения.

Пропускная способность памяти (DRAM/HBM): 84.2% от пикового значения.

Причина простоя варпов (long_scoreboard): 68.5% (варпы простаивают в ожидании завершения чтения глобальной памяти HBM).

Диагноз: ядро находится прямо под наклонным пределом пропускной способности памяти на диаграмме Roofline. Кроме того, проверка l1tex__t_sectors выявляет необъединенные обращения к памяти: потоки одного варпа движутся с шагом по столбцам вместо чтения последовательных адресов строки, тратя впустую 75% каждой загруженной 128-байтной строки кеша HBM.

Шаг 3: применение целевых аппаратных исправлений

1. Исправьте объединение обращений к памяти: обеспечьте чтение соседними потоками (threadIdx.x) соседних элементов вдоль непрерывного внутреннего измерения (n_cols), чтобы 32 чтения FP16 объединялись в одну 64-байтную транзакцию.

2. Векторизованные загрузки (float4 / блок tl.load): загружайте 128 бит (8 значений FP16) за инструкцию вместо 16 бит, сокращая общее число инструкций обращения к памяти в 8×.

3. Редукции через обмен внутри варпа: замените медленные редукции в Shared Memory обменами между регистрами потоков варпа (__shfl_xor_sync в CUDA или tl.sum в Triton), чтобы вычислять среднее значение и дисперсию строки без барьеров синхронизации.

Шаг 4: проверка на Roofline

Повторный запуск ncu подтверждает:

Секторы глобальных загрузок: сокращаются в 4× (использование строк кеша на 100%).

Простои Long Scoreboard: снижаются с 68.5% до < 15%.

Задержка по реальному времени: улучшается в 3.8×, поднимая эффективную пропускную способность выше 85% физического предела HBM у H100.

3. Конвейер развертывания на периферийных устройствах: квантование & запуск LLM на телефоне/ноутбуке

Развертывание LLM с 8 млрд параметров на ноутбуке или смартфоне (8 GB–16 GB общей RAM) требует сжатия весов до 4-битных целых чисел и выполнения через среду C++ без накладных расходов, такую как llama.cpp или Apple MLX.

[HF Safetensors (FP16, 16 GB)] 
       │  1. Convert & Apply Calibration Data (Imatrix)
       ▼
[GGUF Quantized Binary (Q4_K_M, 4.9 GB)]
       │  2. Memory-Map (mmap) into Unified RAM
       ▼
[Metal / Vulkan Compute Shaders (On-the-Fly Dequant + Matrix Mul)]

Шаг 1: преобразование & калибровка (матрица важности)

Прямое 4-битное округление снижает точность рассуждений и программирования. С помощью llama.cpp сначала преобразуйте веса Hugging Face в несжатый файл GGUF, а затем вычислите матрицу важности (imatrix) на небольшой калибровочной выборке текста, чтобы определить, какие каналы весов сильнее всего влияют на активации:

python convert_hf_to_gguf.py ./Llama-3.1-8B --outfile llama-8b-f16.gguf
./llama-imatrix -m llama-8b-f16.gguf -f calibration.txt -o imatrix.dat

Шаг 2: K-квантование со смешанной разрядностью (Q4_K_M)

Квантуйте модель с помощью Q4_K_M (4-bit K-quant Medium): этот формат сохраняет чувствительные слои значений/выходов внимания в 6-битном виде (Q6_K), а плотные слои прямого распространения сжимает до 4 бит (Q4_K) в суперблоках по 256 весов:

./llama-quantize --imatrix imatrix.dat llama-8b-f16.gguf llama-8b-q4_k_m.gguf Q4_K_M

Это уменьшает модель с 16.0 GB до ~4.9 GB при ухудшении перплексии менее чем на 1%.

Шаг 3: выполнение на устройстве с объединенной памятью без копирования

При запуске на iPhone, устройстве Android или MacBook:

Мгновенная загрузка через mmap: среда выполнения отображает файл .gguf объемом 4.9 GB непосредственно из флеш-хранилища в объединенную память (mmap), устраняя дублирование RAM и загружая модель менее чем за 1 секунду.

Объединенные шейдеры деквантования: во время генерации токенов пользовательские вычислительные шейдеры Metal (ggml-metal.metal) или Vulkan (mul_mat_q4_K.comp) загружают упакованные блоки 4-битных весов прямо в регистры GPU, распаковывают их с помощью побитовых сдвигов и на лету умножают на активации FP16, достигая 40–70+ токенов/с на современных GPU ноутбуков и 15–25 токенов/с на флагманских мобильных телефонах.

Больше новостей: ИИ и технологии