GPU আর AI accelerator — হাজার core, SIMT আর tensor unit

১ · CPU বনাম GPU — একটি কাজ দ্রুত, নাকি অনেক কাজ একসাথে

পুরান ঢাকার একটি বিয়েবাড়ির রান্নার কথা ভাবুন। এক দিকে একজন নামী বাবুর্চি — তিনি একাই একটি জটিল পদ খুব দ্রুত রাঁধেন, কোন মশলা কখন লাগবে আগে থেকেই আন্দাজ করেন, হাতের কাছে সব উপকরণ গুছিয়ে রাখেন। অন্য দিকে পাঁচশো জন সহকারী — প্রত্যেকে সাধারণ কাজ পারেন: পেঁয়াজ কাটা, আলু ছোলা। একজন সহকারী বাবুর্চির চেয়ে অনেক ধীর, কিন্তু পাঁচশো জন একসাথে পাঁচশো কেজি পেঁয়াজ কাটলে মোট কাজের হার বাবুর্চির ধারেকাছেও কেউ যেতে পারবে না। আর কারো হাতে পেঁয়াজ ফুরালে তিনি বসে থাকেন না — পাশের ঝুড়ি থেকে আলু তুলে নেন।

CPU সেই বাবুর্চি। তার লক্ষ্য একটি thread-এর latency কমানো — অর্থাৎ একটি কাজ যত দ্রুত সম্ভব শেষ করা। এজন্য প্রতিটি core-এ বিশাল যন্ত্রপাতি: branch predictor (পর্ব ৩৭), out-of-order engine আর reorder buffer (পর্ব ৪০), আর মেমরির অপেক্ষা কমাতে বড় বড় cache (পর্ব ৪৪)। GPU সেই সহকারীর দল। তার লক্ষ্য throughput — প্রতি সেকেন্ডে মোট কত কাজ। প্রতিটি lane সরল, in-order, predictor নেই বললেই চলে; বাঁচানো জায়গায় বসানো হয় আরও অনেক ALU. আর মেমরির অপেক্ষা সে কমায় না, ঢেকে দেয় — একদল thread অপেক্ষায় থাকলে আরেক দল চলে।

২ · Graphics থেকে সাধারণ হিসাবে — ইতিহাস আর GPU-র execution model

Kernel, grid, block, thread

GPU-তে চালানো function-কে বলে kernel. CPU (যাকে বলা হয় host) একটি kernel চালু করার সময় বলে দেয় কতগুলো thread লাগবে, আর GPU (device) সেই kernel-এর একই কোড প্রতিটি thread-এ চালায় — প্রতিটি thread শুধু নিজের নম্বর দেখে ঠিক করে কোন ডেটায় কাজ করবে। Thread-গুলো দুই স্তরে সাজানো:

• Grid: একটি kernel-এর সব thread মিলে একটি grid. Grid অনেকগুলো block-এ ভাগ।
• Thread block (OpenCL-এ workgroup): একসাথে একটি SM-এ বসে এমন thread-এর দল, সাধারণত ১২৮–১০২৪টি। একই block-এর thread-রা দ্রুত on-chip shared memory ভাগ করে আর একটি barrier-এ (__syncthreads()) একে অপরের জন্য অপেক্ষা করতে পারে; আলাদা block-এর thread-রা পারে না।
• Thread: প্রতিটির নিজস্ব register আর নম্বর (threadIdx, blockIdx). Hardware এদের ৩২টি করে warp-এ দলবদ্ধ করে চালায় (অংশ ৩)।

সবচেয়ে সরল kernel — দুটি array যোগ। নিচের CUDA কোড চালানোর জন্য নয় (এই পাতায় GPU নেই), লাইন ধরে বোঝার জন্য:

// device code: one thread adds one element
__global__ void vadd(const float *a, const float *b, float *c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;   // global thread index
    if (i < n) c[i] = a[i] + b[i];                   // guard: last block overhangs
}

// host code (error checks omitted for brevity)
void add_on_gpu(const float *a, const float *b, float *c, int n) {
    size_t bytes = n * sizeof(float);
    float *da, *db, *dc;
    cudaMalloc(&da, bytes); cudaMalloc(&db, bytes); cudaMalloc(&dc, bytes);
    cudaMemcpy(da, a, bytes, cudaMemcpyHostToDevice);   // over PCIe
    cudaMemcpy(db, b, bytes, cudaMemcpyHostToDevice);
    int threads = 256;
    int blocks = (n + threads - 1) / threads;           // ceil(n / 256)
    vadd<<<blocks, threads>>>(da, db, dc, n);            // launch the grid
    cudaMemcpy(c, dc, bytes, cudaMemcpyDeviceToHost);   // waits for the kernel
    cudaFree(da); cudaFree(db); cudaFree(dc);
}

৩ · SIMT আর warp — একটি instruction, অনেক thread, আর branch divergence

