Ingénierie GPU pratique : shaders Metal, profilage et déploiement edge

L'ingénierie GPU pratique couvre la programmation de shaders multiplateforme, le diagnostic par compteurs matériels et le déploiement edge sur appareil. Vous trouverez ci-dessous une analyse approfondie de ces trois flux de travail.

1. Shader de multiplication matricielle par tuiles personnalisé en Apple Metal (MSL)

Sur Apple Silicon (M1–M4 / série A), les kernels de calcul sont écrits en Metal Shading Language (MSL) et exploitent la Threadgroup Memory (l'équivalent Apple de la Shared Memory de CUDA) pour éviter les lectures lentes en mémoire globale.

Vous trouverez ci-dessous un shader de multiplication matricielle par tuiles (C = A × B) en MSL utilisant des tuiles de threadgroup de 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;
    }
}

• Correspondance matérielle clé : [[thread_position_in_grid]] correspond directement à blockIdx * blockDim + threadIdx de CUDA, tandis que threadgroup_barrier(mem_flags::mem_threadgroup) correspond à __syncthreads() de CUDA. Pour les charges de travail FP16 en production, Apple Silicon expose également les intrinsèques matériels simdgroup_float8x8 pour exploiter directement les unités de multiplication matricielle d'Apple.

2. Procédure pas à pas d'optimisation Roofline avec ncu

Supposons que vous écriviez un kernel CUDA ou Triton naïf pour une LayerNorm par ligne sur un tenseur FP16 (4096, 4096), et qu'il s'exécute nettement plus lentement que l'implémentation native de PyTorch.

Étape 1 : capturer les compteurs matériels

Exécutez NVIDIA Nsight Compute (ncu) en ciblant uniquement le nom de votre kernel personnalisé pour éviter la surcharge de profilage sur le code d'initialisation :

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

Étape 2 : diagnostiquer le goulot d'étranglement sur le Roofline

Inspectez les métriques produites :

• Débit de calcul (SM) : 11.4% du pic.

• Débit mémoire (DRAM/HBM) : 84.2% du pic.

• Raison de blocage des warps (long_scoreboard) : 68.5% (les warps sont bloqués en attente du retour des lectures de la mémoire HBM globale).

• Diagnostic : le kernel se situe nettement sous le toit incliné de bande passante mémoire sur le graphique Roofline. De plus, la vérification de l1tex__t_sectors révèle des lectures mémoire non coalescées : les threads d'un même warp parcourent les colonnes au lieu de lire des adresses de lignes contiguës, gaspillant 75% de chaque ligne de cache HBM de 128 octets récupérée.

Étape 3 : appliquer des corrections matérielles ciblées

1. Corriger la coalescence mémoire : assurez-vous que les threads adjacents (threadIdx.x) lisent des éléments adjacents le long de la dimension interne contiguë (n_cols), afin que 32 lectures FP16 fusionnent en une seule transaction coalescée de 64 octets.

2. Chargements vectorisés (float4 / bloc tl.load) : chargez 128 bits (8 valeurs FP16) par instruction au lieu de 16 bits, réduisant ainsi le nombre total d'instructions mémoire par 8×.

3. Réductions par warp-shuffle : remplacez les réductions lentes en Shared Memory par des warp shuffles registre à registre (__shfl_xor_sync en CUDA ou tl.sum en Triton) pour calculer la moyenne et la variance de ligne sans barrières de synchronisation.

Étape 4 : vérifier sur le Roofline

Une nouvelle exécution de ncu confirme :

• Secteurs de chargement globaux : baisse d'un facteur 4× (utilisation des lignes de cache de 100%).

• Blocages Long Scoreboard : baisse de 68.5% à < 15%.

• Latence wall-clock : amélioration d'un facteur 3.8×, poussant la bande passante effective au-dessus de 85% de la limite HBM physique du H100.

3. Pipeline de déploiement edge : quantifier et exécuter un LLM sur mobile/ordinateur portable

Déployer un LLM de 8 milliards de paramètres sur un ordinateur portable ou un smartphone (8 GB–16 GB de RAM partagée) exige de compresser les poids en entiers 4 bits et de les exécuter via un runtime C++ sans surcharge comme llama.cpp ou 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)]

Étape 1 : conversion et calibration (matrice d'importance)

Un arrondi brut en 4 bits dégrade la précision de raisonnement et de codage. Avec llama.cpp, convertissez d'abord les poids Hugging Face en un fichier GGUF non compressé, puis calculez une matrice d'importance (imatrix) sur un petit échantillon de texte de calibration pour mesurer quels canaux de poids ont le plus grand impact d'activation :

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

Étape 2 : K-quantification à bits mixtes (Q4_K_M)

Quantifiez le modèle avec Q4_K_M (K-quant Medium 4 bits), qui conserve les couches sensibles de valeur/sortie d'attention en 6 bits (Q6_K) tout en compressant les couches feed-forward denses en 4 bits (Q4_K) dans des super-blocs de 256 poids :

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

Cela réduit le modèle de 16.0 GB à ~4.9 GB avec moins de 1% de dégradation de la perplexité.

Étape 3 : exécution zéro copie en mémoire unifiée sur l'appareil

Lors de l'exécution sur un iPhone, un appareil Android ou un MacBook :

• Chargement instantané via mmap : le runtime mappe en mémoire (mmap) le fichier .gguf de 4.9 GB directement depuis le stockage flash vers la mémoire unifiée, éliminant la duplication de RAM et chargeant le modèle en moins d'une seconde.

• Shaders de déquantification fusionnés : pendant la génération de tokens, des compute shaders Metal (ggml-metal.metal) ou Vulkan (mul_mat_q4_K.comp) personnalisés chargent les blocs de poids 4 bits compressés directement dans les registres du GPU, les décompressent à l'aide de décalages binaires et les multiplient à la volée par les activations FP16, atteignant 40–70+ tokens/s sur les GPU d'ordinateurs portables modernes et 15–25 tokens/s sur les téléphones mobiles haut de gamme.