Coalesced memory access: Exercises — Hinglish deep dive

6.2.8 · D4 · Hardware › GPU Architecture › Coalesced memory access

Shuru karne se pehle, ek shared model taaki har number reproducible ho. Yahi model architecture hai jo parent note ne fix kiya tha:

Definition Hamara fixed model (jab tak problem kuch aur na kahe)

Warp size $W = 32$ threads (sab same instruction same time pe execute karte hain, yahi SIMT rule hai 6.2.01-GPU-execution-model se).
Transaction (sector) size $T = 32$ bytes — sabse chhota chunk jo DRAM wapas deta hai.
Element size $E = 4$ bytes (ek float32).
Memory latency $L = 300$ ns per transaction; bandwidth $B = 900$ GB/s.

Do formulas jo hum baar baar use karenge. Words mein kya kehte hain: pehla distinct sectors count karta hai jo warp ko controller se fetch karwata hai; doosra fetched bytes ka wo fraction hai jo aap actually rakhte ho. $N_{transactions} = {distinct T -byte aligned sectors touched by any thread of the warp}$ $Efficiency = \frac{bytes used}{bytes fetched} = \frac{W \times E}{N _{transactions} \times T}$ $N_{transactions}$ ko literally padho: har thread ka byte address lo, use $T$ ke multiple se neeche round karo (yahi uska sector name hai), aur dekho kitne alag sector-names hain. Wahi count — threads ki sankhya nahi — wo hai jo memory controller ko fetch karna hi hoga.

Do-extremes figure padna

Picture kya dikhata hai. Left (coalesced) pe, aath black threads sab upar ek single red 32-byte sector ki taraf point kar rahe hain — red box wo ek object hai jo controller fetch karta hai, aur saare 8 useful floats usike andar hain, isliye koi byte waste nahi. Right (strided) pe, chaar black threads har ek apne alag red sector ki taraf point kar raha hai: chaar boxes, ek thread per box, toh har fetched sector ka zyada hissa phenk diya jaata hai. Ise apna mental test banao: coalesced = bahut saare arrows few red boxes pe converge karte hain; strided = ek arrow per red box.

L1 — Recognition

Problem 1.1

Ek warp float v = data[threadIdx.x]; run karta hai. Thread $i$ element $i$ read karta hai. Kya yeh access coalesced hai ya non-coalesced?

Recall Solution

Hum kya check karte hain: kya consecutive threads consecutive addresses hit karte hain? Thread $i$ data[i] read karta hai, jiska byte address $i \times E = 4 i$ hai. Toh thread 0 → byte 0, thread 1 → byte 4, …, thread 31 → byte 124. 32 addresses contiguous block $[0, 127]$ ko bina kisi gap ke fill karte hain. Har address ko $T = 32$ ke multiple se neeche round karo: wo sectors ${0, 32, 64, 96}$ name karte hain — exactly 4 distinct sectors, sab fully used. Definition se, efficiency $= 1$ → coalesced (perfectly).

Problem 1.2

Doosra warp float v = data[threadIdx.x * 1024]; run karta hai. Kya yeh coalesced hai?

Recall Solution

Thread $i$ data[1024 i] read karta hai, byte address $4096 i$ . Thread 0 → byte 0, thread 1 → byte 4096. 32 ke multiples se neeche round karne par 32 alag sector-names milte hain (0, 4096, 8192, …). $N_{transactions} = 32$ , ek per thread. Efficiency $= 128/ (32 \times 32) = 12.5%$ → non-coalesced (worst case).

L2 — Application

Problem 2.1

Coalesced pattern data[threadIdx.x] (32 floats) ke liye compute karo (a) elements per transaction, (b) number of transactions, (c) bytes fetched.

Recall Solution

(a) Ek sector $T = 32$ bytes ka hai aur $T / E = 32/4 = 8$ floats hold karta hai — yahi reason hai ek transaction 8 threads ke contiguous data cover karta hai. (b) Warp ko 32 contiguous floats chahiye; 8 per sector packed, $N = W \div (T / E) = 32/8 = 4$ sectors mein tile ho jaate hain (division valid hai kyunki run contiguous aur aligned shuru hota hai — misaligned case ke liye Problem 2.4 dekho). (c) Bytes fetched $= N \times T = 4 \times 32 = 128$ bytes. Un 128 bytes mein se har ek use hota hai, isliye kuch waste nahi.

