Metal Shader, ncu Roofline மற்றும் Edge Deployment: நடைமுறை GPU பொறியியல்

நடைமுறை GPU பொறியியல் (GPU engineering) என்பது தளங்கள் கடந்த shader நிரலாக்கம், வன்பொருள் கணிப்பி (hardware counter) கண்டறிதல், சாதனத்திலேயே நிகழும் edge பயன்படுத்தல் (deployment) ஆகிய மூன்றையும் உள்ளடக்கியது. இந்த மூன்று பணிப்பாய்வுகளையும் (workflows) கீழே ஆழமாகப் பார்ப்போம்.

1. தனிப்பயன் Apple Metal (MSL) Tiled Matrix Multiplication Shader

Apple Silicon (M1–M4 / A-series) சில்லுகளில், compute kernel-கள் Metal Shading Language (MSL) இல் எழுதப்படுகின்றன. மெதுவான global memory வாசிப்புகளைத் தவிர்க்க, இவை Threadgroup Memory (CUDA Shared Memory க்கு Apple இன் சமமான அம்சம்) ஐப் பயன்படுத்துகின்றன.

கீழே 16 × 16 threadgroup tile-களைப் பயன்படுத்தி MSL இல் எழுதப்பட்ட tiled Matrix Multiplication (C = A × B) shader உள்ளது.

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

• முக்கிய வன்பொருள் பொருத்தம் (Key Hardware Mapping): [[thread_position_in_grid]] ஆனது CUDA இன் blockIdx * blockDim + threadIdx உடன் நேரடியாகப் பொருந்துகிறது; threadgroup_barrier(mem_flags::mem_threadgroup) ஆனது CUDA இன் __syncthreads() உடன் பொருந்துகிறது. உற்பத்திநிலை FP16 வேலைச்சுமைகளுக்கு, Apple இன் matrix-multiply அலகுகளை நேரடியாகப் பயன்படுத்த Apple Silicon ஆனது simdgroup_float8x8 hardware intrinsics ஐயும் வழங்குகிறது.

2. படிப்படியான ncu Roofline மேம்படுத்தல் விளக்கம் (Step-by-Step ncu Roofline Optimisation Walkthrough)

(4096, 4096) FP16 tensor மீது Row-Wise LayerNorm க்காக ஒரு எளிமையான (naive) CUDA அல்லது Triton kernel ஐ நீங்கள் எழுதுகிறீர்கள் என்று வைத்துக்கொள்ளுங்கள்; அது PyTorch இன் இயல்பான (native) செயலாக்கத்தை விடக் குறிப்பிடத்தக்க அளவு மெதுவாக ஓடுகிறது.

படி 1: வன்பொருள் கணிப்பிகளைப் பிடித்தல் (Step 1: Capture Hardware Counters)

அமைப்புக் குறியீட்டில் (setup code) profiling சுமை ஏற்படாமல் இருக்க, உங்கள் தனிப்பயன் kernel பெயரை மட்டும் இலக்காக வைத்து 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 இல் தடைக்கற்களைக் கண்டறிதல் (Step 2: Diagnose the Bottleneck on the Roofline)

வெளியீட்டு metrics ஐப் பாருங்கள்:

• Compute (SM) Throughput: உச்சத்தில் 11.4%.

• Memory (DRAM/HBM) Throughput: உச்சத்தில் 84.2%.

• Warp Stall Reason (long_scoreboard): 68.5% (global HBM memory வாசிப்புகள் திரும்பி வரும் வரை warp-கள் காத்திருக்கின்றன).

• கண்டறிதல் (Diagnosis): Roofline வரைபடத்தில், சரிவான Memory-Bandwidth Roof க்குக் கீழே kernel நேராக அமர்ந்துள்ளது. மேலும், l1tex__t_sectors ஐப் பார்க்கும்போது, ஒருங்கிணைக்கப்படாத (uncoalesced) memory வாசிப்புகள் தெரிகின்றன - ஒரே warp இல் உள்ள threads தொடர்ச்சியான வரிசை முகவரிகளை வாசிப்பதற்குப் பதிலாக நிரல்கள் (columns) வழியே தாவிச் செல்கின்றன; இதனால் பெறப்படும் ஒவ்வொரு 128-byte HBM cache line இலும் 75% வீணாகிறது.

படி 3: இலக்கு வைத்த வன்பொருள் திருத்தங்களைப் பயன்படுத்துதல் (Step 3: Apply Targeted Hardware Fixes)

1. Memory Coalescing ஐச் சரிசெய்தல்: அருகருகே உள்ள threads (threadIdx.x) தொடர்ச்சியான உள் பரிமாணத்தில் (n_cols) அருகருகே உள்ள உறுப்புகளை வாசிக்குமாறு செய்யுங்கள்; அப்போது 32 FP16 வாசிப்புகள் ஒரே 64-byte ஒருங்கிணைந்த பரிமாற்றமாக (coalesced transaction) இணைகின்றன.

