ප්‍රායෝගික GPU ඉංජිනේරුකරණය: Metal ෂේඩර්, පැතිකඩ විශ්ලේෂණය සහ එජ් යෙදවීම

ප්‍රායෝගික GPU ඉංජිනේරුකරණයට විවිධ වේදිකා සඳහා ෂේඩර් ක්‍රමලේඛනය, දෘඪාංග ගණක රෝග විනිශ්චය සහ උපාංගය මත එජ් යෙදවීම ඇතුළත් වේ. පහත දැක්වෙන්නේ මෙම කාර්ය ප්‍රවාහ තුනම පිළිබඳ සවිස්තර විමසුමකි.

1. අභිරුචි Apple Metal (MSL) ටයිල් කළ න්‍යාස ගුණන ෂේඩරය

Apple Silicon (M1–M4 / A-series) මත ගණන කර්නල් Metal Shading Language (MSL) තුළ ලියන අතර මන්දගාමී ගෝලීය මතක කියවීම් වළක්වා ගැනීමට Threadgroup Memory (CUDA Shared Memory සඳහා Apple හි සමාන ක්‍රමය) භාවිත කරයි.

16 × 16 ත්‍රෙඩ් කණ්ඩායම් ටයිල් භාවිත කරන, MSL තුළ ලියූ ටයිල් කළ න්‍යාස ගුණන (C = A × B) ෂේඩරයක් පහත දැක්වේ.

#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]] සෘජුවම CUDA හි blockIdx * blockDim + threadIdx වෙත අනුරූප වන අතර threadgroup_barrier(mem_flags::mem_threadgroup), CUDA හි __syncthreads() වෙත අනුරූප වේ. නිෂ්පාදන භාවිතයේ FP16 කාර්යභාර සඳහා Apple Silicon විසින් Apple හි න්‍යාස ගුණන ඒකක සෘජුව භාවිත කිරීමට simdgroup_float8x8 දෘඪාංග අන්තර්නිහිත ක්‍රියා ද සපයයි.

2. පියවරෙන් පියවර ncu Roofline ප්‍රශස්තකරණ මාර්ගෝපදේශය

ඔබ (4096, 4096) FP16 ටෙන්සරයක පේළි අනුව LayerNorm සඳහා සරල CUDA හෝ Triton කර්නලයක් ලියා ඇතැයි සිතන්න. එය 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%.

• Warp නැවතීමේ හේතුව (long_scoreboard): 68.5% (ගෝලීය HBM මතක කියවීම් ආපසු ලැබෙන තුරු warp නතර වී බලා සිටී).

• විනිශ්චය: කර්නලය Roofline ප්‍රස්තාරයේ ආනත මතක කලාප පළල සීමාවට හරියටම යටින් පිහිටයි. තවද l1tex__t_sectors පරීක්ෂා කිරීමෙන් ඒකාබද්ධ නොවූ මතක කියවීම් හෙළි වේ. එකම warp තුළ ඇති ත්‍රෙඩ් යාබද පේළි ලිපින කියවීම වෙනුවට තීරු දිගේ පියවර තබමින් ගමන් කරන බැවින් ලබාගන්නා සෑම 128-byte HBM හැඹිලි රේඛාවකම 75% නාස්ති වේ.

පියවර 3: ඉලක්කගත දෘඪාංග පිළියම් යොදන්න

1. මතක ඒකාබද්ධකරණය නිවැරදි කරන්න: යාබද ත්‍රෙඩ් (threadIdx.x) අඛණ්ඩ අභ්‍යන්තර මානය (n_cols) ඔස්සේ යාබද මූලද්‍රව්‍ය කියවන බව තහවුරු කරන්න. එවිට FP16 කියවීම් 32 ක් එක් 64-byte ඒකාබද්ධ ගනුදෙනුවකට එක් වේ.

2. දෛශිකගත පූරණ (float4 / tl.load block): එක් උපදෙසකට බිටු 16 වෙනුවට බිටු 128 (FP16 අගයන් 8) පූරණය කිරීමෙන් මුළු මතක උපදෙස් ගණන 8× කින් අඩු වේ.

3. Warp-Shuffle අඩුකිරීම්: සමමුහුර්තකරණ බාධක නොමැතිව පේළියේ මධ්‍යන්‍යය සහ විචලතාව ගණනය කිරීමට මන්දගාමී Shared Memory අඩුකිරීම් වෙනුවට රෙජිස්ටරයෙන් රෙජිස්ටරයට warp shuffle (CUDA හි __shfl_xor_sync හෝ Triton හි tl.sum) භාවිත කරන්න.

