Rekayasa GPU praktis: shader Metal, pemrofilan, dan penerapan di edge

Rekayasa GPU praktis mencakup pemrograman shader lintas platform, diagnostik penghitung perangkat keras, dan penerapan edge di perangkat. Berikut ulasan mendalam tentang ketiga alur kerja tersebut.

1. Shader Perkalian Matriks Berubin (Tiled) Apple Metal (MSL) Kustom

Pada Apple Silicon (M1–M4 / seri A), kernel komputasi ditulis dalam Metal Shading Language (MSL) dan memanfaatkan Threadgroup Memory (padanan Shared Memory CUDA di Apple) untuk menghindari pembacaan memori global yang lambat.

Berikut shader Perkalian Matriks berubin (C = A × B) dalam MSL menggunakan tile threadgroup 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;
    }
}

• Pemetaan Perangkat Keras Utama: [[thread_position_in_grid]] dipetakan langsung ke blockIdx * blockDim + threadIdx milik CUDA, sedangkan threadgroup_barrier(mem_flags::mem_threadgroup) dipetakan ke __syncthreads() milik CUDA. Untuk beban kerja FP16 produksi, Apple Silicon juga menyediakan intrinsik perangkat keras simdgroup_float8x8 untuk langsung memanfaatkan unit perkalian matriks Apple.

2. Panduan Langkah demi Langkah Optimasi Roofline dengan ncu

Misalkan Anda menulis kernel CUDA atau Triton yang naif untuk Row-Wise LayerNorm pada tensor FP16 (4096, 4096), dan kernel itu berjalan jauh lebih lambat daripada implementasi bawaan PyTorch.

Langkah 1: Tangkap Penghitung Perangkat Keras

Jalankan NVIDIA Nsight Compute (ncu) dengan menargetkan hanya nama kernel kustom Anda agar tidak ada overhead pemrofilan pada kode penyiapan:

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

Langkah 2: Diagnosis Hambatan pada Roofline

Periksa metrik keluaran:

• Throughput Komputasi (SM): 11,4% dari puncak.

• Throughput Memori (DRAM/HBM): 84,2% dari puncak.

• Alasan Warp Stall (long_scoreboard): 68,5% (warp tertahan menunggu pembacaan memori HBM global selesai).

• Diagnosis: Kernel berada tepat di bawah Memory-Bandwidth Roof yang miring pada grafik Roofline. Selain itu, pemeriksaan l1tex__t_sectors menunjukkan pembacaan memori yang tidak coalesced—thread dalam warp yang sama melangkah menuruni kolom alih-alih membaca alamat baris yang berurutan, sehingga membuang 75% dari setiap cache line HBM 128 byte yang diambil.

Langkah 3: Terapkan Perbaikan Perangkat Keras yang Terarah

1. Perbaiki Memory Coalescing: Pastikan thread yang bersebelahan (threadIdx.x) membaca elemen yang bersebelahan sepanjang dimensi dalam yang berurutan (n_cols) sehingga 32 pembacaan FP16 menyatu menjadi satu transaksi coalesced 64 byte.

2. Vectorised Loads (float4 / tl.load block): Muat 128 bit (8 nilai FP16) per instruksi, bukan 16 bit, sehingga jumlah total instruksi memori berkurang 8×.

3. Reduksi Warp-Shuffle: Ganti reduksi Shared Memory yang lambat dengan warp shuffle antar register (__shfl_xor_sync di CUDA atau tl.sum di Triton) untuk menghitung rata-rata dan varians baris tanpa barrier sinkronisasi.

Langkah 4: Verifikasi pada Roofline

Menjalankan ulang ncu mengonfirmasi:

• Global Load Sectors: Turun 4× (pemanfaatan cache line 100%).

• Long Scoreboard Stalls: Turun dari 68,5% menjadi < 15%.

• Latensi Wall-Clock: Membaik 3,8×, mendorong bandwidth efektif di atas 85% dari batas HBM fisik H100.

3. Pipeline Penerapan Edge: Kuantisasi & Menjalankan LLM di Seluler/Laptop

Menerapkan LLM 8B parameter ke laptop atau ponsel pintar (RAM bersama 8 GB–16 GB) mengharuskan kompresi bobot menjadi integer 4-bit dan eksekusi melalui runtime C++ tanpa overhead seperti llama.cpp atau 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)]

Langkah 1: Konversi & Kalibrasi (Importance Matrix)

Pembulatan 4-bit mentah menurunkan akurasi penalaran dan pengodean. Dengan llama.cpp, pertama konversikan bobot Hugging Face ke berkas GGUF tanpa kompresi, lalu hitung Importance Matrix (imatrix) pada sampel teks kalibrasi kecil untuk mengukur kanal bobot mana yang memiliki dampak aktivasi terbesar:

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

Langkah 2: K-Quantization Campuran-Bit (Q4_K_M)

Kuantisasi model menggunakan Q4_K_M (4-bit K-quant Medium), yang mempertahankan layer value/output attention yang sensitif pada 6-bit (Q6_K) sambil mengompresi layer feed-forward padat menjadi 4-bit (Q4_K) dalam super-block berisi 256 bobot:

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

Ini mengecilkan model dari 16,0 GB menjadi ~4,9 GB dengan degradasi perplexity kurang dari 1%.

Langkah 3: Eksekusi Unified Memory Zero-Copy di Perangkat

Saat berjalan di iPhone, perangkat Android, atau MacBook:

• Pemuatan Instan melalui mmap: Runtime memetakan memori (mmap) berkas .gguf 4,9 GB langsung dari penyimpanan flash ke Unified Memory—menghilangkan duplikasi RAM dan memuat model dalam waktu kurang dari 1 detik.

• Shader Dekuantisasi Terpadu: Selama pembuatan token, compute shader kustom Metal (ggml-metal.metal) atau Vulkan (mul_mat_q4_K.comp) memuat blok bobot 4-bit yang dikemas langsung ke register GPU, membukanya dengan bitwise shift, dan mengalikannya dengan aktivasi FP16 secara langsung—mencapai 40–70+ token/detik pada GPU laptop modern dan 15–25 token/detik pada ponsel andalan seluler.