Coalesced memory access — Hinglish notes

6.2.8 · Hardware › GPU Architecture

Coalesced Memory Access Kya Hai?

Coalesced memory access tab hoti hai jab ek GPU warp ke threads consecutive memory addresses access karte hain jinhe minimum number of memory transactions se serve kiya ja sake. GPUs fixed-size transactions (ek fixed number of bytes ke bursts) mein memory fetch karti hain. Jab ek warp ke threads ka data kam se kam possible transactions mein aa jaye, toh access coalesced hai. Jab requests bahut saari transactions mein scatter ho jaayein, memory efficiency gir jaati hai.

GPUs Kyun Care Karti Hain: Hardware Reality

Memory Transaction Granularity (Architecture-Dependent!)

GPU memory controllers fixed-size transactions mein data fetch karte hain. NVIDIA GPUs par representative values:

Level	Fermi/Kepler/Pascal	Volta/Turing/Ampere+
L1 cache line (fill granularity)	128 bytes	128 bytes
L2 cache line	32 bytes (sector) within 128-byte line	32-byte sectors within 128-byte line
Global memory (DRAM) transaction	128-byte transactions (segmented into 32-byte parts)	32-byte sectors independently service kar sakta hai

Key architectural distinction:

Older GPUs (Fermi/Kepler/Pascal): L1 se ek global load poora 128-byte cache line pull karta hai.
Newer GPUs (Volta+): memory subsystem sectored hai, isliye yeh 32-byte sectors ko independently service kar sakta hai, partial accesses par kam bandwidth waste hoti hai.

Jab thread 0 address 0x1000 read karta hai, hardware sirf 4 bytes fetch nahi karta — woh poora transaction/sector fetch karta hai jo us address ko contain karta hai (jaise ki ek 32-byte sector 0x1000–0x101F, ya architecture aur cache configuration ke hisaab se ek full 128-byte line).

KYUN? DRAM chips rows aur columns mein organized hoti hain. Ek byte padhne ka cost almost ek poora burst padhne jaisa hi hota hai. Memory controller bursts fetch karne ke liye optimized hai, individual bytes ke liye nahi. Coalescing is hardware reality ko exploit karti hai exact burst size chahe jo bhi ho.

Warp Execution Model

Ek warp ke saare 32 threads same instruction simultaneously execute karte hain (SIMT). Jab woh instruction ek memory load ho, warp scheduler saare 32 addresses collect karta hai aur unhe cover karne ke liye minimum number of memory transactions issue karta hai.

Best case: Saare 32 addresses possible transactions ke smallest set mein aa jaayein → minimum transactions (e.g., 1 × 128-byte line, ya 4 × 32-byte sectors). Worst case: Saare 32 addresses scatter ho jaayein → up to 32 transactions.

Performance Impact Derive Karna

Cost ko quantify karte hain. Neeche di gayi numbers ke liye hum ek model architecture fix karte hain: ek Volta-style sectored memory subsystem jahan global-memory transaction (sector) size $T = 32$ bytes hai. (Fermi/Kepler/Pascal par, $T = 32$ ko 128-byte line se replace karo aur arithmetic redo karo — result ki shape identical hai.)

Assume karo:

Warp size: $W = 32$ threads
Memory transaction (sector) size: $T = 32$ bytes (256 bits)
Element size: $E = 4$ bytes (float32)
Memory bandwidth: $B = 900$ GB/s
Memory latency: $L = 300$ ns

Coalesced Access (Optimal Pattern)

Threads consecutive floats access karte hain: arr[threadIdx.x] jahan threadIdx.x = 0, 1, 2, ..., 31.

Step 1: Ek transaction mein kitne elements?

$Elements per transaction = \frac{T}{E} = \frac{32}{4} = 8 floats$

Step 2: Warp ke liye kitni transactions chahiye?

$Transactions = \frac{W}{Elements per transaction} = \frac{32}{8} = 4 transactions$

4 KYUN? 32-byte sectors ke saath, har transaction 8 consecutive floats fetch karta hai. 32 threads ko 32 floats chahiye, isliye $32 \div 8 = 4$ transactions. (128-byte-line architecture par, wohi 128 bytes 1 transaction hain — same data, coarser granularity.)

Step 3: Kitne bytes fetch hue?

$Bytes = 4 \times 32 = 128 bytes$

Step 4: Fetch karne mein kitna time?

Bandwidth-limited time (sustained transfers ke liye latency ignore karte hue):

$t_{coalesced} = \frac{128 bytes}{900 \times 1 0 ^{9} bytes/s} = 0.142 ns$