প্রোগ্রামারের চোখে প্রতিটি thread স্বাধীন — নিজের PC, নিজের if. কিন্তু hardware-এ ৩২টি thread-এর একটি warp একটি instruction fetch-decode করে আর ৩২টি lane-এ একসাথে চালায়, প্রতিটি lane নিজের register নিয়ে। NVIDIA এই ধরনকে বলে SIMT (single instruction, multiple threads). এটি পর্ব ৪২-এর SIMD-এরই আত্মীয় — পার্থক্য হলো, SIMD-এ প্রোগ্রামার বা compiler নিজে lane আর mask সামলায়, আর SIMT-এ প্রতিটি lane-কে আলাদা thread হিসেবে লেখা হয়, mask সামলায় hardware. Warp-এর আকার নির্মাতাভেদে আলাদা: NVIDIA-র warp ৩২ thread; AMD-এর পুরোনো GCN নকশায় wavefront ৬৪, নতুন RDNA নকশায় ৩২ বা ৬৪ — প্রকাশিত বিবরণ অনুযায়ী।

সমস্যা বাধে শাখায়। Warp-এর কিছু thread if-এর ভেতরে যেতে চায়, কিছু else-এ — অথচ instruction একটিই। সমাধান: hardware একটি active mask রাখে (প্রতি lane-এ এক bit, পর্ব ৪২-এর vector mask-এর মতো)। প্রথমে if-পথ চলে শুধু সেই lane-গুলোতে, বাকিরা চুপ; তারপর else-পথ উল্টো mask-এ; শেষে দুই পথ যেখানে মেলে (reconvergence point), সেখানে আবার সবাই একসাথে। কোথায় মিলবে আর কোন পথ বাকি, তা মনে রাখতে একটি reconvergence stack — ক্লাসিক নকশায় প্রতিটি entry-তে (PC, mask, মেলার PC). একে বলে branch divergence. নিচে ৮ lane-এর একটি ছোট warp (বোঝার সুবিধায়; আসলে ৩২):

লুপেও একই ঘটনা: প্রতিটি thread-এর ঘোরার সংখ্যা আলাদা হলে যে thread আগে শেষ, তার lane বসে থাকে, যতক্ষণ না warp-এর সবচেয়ে ধীর thread শেষ করে। নিচের simulator ৮ lane-এর একটি warp-এ একটি if/else আর একটি লুপ (প্রতিটি thread-এর আলাদা trip count) চালায় — প্রতিটি issue-তে PC, instruction, mask, সক্রিয় lane আর stack-এর গভীরতা ছাপে। Run চাপুন:

// SIMT warp of 8 threads: active mask + reconvergence stack
const W = 8, trip = [2, 0, 1, 3, 0, 2, 1, 0];
const v = new Array(W).fill(0), n = new Array(W).fill(0);
const P = [ // br: lanes where cond is true jump to tgt; rpc = reconvergence pc
  { s: "v = t; n = trip[t]", f: t => { v[t] = t; n[t] = trip[t]; } },
  { s: "if (t % 2) goto 4", br: t => t % 2 === 1, tgt: 4, rpc: 6 },
  { s: "v = v * 10   (even)", f: t => { v[t] *= 10; } },
  { s: "goto 6", br: () => true, tgt: 6, rpc: 6 },
  { s: "v = v + 100  (odd)", f: t => { v[t] += 100; } },
  { s: "v = v - 1    (odd)", f: t => { v[t] -= 1; } },
  { s: "if (n == 0) goto 9", br: t => n[t] === 0, tgt: 9, rpc: 9 },
  { s: "v = v + 1; n = n - 1", f: t => { v[t] += 1; n[t] -= 1; } },
  { s: "goto 6", br: () => true, tgt: 6, rpc: 9 },
  { s: "out[t] = v", f: () => {} }];
const stack = [{ pc: 0, mask: (1 << W) - 1, rpc: -1 }];
const bits = m => [...Array(W)].map((_, t) => (m >> t) & 1).join("");
let issued = 0, busy = 0;
while (stack.length) {
  const top = stack[stack.length - 1];
  if (top.pc === top.rpc || top.pc >= P.length) { stack.pop(); continue; }
  const pc = top.pc, op = P[pc], lanes = [...Array(W).keys()].filter(t => (top.mask >> t) & 1);
  const act = lanes.length; issued++; busy += act;
  console.log(String(pc).padStart(2) + " " + op.s.padEnd(21) + bits(top.mask) + " " + act + "/" + W + " depth " + stack.length);
  if (!op.br) { lanes.forEach(op.f); top.pc++; continue; }
  let T = 0; lanes.forEach(t => { if (op.br(t)) T |= 1 << t; });
  const N = top.mask & ~T;
  if (!N) top.pc = op.tgt; else if (!T) top.pc++;
  else if (top.rpc === op.rpc) { top.mask = N; top.pc++; }   // taken lanes just wait at rpc
  else { top.pc = op.rpc;                                     // this entry = reconvergence point
    stack.push({ pc: op.tgt, mask: T, rpc: op.rpc }, { pc: pc + 1, mask: N, rpc: op.rpc }); }
}
console.log("out  = [" + v.join(", ") + "]");
console.log("SIMD efficiency = " + busy + " / (" + issued + " x " + W + ") = " + (100 * busy / (issued * W)).toFixed(1) + "%");

 0 v = t; n = trip[t]   11111111 8/8 depth 1
 1 if (t % 2) goto 4    11111111 8/8 depth 1
 2 v = v * 10   (even)  10101010 4/8 depth 3
 3 goto 6               10101010 4/8 depth 3
 4 v = v + 100  (odd)   01010101 4/8 depth 2
 5 v = v - 1    (odd)   01010101 4/8 depth 2
 6 if (n == 0) goto 9   11111111 8/8 depth 1
 7 v = v + 1; n = n - 1 10110110 5/8 depth 3
 8 goto 6               10110110 5/8 depth 3
 6 if (n == 0) goto 9   10110110 5/8 depth 3
 7 v = v + 1; n = n - 1 10010100 3/8 depth 3
 8 goto 6               10010100 3/8 depth 3
 6 if (n == 0) goto 9   10010100 3/8 depth 3
 7 v = v + 1; n = n - 1 00010000 1/8 depth 3
 8 goto 6               00010000 1/8 depth 3
 6 if (n == 0) goto 9   00010000 1/8 depth 3
 9 out[t] = v           11111111 8/8 depth 1
