Parallel computing and GPUs

In short

Parallel computing solves one problem by running many operations at the same time on multiple cores, processors or accelerators, instead of one instruction stream on one core. It became unavoidable once clock speeds stopped rising (power wall) around 2004, so performance now comes from more cores rather than faster cores. A GPU is the extreme form: thousands of simple arithmetic units run the same instruction on different data (SIMT), giving tens of TFLOPS for graphics, deep learning, simulation and crypto workloads. Universities examine it in "Parallel and Distributed Computing" / "GPU Computing" electives and in computer architecture papers (Flynn's taxonomy, Amdahl's law, coherence, interconnects); GATE asks speedup, dependence and synchronisation questions. It is interviewed hard at NVIDIA, AMD, Intel, Qualcomm, Samsung R&D, Texas Instruments, Mercedes-Benz R&D India, Ola Krutrim, Zoho, Mathworks and every HPC/AI-infrastructure team.

Explained

Why parallelism, and what limits it

• Drivers: the power wall (dynamic power ≈ C·V²·f, so raising f is costly), the ILP wall (limited instruction-level parallelism to extract) and the memory wall (DRAM latency improves far slower than compute).
• Levels of parallelism: bit-level → instruction-level (pipelining, superscalar) → data-level (SIMD/vector) → thread-level (multicore, SMT) → task/process-level (clusters).
• Granularity: fine-grained (few instructions per sync, high communication) vs coarse-grained (large independent chunks, better for clusters).
• Bernstein's conditions — two statements S1, S2 (S1 first in program order) can run in parallel if: In(S1) ∩ Out(S2) = ∅ (no anti-dependence), In(S2) ∩ Out(S1) = ∅ (no flow dependence), Out(S1) ∩ Out(S2) = ∅ (no output dependence).
• Dependence types in loops: flow (true), anti, output; loop-carried dependences block vectorisation. a[i] = a[i-1] + 1 is serial; a[i] = b[i] + c[i] is fully parallel.
• Overheads that eat speedup: communication, synchronisation (barriers), load imbalance, redundant work, and the serial fraction.

Classification and architectures

• Flynn's taxonomy (1966) — by instruction and data streams:
  • SISD: single core, classical von Neumann machine.
  • SIMD: one instruction on many data — vector units (SSE/AVX on x86, NEON on ARM), GPUs.
  • MISD: rare; fault-tolerant/pipelined systolic designs are cited as examples.
  • MIMD: independent cores running different instructions — multicore CPUs, clusters.
• SPMD (single program, multiple data) is the usual *programming* style on MIMD hardware; SIMT (single instruction, multiple threads) is NVIDIA's GPU model — SIMD hardware with a per-thread programming view.
• Memory organisation:
  • Shared memory / UMA (SMP): all cores see one address space at equal cost; easy to program, limited scaling.
  • NUMA: each socket has local DRAM; remote access costs more — bind threads and memory (first-touch) to the same node.
  • Distributed memory: nodes with private memory communicating by messages (MPI); scales to thousands of nodes.
  • Hybrid / PGAS: MPI across nodes + OpenMP/CUDA inside a node; PGAS (UPC, Chapel) gives a global address space with locality awareness.
• Multicore terms: chip multiprocessor, SMT/hyper-threading (one core, several hardware thread contexts), heterogeneous big.LITTLE (Snapdragon, MediaTek), accelerators (GPU, TPU, NPU, FPGA).

Coherence, consistency and synchronisation

• Cache coherence: with private caches, the same line can sit in several caches; coherence keeps a single value visible. Write-invalidate (common) vs write-update.
• Snooping protocols: MSI (Modified, Shared, Invalid), MESI (adds Exclusive — avoids a bus transaction on write to a clean, privately held line), MOESI/MESIF (adds Owned/Forward for cache-to-cache supply). Snooping needs a broadcast medium, so it stops scaling at a few dozen cores.
• Directory-based coherence: a directory tracks sharers per line; messages are point-to-point, so it scales — used in large NUMA machines.
• False sharing: two threads write different variables in the *same* cache line; the line ping-pongs and performance collapses. Fix by padding/aligning to 64 bytes.
• Memory consistency is a different question — the allowed *order* of operations. Sequential consistency is the intuitive model; real CPUs use relaxed models (x86-TSO allows store buffering; ARM/POWER are weaker), so portable code needs fences/barriers and the atomics of the language memory model (C++11 std::atomic; volatile in the Java JMM — C++ volatile gives no ordering or atomicity).
• Synchronisation primitives: atomic read-modify-write (test-and-set, compare-and-swap), spinlocks, mutexes, semaphores, monitors, barriers, reader-writer locks; lock-free and wait-free structures use CAS.
• Hazards of concurrency: race condition, deadlock (mutual exclusion + hold-and-wait + no preemption + circular wait), livelock, starvation, priority inversion, ABA problem in CAS code.