Lekin single accesses ke liye memory latency dominant hoti hai: $t_{total} \approx 300$ ns (latency) + $0.142$ ns (transfer) $\approx 300$ ns.

Strided Access (Non-Coalesced)

Threads itne bade stride se access karte hain ki har ek alag transaction mein land kare, jaise arr[threadIdx.x * 8] (stride = 32 bytes = ek sector) ya usse bada:

Thread 0 → address 0x000 (sector 0) Thread 1 → address 0x0020 (sector 1) Thread 2 → address 0x0040 (sector 2) ...

Har thread ka address ek alag transaction/sector mein hota hai.

Transactions needed: 32 (ek per thread, har ek alag 32-byte sector mein).

Bytes fetched: $32 \times 32 = 1024$ bytes (lekin actually sirf $32 \times 4 = 128$ bytes use hote hain).

$t_{strided} = 32 \times 300 ns = 9600 ns = 9.6 μ s$

Coalescing se Speedup:

$Speedup = \frac{9600}{300} = 32 \times$

KYUN? Non-coalesced access 32 separate memory transactions force karta hai, har ek ke saath full latency overhead. Tum zaroorat se 8× zyada data fetch karte ho aur 32× zyada latency cost pay karte ho.

Access Patterns: Derivation ke Saath Examples

Example 1: Sequential Access (Perfect Coalescing)

__global__ void sequential(float *data) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    float val = data[idx];  // Thread i reads data[i]
}

Address analysis:

Thread 0: data[0] → 0x000
Thread 1: data[1] → 0x0004 (4 bytes baad)
Thread 31: data[31] → 0x007C (thread 0 se 124 bytes)

Saare 32 addresses exactly 128 contiguous bytes span karte hain: $[0 x 0000, 0 x 007 F]$ .

Transactions (architecture-dependent, lekin hamesha minimum possible):

128-byte-line architecture (Fermi/Kepler/Pascal): yeh 128 bytes ek aligned cache line hain → 1 transaction.
32-byte-sector architecture (Volta+): 128 bytes = chaar aligned sectors → 4 transactions.

Dono cases mein, 128 useful bytes padhne ke liye yeh minimum hai, aur fetched bytes ka 100% use hota hai. Iska matlab yahi hai "coalesced" — ek fixed transaction count nahi.

KYUN kaam karta hai: Memory addresses stride 4 ke saath monotonically increase karte hain (float ka size). Controller saare 32 requests ko minimum transactions aur zero wasted bytes se cover karta hai.

Example 2: Strided Access (Worst Case)

__global__ void strided(float *data) {
    int idx = threadIdx.x * 1024;  // Stride = 1024 floats = 4096 bytes
    float val = data[idx];
}

Address analysis:

Thread 0: 0x0000
Thread 1: 0x1000 (4096 bytes door)
Har thread ka address ek alag transaction/cache line mein hai.

Transactions: 32 (ek per thread), kisi bhi architecture par, kyunki koi bhi do threads ek transaction share nahi karte.

Bytes fetched (32-byte-sector arch): $32 \times 32 = 1024$ bytes. Bytes used: $32 \times 4 = 128$ bytes. Waste: $1024 - 128 = 896$ bytes (sectored GPU par 87.5% waste; 128-byte-line GPU par toh aur bhi bura — 96.9% — kyunki use per thread ek full line pull karni padti hai).

KYUN fail hota hai: Stride transaction size se bada hai. Koi bhi do threads ek transaction share nahi karte, isliye har thread ek alag transaction trigger karta hai.

Example 3: Structure of Arrays vs. Array of Structures

Array of Structures (AoS) – Non-coalesced:

struct Particle {
    float x, y, z;  // 12 bytes
    float vx, vy, vz;  // 12 bytes
};
__global__ void aos(Particle *particles) {
    int idx = threadIdx.x;
    float x = particles[idx].x;  // Access x-coordinate
}

Address analysis (struct = 24 bytes):

Thread 0: particles[0].x → 0x0000
Thread 1: particles[1].x → 0x0018 (24 bytes baad)
Thread 31: particles[31].x → 0x02E8 (744 bytes)

32 accessed x-values $32 \times 24 = 768$ bytes ke address space mein spread hain, lekin har 24 bytes mein sirf 4 useful bytes read hoti hain.

Transactions (32-byte-sector arch): touched bytes $⌈ 768/32 ⌉ = 24$ sectors span karti hain, isliye ~24 transactions, jinmein zyada bytes waste hain.

Structure of Arrays (SoA) – Coalesced:

