Tensor cores and matrix operations: Exercises — Hinglish deep dive

6.2.12 · D4 · Hardware › GPU Architecture › Tensor cores and matrix operations

Intuition Is page ko kaise use karein

Ye Tensor cores and matrix operations ke liye graded drills hain. Har problem pe ek level tag hai (L1 → L5). Pehle solution fold karke try karo, phir unfold karo. Numbers aise choose kiye hain ki tum inhe haath se check kar sako. Agar kisi step mein koi aisa word aata hai jo tum abhi nahi jaante, to woh pehli baar aane par define kar diya jayega.

Shuru karne se pehle, teen plain-word reminders jo hum baar baar use karenge:

Matrix bas ek numbers ki grid hai jisme rows (aadhe mein jaati hain) aur columns (neeche jaate hain) hote hain. Hum iska size rows × columns likhte hain.
FMA ("fused multiply-add") ek aisi operation hai jo a·b + c karti hai — do numbers multiply karo, phir teesra add karo — ek hi step mein. Yeh matrix multiply ka atom hai.
FLOP ek floating-point arithmetic operation hai. Ek FMA 2 FLOPs count hoti hai (ek multiply + ek add).

Figure — Tensor cores and matrix operations

Upar ki picture haara anchor hai: ek output cell C[i][j] fill karne ke liye, tum A ki row i aur B ke column j ke saath slide karte ho, milne wale pairs ko multiply karte ho, aur unhe add karte jaate ho. Ise dhyan mein rakho — har exercise isi slide-and-sum ki variation hai.

Level 1 — Recognition

L1.1 — Operation ka naam lena

Problem. Ek Tensor Core chhote square tiles pe D = A × B + C execute karta hai. Plain words mein, wo do arithmetic actions kya hain jo us single instruction mein packed hain, aur accumulator kya hai?

Recall Solution

Do actions hain: ek matrix multiply (A × B) aur phir ek matrix add (+ C). Yeh combined action ek matrix multiply-accumulate (MMA) hai. Accumulator C hai (aur result D mein overwrite hota hai): yeh woh running total hai jisme hum fresh partial products add karte rehte hain. Ek alag accumulator rakhna hi allow karta hai ki hum kai tiles sum kar sakein bina kuch khoe.

L1.2 — Tile shapes padhna

Problem. Ek Ampere Tensor Core M×N×K notation mein 16×8×16 tile advertise karta hai. M, N, K mein se kaunsa shared inner dimension hai jo sum hota hai, aur woh kaunsa output tile produce karta hai?

Recall Solution

C = A × B mein jahan A M×K hai aur B K×N hai, letter K woh shared dimension hai jo A aur B dono mein appear karta hai — yahi woh hai jise hum sum karte hain (row/column slide ki length). Toh yahan K = 16 inner dimension hai. Output tile C ka size M×N = 16×8 hai.

Level 2 — Application

L2.1 — FMAs aur FLOPs count karna

Problem. Tum do 512 × 512 matrices multiply karte ho. (a) Naive element-by-element algorithm kitne FMAs perform karta hai? (b) Woh kitne FLOPs hain?

Recall Solution

Har output cell C[i][j] ko K FMAs chahiye (sum ke har term ke liye ek). Yahan matrices M×K times K×N hain jahan M = N = K = 512.

Output cells ki sankhya: M·N = 512·512.
FMAs per cell: K = 512.
(a) Total FMAs = M·N·K = 512³ = 134{,}217{,}728 ≈ 1.34 × 10⁸ FMAs.
(b) Har FMA = 2 FLOPs, toh 2 × 512³ = 268{,}435{,}456 ≈ 2.68 × 10⁸ FLOPs.

L2.2 — Tensor Core instructions count karna (4×4 tiles)

Problem. 256 × 256 matrices ko Volta-style 4×4 hardware tiles use karke multiply karo. Kitne mma.sync tile-instructions issue hote hain, aur yeh scalar FMAs se kitne factor kam hain?

Recall Solution

Hum har dimension ko TILE = 4 size ke tiles mein kaatenge. Har axis mein 256 / 4 = 64 tiles hain.

Output grid 64 × 64 tiles ka hai.
Har output tile ke liye hum inner dimension pe 64 tile-steps mein slide karte hain.
Tile-instructions = 64 × 64 × 64 = 64³ = 262{,}144.

Scalar FMAs hote 256³ = 16{,}777{,}216. Ratio: $\frac{25 6 ^{3}}{6 4 ^{3}} = 4^{3} = 64.$ Har tile op internally 4×4×4 = 64 FMAs karta hai, isliye 64× kam instructions bilkul wahi kaam hai jo ek Tensor Core op ke andar packed hai. 262,144 instructions, ek 64× reduction.