Interconnection networks (examined with diameter and bisection width)

• Metrics: degree (links per node), diameter (worst-case hop count), bisection width (minimum links cut to halve the network — a bandwidth proxy), cost (number of links), latency = startup + hops × per-hop.
• Topologies: bus (diameter 1, bisection 1, does not scale) · crossbar (full, p² switches) · ring (diameter ⌊p/2⌋) · 2-D mesh (diameter 2(√p − 1)) · torus (wrap-around, halves mesh diameter) · hypercube of p = 2^d nodes (degree d, diameter d, bisection p/2) · fat tree (used in InfiniBand clusters) · dragonfly.
• Real links: PCIe (Gen4 x16 ≈ 32 GB/s each way), NVLink (hundreds of GB/s GPU-to-GPU), NVSwitch, InfiniBand HDR/NDR for cluster fabric, Ethernet/RoCE.
• Routing/switching: store-and-forward vs wormhole/cut-through; deterministic (dimension-order) vs adaptive routing; deadlock avoided by virtual channels.

GPU architecture and the CUDA model

• Design difference: a CPU spends area on large caches, branch prediction and out-of-order logic to minimise *latency* for few threads; a GPU spends area on ALUs and runs tens of thousands of threads to maximise *throughput*, hiding memory latency by switching warps.
• Hardware hierarchy: GPU → several SMs (streaming multiprocessors) → CUDA cores (FP32 lanes), load/store units, SFUs, tensor cores (matrix-multiply-accumulate for FP16/BF16/TF32/INT8 — the basis of deep-learning throughput), register file, shared memory/L1, warp schedulers; shared L2 and HBM2e/HBM3 or GDDR device memory.
• Software hierarchy: kernel → grid → blocks (cooperative thread arrays) → warps of 32 threads → threads. A block runs entirely on one SM; blocks are scheduled in any order, which is what makes a kernel scale across GPUs of different sizes.
• Memory spaces: registers (per thread, fastest) · local memory (register spill, in DRAM) · shared memory (per block, software-managed scratchpad, ~100× lower latency than uncached global) · global memory (HBM/GDDR, visible to all) · constant and texture/read-only caches. Unified/managed memory migrates pages automatically.
• Execution rules that decide performance:
  • Coalescing: threads of a warp should touch consecutive addresses so one request fills whole 32-byte sectors / 128-byte lines. Strided or random access multiplies traffic.
  • Warp divergence: an if that splits a warp makes both paths execute serially with lanes masked; branch on warp-aligned conditions.
  • Shared-memory bank conflicts: shared memory has 32 banks of 4 bytes; two threads hitting different rows of the same bank serialise. Pad arrays (e.g. [32][33]) to break conflicts.
  • Occupancy: active warps per SM ÷ maximum warps per SM; limited by registers per thread, shared memory per block, threads per block and blocks per SM. High occupancy helps latency hiding but is not the goal in itself.
  • Synchronisation: __syncthreads() (block barrier), warp-level primitives (__shfl_sync, __ballot_sync), atomics, cooperative groups; there is no global barrier inside a kernel — launch a second kernel instead.
  • Streams and overlap: asynchronous copies + kernels in multiple streams overlap PCIe transfer with compute; pinned (page-locked) host memory is needed for true async copy.
• Programming stack: CUDA C/C++ (nvcc), OpenCL (vendor-neutral), SYCL/oneAPI, AMD HIP/ROCm, OpenACC/OpenMP target directives, Python via Numba/CuPy/PyTorch/TensorFlow; libraries cuBLAS, cuFFT, cuDNN, cuSPARSE, Thrust, NCCL; profiling with Nsight Compute/Systems.