struct Particles {
    float *x;  // All x-coordinates together
    float *y, *z;
    float *vx, *vy, *vz;
};
__global__ void soa(Particles p) {
    int idx = threadIdx.x;
    float x = p.x[idx];  // Consecutive access
}

Address analysis: Example 1 (sequential) jaisa hi. Thread $i$ , p.x[i] read karta hai, stride = 4 bytes.

Transactions: 4 (32-byte sectors) ya 1 (128-byte line) — minimum, 100% efficiency.

KYUN SoA jeet ta hai: Fields ko alag karne se threads homogeneous data ko sequence mein access kar paate hain, aur har thread ke zaroorat ke data ke beech koi wasted bytes nahi hote.

Worked example Misaligned Access

__global__ void misaligned(float *data) {
   int idx = threadIdx.x;
   float val = data[idx + 1];  // Offset by 1
}

Kya hota hai:

Thread 0: data[1] → 0x0004
Thread 31: data[32] → 0x0080 (base se 128 bytes)

Addresses $[0 x 0004, 0 x 0083]$ span karti hain. Ek 128-byte-line architecture par yeh do aligned 128-byte lines (0x0000–0x007F aur 0x0080–0x00FF) cross karta hai → 1 ki jagah 2 transactions. Ek 32-byte-sector architecture par yeh ek extra sector touch karta hai → 4 ki jagah 5 sectors.

Extra transactions KYUN: Offset pattern ko shift kar deta hai taaki woh ek transaction boundary par straddle kare. Alignment matter karta hai: sirf ek ek-element offset bhi ek extra transaction add kar sakta hai.

Common Mistakes aur Steel-Manning

Common mistake Mistake 1: "Coalescing sirf global memory ke liye matter karti hai"

Kyun sahi lagta hai: Documentation global memory coalescing par emphasize karti hai kyunki wahan biggest wins hote hain. Shared memory on-chip aur fast hai, isliye students assume karte hain ki wahan access patterns matter nahi karte.

Reality: Shared memory mein memory banks hote hain (modern GPUs par typically 32 banks). Agar ek warp ke multiple threads same bank lekin alag addresses access karein, toh bank conflicts hote hain, jo access ko serialize kar dete hain. Shared memory ka sequential access aksar conflict-free patterns par map hota hai, lekin strided access (stride = bank count ka multiple) worst-case conflicts cause karta hai.

Fix: Locality principles har jagah apply hote hain. Shared memory ke liye, conflict-free access aim karo: stride-1 ya broadcast (sab threads same address access karein). Bank conflicts avoid karne ke liye arrays pad karo.

Common mistake Mistake 2: "Agar mera data naturally sequential nahi hai, toh main coalesce nahi kar sakta"

Kyun sahi lagta hai: Bahut saare algorithms mein inherent random access patterns hote hain (hash tables, trees, sparse matrices). Agar algorithm access order dictate kare, toh tum stuck ho, hai na?

Reality: Tum aksar computation reorder kar sakte ho ya data restructure kar sakte ho taaki coalescing enable ho:

Tiling: Kaam ko blocks mein tod do jahan har warp ek contiguous tile process kare.
Sort: Memory address ke hisaab se sort karo access karne se pehle (agar access pattern known ho).
SoA transformation: Jaisa upar dikhaya, data layout changes scattered accesses ko sequential mein badal sakti hain.
Texture cache: Truly random patterns ke liye, texture memory (2D spatial locality ke saath) raw global memory se behtar perform kar sakti hai.

Fix: Coalescing ek target hai, constraint nahi. Pehle profile karo. Agar memory bottleneck hai aur coalescing kharab hai, toh data layout changes ya algorithmic tweaks explore karo.

Common mistake Mistake 3: "Maine

__align__(128) use kiya isliye meri access coalesced hai" Kyun sahi lagta hai: Memory transactions ke liye alignment matter karti hai, isliye data ko cache-line boundaries par align karne se coalescing guarantee honi chahiye.

Reality: Alignment necessary hai lekin sufficient nahi. Coalescing ke liye chahiye:

Base address ki alignment (__align__ se solve hota hai)
Ek warp ke threads ka sequential access pattern (alignment se akele solve nahi hota) Agar threads abhi bhi large strides ya random offsets se access karen, toh alignment help nahi karta. Alignment boundary-crossing penalties prevent karta hai; coalescing ke liye sahi access pattern bhi chahiye.

Fix: Apne arrays ko transaction boundaries par align karo aur ensure karo ki ek warp ke threads consecutive elements access karein. Dono zaroori hain.

Active Recall Flashcards

#flashcards/hardware

GPU programming mein coalesced memory access kya hai?

