मराठी बातम्या
ताज्या
देशाची पुन्हा दिवाळखोरीकडे वाटचाल सुरू आहे, असा इशारा माजी मंत्री रजिथा सेनारत्ने यांनी दिला.माजी राष्ट्रपती थंबुटेगामा येथील गरजू कुटुंबासाठी घरासाठी निधी देतातअब्जावधी रुपयांची कॅन्सर उपचार यंत्रे निष्क्रिय असल्याने रुग्णांना अधिक विलंब होतोसूर्याला एकाच वर्षात प्रत्येकी 2 अब्ज भारतीय रुपये कमावणाऱ्या तीन चित्रपटांची शक्यतापाकिस्तानमधील तारबेला जलविद्युत प्रकल्पात रिटेनिंग वॉल कोसळल्याची माहिती आहेमंत्री म्हणतात की आंतरराष्ट्रीय प्रसारक श्रीलंकेच्या पर्यटनासाठी समर्थन देतातनोचियागामा हत्येत कथित मदत केल्याप्रकरणी पोलिसांनी एका व्यक्तीला अटक केलीसिंहराजामध्ये कथित वन्यजीव नमुने गोळा केल्याप्रकरणी चार चिनी नागरिकांना कोठडी
एआय आणि तंत्रज्ञान

व्यावहारिक GPU अभियांत्रिकी: Metal शेडर, प्रोफाइलिंग आणि एज डिप्लॉयमेंट

वायव्य अमेरिकेतील शहराच्या दिव्यांचे अवकाशातून दिसणारे उदाहरणात्मक दृश्य; ही AI कॉम्प्युटिंग प्रणाली नाही.
वायव्य अमेरिकेतील शहराच्या दिव्यांचे अवकाशातून दिसणारे उदाहरणात्मक दृश्य; ही AI कॉम्प्युटिंग प्रणाली नाही.

व्यावहारिक GPU अभियांत्रिकीमध्ये क्रॉस-प्लॅटफॉर्म शेडर प्रोग्रामिंग, हार्डवेअर काउंटर डायग्नोस्टिक्स आणि ऑन-डिव्हाइस एज डिप्लॉयमेंट यांचा समावेश होतो. खाली तिन्ही वर्कफ्लोचा सखोल आढावा दिला आहे.

1. कस्टम Apple Metal (MSL) टाइल्ड मॅट्रिक्स गुणाकार शेडर

Apple Silicon (M1–M4 / A-series) वर, कॉम्प्युट कर्नल Metal Shading Language (MSL) मध्ये लिहिले जातात आणि संथ ग्लोबल मेमरी रीड टाळण्यासाठी Threadgroup Memory (CUDA Shared Memory चे Apple समतुल्य) वापरतात.

खाली 16 × 16 threadgroup टाइल वापरणारे MSL मधील टाइल्ड Matrix Multiplication (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 टेन्सरवर Row-Wise 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 Stall Reason (long_scoreboard): 68.5% (warp ग्लोबल HBM मेमरी रीड परत येण्याची वाट पाहत थांबलेले आहेत).
  • निदान: कर्नल Roofline चार्टवर उतरत्या Memory-Bandwidth Roof च्या अगदी खाली आहे. शिवाय, l1tex__t_sectors तपासल्यास अनकोअलेस्ड मेमरी रीड दिसतात—एकाच warp मधील थ्रेड सलग पंक्तीचे पत्ते वाचण्याऐवजी स्तंभांमधून खाली स्ट्राइड करतात, ज्यामुळे आणलेल्या प्रत्येक 128-बाइट HBM कॅश लाइनपैकी 75% वाया जाते.

पायरी 3: लक्ष्यित हार्डवेअर सुधारणा लागू करा

1. मेमरी कोअलेसिंग दुरुस्त करा: लगतचे थ्रेड (threadIdx.x) सलग अंतर्गत परिमाणावर (n_cols) लगतचे घटक वाचतील याची खात्री करा, म्हणजे 32 FP16 रीड एकाच 64-बाइट कोअलेस्ड ट्रान्झॅक्शनमध्ये विलीन होतील.

2. व्हेक्टराइज्ड लोड (float4 / tl.load ब्लॉक): 16 बिटऐवजी प्रति सूचना 128 बिट (8 FP16 मूल्ये) लोड करा, ज्यामुळे एकूण मेमरी सूचनांची संख्या 8× कमी होते.

3. Warp-Shuffle रिडक्शन: संथ Shared Memory रिडक्शनऐवजी रजिस्टर-टू-रजिस्टर warp शफल (CUDA मध्ये __shfl_xor_sync किंवा Triton मध्ये tl.sum) वापरा, म्हणजे सिंक्रोनायझेशन बॅरियरशिवाय पंक्तीची सरासरी आणि व्हेरिअन्स काढता येतात.

पायरी 4: Roofline वर पडताळणी करा

ncu पुन्हा चालवल्यावर पुढील गोष्टी निश्चित होतात:

  • Global Load Sectors: 4× ने कमी होतात (100% कॅश-लाइन वापर).
  • Long Scoreboard Stalls: 68.5% वरून < 15% पर्यंत कमी होतात.
  • Wall-Clock Latency: 3.8× ने सुधारते, आणि प्रभावी बँडविड्थ H100 च्या भौतिक HBM मर्यादेच्या 85% च्या वर जाते.

3. एज डिप्लॉयमेंट पाइपलाइन: मोबाइल/लॅपटॉपवर LLM क्वांटाइज करणे आणि चालवणे

8B-पॅरामीटर LLM लॅपटॉप किंवा स्मार्टफोनवर (8 GB–16 GB शेअर्ड RAM) डिप्लॉय करण्यासाठी वेट 4-बिट इंटिजरमध्ये संकुचित करावे लागतात आणि 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-बिट राउंडिंग तर्क आणि कोडिंगची अचूकता कमी करते. 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) वापरून मॉडेल क्वांटाइज करा, जे संवेदनशील attention value/output लेयर 6-बिट (Q6_K) वर ठेवते, तर दाट फीड-फॉरवर्ड लेयर 256 वेटच्या सुपर-ब्लॉकमध्ये 4-बिट (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: उपकरणावर झिरो-कॉपी Unified Memory एक्झिक्युशन

iPhone, Android उपकरण किंवा MacBook वर चालवताना:

  • mmap द्वारे तात्काळ लोड: रनटाइम 4.9 GB .gguf फाइल फ्लॅश स्टोरेजमधून थेट Unified Memory मध्ये मेमरी-मॅप (mmap) करते—RAM ची द्विरुक्ती टाळते आणि मॉडेल 1 सेकंदापेक्षा कमी वेळात लोड करते.
  • Fused Dequantization Shaders: टोकन जनरेशन दरम्यान, कस्टम Metal (ggml-metal.metal) किंवा Vulkan (mul_mat_q4_K.comp) कॉम्प्युट शेडर पॅक केलेले 4-बिट वेट ब्लॉक थेट GPU रजिस्टरमध्ये लोड करतात, बिटवाइज शिफ्टने अनपॅक करतात आणि त्यांना FP16 अ‍ॅक्टिव्हेशनशी तत्काळ गुणतात—ज्यामुळे आधुनिक लॅपटॉप GPU वर 40–70+ टोकन/सेकंद आणि फ्लॅगशिप मोबाइल फोनवर 15–25 टोकन/सेकंद मिळतात.

अधिक बातम्या: एआय आणि तंत्रज्ञान