Problem 2.2

Strided pattern data[threadIdx.x * 8] (byte stride $= 8 \times 4 = 32$ bytes $= T$ ) ke liye memory efficiency compute karo.

Recall Solution

Har thread kahan land karta hai: thread $i$ → byte $32 i$ . 32 ke multiples se neeche round karne par sector-names ${0, 32, 64, \dots}$ milte hain — 32 distinct names. Toh $N = 32$ transactions. Bytes used $= W \times E = 32 \times 4 = 128$ . Bytes fetched $= N \times T = 32 \times 32 = 1024$ . $Efficiency = \frac{128}{1024} = 0.125 = 12.5%.$ Har 8 fetched bytes mein se sirf 1 rakha jaata hai.

Problem 2.3

Latency $L = 300$ ns per transaction use karke, (a) coalesced warp (4 transactions ek saath issue) aur (b) strided warp (32 transactions har ek alag sector mein) ke liye fetch time estimate karo. Phir speedup do. Wo hardware mechanism batao jo coalesced transactions ko overlap karne deta hai.

Recall Solution

(a) Coalesced: 4 sectors ek single warp-load instruction ke one 128-byte line ke aligned quarters hain. Memory controller ke paas multiple outstanding-request slots (miss-status/handling registers) hote hain, isliye wo pehle ke return ka intezaar kiye bina chaaron sector fetches back-to-back launch karta hai. Unke 300 ns latencies isliye ek mein overlap ho jaate hain $\approx L = 300$ ns — alag alag add nahi hote: $t_{coal} \approx 300$ ns. Overlap tabhi kaam karta hai kyunki requests ek instruction se simultaneously issue hoti hain aur inhe hold karne ke liye enough outstanding slots hain. (b) Strided: 32 sectors, 32 unrelated locations mein. Worst case mein wo outstanding-slot budget se zyada ho jaate hain aur serialize ho jaate hain, har ek full latency pay karta hai: $t_{strided} = 32 \times 300 = 9600$ ns $= 9.6 μ s$ . Speedup $= 9600/300 = 32 \times$ .

Problem 2.4 — Misaligned base address

Same contiguous pattern data[threadIdx.x], lekin array ka pehla byte address 0x0010 (ek sector mein 16 bytes andar) pe hai, 32-byte boundary pe nahi. Ab kitne transactions hain, aur efficiency kya hai?

Recall Solution

32 floats bytes $[16, 143]$ occupy karte hain. Har address ko 32 ke multiple se neeche round karo: touched sector-names hain ${0, 32, 64, 96, 128}$ — 5 sectors, kyunki block ab dono ends pe boundary straddle karta hai (sector 0 ke upper half aur sector 128 ke lower half dono mein dip karta hai). $N = 5$ , bytes used $= 128$ , bytes fetched $= 5 \times 32 = 160$ . $Efficiency = \frac{128}{160} = 0.8 = 80%.$ Lesson: misalignment ek extra transaction cost karta hai aur efficiency 100% se neeche le jaata hai chahe access contiguous ho. Yahi reason hai CUDA allocators sector-aligned pointers return karte hain, aur cudaMalloc alignment guarantee karta hai — apna base address align karo aur naive division $N = W / (T / E)$ phir exact hoga.

L3 — Analysis

Problem 3.1

Ek kernel ek matrix read karta hai jo row-major A[row * N + col] mein stored hai. Do mappings compare karo:

Mapping X: threadIdx.x → col (threads ek row ke across sweep karte hain).
Mapping Y: threadIdx.x → row (threads ek column ke down sweep karte hain).

Kaun sa mapping coalesce karta hai, aur kyun?

Recall Solution

Row-major matlab element A[row * N + col] byte $(r o w \cdot N + co l) \cdot E$ pe hai. Neighbouring columns 4 bytes apart hain; neighbouring rows $N \cdot 4$ bytes apart hain.