out  = [2, 100, 21, 105, 40, 106, 61, 106]
SIMD efficiency = 75 / (17 x 8) = 55.1%

৪ · Warp বদলে অপেক্ষা লুকানো — occupancy-র হিসাব

GPU-র global memory থেকে একটি load-এ কয়েকশো cycle লাগতে পারে। একটি warp load চালিয়ে ডেটার অপেক্ষায় থাকলে SM-এর warp scheduler পরের cycle-এই আরেকটি প্রস্তুত warp-এর instruction issue করে। এটি পর্ব ৪৩-এর fine-grained multithreading, কিন্তু context switch একেবারে বিনা খরচে — কারণ প্রতিটি জীবিত warp-এর register আগে থেকেই SM-এর বিশাল register file-এ বসানো, কিছু সরাতে হয় না। একে বলে latency hiding. সরল হিসাব: একটি warp প্রতি ৪টি হিসাবের instruction-এর পর একটি load দেয়, যার latency ৪০০ cycle; SM প্রতি cycle-এ একটি warp-instruction issue করলে, ৪০০ cycle ব্যস্ত রাখতে মোটামুটি ৪০০ / ৪ = ১০০টি warp-এর কাজ লাগবে — অথবা প্রতিটি warp-এর ভেতরে কয়েকটি স্বাধীন load, যাতে কম warp-এ চলে।

কিন্তু SM-এ একসাথে কতগুলো warp বসতে পারবে, তা সীমিত তিনটি জিনিসে: register file-এর মাপ, shared memory-র মাপ, আর thread/block-এর সর্বোচ্চ সংখ্যা। জীবিত warp / সর্বোচ্চ সম্ভব warp — এই অনুপাতের নাম occupancy. নিচের হিসাবগুলোতে SM-এর সীমা শেখার জন্য ধরে নেওয়া — ৬৫,৫৩৬টি 32-bit register, ৬৪ KB shared memory, সর্বোচ্চ ২,০৪৮ thread (৬৪টি warp) আর ৩২টি block. (অনেক বাস্তব GPU-র সংখ্যা এর কাছাকাছি, কিন্তু প্রজন্মভেদে আলাদা, আর hardware register বরাদ্দ করে নির্দিষ্ট খণ্ডে — তাই আসল হিসাব নির্মাতার occupancy calculator-এ মেলাতে হয়।)

নিচের calculator একই হিসাব কয়েকটি kernel-এর জন্য করে আর বলে কোন সম্পদ সীমা টানছে:

// Occupancy calculator. SM limits below are ASSUMED teaching values,
// not the datasheet of any real GPU (real GPUs also round allocations up).
const SM = { regs: 65536, smem: 64 * 1024, maxThreads: 2048, maxBlocks: 32, warp: 32 };
const maxWarps = SM.maxThreads / SM.warp;
function occupancy(threadsPerBlock, regsPerThread, smemPerBlock) {
  const byRegs = Math.floor(SM.regs / (threadsPerBlock * regsPerThread));
  const bySmem = smemPerBlock ? Math.floor(SM.smem / smemPerBlock) : Infinity;
  const byThreads = Math.floor(SM.maxThreads / threadsPerBlock);
  const limits = { registers: byRegs, "shared mem": bySmem, threads: byThreads, "block cap": SM.maxBlocks };
  const blocks = Math.min(...Object.values(limits));
  const who = Object.keys(limits).filter(k => limits[k] === blocks).join(" + ");
  const warps = blocks * threadsPerBlock / SM.warp;
  return { byRegs, bySmem, byThreads, blocks, warps, who, occ: (100 * warps / maxWarps).toFixed(1) };
}
console.log("max resident warps per SM = " + maxWarps);
console.log("block regs/thr smem(KB) | regs smem thr -> blocks warps  occ   limiter");
const cases = [[256, 64, 16], [256, 32, 16], [256, 32, 8], [256, 40, 0], [128, 128, 0], [1024, 32, 48]];
for (const [tpb, r, kb] of cases) {
  const o = occupancy(tpb, r, kb * 1024);
  const sm = o.bySmem === Infinity ? "-" : o.bySmem;
  console.log(String(tpb).padStart(5) + String(r).padStart(9) + String(kb).padStart(9) + "   |" +
    String(o.byRegs).padStart(5) + String(sm).padStart(5) + String(o.byThreads).padStart(4) + " ->" +
    String(o.blocks).padStart(6) + String(o.warps).padStart(6) + (o.occ + "%").padStart(7) + "  " + o.who);
}