Jab ek warp ke threads consecutive memory addresses access karte hain jo minimum number of memory transactions se serve ho sakein, memory controller ke burst-fetching behavior ko exploit karte hue.

GPU memory controller individual bytes ki jagah bursts mein data kyun fetch karta hai?

DRAM architecture mein ek row padhna almost utna hi sasta hai jitna ek single cell padhna. Memory controllers burst transfers ke liye optimized hain. Bursts fetch karne se fixed latency cost bahut saare bytes mein amortize ho jaati hai.

Kya GPU transaction/cache-line sizes saare NVIDIA GPUs par universal hain?

Nahi. Yeh compute capability par depend karti hain. Fermi/Kepler/Pascal full 128-byte L1 lines pull karte hain; Volta/Turing/Ampere 32-byte sectors use karte hain aur partial lines service kar sakte hain. Hamesha efficiency ratio ke saath reason karo, fixed count ke saath nahi.

Agar 32 threads itne bade stride se floats access karein ki har ek alag transaction mein land kare, toh kitni transactions chahiye hongi?

Up to 32 — ek per thread. Koi coalescing nahi hoti kyunki koi bhi do threads ek transaction share nahi karte.

GPU access patterns ke liye memory efficiency formula kya hai, aur yeh architecture-independent kyun hai?

Efficiency = (Bytes used) / (Bytes fetched) = (Warp size × Element size) / (Num transactions × Transaction size). Yeh ek ratio hai, isliye yeh absolute transaction size se independent bandwidth waste capture karta hai.

GPUs par Structure of Arrays (SoA) typically Array of Structures (AoS) se behtar kyun perform karta hai?

SoA saare elements ke liye same field ko contiguously group karta hai, isliye ek field access karne wale threads consecutive memory read karte hain 100% efficiency ke saath. AoS fields interleave karta hai, jo large strides aur wasted bytes per transaction force karta hai.

1-element offset (jaise data[threadIdx.x + 1]) coalescing ko kaise affect karta hai?

Yeh ek extra transaction add kar sakta hai. Offset access pattern ko shift kar deta hai taaki woh ek transaction boundary par straddle kare, warp ko ek extra line/sector touch karne par majboor kar ke.

Connections

6.2.01-GPU-execution-model – Warps lockstep mein execute karte hain, isliye saare 32 threads ke memory requests simultaneously collect hote hain. Coalescing SIMT model par depend karta hai.
6.2.04-Shared-memory-and-synchronization – Locality shared memory par bank conflicts ke zariye apply hoti hai. Strided shared memory access serialization cause karta hai.
6.2.09-Memory-hierarchy – Transaction aur cache-line sizes (architecture-dependent) coalescing granularity determine karti hain. L1 (128 bytes) aur L2 (128-byte lines ke andar 32-byte sectors) transaction counts affect karte hain.
6.3.02-Roofline-model – Coalescing directly achieved memory bandwidth ko impact karta hai. Non-coalesced code hardware roofline se bahut neeche operate karta hai.
5.1.05-Cache-coherence-protocols – GPUs traditional snooping coherence use nahi karte, lekin L1/L2 caches abhi bhi spatial locality exploit karte hain, jo coalescing maximize karta hai.

Recall Ek 12-Saal-Ke Bachche Ko Explain Karo

Socho tum aur tumhare 31 dost ek line mein library mein khade ho, aur tumhe sab ko books chahiye. Librarian ek badi cart le ja sakta hai jo ek baar mein kaafi books carry kar sake. Agar tum sab ek hi shelf par aas-paas rakhi books maango (book 1, book 2, book 3, ...), toh librarian cart ko chand quick grabs mein bhar leta hai aur sab ko unki book de deta hai. Aasaan!

Lekin agar tum mein se har koi library mein bilkul alag shelf se book maange (shelf A se book 1, shelf Z se book 2000, shelf Q se book 5000, ...), toh librarian ko har insaan ke liye alag trip karni padegi — 32 trips tak. Yeh bahut zyada time leta hai, chahe tum phir bhi ek hi book le rahe ho.

Coalesced memory access woh hai jab tumhare GPU ka "librarian" (memory controller) tumhare 32 threads ka saara zaroori data kam se kam trips mein fetch kar sake kyunki tumne memory mein ek doosre ke paas ki cheezein maangi theen. Jab tumhara code saara jagah scattered data maangta hai, GPU ko bahut saari trips karni padti hain, aur tumhara program slow ho jaata hai. Trick yeh hai ki apna data organize karo aur use order mein access karo, taaki GPU sabkuch kam se kam trips mein grab kar sake!

Generated: 2026-07-01 | Context: GPU memory performance optimization