Uhandisi wa GPU kwa vitendo: shader za Metal, uchanganuzi wa utendaji na usambazaji kwenye vifaa vya pembeni

Uhandisi wa GPU kwa vitendo unahusisha uandishi wa shader kwa majukwaa tofauti, uchunguzi kwa vihesabu vya maunzi na usambazaji kwenye vifaa vya pembeni. Hapa chini kuna uchambuzi wa kina wa mitiririko yote mitatu ya kazi.

1. Shader Maalumu ya Apple Metal (MSL) ya Kuzidisha Matriki kwa Vigae

Kwenye Apple Silicon (M1–M4 / A-series), kerneli za ukokotoaji huandikwa katika Metal Shading Language (MSL) na kutumia Threadgroup Memory (sawa ya Apple ya CUDA Shared Memory) ili kuepuka usomaji wa polepole wa kumbukumbu ya jumla.

Hapa chini kuna shader ya Kuzidisha Matriki kwa vigae (C = A × B) katika MSL kwa kutumia vigae vya threadgroup vya 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;
    }
}

• Ulinganishaji Muhimu wa Maunzi: [[thread_position_in_grid]] hulingana moja kwa moja na blockIdx * blockDim + threadIdx ya CUDA, huku threadgroup_barrier(mem_flags::mem_threadgroup) ikilingana na __syncthreads() ya CUDA. Kwa mizigo ya kazi ya FP16 katika uzalishaji, Apple Silicon pia hutoa maagizo asilia ya maunzi ya simdgroup_float8x8 ili kutumia moja kwa moja vitengo vya Apple vya kuzidisha matriki.

2. Mwongozo wa Hatua kwa Hatua wa Uboreshaji wa Roofline kwa ncu

Tuseme unaandika kerneli rahisi ya CUDA au Triton kwa Row-Wise LayerNorm kwenye tensori ya FP16 ya (4096, 4096), na inaendeshwa polepole zaidi kwa kiasi kinachoonekana kuliko utekelezaji asilia wa PyTorch.

Hatua 1: Nasa Vihesabu vya Maunzi

Endesha NVIDIA Nsight Compute (ncu) ukilenga jina la kerneli yako maalumu pekee ili kuepuka mzigo wa ziada wa uchanganuzi wa utendaji kwenye msimbo wa maandalizi:

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

Hatua 2: Tambua Kizuizi kwenye Roofline

Kagua vipimo vya matokeo:

• Kiwango cha Ukokotoaji (SM): 11.4% ya kilele.

• Kiwango cha Uhamishaji wa Kumbukumbu (DRAM/HBM): 84.2% ya kilele.

• Sababu ya Kusimama kwa Warp (long_scoreboard): 68.5% (warp zimesimama zikisubiri usomaji wa kumbukumbu ya jumla ya HBM urudi).

• Utambuzi: Kerneli iko moja kwa moja chini ya paa lenye mteremko la kipimo-data cha kumbukumbu kwenye chati ya Roofline. Zaidi ya hayo, kukagua l1tex__t_sectors kunaonyesha usomaji wa kumbukumbu ambao haujaunganishwa: thread katika warp moja zinaruka hatua zikishuka kwenye safu wima badala ya kusoma anwani mfululizo za safu mlalo, zikipoteza 75% ya kila mstari wa kache wa HBM wa baiti 128 unaoletwa.

Hatua 3: Tumia Marekebisho Yanayolengwa ya Maunzi

1. Rekebisha Uunganishaji wa Ufikiaji wa Kumbukumbu: Hakikisha thread jirani (threadIdx.x) zinasoma elementi jirani kwenye kipimo cha ndani kinachoendelea (n_cols) ili usomaji 32 wa FP16 uungane kuwa muamala mmoja uliounganishwa wa baiti 64.

2. Upakiaji wa Vekta (float4 / bloku ya tl.load): Pakia biti 128 (thamani 8 za FP16) kwa kila agizo badala ya biti 16, ukipunguza jumla ya maagizo ya kumbukumbu kwa mara 8×.