max resident warps per SM = 64
block regs/thr smem(KB) | regs smem thr -> blocks warps  occ   limiter
  256       64       16   |    4    4   8 ->     4    32  50.0%  registers + shared mem
  256       32       16   |    8    4   8 ->     4    32  50.0%  shared mem
  256       32        8   |    8    8   8 ->     8    64 100.0%  registers + shared mem + threads
  256       40        0   |    6    -   8 ->     6    48  75.0%  registers
  128      128        0   |    4    -  16 ->     4    16  25.0%  registers
 1024       32       48   |    2    1   2 ->     1    32  50.0%  shared mem

৫ · GPU-র মেমরি — register, shared memory, coalescing আর bank conflict

GPU-র মেমরি হায়ারার্কি CPU-র মতোই স্তরে স্তরে (পর্ব ৪৪), কিন্তু ভাগাভাগি আলাদা:

Memory coalescing — পাশাপাশি thread, পাশাপাশি ঠিকানা

একটি warp-এর ৩২টি thread একটি load instruction একসাথে চালায়, তাই ৩২টি ঠিকানা একসাথে আসে। Memory system সেগুলো মিলিয়ে দেখে কয়টি আলাদা মেমরি-খণ্ড (segment) ছুঁতে হবে, আর প্রতিটি খণ্ডের জন্য একটি transaction পাঠায়। ঠিকানাগুলো পাশাপাশি হলে সব একটি বা দুটি খণ্ডে পড়ে — একে বলে memory coalescing. ছড়ানো হলে (strided access) প্রায় প্রতিটি thread-এর জন্য আলাদা transaction, আর আনা ডেটার বেশিরভাগ অপচয়। নিচের হিসাবে পাঠ্যবইয়ের মডেল ধরে নেওয়া: ৩২ thread, প্রতিটি 4-byte উপাদান, thread t পড়ে base + t × stride × 4, আর মেমরি ১২৮-byte খণ্ডে আসে (base 0 মানে খণ্ডের শুরুতে align করা):

নিচের calculator যেকোনো stride, উপাদানের মাপ আর base ঠিকানার জন্য খণ্ড গোনে — আর একই access ৩২-byte “sector”-এ গুনে দেখায় (নতুন অনেক GPU আরও সূক্ষ্ম ৩২-byte একক ব্যবহার করে বলে প্রকাশিত):

// Memory coalescing: which 128-byte segments does one warp touch?
// Assumption (textbook model): 32 threads, thread t reads base + t*stride*size
const SEG = 128, WARP = 32;
function warpLoad(stride, size, base) {
  const segs = new Set();
  for (let t = 0; t < WARP; t++) {
    const first = base + t * stride * size, last = first + size - 1;
    for (let s = Math.floor(first / SEG); s <= Math.floor(last / SEG); s++) segs.add(s);
  }
  const useful = WARP * size, moved = segs.size * SEG;
  return { n: segs.size, moved, eff: (100 * useful / moved).toFixed(1) + "%" };
}
const cases = [ // [stride in elements, element bytes, base address]
  [1, 4, 0], [2, 4, 0], [4, 4, 0], [32, 4, 0], [1, 4, 4], [1, 8, 0], [1, 1, 0]];
console.log("stride size base | segments  bytes moved  efficiency");
for (const [st, sz, b] of cases) {
  const r = warpLoad(st, sz, b);
  console.log(String(st).padStart(6) + String(sz).padStart(5) + String(b).padStart(5) + " |" +
    String(r.n).padStart(9) + String(r.moved).padStart(13) + r.eff.padStart(12));
}
// Same accesses counted in 32-byte sectors (finer granularity, newer GPUs)
const sectors = (stride, size, base) => {
  const s = new Set();
  for (let t = 0; t < WARP; t++) s.add(Math.floor((base + t * stride * size) / 32));
  return s.size;
};
console.log("32-byte sectors, 4-byte elements: stride 1 -> " + sectors(1, 4, 0) +
  ", stride 2 -> " + sectors(2, 4, 0) + ", stride 32 -> " + sectors(32, 4, 0));

stride size base | segments  bytes moved  efficiency
     1    4    0 |        1          128      100.0%
     2    4    0 |        2          256       50.0%
     4    4    0 |        4          512       25.0%
    32    4    0 |       32         4096        3.1%
     1    4    4 |        2          256       50.0%
     1    8    0 |        2          256      100.0%
     1    1    0 |        1          128       25.0%