L2.3 — Peak throughput timing

Problem. Ek GPU Tensor Cores pe 100 TFLOP/s sustain karta hai. Ek 1024 × 1024 × 1024 matrix multiply ke liye pure compute kitna time (milliseconds mein) lega, memory ignore karo?

Recall Solution

FLOPs needed = 2 × 1024³ = 2 × 1{,}073{,}741{,}824 = 2{,}147{,}483{,}648 ≈ 2.147 × 10⁹.
Rate = 100 TFLOP/s = 100 × 10¹² = 10¹⁴ FLOP/s.
Time = FLOPs ÷ rate = 2.147 × 10⁹ / 10¹⁴ = 2.147 × 10⁻⁵ s. $t \approx 2.147 \times 1 0^{- 5} s = 0.0215 ms .$ Toh pure arithmetic mein lagbhag 0.021 ms. Real time zyada hoga kyunki memory ko yeh feed karna padta hai — dekho 9.1.5-Roofline-model.

Level 3 — Analysis

L3.1 — Ek tile multiply haath se

Problem. TILE = 2 ke saath (ek mini Tensor Core), D = A × B + C compute karo jahan $A = [1324], B = [5768], C = [1001] .$

Recall Solution

Slide-and-sum rule D[i][j] = Σₖ A[i][k]·B[k][j] + C[i][j] use karo, jahan k 0,1 run karta hai.

D[0][0] = (1·5 + 2·7) + 1 = (5+14)+1 = 20
D[0][1] = (1·6 + 2·8) + 0 = (6+16)+0 = 22
D[1][0] = (3·5 + 4·7) + 0 = (15+28)+0 = 43
D[1][1] = (3·6 + 4·8) + 1 = (18+32)+1 = 51 $D = [20432251] .$ Notice karo ki accumulator C yahan sirf diagonal ko thoda sa nudge karta hai — yahi exactly woh hai jo ek identity matrix karta hai.

L3.2 — K ke upar accumulation sahi kyun rehta hai

Problem. Tum ek K = 8 multiply ko do 4×4 tile-steps mein split karte ho (k = 0..3, phir k = 4..7), doosra result pehle ke upar add karte hue. Woh algebraic property explain karo jo ise ek saath 8 karne ke barabar banati hai, aur numerically ek cell ke liye confirm karo: row A[0] = [1,1,1,1,1,1,1,1] times column B[:,0] = [2,2,2,2,2,2,2,2].

Recall Solution

Rule yeh hai ki ek sum ko partial sums mein split karke re-add kiya ja sakta hai — yeh associativity of addition hai: $\sum_{k = 0}^{7} x_{k} = (\sum_{k = 0}^{3} x_{k}) + (\sum_{k = 4}^{7} x_{k}) .$ Accumulator C pehla partial sum hold karta hai, phir doosra tile op apna partial sum uske upar add karta hai. Kuch bhi lost nahi hota kyunki addition care nahi karta ki hum terms ko kaise group karte hain. Numerically har term 1·2 = 2 hai.

Pehla tile (k=0..3): 2+2+2+2 = 8.
Doosra tile (k=4..7): 2+2+2+2 = 8.
Accumulated total: 8 + 8 = 16, bilkul wahi jitna ek baar mein 2·8 = 16 karne se aata. ✓

Level 4 — Synthesis

L4.1 — Mixed-precision error budget

Problem. TF32 mein 10-bit mantissa hai; FP32 mein 23 bits hain. Ek m-bit mantissa wale format ki relative rounding error lagbhag 2⁻ᵐ hoti hai. (a) TF32 aur FP32 ki per-number relative errors batao. (b) TF32 kitne factor zyada coarse hai?

Recall Solution

Mantissa floating-point number ka "significant digits" wala hissa hai; zyada mantissa bits = finer resolution. Relative error roughly 2⁻ᵐ hai.

(a) TF32: 2⁻¹⁰ = 1/1024 ≈ 9.77 × 10⁻⁴ (lagbhag 0.098%). FP32: 2⁻²³ ≈ 1.19 × 10⁻⁷.
(b) Ratio = 2⁻¹⁰ / 2⁻²³ = 2¹³ = 8192. Toh TF32 mein stored har number lagbhag 8192× coarser hai, phir bhi kyunki Tensor Cores FP32 mein accumulate karte hain, summed result training ke liye kaafi accurate rehta hai. Dekho 8.3.4-Mixed-precision-training.

