عملی GPU انجینئرنگ: Metal شیڈرز، پروفائلنگ اور ایج تعیناتی

عملی GPU انجینئرنگ میں مختلف پلیٹ فارموں کے لیے شیڈر پروگرامنگ، ہارڈویئر کاؤنٹرز کی تشخیص، اور ڈیوائس پر ایج تعیناتی شامل ہیں۔ ذیل میں ان تینوں طریقہ ہائے کار کا تفصیلی جائزہ ہے۔

1. مخصوص Apple Metal (MSL) ٹائل شدہ میٹرکس ضرب کا شیڈر

Apple Silicon (M1–M4 / A-series) پر کمپیوٹ کرنلز Metal Shading Language (MSL) میں لکھے جاتے ہیں اور سست گلوبل میموری ریڈز سے بچنے کے لیے Threadgroup Memory (CUDA Shared Memory کا Apple متبادل) استعمال کرتے ہیں۔

ذیل میں MSL میں ٹائل شدہ میٹرکس ضرب (C = A × B) کا شیڈر ہے، جو 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;
    }
}

• اہم ہارڈویئر میپنگ: [[thread_position_in_grid]] براہ راست CUDA کے blockIdx * blockDim + threadIdx کے مطابق ہے، جبکہ threadgroup_barrier(mem_flags::mem_threadgroup)، CUDA کے __syncthreads() کے مطابق ہے۔ پروڈکشن FP16 ورک لوڈز کے لیے Apple Silicon، simdgroup_float8x8 ہارڈویئر انٹرنسکس بھی فراہم کرتا ہے تاکہ Apple کے میٹرکس ضرب یونٹس کو براہ راست استعمال کیا جا سکے۔

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%۔

• وارپ رکنے کی وجہ (long_scoreboard): 68.5% (وارپس گلوبل HBM میموری ریڈز کے واپس آنے کے انتظار میں رکے ہوئے ہیں)۔

• تشخیص: کرنل Roofline چارٹ پر ڈھلوان میموری بینڈوڈتھ کی حد کے عین نیچے ہے۔ مزید یہ کہ l1tex__t_sectors کی جانچ سے غیر یکجا میموری ریڈز سامنے آتے ہیں۔ ایک ہی وارپ کے تھریڈز مسلسل قطار کے ایڈریس پڑھنے کے بجائے کالموں میں قدم بہ قدم نیچے جا رہے ہیں، جس سے لائی گئی ہر 128-byte HBM کیش لائن کا 75% ضائع ہو رہا ہے۔

مرحلہ 3: ہدفی ہارڈویئر اصلاحات لاگو کریں

1. میموری یکجائی درست کریں: یقینی بنائیں کہ متصل تھریڈز (threadIdx.x) مسلسل اندرونی بُعد (n_cols) میں متصل عناصر پڑھیں تاکہ 32 FP16 ریڈز ایک ہی 64-byte یکجا ٹرانزیکشن میں ضم ہو جائیں۔

2. ویکٹر شدہ لوڈز (float4 / tl.load block): ہر ہدایت میں 16 بٹس کے بجائے 128 بٹس (8 FP16 اقدار) لوڈ کریں، جس سے میموری ہدایات کی کل تعداد 8× کم ہو جاتی ہے۔

3. وارپ شفل ریڈکشنز: سست Shared Memory ریڈکشنز کی جگہ رجسٹر سے رجسٹر وارپ شفلز (CUDA میں __shfl_xor_sync یا Triton میں tl.sum) استعمال کریں تاکہ ہم وقت سازی کی رکاوٹوں کے بغیر قطار کا اوسط اور تغیر حساب کیا جا سکے۔

مرحلہ 4: Roofline پر تصدیق کریں

ncu دوبارہ چلانے سے تصدیق ہوتی ہے:

• گلوبل لوڈ سیکٹرز: 4× کم ہوتے ہیں (کیش لائن کا 100% استعمال)۔

• لانگ اسکوربورڈ اسٹالز: 68.5% سے کم ہو کر < 15% رہ جاتے ہیں۔

• وال کلاک تاخیر: 3.8× بہتر ہوتی ہے، جس سے مؤثر بینڈوڈتھ H100 کی طبعی HBM حد کے 85% سے اوپر پہنچ جاتی ہے۔

3. ایج تعیناتی کی پائپ لائن: موبائل/لیپ ٹاپ پر LLM کی کوانٹائزیشن اور اجرا

لیپ ٹاپ یا اسمارٹ فون (8 GB–16 GB مشترکہ RAM) پر 8B پیرامیٹرز والا LLM تعینات کرنے کے لیے ویٹس کو 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-Quantization (Q4_K_M)

Q4_K_M (4-bit K-quant Medium) استعمال کرکے ماڈل کی کوانٹائزیشن کریں، جو حساس اٹینشن ویلیو/آؤٹ پٹ تہوں کو 6-bit (Q6_K) پر رکھتا ہے، جبکہ ڈینس فیڈ فارورڈ تہوں کو 256 ویٹس کے سپر بلاکس میں 4-bit (Q4_K) تک کمپریس کرتا ہے:

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

اس سے ماڈل 16.0 GB سے گھٹ کر ~4.9 GB ہو جاتا ہے اور پرپلیکسیٹی میں خرابی 1% سے کم رہتی ہے۔

مرحلہ 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 ایکٹیویشنز سے ضرب دیتے ہیں۔ یوں جدید لیپ ٹاپ GPUs پر 40–70+ tokens/sec اور اعلیٰ درجے کے موبائل فونز پر 15–25 tokens/sec حاصل ہوتے ہیں۔