32-byte sectors, 4-byte elements: stride 1 -> 4, stride 2 -> 8, stride 32 -> 32

Shared memory আর bank conflict

Shared memory একসাথে অনেক thread-কে ডেটা দিতে পারে কারণ সেটি কয়েকটি স্বাধীন bank-এ ভাগ — অনেক NVIDIA GPU-তে ৩২টি bank, প্রতিটি 4 byte চওড়া; পরপর 4-byte word পরপর bank-এ: bank = (ঠিকানা / 4) mod 32. একটি warp-এর ৩২টি thread ৩২টি আলাদা bank ছুঁলে এক cycle-এ সব। দুটি thread একই bank-এর আলাদা word চাইলে সেগুলো পরপর পরিবেশন হয় — একে বলে bank conflict. (একই word হলে সমস্যা নেই — broadcast হয়।) Word-এর stride 1-এ conflict নেই; stride 2-এ ২-way (দুই গুণ সময়); stride 32-এ সব thread একই bank-এ — ৩২-way. প্রচলিত সমাধান padding: ৩২ × ৩২-এর একটি tile [32][33] হিসেবে ঘোষণা করলে কলাম ধরে পড়ার stride হয় ৩৩, আর ৩৩ mod ৩২ = ১ — conflict উধাও।

Tiled matrix multiply — সব মিলিয়ে

n × n matrix গুণে C-এর প্রতিটি উপাদানের জন্য A-এর একটি সারি আর B-এর একটি কলাম লাগে — ২n বার global memory থেকে পড়া। কিন্তু একই সারি আর কলাম বহু thread-এর দরকার। পর্ব ৪৮-এর blocking-এর ধারণাই এখানে: একটি block A আর B-এর একটি করে ১৬ × ১৬ tile একবার shared memory-তে আনে, তারপর সব thread সেখান থেকে ১৬ বার করে পড়ে:

#define T 16                                   // tile = 16 x 16 = one block
__global__ void matmul(const float *A, const float *B, float *C, int n) {
    __shared__ float As[T][T], Bs[T][T];       // per-block scratchpad tiles
    int tx = threadIdx.x, ty = threadIdx.y;
    int row = blockIdx.y * T + ty;
    int col = blockIdx.x * T + tx;
    float sum = 0.0f;                          // lives in a register
    for (int k0 = 0; k0 < n; k0 += T) {        // walk tiles along k
        As[ty][tx] = A[row * n + (k0 + tx)];   // neighbours read neighbours
        Bs[ty][tx] = B[(k0 + ty) * n + col];   //   -> coalesced loads
        __syncthreads();                       // whole tile loaded?
        for (int k = 0; k < T; k++)
            sum += As[ty][k] * Bs[k][tx];      // 16 MACs from shared memory
        __syncthreads();                       // done before tile is reused
    }
    C[row * n + col] = sum;
}
// launch: dim3 block(T, T); dim3 grid(n / T, n / T);
//         matmul<<<grid, block>>>(dA, dB, dC, n);

৬ · CPU থেকে GPU-তে ডেটা পাঠানো — PCIe-র খরচ আর কখন offload লাভজনক

আলাদা graphics card-এর GPU-র নিজের মেমরি, CPU-র নিজের। ডেটা যায় PCIe bus দিয়ে (পর্ব ৫১), যার bandwidth GPU-র নিজের মেমরির তুলনায় অনেক কম — প্রায়ই দশ গুণ বা তারও বেশি কম। আর প্রতিটি kernel চালু করতেও কয়েক মাইক্রোসেকেন্ডের মতো একটি স্থির খরচ (launch overhead) আছে। তাই প্রশ্নটি সবসময়: হিসাবের লাভ কি পাঠানোর খরচ পুষিয়ে দেবে? নিচের সংখ্যাগুলো শুধু হিসাবের জন্য ধরে নেওয়া, কোনো নির্দিষ্ট পণ্যের নয়: PCIe কার্যকর ২৫ GB/s; CPU টানা ১০০ GFLOP/s আর মেমরি ৫০ GB/s; GPU টানা ১০ TFLOP/s আর মেমরি ১ TB/s.

এই অনুপাত — প্রতি byte আনা-নেওয়ায় কতটি FLOP — এর নাম arithmetic intensity. Vector add-এ ১ / ১২ FLOP/byte, n = ১০০০-এর matmul-এ ২ × ১০^৯ / (১.২ × ১০^৭) ≈ ১৬৭. পর্ব ৫৪-এর roofline model এই অনুপাত দিয়েই বলে একটি kernel মেমরি-সীমিত না হিসাব-সীমিত। Laptop আর ফোনের chip-এ, যেখানে CPU আর GPU একই মেমরি ভাগ করে (integrated GPU, অনেক SoC), PCIe-র এই খরচ নেই — তবে মেমরির মোট bandwidth-ও তুলনায় কম।