L4.2 — Precision choose karna, justification ke saath

Problem. Tumhe ek pre-trained, frozen image classifier pe inference run karna hai aur maximum throughput chahiye. Turing Tensor Cores FP16 aur INT8 support karte hain. Agar INT8 same matmul ke liye FP16 se 4× throughput deta hai, aur INT8 quantization se accuracy drop 1% se kam hai, to tum kya pick karoge aur kyun?

Recall Solution

INT8 pick karo. Reasoning chain:

Weights frozen hain — koi gradients protect nahi karne, isliye floating point ka wide dynamic range zaruri nahi.
INT8 yahan FP16 se 4× throughput deta hai, aur FP16 ka FP32 se already bada lead hai usse 4× zyada.
Accuracy cost (<1%) classification inference ke liye acceptable hai. Accumulator INT32 hai taaki kai int8 × int8 products sum karte waqt overflow na ho. Yeh exactly 8.4.2-Quantization-techniques wala tradeoff hai: woh precision chhod do jo tumhe ab chahiye nahi, speed kharido. Agar yeh training hota, to tum BF16 use karte (L4.1 note dekho).

Level 5 — Mastery

L5.1 — Ek full tile tak padding karna

Problem. Ek layer 1000 × 1000 output produce karti hai lekin Tensor Cores chahte hain ki dimensions 16 ke multiples hon. (a) Sabse chhota padded size kya hai? (b) Padding kitne extra "wasted" output cells create karta hai, aur padded work ka kitna fraction waste hota hai?

Recall Solution

(a) Har dimension ko 16 ke agle multiple tak round up karo. 1000 / 16 = 62.5, toh humein 63 tiles per axis chahiye → 63 × 16 = 1008. Padded size 1008 × 1008.
(b) Padded cells = 1008² = 1{,}016{,}064. Real cells = 1000² = 1{,}000{,}000. Wasted cells = 1{,}016{,}064 − 1{,}000{,}000 = 16{,}064.
Fraction wasted = 16{,}064 / 1{,}016{,}064 ≈ 0.01581 = 1.58%. Ek chhota padding tax puri Tensor Core utilization dilata hai — iske layak hai kyunki ek partial tile fast path se bilkul bahar gir jaata hai (the "performance cliff").

L5.2 — End-to-end throughput comparison

Problem. Do 2048 × 2048 × 2048 matmuls ek hi GPU pe run hote hain. Path A CUDA cores use karta hai 12 TFLOP/s pe. Path B Tensor Cores use karta hai 120 TFLOP/s pe. (a) Dono ke liye compute-only time. (b) Speedup. (c) Agar Path B actually memory-bound hai aur apni peak ka sirf 40% hi reach kar paata hai, to Path A ke upar realistic speedup kya hai?

Recall Solution

FLOPs = 2 × 2048³ = 2 × 8{,}589{,}934{,}592 = 17{,}179{,}869{,}184 ≈ 1.718 × 10¹⁰.
(a) Path A: 1.718×10¹⁰ / (12×10¹²) = 1.432×10⁻³ s ≈ 1.432 ms. Path B (peak): 1.718×10¹⁰ / (120×10¹²) = 1.432×10⁻⁴ s ≈ 0.1432 ms.
(b) Ideal speedup = 120 / 12 = 10× (also 1.432ms / 0.1432ms = 10).
(c) Peak ka 40% pe, Path B effectively 0.40 × 120 = 48 TFLOP/s pe run karta hai. Time = 1.718×10¹⁰ / (48×10¹²) = 3.579×10⁻⁴ s ≈ 0.358 ms. Realistic speedup = 1.432 / 0.358 = 4×. Lesson: peak FLOP/s ek ceiling hai, promise nahi. Tum ise reach karo ya na karo yeh depend karta hai cores ko feed karne pe — dekho 6.2.8-Memory-hierarchy-in-GPUs aur 9.1.5-Roofline-model.

Recall Quick self-test (cloze)

Ek Tensor Core instruction D = ==A × B + C== ek tile pe compute karta hai. K-multiply ko partial sums mein split karna valid hai kyunki associativity of addition hai. Ek single FMA 2 FLOPs count hoti hai. Teeno axes pe matmul tiling karna instruction count ko (N/TILE)³ se reduce karta hai. Mixed precision low-precision inputs use karta hai lekin FP32 (high-precision) accumulation karta hai.

Related: 6.2.1-GPU-vs-CPU-architecture · 6.2.10-Warp-scheduling-and-execution · 7.3.6-cuBLAS-and-cuDNN · 6.2.12 Tensor cores and matrix operations (Hinglish)