3. Ukusanyaji wa Warp-Shuffle: Badilisha ukusanyaji wa polepole wa Shared Memory kwa ubadilishanaji wa warp kutoka rejista hadi rejista (__shfl_xor_sync katika CUDA au tl.sum katika Triton) ili kukokotoa wastani na variansi ya safu mlalo bila vizuizi vya usawazishaji.

Hatua 4: Thibitisha kwenye Roofline

Kuendesha ncu tena kunathibitisha:

• Sekta za Upakiaji wa Jumla: Hupungua kwa mara 4× (matumizi ya 100% ya mstari wa kache).

• Kusimama kwa Long Scoreboard: Hupungua kutoka 68.5% hadi < 15%.

• Muda Halisi wa Kusubiri: Huboreka kwa mara 3.8×, na kusukuma kipimo-data halisi zaidi ya 85% ya kikomo halisi cha HBM cha H100.

3. Mlolongo wa Usambazaji kwenye Vifaa vya Pembeni: Kukwantisha & Kuendesha LLM kwenye Simu/Kompyuta Mpakato

Kusambaza LLM yenye vigezo 8B kwenye kompyuta mpakato au simu mahiri (RAM inayoshirikiwa ya 8 GB–16 GB) kunahitaji kubana uzani kuwa namba kamili za biti 4 na kutekeleza kupitia mazingira ya C++ yasiyo na mzigo wa ziada kama llama.cpp au 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)]

Hatua 1: Badilisha & Rekebisha kwa Ulinganifu (Matriki ya Umuhimu)

Kukadiria moja kwa moja hadi biti 4 hupunguza usahihi wa kufikiri na kuandika msimbo. Kwa kutumia llama.cpp, kwanza badilisha uzani wa Hugging Face kuwa faili ya GGUF isiyobanwa, kisha kokotoa Matriki ya Umuhimu (imatrix) kwenye sampuli ndogo ya maandishi ya urekebishaji ili kupima ni chaneli zipi za uzani zenye athari kubwa zaidi kwenye uamilishaji:

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

Hatua 2: K-Quantization ya Biti Mchanganyiko (Q4_K_M)

Kwantisha modeli kwa kutumia Q4_K_M (4-bit K-quant Medium), ambayo huweka tabaka nyeti za thamani/matokeo za attention kwenye biti 6 (Q6_K) huku ikibana tabaka zenye msongamano za feed-forward kuwa biti 4 (Q4_K) katika super-block zenye uzani 256:

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

Hii hupunguza modeli kutoka 16.0 GB hadi ~4.9 GB huku uharibifu wa perplexity ukiwa chini ya 1%.

Hatua 3: Utekelezaji wa Kumbukumbu Iliyounganishwa Bila Kunakili kwenye Kifaa

Unapoendesha kwenye iPhone, kifaa cha Android au MacBook:

• Upakiaji wa Papo Hapo kupitia mmap: Mazingira ya utekelezaji huunganisha faili ya .gguf ya 4.9 GB kwenye kumbukumbu (mmap) moja kwa moja kutoka hifadhi ya flash hadi Kumbukumbu Iliyounganishwa, yakiondoa urudufu wa RAM na kupakia modeli kwa chini ya sekunde 1.

• Shader Zilizounganishwa za Kuondoa Ukwantishaji: Wakati wa uzalishaji wa tokeni, shader maalumu za ukokotoaji za Metal (ggml-metal.metal) au Vulkan (mul_mat_q4_K.comp) hupakia bloku zilizofungashwa za uzani wa biti 4 moja kwa moja kwenye rejista za GPU, huzifungua kwa kutumia uhamishaji wa biti, na kuzizidisha kwa uamilishaji wa FP16 papo hapo, zikifikia tokeni 40–70+ kwa sekunde kwenye GPU za kisasa za kompyuta mpakato na tokeni 15–25 kwa sekunde kwenye simu za kiwango cha juu.