৭ · AI accelerator — tensor core, systolic array আর কম precision

একটি neural network-এর প্রতিটি স্তর আসলে একটি বড় matrix গুণ: input-এর সারি × weight-এর matrix, তারপর একটি সরল function. Model train করা আর তাকে ব্যবহার করে উত্তর বানানো (inference) — দুটিতেই সময়ের সিংহভাগ যায় ঘন matrix গুণ আর যোগে, অর্থাৎ কোটি কোটি MAC (multiply-accumulate). Matrix গুণের বিশেষ গুণ: n × n-এ ডেটা n², হিসাব n³ — প্রতিটি সংখ্যা n বার ব্যবহার হয়। যে hardware এই পুনর্ব্যবহার chip-এর ভেতরে ধরে রাখতে পারে, সে প্রতি byte-এ অনেক বেশি হিসাব করতে পারে।

Tensor core আর matrix unit

সাধারণ GPU lane প্রতি instruction-এ একটি FMA করে। Tensor core (NVIDIA-র Volta, ২০১৭ থেকে) একটি instruction-এ একটি ছোট matrix-খণ্ডের পুরো গুণ-যোগ করে — যেমন প্রথম প্রজন্মে ৪ × ৪ FP16 matrix-এর গুণ, FP32-এ জমা, প্রকাশিত বিবরণ অনুযায়ী। Warp-এর ৩২টি thread মিলে বড় একটি খণ্ড (যেমন ১৬ × ১৬) নেয়। Instruction fetch-decode আর register-এর পড়া-লেখার খরচ এক instruction-এ বহু MAC-এ ভাগ হয়ে যায় — পর্ব ৪২-এর SIMD-যুক্তির আরেক ধাপ। AMD-এর matrix core, Intel-এর AMX, ARM-এর SME — বিভিন্ন নির্মাতা একই ধারণার নিজস্ব রূপ এনেছে।

Systolic array — ডেটা হাঁটে, হিসাব হয় পথে

১৯৭৮ সালের দিকে H. T. Kung আর Charles Leiserson systolic array-এর ধারণা দেন, আর ১৯৮২ সালে Kung-এর “Why Systolic Architectures?” প্রবন্ধ একে পরিচিত করে। নামটি এসেছে হৃৎপিণ্ডের systole থেকে — প্রতিটি “স্পন্দনে” (clock) ডেটা এক ঘর এগোয়। একটি grid-এ সাজানো অনেকগুলো সরল processing element (PE), প্রতিটি শুধু গুণ-যোগ করে আর ডেটা প্রতিবেশীকে দিয়ে দেয়। একটি সংখ্যা মেমরি থেকে একবার ঢোকে, তারপর পথে বহু PE তাকে ব্যবহার করে — মেমরিতে ফেরত না গিয়ে।

এখানে একটি output-stationary নকশা: n × n PE, PE(i, j) জমায় C[i][j]. A-এর সারি i বাঁ দিক থেকে ঢোকে, i cycle দেরিতে; B-এর কলাম j ওপর থেকে নামে, j cycle দেরিতে — এই “তির্যক” (skewed) খাওয়ানোর ফলে PE(i, j)-তে a[i][k] আর b[k][j] একই cycle-এ পৌঁছায়, cycle t = i + j + k (0 থেকে গোনা)। শেষ MAC হয় t = ৩(n − ১)-এ, তাই মোট ৩n − ২ cycle. A = [[1,2,3],[4,5,6],[7,8,9]], B = [[9,8,7],[6,5,4],[3,2,1]]:

নিচের simulator আসল ডেটা-চলাচল মডেল করে — প্রতি cycle-এ a-register ডানে, b-register নিচে এক ঘর সরে, তারপর প্রতিটি PE গুণ-যোগ করে — আর শেষে সাধারণ matmul-এর সাথে মেলায়:

// Output-stationary systolic array: n x n PEs, A flows right, B flows down
const A = [[1, 2, 3], [4, 5, 6], [7, 8, 9]];
const B = [[9, 8, 7], [6, 5, 4], [3, 2, 1]];
const n = A.length, Z = () => A.map(r => r.map(() => 0));
let a = Z(), b = Z(), acc = Z(), cycles = 0, macs = 0;
for (let t = 0; t <= 3 * n - 3; t++) {          // last MAC happens at t = 3n - 3
  for (let i = 0; i < n; i++) for (let j = n - 1; j > 0; j--) a[i][j] = a[i][j - 1];
  for (let j = 0; j < n; j++) for (let i = n - 1; i > 0; i--) b[i][j] = b[i - 1][j];
  for (let i = 0; i < n; i++) {                   // skewed feed: row i starts i cycles late
    const k = t - i; a[i][0] = k >= 0 && k < n ? A[i][k] : 0;
  }
  for (let j = 0; j < n; j++) {
    const k = t - j; b[0][j] = k >= 0 && k < n ? B[k][j] : 0;
  }
  let busy = 0;
  for (let i = 0; i < n; i++) for (let j = 0; j < n; j++) {
    const k = t - i - j;                          // PE(i,j) sees a[i][k] and b[k][j] now
    if (k >= 0 && k < n) { acc[i][j] += a[i][j] * b[i][j]; busy++; }
  }
  cycles++; macs += busy;
  console.log("cycle " + cycles + " (" + busy + " PEs busy): " +
    acc.map(r => "[" + r.map(x => String(x).padStart(3)).join(" ") + "]").join(" "));
}
const ref = A.map((r, i) => B[0].map((_, j) => r.reduce((s, x, k) => s + x * B[k][j], 0)));
const ok = ref.every((r, i) => r.every((x, j) => x === acc[i][j]));
console.log("total cycles = " + cycles + " (3n - 2), MACs = " + macs + " (n^3)");
console.log("matches normal matmul: " + ok + ", utilisation = " +
  (100 * macs / (cycles * n * n)).toFixed(1) + "%");

