Uygulamalı GPU mühendisliği: Metal gölgelendiricileri, profilleme ve uç cihazlara dağıtım

Uygulamalı GPU mühendisliği, platformlar arası gölgelendirici programlamayı, donanım sayacı tanılamasını ve cihaz üzerinde uç dağıtımını kapsar. Aşağıda bu üç iş akışı ayrıntılı olarak ele alınmaktadır.

1. Özel Apple Metal (MSL) Karolu Matris Çarpımı Gölgelendiricisi

Apple Silicon (M1–M4 / A-series) üzerinde hesaplama çekirdekleri Metal Shading Language (MSL) ile yazılır ve yavaş genel bellek okumalarından kaçınmak için Threadgroup Memory (Apple'ın CUDA Shared Memory karşılığı) kullanır.

Aşağıda 16 × 16 iş parçacığı grubu karoları kullanan, MSL ile yazılmış bir karolu Matris Çarpımı (C = A × B) gölgelendiricisi yer almaktadır.

#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;
    }
}

• Temel Donanım Eşlemesi: [[thread_position_in_grid]], doğrudan CUDA'nın blockIdx * blockDim + threadIdx ifadesine; threadgroup_barrier(mem_flags::mem_threadgroup) ise CUDA'nın __syncthreads() işlevine karşılık gelir. Üretim ortamındaki FP16 iş yükleri için Apple Silicon, Apple'ın matris çarpma birimlerine doğrudan erişmek üzere simdgroup_float8x8 donanım içsel işlevlerini de sunar.

2. Adım Adım ncu Roofline Optimizasyonu

(4096, 4096) boyutlu bir FP16 tensöründe Satır Bazlı LayerNorm için basit bir CUDA veya Triton çekirdeği yazdığınızı ve bunun PyTorch'un yerel uygulamasından belirgin biçimde yavaş çalıştığını varsayalım.

Adım 1: Donanım Sayaçlarını Kaydedin

Kurulum kodunda profilleme ek yükünden kaçınmak için NVIDIA Nsight Compute (ncu) aracını yalnızca özel çekirdeğinizin adını hedefleyerek çalıştırın:

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

Adım 2: Roofline Üzerinde Darboğazı Teşhis Edin

Çıktı metriklerini inceleyin:

• Hesaplama (SM) İş Hacmi: Tepe değerin %11.4'ü.

• Bellek (DRAM/HBM) İş Hacmi: Tepe değerin %84.2'si.

• Warp Duraklama Nedeni (long_scoreboard): %68.5 (warp'lar genel HBM bellek okumalarının dönmesini beklerken duraklar).

• Teşhis: Çekirdek, Roofline grafiğinde eğimli Bellek Bant Genişliği Tavanı'nın tam altında yer alır. Ayrıca l1tex__t_sectors incelendiğinde birleştirilmemiş bellek okumaları görülür; aynı warp içindeki iş parçacıkları ardışık satır adreslerini okumak yerine sütunlar boyunca adımlamakta ve getirilen her 128 baytlık HBM önbellek satırının %75'ini boşa harcamaktadır.

Adım 3: Hedefe Yönelik Donanım Düzeltmelerini Uygulayın

1. Bellek Erişimlerini Birleştirmeyi Düzeltin: Bitişik iş parçacıklarının (threadIdx.x), ardışık iç boyut (n_cols) boyunca bitişik elemanları okumasını sağlayın; böylece 32 FP16 okuması tek bir 64 baytlık birleştirilmiş işlemde birleşir.

2. Vektörleştirilmiş Yüklemeler (float4 / tl.load bloğu): Komut başına 16 bit yerine 128 bit (8 FP16 değeri) yükleyerek toplam bellek komutu sayısını 8× azaltın.

3. Warp-Shuffle İndirgemeleri: Satır ortalamasını ve varyansını eşzamanlama bariyerleri olmadan hesaplamak için yavaş Shared Memory indirgemeleri yerine yazmaçtan yazmaca warp shuffle işlemlerini (CUDA'da __shfl_xor_sync veya Triton'da tl.sum) kullanın.

Adım 4: Roofline Üzerinde Doğrulayın

ncu tekrar çalıştırıldığında şu sonuçlar doğrulanır:

• Genel Yükleme Sektörleri: 4× azalır (%100 önbellek satırı kullanımı).

• Long Scoreboard Duraklamaları: %68.5'ten < %15'e düşer.

• Gerçek Süre Gecikmesi: 3.8× iyileşerek etkin bant genişliğini H100'ün fiziksel HBM sınırının %85'inin üzerine çıkarır.

3. Uç Dağıtım Hattı: Bir LLM'yi Nicemleme & Mobil/Dizüstü Bilgisayarda Çalıştırma

8 milyar parametreli bir LLM'yi dizüstü bilgisayara veya akıllı telefona (8 GB–16 GB paylaşılan RAM) dağıtmak, ağırlıkları 4 bit tamsayılara sıkıştırmayı ve llama.cpp veya Apple MLX gibi sıfır ek yüklü bir C++ çalışma zamanı üzerinden yürütmeyi gerektirir.

[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)]

Adım 1: Dönüştürün & Kalibre Edin (Önem Matrisi)

Doğrudan 4 bit yuvarlama, akıl yürütme ve kodlama doğruluğunu düşürür. llama.cpp kullanarak önce Hugging Face ağırlıklarını sıkıştırılmamış bir GGUF dosyasına dönüştürün, ardından hangi ağırlık kanallarının aktivasyon üzerinde en büyük etkiye sahip olduğunu ölçmek için küçük bir kalibrasyon metni örneğinde Önem Matrisi (imatrix) hesaplayın:

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

Adım 2: Karma Bitli K-Nicemleme (Q4_K_M)

Hassas dikkat değer/çıktı katmanlarını 6 bitte (Q6_K) tutarken yoğun ileri beslemeli katmanları 256 ağırlıktan oluşan süper bloklarda 4 bite (Q4_K) sıkıştıran Q4_K_M (4-bit K-quant Medium) kullanarak modeli nicemleyin:

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

Bu işlem modeli, %1'den az perpleksite (perplexity) bozulmasıyla 16.0 GB'tan ~4.9 GB'a küçültür.

Adım 3: Cihaz Üzerinde Sıfır Kopyalı Birleşik Bellek Yürütmesi

iPhone, Android cihazı veya MacBook üzerinde çalıştırırken:

• mmap ile Anında Yükleme: Çalışma zamanı, 4.9 GB boyutlu .gguf dosyasını doğrudan flaş depolamadan Birleşik Belleğe bellek eşlemesiyle (mmap) aktarır; RAM çoğaltmasını ortadan kaldırır ve modeli 1 saniyeden kısa sürede yükler.

• Birleştirilmiş Ters Nicemleme Gölgelendiricileri: Token üretimi sırasında özel Metal (ggml-metal.metal) veya Vulkan (mul_mat_q4_K.comp) hesaplama gölgelendiricileri, paketlenmiş 4 bit ağırlık bloklarını doğrudan GPU yazmaçlarına yükler, bit kaydırmalarıyla açar ve anında FP16 aktivasyonlarıyla çarpar; böylece modern dizüstü GPU'larında 40–70+ token/sn, amiral gemisi cep telefonlarında ise 15–25 token/sn elde edilir.