2. Vectorised Loads (float4 / tl.load block): ஒவ்வொரு instruction இலும் 16 bits க்குப் பதிலாக 128 bits (8 FP16 மதிப்புகள்) ஐ ஏற்றுங்கள்; இதனால் மொத்த memory instruction எண்ணிக்கை 8× குறைகிறது.

3. Warp-Shuffle Reductions: மெதுவான Shared Memory reductions ஐ, register-இலிருந்து register க்குச் செல்லும் warp shuffles (CUDA இல் __shfl_xor_sync அல்லது Triton இல் tl.sum) கொண்டு மாற்றுங்கள்; இதனால் ஒத்திசைவுத் தடைகள் (synchronization barriers) இல்லாமல் வரிசையின் சராசரி (mean) மற்றும் மாறுபாடு (variance) கணிக்கப்படுகிறது.

படி 4: Roofline இல் சரிபார்த்தல் (Step 4: Verify on the Roofline)

ncu ஐ மீண்டும் இயக்கும்போது பின்வருவன உறுதியாகின்றன:

• Global Load Sectors: 4× குறைகிறது (cache-line பயன்பாடு 100%).

• Long Scoreboard Stalls: 68.5% இலிருந்து 15% க்குக் குறைவாகக் குறைகிறது.

• Wall-Clock Latency: 3.8× மேம்படுகிறது; இதனால் H100 இன் இயற்பியல் HBM வரம்பில் 85% க்கும் மேலான பயனுறு அலைவரிசை (effective bandwidth) கிடைக்கிறது.

3. Edge Deployment Pipeline: மொபைல்/மடிக்கணினியில் ஒரு LLM ஐ Quantize செய்து இயக்குதல்

8B-parameter LLM ஒன்றை மடிக்கணினி அல்லது திறன்பேசியில் (8 GB–16 GB பகிர்ந்த RAM) பயன்படுத்துவதற்கு, weights ஐ 4-bit integers ஆகச் சுருக்கி, llama.cpp அல்லது Apple MLX போன்ற zero-overhead C++ runtime மூலம் இயக்க வேண்டும்.

[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) (Step 1: Convert & Calibrate)

வெறும் 4-bit rounding ஆனது reasoning மற்றும் coding துல்லியத்தைக் குறைக்கிறது. llama.cpp ஐப் பயன்படுத்தி, முதலில் Hugging Face weights ஐச் சுருக்கப்படாத GGUF கோப்பாக மாற்றுங்கள்; பின்னர் எந்த weight channels க்கு activation தாக்கம் மிக அதிகம் என்பதை அளவிட, சிறிய அளவுத்திருத்த உரை மாதிரியில் (calibration text sample) ஒரு 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: கலப்பு-bit K-Quantization (Q4_K_M) (Step 2: Mixed-Bit K-Quantization)

Q4_K_M (4-bit K-quant Medium) ஐப் பயன்படுத்தி மாதிரியை quantize செய்யுங்கள்; இது உணர்திறன் மிக்க attention value/output அடுக்குகளை 6-bit (Q6_K) இல் வைத்திருக்கிறது, அடர்த்தியான feed-forward அடுக்குகளை 256 weights கொண்ட super-blocks இல் 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 ஆகச் சுருக்குகிறது; perplexity சிதைவு 1% க்கும் குறைவு.

படி 3: சாதனத்தில் Zero-Copy Unified Memory செயலாக்கம் (Step 3: Zero-Copy Unified Memory Execution on Device)

iPhone, Android சாதனம் அல்லது MacBook இல் இயக்கும்போது:

• mmap மூலம் உடனடி ஏற்றம் (Instant Load via mmap): runtime ஆனது 4.9 GB .gguf கோப்பை flash சேமிப்பகத்திலிருந்து நேரடியாக Unified Memory க்குள் memory-map (mmap) செய்கிறது; இதனால் RAM இரட்டிப்பாதல் தவிர்க்கப்படுகிறது, மாதிரி 1 வினாடிக்குள் ஏற்றப்படுகிறது.

• இணைக்கப்பட்ட Dequantization Shaders (Fused Dequantization Shaders): token உருவாக்கத்தின்போது, தனிப்பயன் Metal (ggml-metal.metal) அல்லது Vulkan (mul_mat_q4_K.comp) compute shaders ஆனவை பொதிந்த 4-bit weight blocks ஐ நேரடியாக GPU registers க்குள் ஏற்றி, bitwise shifts மூலம் பிரித்து, FP16 activations உடன் உடனுக்குடன் பெருக்குகின்றன - இதனால் நவீன மடிக்கணினி GPU களில் 40–70+ tokens/sec உம், முதன்மை (flagship) கைப்பேசிகளில் 15–25 tokens/sec உம் கிடைக்கின்றன.