cycle 1 (1 PEs busy): [  9   0   0] [  0   0   0] [  0   0   0]
cycle 2 (3 PEs busy): [ 21   8   0] [ 36   0   0] [  0   0   0]
cycle 3 (6 PEs busy): [ 30  18   7] [ 66  32   0] [ 63   0   0]
cycle 4 (7 PEs busy): [ 30  24  15] [ 84  57  28] [111  56   0]
cycle 5 (6 PEs busy): [ 30  24  18] [ 84  69  48] [138  96  49]
cycle 6 (3 PEs busy): [ 30  24  18] [ 84  69  54] [138 114  81]
cycle 7 (1 PEs busy): [ 30  24  18] [ 84  69  54] [138 114  90]
total cycles = 7 (3n - 2), MACs = 27 (n^3)
matches normal matmul: true, utilisation = 42.9%

কম precision — কম bit, বেশি গুণ

Neural network ছোট ছোট গোলমাল সহ্য করতে পারে, তাই 32-bit float সবসময় লাগে না। পর্ব ২৩-এ যে format-গুলো দেখেছেন: FP16 (৫ bit exponent, ১০ bit fraction), BF16 (FP32-এর সমান ৮ bit exponent, মাত্র ৭ bit fraction — পরিসর বড়, নির্ভুলতা কম), FP8 (E4M3 আর E5M2 রূপ), আর integer INT8, এমনকি ৪-bit. একটি 8-bit গুণক একটি 32-bit গুণকের চেয়ে অনেক ছোট — গুণকের জায়গা মোটামুটি bit-সংখ্যার বর্গের মতো বাড়ে (পর্ব ২১) — তাই একই silicon-এ অনেক বেশি MAC, আর প্রতিটি সংখ্যা আনতে কম byte. সাধারণ কৌশল mixed precision: গুণ কম precision-এ, যোগফল জমানো FP32-এ। Float model-কে INT8-এর মতো integer-এ রূপান্তর করাকে বলে quantization — নির্ভুলতা কতটা কমে, তা model আর কৌশলভেদে আলাদা, মেপে দেখতে হয়।

TOPS বনাম আসল কাজ

AI chip-এর বিজ্ঞাপনে প্রায়ই TOPS (প্রতি সেকেন্ডে ট্রিলিয়ন operation) লেখা থাকে। হিসাবটি সরল: MAC-এর সংখ্যা × ২ (গুণ আর যোগ আলাদা গোনা) × clock. একটি কাল্পনিক NPU: ৪,০৯৬টি INT8 MAC, ১ GHz → ৪,০৯৬ × ২ × ১০^৯ = ৮.১৯২ TOPS চূড়ান্ত (peak). এটি তখনই সম্ভব যখন প্রতিটি MAC প্রতিটি cycle-এ ব্যস্ত — বাস্তবে মেমরি থেকে ডেটা আসতে দেরি, ছোট বা অদ্ভুত আকারের স্তর, আর matrix ছাড়া অন্য কাজের জন্য ব্যবহার অনেক কম হতে পারে; ৩০% হলে ২.৪৬ TOPS. তাছাড়া কোন precision-এ (INT8 না INT4), আর sparsity (শূন্য বাদ দেওয়া) ধরে সংখ্যা দ্বিগুণ দেখানো হয়েছে কি না — এসব না জেনে দুটি chip-এর TOPS তুলনা অর্থহীন।

ফোনে আর সাম্প্রতিক laptop-এর SoC-তে আলাদা NPU (neural processing unit) থাকে — ছবি থেকে পেছনের দৃশ্য ঝাপসা করা, কথা থেকে লেখা, ছোট ভাষা-model — কম power-এ চালানোর জন্য। এগুলোর ভেতরের নকশা নির্মাতাভেদে আলাদা আর প্রায়ই পুরো প্রকাশিত নয়; সাধারণত ছোট systolic বা MAC-array, বড় on-chip SRAM আর INT8/FP16-এর মতো কম precision থাকে বলে জানা যায়।

AI inference-এর আসল ছাদ — মেমরির bandwidth