__global__ void saxpy(int n, float a, const float* x, float* y) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;   // global index
    if (i < n) y[i] = a * x[i] + y[i];               // coalesced, 1 FMA
}
// host: saxpy<<<(n + 255) / 256, 256>>>(n, 2.0f, d_x, d_y);

Parallel algorithm patterns and performance analysis

• Patterns: map (element-wise) · reduction (tree, log p depth) · scan/prefix sum (Hillis–Steele O(n log n) work, Blelloch work-efficient O(n)) · stencil (shared-memory tiling, halo exchange) · histogram (privatisation + atomics) · sort (bitonic, radix) · tiled matrix multiply (load tiles into shared memory to reuse data) · sparse SpMV · graph BFS (frontier-based).
• Decomposition strategies: data (domain) decomposition, functional decomposition, recursive (divide and conquer), pipeline, master–worker.
• Load balancing: static block/cyclic distribution; dynamic self-scheduling and guided scheduling (OpenMP schedule(dynamic, chunk) / schedule(guided, chunk)), work stealing (Cilk, TBB).
• Scaling: strong scaling — fixed problem, more processors (Amdahl); weak scaling — problem grows with processors (Gustafson). Isoefficiency function gives how fast the problem must grow to hold efficiency constant; Karp–Flatt metric estimates the experimental serial fraction and detects overhead growth.
• Roofline model: plot attainable GFLOP/s against arithmetic intensity (FLOP per byte of DRAM traffic); below the machine balance point a kernel is memory-bound, above it compute-bound. Most real kernels (SAXPY, stencils, SpMV) are memory-bound, which is why coalescing and reuse matter more than FLOPs.
• Little's law for latency hiding: bytes in flight = bandwidth × latency; a GPU needs thousands of outstanding requests, which is why it needs many warps.

Key points

• Speedup - serial time ÷ parallel time; efficiency = speedup ÷ processors; superlinear speedup usually means cache effects, not magic.
• Amdahl's law - the serial fraction s caps speedup at 1/s however many processors you add.
• Gustafson's law - if the problem grows with the machine, speedup grows almost linearly; explains why supercomputers are useful despite Amdahl.
• Flynn's taxonomy - SISD, SIMD, MISD, MIMD by instruction and data streams; SIMT is the GPU variant of SIMD.
• UMA vs NUMA - equal-cost shared memory vs socket-local memory; NUMA needs thread and memory affinity.
• Cache coherence - MSI/MESI/MOESI by snooping for small systems, directory protocols for large ones.
• False sharing - independent variables in one 64-byte line destroy scaling; pad them.
• Memory consistency - ordering rules; relaxed models need explicit fences, coherence alone is not enough.
• Warp - 32 threads issued together; divergence serialises the taken paths.
• Coalescing - consecutive addresses per warp; a stride of 32 floats can waste 8× of the bandwidth fetched.
• Shared memory - per-block scratchpad, 32 banks; conflicts serialise, padding fixes them.
• Occupancy - active warps ÷ max warps per SM; limited by registers and shared memory, not an end in itself.
• Arithmetic intensity - FLOP/byte; compare with machine balance (peak FLOPS ÷ peak bandwidth) to label a kernel memory- or compute-bound.
• Tensor cores - fixed-function matrix-multiply units giving mixed-precision throughput many times the FP32 rate.
• No global barrier inside a kernel - inter-block synchronisation means a new kernel launch (or cooperative-groups grid sync).
• Host-device transfer is often the bottleneck - minimise copies, use pinned memory and streams to overlap.
• Bisection width and diameter are the two network numbers exams ask for; hypercube of 2^d nodes has diameter d and bisection 2^d−1.
• Strong vs weak scaling must be stated with every speedup number, along with the processor count.

Formulas and examples

Example 1 - Amdahl vs Gustafson. A program is 95% parallelisable, so s = 0.05.

• p = 16: S = 1/(0.05 + 0.95/16) = 1/(0.05 + 0.059375) = 1/0.109375 = 9.14; E = 9.14/16 = 57.1%.
• p = 64: S = 1/(0.05 + 0.0148438) = 1/0.0648438 = 15.42; E = 24.1% — four times the hardware bought 1.7× the speed.
• p → ∞: S_max = 1/0.05 = 20.
• Gustafson (problem scaled with p), p = 16: S = 0.05 + 16 × 0.95 = 15.25 — the same code looks far better under weak scaling.
• Karp–Flatt check on the p = 16 measurement: e = (1/9.14 − 1/16)/(1 − 1/16) = (0.10941 − 0.0625)/0.9375 = 0.050 ✓ consistent with s = 0.05 (if e *rose* with p, the loss would be overhead, not serial code).