Mapping X (thread $i$ → col = i, fixed row): addresses $4 i$ apart → contiguous → distinct sectors ${0, 32, 64, 96}$ → coalesced, 4 transactions.
Mapping Y (thread $i$ → row = i, fixed col): addresses $4 N i$ apart. Kisi bhi realistic $N$ ke liye yeh stride $≫ T$ → 32 distinct sector-names → non-coalesced, 32 transactions.

Rule: threadIdx.x ko us axis pe map karo jo memory mein contiguous ho (row-major mein last index).

Problem 3.2

Ek 128-byte-line GPU (Fermi/Kepler/Pascal) pe, coalesced 32-float read 1 transaction hai; ek 32-byte-sector GPU (Volta+) pe yeh 4 transactions hai. Dono same 128 bytes 100% efficiency se read karte hain. Strided data[threadIdx.x * 1024] pattern pe kaun sa architecture kam bytes waste karta hai, aur sectoring wahan kyun help karta hai?

Recall Solution

Strided, 128-byte-line GPU: har thread ek full 128-byte line kheechta hai jo barely use hoti hai. Fetched $= 32 \times 128 = 4096$ bytes, used $= 128$ . Efficiency $= 128/4096 = 3.125%$ (96.9% waste). Strided, 32-byte-sector GPU: har thread sirf apna 32-byte sector kheechta hai. Fetched $= 32 \times 32 = 1024$ , used $= 128$ . Efficiency $= 128/1024 = 12.5%$ (87.5% waste). Sectoring isliye jeetta hai kyunki yeh finer granularity pe fetch karta hai — ek scattered access sirf 32 bytes drag karta hai, 128 nahi. Sectors kahan rehte hain yeh jaanne ke liye 6.2.09-Memory-hierarchy dekho.

L4 — Synthesis

Problem 4.1 — Array-of-Structs vs Struct-of-Arrays

Ek particle struct P { float x, y, z, w; } (16 bytes) hai. Ek warp sirf 32 particles ka x field read karta hai.

AoS: P *p; float xi = p[threadIdx.x].x; — 32 x values 16 bytes apart hain.
SoA: float *px; float xi = px[threadIdx.x]; — 32 x values 4 bytes apart hain.

32-byte-sector model ke liye har ek ki efficiency compute karo. Kaun sa layout choose karoge?

Recall Solution

SoA: neighbouring x 4 bytes apart hain → contiguous 128-byte block → distinct sectors ${0, 32, 64, 96}$ → $N = 4$ . Efficiency $= (32 \times 4) / (4 \times 32) = 128/128 = 100%$ . AoS: thread $i$ ka x byte $16 i$ pe hai. Har ek ko 32 ke multiple se neeche round karo sector name ke liye: $thread i 012345 byte 16 i 01632486480 sector ⌊ 16 i /32 ⌋ \cdot 32 0032326464$ Pattern clear hai: do consecutive threads ek sector share karte hain (bytes $0$ – $15$ aur $16$ – $31$ dono sector 0 mein hain), isliye 32 threads $32/2 = 16$ distinct sectors name karte hain → $N = 16$ . Efficiency $= (32 \times 4) / (16 \times 32) = 128/512 = 25%$ . SoA choose karo — yeh ek strided field-access ko contiguous stream mein badal deta hai. Yeh canonical GPU data-layout transform hai.

Problem 4.2

4.1 continue karte hue, agar kernel 32 particles ke chaaron fields x, y, z, w read karta hai, toh kya AoS achanak theek ho jaata hai? Reason through karo.

Recall Solution

Chaaron fields read karne ka matlab warp p[0..31] ke saare $32 \times 16 = 512$ contiguous bytes touch karta hai. Yeh $512/32 = 16$ sectors hain, sab fully used: efficiency $= 512/ (16 \times 32) = 512/512 = 100%$ . Toh AoS tab efficient hota hai jab aap poora struct consume karo; yeh wasteful tab hota hai jab aap sirf ek field cherry-pick karo. Layout choice access pattern pe depend karti hai, abstractly AoS/SoA pe nahi.