පියවර 4: Roofline මත තහවුරු කරන්න

ncu නැවත ධාවනය කිරීමෙන් තහවුරු වන්නේ:

• ගෝලීය පූරණ අංශ: 4× කින් අඩු වේ (හැඹිලි රේඛාවේ 100% භාවිතය).

• Long Scoreboard නැවතීම්: 68.5% සිට < 15% දක්වා අඩු වේ.

• සැබෑ ගතවන කාල ප්‍රමාදය: 3.8× කින් වැඩිදියුණු වී ඵලදායී කලාප පළල H100 හි භෞතික HBM සීමාවෙන් 85% ඉක්මවා ගෙන යයි.

3. එජ් යෙදවුම් පයිප්ලයිනය: ජංගම/ලැප්ටොප් මත LLM ක්වොන්ටීකරණය හා ධාවනය

පරාමිති 8B ක් ඇති LLM එකක් ලැප්ටොප් හෝ ස්මාර්ට්ෆෝනයක් මත (8 GB–16 GB හවුල් RAM) යෙදවීමට බර 4-bit පූර්ණ සංඛ්‍යාවලට සම්පීඩනය කළ යුතු අතර llama.cpp හෝ Apple MLX වැනි අතිරේක බර රහිත C++ ධාවන පරිසරයක් හරහා ක්‍රියාත්මක කළ යුතුය.

[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: පරිවර්තනය සහ ක්‍රමාංකනය (Importance Matrix)

සරල 4-bit වටකිරීම තර්කන හා කේතකරණ නිරවද්‍යතාව අඩු කරයි. llama.cpp භාවිතයෙන් පළමුව Hugging Face බර සම්පීඩනය නොකළ GGUF ගොනුවකට පරිවර්තනය කරන්න. ඉන්පසු සක්‍රියකරණයට වැඩිම බලපෑමක් ඇති බර නාලිකා මැනීමට කුඩා ක්‍රමාංකන පෙළ නියැදියක් මත Importance Matrix (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) භාවිතයෙන් මාදිලිය ක්වොන්ටීකරණය කරන්න. එය සංවේදී attention value/output ස්තර 6-bit (Q6_K) මට්ටමේ තබා ගනිමින් ඝන feed-forward ස්තර බර 256 බැගින් වූ සුපිරි බ්ලොක් තුළ 4-bit (Q4_K) දක්වා සම්පීඩනය කරයි:

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

මෙය perplexity පිරිහීම 1% ට අඩුවෙන් තබා ගනිමින් මාදිලිය 16.0 GB සිට ~4.9 GB දක්වා කුඩා කරයි.

පියවර 3: උපාංගය මත පිටපත් රහිත ඒකාබද්ධ මතක ක්‍රියාත්මක කිරීම

iPhone, Android උපාංගයක් හෝ MacBook මත ධාවනය කරන විට:

• mmap හරහා ක්ෂණික පූරණය: ධාවන පරිසරය 4.9 GB .gguf ගොනුව ෆ්ලෑෂ් ආචයනයෙන් සෘජුවම Unified Memory වෙත මතක අනුරූපණය (mmap) කරයි. එමඟින් RAM අනුපිටපත් ඉවත් කර තත්පර 1 ට අඩු කාලයකින් මාදිලිය පූරණය කරයි.

• ඒකාබද්ධ ප්‍රතික්වොන්ටීකරණ ෂේඩර්: ටෝකන ජනනය අතරතුර අභිරුචි Metal (ggml-metal.metal) හෝ Vulkan (mul_mat_q4_K.comp) ගණන ෂේඩර් ඇසුරුම් කළ 4-bit බර බ්ලොක් සෘජුවම GPU රෙජිස්ටර වෙත පූරණය කරයි, බිටු මාරු කිරීම් මඟින් ඒවා දිගහරියි, එවෙලේම FP16 සක්‍රියකරණ සමඟ ගුණ කරයි. එමඟින් නවීන ලැප්ටොප් GPU මත 40–70+ tokens/sec සහ ප්‍රමුඛ පෙළේ ජංගම දුරකථන මත 15–25 tokens/sec ලබාගනී.