Example 2 - GPU peak and machine balance. A GPU has 80 SMs × 64 FP32 lanes = 5,120 cores at 1.41 GHz, with 900 GB/s HBM2.

• Peak FP32 = 5,120 × 1.41 × 10⁹ × 2 = 1.4438 × 10¹³ = 14.44 TFLOPS.
• Machine balance = 14,438 GFLOP/s ÷ 900 GB/s = 16.0 FLOP/byte. Any kernel below 16 FLOP/byte is memory-bound.

Example 3 - SAXPY is memory-bound. y[i] = a*x[i] + y[i] on N = 2²⁰ × 64 = 67,108,864 floats.

• Per element: read x and y (8 B) + write y (4 B) = 12 B; FLOPs = 2 (one FMA) → AI = 2/12 = 0.167 FLOP/byte.
• Attainable = min(14,438, 900 × 0.1667) = 150 GFLOP/s ≈ 1.0% of peak — correct behaviour, not a bug.
• Time = 67,108,864 × 12 B ÷ 900 GB/s = 805.3 MB ÷ 900 GB/s ≈ 0.89 ms.
• Adding + 3.0f to the expression changes AI only to 0.25 FLOP/byte; fusing kernels to cut traffic is the real optimisation.

Example 4 - Occupancy from registers. SM limits: 2,048 threads (64 warps), 65,536 32-bit registers. Kernel: 256 threads/block, 48 registers/thread.

• Registers per block = 256 × 48 = 12,288 → blocks per SM = ⌊65,536/12,288⌋ = 5 → threads = 1,280 = 40 warps → occupancy = 1,280/2,048 = 62.5%.
• Cut to 32 registers/thread: 256 × 32 = 8,192 → ⌊65,536/8,192⌋ = 8 blocks = 2,048 threads = 100% occupancy (use __launch_bounds__ or -maxrregcount).

Example 5 - Coalescing cost. A warp of 32 threads reads 4-byte floats.

• Contiguous: 32 × 4 = 128 B → one 128-byte line, 4 sectors, 100% useful.
• Stride 32 floats (128 B apart, i.e. column access of a row-major matrix): 32 separate 32-byte sectors = 1,024 B fetched for 128 B used → 12.5% efficiency, 8× wasted bandwidth. Transposing through shared memory restores it.

Example 6 - Reduction of N = 1,048,576 with 1,024 threads. Steps ≈ N/p + log₂p = 1,024 + 10 = 1,034 against 1,048,576 serial additions → speedup ≈ 1,014, efficiency ≈ 99% (ignoring memory limits) — reductions are work-efficient; the real cost is the memory sweep.

Where it is used