L5 — Mastery

Problem 5.1 — Shared memory se Transpose

Ek naive matrix transpose B[col*N+row] = A[row*N+col] mein A ki read coalesced hai lekin B ki write strided hai (neighbouring threads $N \cdot 4$ bytes apart likhte hain). Explain karo kaise ek shared-memory tile (6.2.04-Shared-memory-and-synchronization se) write ko fix karta hai, aur global write ki efficiency pehle aur baad mein batao.

Recall Solution

Pehle: write B[col*N+row] ka stride $4 N$ per thread hai → 32 distinct sectors, ek per thread → efficiency $128/ (32 \times 32) = 12.5%$ . Fix: har block A ka ek tile shared-memory buffer tile[ ][ ] mein coalesced (row-contiguous) read karta hai, __syncthreads(), phir tile ko B mein wapas write karta hai — tile transposed read karte hue lekin global memory row-contiguous likhte hue. Kyunki transpose ab on-chip shared memory ke andar hota hai (jisme koi coalescing penalty nahi), dono global read aur global write contiguous ho jaate hain. Baad mein: global write contiguous hai → 4 sectors, sab used → efficiency $128/128 = 100%$ . Strided cost fast shared memory mein pay hoti hai, slow DRAM mein nahi.

Problem 5.2 — Latency-bound vs bandwidth-bound

Aap ek kernel ko 12.5% efficiency aur 32 transactions per warp pe measure karte ho. Aap refactor karke coalesced layout achieve karte ho — 4 transactions per warp, 100% efficiency. Agar har transaction 300 ns pe latency-bound hai aur kernel 1,000,000 warp-loads issue karta hai, toh pehle aur baad mein wall-clock memory time estimate karo, aur speedup do. Phir explain karo kab yeh latency model apply karna band ho jaata hai.

Recall Solution

Pehle (latency-bound, serialized): $32 transactions \times 300 ns \times 1 0^{6} = 9.6 \times 1 0^{9}$ ns $= 9.6$ s. Baad mein: 4 sectors ek latency mein pipeline ho jaate hain $\approx 300$ ns per warp: $300 ns \times 1 0^{6} = 3.0 \times 1 0^{8}$ ns $= 0.3$ s. Speedup $= 9.6/0.3 = 32 \times$ . Latency model kab apply karna band hota hai: do regimes tab milte hain jab data volume itna bada ho ki transfer time latency ko overtake kar le. Dono times compare karo: latency-bound time hai $N_{transactions} \times L$ ; bandwidth-bound time hai $(bytes fetched) / B$ . Access bandwidth-bound tab hota hai jab $\frac{bytes fetched}{B} ≳ N_{transactions} \times L \times (concurrency hiding factor) .$ Practice mein, bahut concurrent warps launch karne se scheduler 300 ns latency ko doosre warps ke kaam ke peeche hide kar leta hai; ek baar enough warps in-flight hon GPU $B$ saturate karta hai aur aapko latency nahi bandwidth se model karna chahiye. Coalescing dono regimes mein help karta hai: yeh $N_{transactions}$ kum karta hai (latency-bound) aur $B$ ka wo fraction badhata hai jo useful bytes pe spend hota hai (bandwidth-bound). Apna kernel 6.3.02-Roofline-model pe locate karo taaki pata chale kaun sa term dominate karta hai.

Recall Quick self-check clozes

32 threads ka ek warp jo data[threadIdx.x] (aligned) read karta hai use 32-byte-sector GPU pe 4 transactions chahiye. $N_{transactions}$ distinct 32-byte aligned sectors ki sankhya ke barabar hai jo kisi bhi thread ne touch kiye hoon ke barabar hota hai. Coalescing ko memory efficiency (bytes used / bytes fetched) se judge kiya jaata hai, absolute transaction count se nahi. Ek contiguous read jo misaligned base address se shuru ho use one extra transaction chahiye. Wo layout transform jo single-field reads ko coalesced banata hai wo hai Array-of-Structs → Struct-of-Arrays (SoA).

Ise Hinglish mein padho: 6.2.08 Coalesced memory access (Hinglish).