একটি বড় ভাষা-model একটি একটি করে token (শব্দাংশ) বানায়। একজন ব্যবহারকারীর জন্য (batch = ১) প্রতিটি token-এ model-এর প্রায় সব weight একবার মেমরি থেকে পড়তে হয়, আর প্রতিটি weight-এ মাত্র একটি গুণ-যোগ। তাই সময়ের নিচের সীমা হিসাব দিয়ে নয়, bytes দিয়ে:

এজন্যই AI accelerator-এ HBM-এর মতো দামি, চওড়া মেমরি (পর্ব ৪৯), আর quantization এত জনপ্রিয় — bit কমালে সরাসরি token/s বাড়ে। অনেক ব্যবহারকারীর অনুরোধ একসাথে (batch) চালালে একবার পড়া weight অনেক token-এ কাজে লাগে, arithmetic intensity বাড়ে — server-এ throughput বাড়ানোর মূল কৌশল এটাই। (বাস্তবে attention-এর জন্য জমানো তথ্য, KV cache, আরও bytes যোগ করে; এখানে হিসাব শুধু weight-এর।)

৮ · GPU, FPGA না ASIC — নমনীয়তা, energy আর খরচ

Energy-র দিকটি সবচেয়ে গভীর। আধুনিক chip-এ একটি গুণ-যোগের চেয়ে ডেটা সরানো অনেক দামি — বাইরের DRAM থেকে একটি সংখ্যা আনতে খরচ হয় একটি গুণের চেয়ে শত গুণ বা তারও বেশি energy, বহুল উদ্ধৃত হিসাব অনুযায়ী। তাই systolic array, shared memory-র tile আর কম precision — তিনটিই মূলত ডেটা-চলাচল কমানোর কৌশল। আর পর্ব ০২-এর power wall আর dark silicon মনে রাখলে বোঝা যায় কেন বিশেষায়িত accelerator বাড়ছে: সব transistor একসাথে চালানো যায় না, তাই নির্দিষ্ট কাজে দক্ষ unit বসিয়ে দরকারের সময় চালানোই লাভজনক। দামের দিক থেকে বড় AI accelerator আর তার HBM মেমরি খুব ব্যয়বহুল, data center-এর বিদ্যুৎ আর ঠান্ডা রাখার খরচও বড় — তাই performance per watt আর প্রতি টাকায় কাজ আজ গতির মতোই গুরুত্বপূর্ণ মাপকাঠি। পরের পর্বে power আর energy-র হিসাব বিস্তারিত।

৯ · AI-এর যুগে — GPU-র প্রশ্নে AI, আর যাচাইয়ের অভ্যাস

মজার ব্যাপার, যে GPU-তে AI চলে, তার সম্পর্কে AI-কে জিজ্ঞেস করলে প্রায়ই আত্মবিশ্বাসী ভুল আসে। এই বিষয়ে সাধারণ ভুলগুলো:

• নির্দিষ্ট GPU-র SM-সংখ্যা, register, shared memory, TOPS বা bandwidth বানিয়ে বলা, বা এক প্রজন্মের সংখ্যা আরেক প্রজন্মে বসানো।
• Divergence-এর সময়কে দুই পথের যোগফলের বদলে সর্বোচ্চটি ধরা; occupancy-তে shared memory-র সীমা ভুলে যাওয়া; coalescing-এ misalignment বা উপাদানের মাপ উপেক্ষা।
• CUDA kernel-এ if (i < n) পাহারা বা __syncthreads() বাদ দেওয়া, বা শাখার ভেতরে barrier বসানো (সব thread না পৌঁছালে আটকে যেতে পারে)।
• PCIe-তে ডেটা পাঠানোর খরচ বাদ দিয়ে “GPU-তে ১০০ গুণ দ্রুত” দাবি; AMD-র wavefront আর NVIDIA-র warp-এর আকার গুলিয়ে ফেলা।
• TOPS-কে আসল গতি ধরা, precision বা sparsity উল্লেখ না করা; token/s-এর হিসাবে মেমরির bandwidth না দেখা।

AI-এর দেওয়া GPU-হিসাব যাচাইয়ের চেকলিস্ট

• Divergence-এর দক্ষতা নিজে গুনেছি — দুই পথের instruction-সংখ্যা যোগ করে, সক্রিয় lane গুনে
• Coalescing-এর খণ্ড ঠিকানার সূত্র থেকে নিজে বা ছোট simulator-এ বের করেছি, alignment আর উপাদানের মাপসহ
• Occupancy তিনটি সীমা (register, shared memory, thread) দিয়ে হিসাব করে নির্মাতার occupancy calculator বা profiler-এর সাথে মিলিয়েছি
• GPU-র যেকোনো spec নির্মাতার নিজস্ব documentation বা whitepaper-এ মিলিয়েছি, AI-এর কথায় নয়
• “কত গুণ দ্রুত” দাবিতে ডেটা পাঠানো, launch খরচ আর CPU-র একটি ভালো রূপের সাথে তুলনা আছে কি না দেখেছি — আর সম্ভব হলে আসল GPU-তে মেপেছি