• AI and deep learning: GPU clusters train and serve models — NVIDIA/AMD GPU software roles in Bengaluru, Pune and Hyderabad; PyTorch/CUDA kernel engineering at Ola Krutrim, Sarvam AI, Fractal and the big IT services' AI labs.
• Indian HPC: the National Supercomputing Mission (MeitY and DST, implemented by C-DAC with IISc) runs the PARAM series — PARAM Shivay (IIT BHU), PARAM Siddhi-AI (C-DAC, NVIDIA DGX A100-based, India's largest AI supercomputer when commissioned), AIRAWAT and the newer PARAM Rudra systems; used for weather (IITM/NCMRWF), genomics, CFD and materials. Capacities and GPU counts under the IndiaAI Mission's empanelment keep changing — quote current MeitY figures rather than a remembered number.
• Semiconductor and systems R&D: GPU/NPU architecture, driver, compiler and verification jobs at NVIDIA, AMD, Intel, Qualcomm, ARM, Samsung R&D, Texas Instruments, Micron, MediaTek and Indian fabless startups.
• Engineering simulation and finance: CFD/FEA (ANSYS, Simcenter) on GPU clusters at Mercedes-Benz R&D India, Tata Motors, ISRO/HAL suppliers; option pricing and risk Monte Carlo on GPUs at trading firms and bank quant teams.
• Graphics, games and media: game engines, real-time ray tracing and video transcoding — studios and product teams in Bengaluru, Hyderabad and Pune.
• Everyday devices: Adreno/Mali GPUs in phones (compute shaders, camera ISP pipelines), automotive ADAS SoCs, and cloud GPU rentals from Indian providers for startups that cannot buy hardware.

Interview questions

• Q: State Amdahl's law and its practical message. A: S(p) = 1/(s + (1 − s)/p); the serial fraction caps speedup at 1/s, so 5% serial code limits you to 20× no matter how many cores you add.
• Q: How does Gustafson's law escape that limit? A: It assumes the problem size grows with the machine, so the serial part stays a fixed absolute cost and speedup grows nearly linearly — valid for weak scaling, not for a fixed-size problem.
• Q: Difference between SIMD and SIMT? A: SIMD exposes a vector of fixed width to the programmer; SIMT exposes independent scalar threads that hardware groups into warps and executes in lockstep, so divergence is legal but costs performance.
• Q: What is false sharing and how do you detect it? A: Threads writing distinct variables inside one cache line cause the line to bounce between caches; it shows up as poor scaling with high coherence traffic in perf/VTune, and is fixed by padding to a cache-line boundary.
• Q: Why is MESI better than MSI? A: The Exclusive state records that a clean line is held by only one cache, so writing it needs no invalidation broadcast — saving bus traffic in the common private-data case.
• Q: Is cache coherence the same as memory consistency? A: No — coherence makes one address appear to have a single value; consistency defines the allowed ordering between different addresses, and relaxed models need explicit fences.
• Q: What is memory coalescing and how much can it cost? A: Having a warp access consecutive addresses so one transaction serves all lanes; a 128-byte stride can turn a 128-byte need into 1,024 bytes fetched — 8× the bandwidth.
• Q: What is warp divergence and how do you avoid it? A: When lanes of a warp take different branches, the paths run serially with masking; avoid it by making the branch condition uniform across a warp or by sorting/binning data.
• Q: Why does shared-memory tiling speed up matrix multiplication? A: Each tile loaded once from global memory is reused by many threads, raising arithmetic intensity from about 0.25 to O(tile width) FLOP/byte and moving the kernel from memory-bound to compute-bound.
• Q: Does 100% occupancy guarantee peak performance? A: No — occupancy only helps hide latency; a kernel can be bandwidth-limited or instruction-limited, and using more registers per thread with lower occupancy is often faster.
• Q: How do blocks synchronise with each other inside one kernel? A: They normally cannot — __syncthreads() is per block; global synchronisation needs a second kernel launch or a cooperative-groups grid barrier with a resident grid.
• Q: A GPU peaks at 14 TFLOPS but your kernel reaches 150 GFLOP/s. Is it broken? A: Not necessarily — compute its arithmetic intensity; at 0.17 FLOP/byte on 900 GB/s the roofline caps it near 150 GFLOP/s, so it is already at the memory limit.
• Q: When is a GPU the wrong choice? A: For small data sets where PCIe transfer dominates, or for branchy, pointer-chasing, latency-sensitive and heavily serial code — a CPU wins there.

Common mistakes

• Quoting a speedup without the processor count, the problem size, or whether it is strong or weak scaling.
• Applying Amdahl's law to a scaled problem (or Gustafson's to a fixed one) and drawing the wrong conclusion.
• Forgetting that efficiency must drop as p rises for a fixed problem, and treating E < 100% as a coding bug.
• Assuming cache coherence removes the need for locks and fences — it removes neither races nor ordering problems.
• Writing CUDA kernels that index with threadIdx.x only, ignoring blockIdx.x, or omitting the if (i < n) guard, so part of the array is never processed or memory is written out of bounds.
• Measuring GPU time with CPU timers without cudaDeviceSynchronize(), or excluding host-device copy time from the comparison with the CPU version.
• Optimising FLOPs in a kernel that is memory-bound, instead of fixing access patterns, reuse and kernel fusion.
• Treating a race-free run as proof of correctness — non-determinism, missing __syncthreads() and unprotected atomics may surface only under different scheduling.

Related topics

• Computer organisation and architecture · Distributed systems basics · Operating systems: processes, threads and synchronisation · Computer graphics basics · Algorithm design and analysis · Machine learning systems and accelerators