V I S V E S VA R AYA T E C H N O L O G I C A L U N I V E R S I T Y | B . E .

C O M P U T E R
SCIENCE & ENGINEERING

BCS702 – Parallel Computing

Module-wise Notes (Semester VII)

COURSE MODULES TEXTBOOK

CODE 5 Grama et
BCS702 al., 2e

Prepared by: Ravi N

Asst. Professor, Department of CSE
SSSE, Tumkur

Based on the prescribed textbook: Introduction to Parallel Computing (2nd Ed.) by Ananth
Grama, Anshul Gupta, George Karypis & Vipin Kumar, Addison-Wesley / Pearson, 2003.
Content follows VTU syllabus topic order.

TABLE OF CONTENTS

1. Module 1 – Introduction to Parallel Programming & Hardware/

Software

1. Introduction to Parallel Programming

2. Motivating Parallelism

3. Scope of Parallel Computing

4. Classifications of Parallel Computers (Flynn's Taxonomy)

 5. SIMD Systems

6. MIMD Systems

7. Interconnection Networks

8. Cache Coherence

9. Shared-Memory vs. Distributed-Memory

10. Coordinating Processes/Threads

2. Module 2 – GPU Programming, MIMD Performance & Scalability

1. Introduction to GPU Programming

2. Programming Hybrid Systems

3. Speedup and Efficiency in MIMD Systems

4. Amdahl's Law

5. Scalability in MIMD Systems

6. Taking Timings of MIMD Programs

7. GPU Performance

3. Module 3 – Distributed Memory Programming with MPI

1. Principles of MPI Programming

2. Key MPI Functions

3. The Trapezoidal Rule in MPI

4. Dealing with I/O

5. Collective Communication

6. MPI Derived Datatypes

7. Performance Evaluation of MPI Programs

8. A Parallel Sorting Algorithm

4. Module 4 – Shared-Memory Programming with OpenMP

1. OpenMP Pragmas and Directives

2. The Trapezoidal Rule with OpenMP

3. Scope of Variables

4. The Reduction Clause

5. Loop-Carried Dependency

6. Scheduling
 7. Producers and Consumers

8. Caches, Cache Coherence & False Sharing

9. Tasking

10. Thread Safety

5. Module 5 – GPU Programming with CUDA

1. GPUs and GPGPU

2. GPU Architectures

3. Heterogeneous Computing

4. Threads, Blocks, and Grids

5. NVIDIA Compute Capabilities

6. Vector Addition in CUDA

7. Returning Results from CUDA Kernels

8. CUDA Trapezoidal Rule

MODULE 1

Introduction to Parallel Programming &

Parallel Hardware/Software
Classifications of parallel computers · SIMD systems · MIMD systems · Interconnection
networks · Cache coherence · Shared-memory vs. distributed-memory · Coordinating
processes/threads

1.1 Introduction to Parallel Programming

A parallel computer is a collection of processing elements that cooperate to solve large

problems quickly. Rather than executing a single instruction stream on one processor,
parallel computing divides a task into sub-tasks that execute simultaneously on multiple
processors, reducing overall execution time.
 Why parallel computing? Moore's Law (circuit complexity doubles every ~18
months) has historically delivered faster single-core processors, but physical limits
— heat dissipation, memory latency, power consumption — now prevent indefinite
single-core speedup. Parallel computing is the principal way to continue improving
computational throughput.

Key Terms

Task
An arbitrarily defined unit of computational work.

Process / Thread
A unit of execution that runs concurrently with other units on the same or different
processors.

Speedup (S)
Ratio of sequential execution time to parallel execution time: S = Tserial / Tparallel.

Efficiency (E)
Speedup per processor: E = S / p, where p is the number of processors.

1.2 Motivating Parallelism

There are three classical arguments for parallelism (from the textbook, Chapter 1):

(a) The Computational Power Argument — from Transistors to FLOPS

Gordon Moore's 1965 observation (Moore's Law) predicted that transistor count per chip
doubles approximately every 18 months. This translates to exponential growth in raw
compute power (FLOPS). However, translating transistor count into actual throughput
requires architectural support — pipelining, superscalar execution, and ultimately explicit
parallelism.
(b) The Memory/Disk Speed Argument

Processor clock speeds improved at ~40% per year over the past decade, but DRAM
access times improved at only ~10% per year. This growing memory wall means a
processor often sits idle waiting for data. Parallel platforms alleviate this by providing:

• Larger aggregate caches (linear in number of processors).

• Higher aggregate bandwidth to memory (also linear in processors).

Remember: Memory latency = time to get first byte. Memory bandwidth = rate at
which data flows. They require different solutions: caching addresses latency; wider
buses address bandwidth.

(c) The Data Communication Argument

Many real-world problems involve data that is naturally distributed across locations (e.g.,
web servers, sensor networks, genome databases). Parallel/distributed computing is
necessary not only for speed but because centralising all data is infeasible.

1.3 Scope of Parallel Computing

Application Domains of Parallel Computing

Domain Example Applications

Engineering & Airfoil optimisation, IC circuit layout, MEMS/NEMS design,

Design structural analysis

Scientific Human genome sequencing, weather modelling, astrophysics

Applications simulations, drug discovery

Commercial Transaction processing, web servers, data mining, e-commerce

Applications databases

Computer Systems Cryptography (integer factoring), network intrusion detection,

embedded automotive control
 Table 1.1: Major application areas of parallel computing (Grama et al., Chapter 1).

1.4 Classifications of Parallel Computers — Flynn's Taxonomy

Flynn (1966) classified computers according to the number of concurrent instruction

streams and data streams:

Class Full Name Description Example

SISD Single Traditional uniprocessor — one Classic desktop

Instruction, instruction acts on one data item at a CPU
Single Data time.

SIMD Single One control unit broadcasts the same GPU cores,
Instruction, instruction to many processing units, MMX/SSE,
Multiple Data each operating on different data. CM-2

MISD Multiple Multiple instruction streams on the Theoretical /

Instruction, same data stream. Rarely used in fault-tolerant
Single Data practice. systems

MIMD Multiple Each processor executes its own Multicore CPUs,

Instruction, independent instruction stream on its clusters, IBM SP
Multiple Data own data. Most modern parallel
computers.

Table 1.2: Flynn's Taxonomy of Computer Architectures.

1.5 SIMD Systems

In an SIMD architecture (Figure 2.3a from the textbook), a single control unit dispatches
the same instruction to all processing units simultaneously. Each unit applies that
instruction to its own local data.
 Figure 1.1: Typical SIMD architecture — one control unit, multiple processing units (Grama et al., Fig.
2.3a).

Characteristics of SIMD

• All processing elements execute the same instruction each clock cycle.

• Execution is synchronous across all units.

• Only one copy of the program is stored (less memory overhead).

• An activity mask (bitmask) can selectively disable certain processors, enabling

conditional operations.
 • Conditional execution (if-else) causes some processors to be idle, reducing utilisation
— a key drawback.

Remember: SIMD is efficient for regular, structured computations (array/vector

operations, image processing) but performs poorly on data-dependent conditional code.

SIMD Conditional Execution — Example

For the statement:

if (B == 0)
C = A;
else
C = A / B;

On an SIMD machine with 4 processors: Step 1 — processors where B=0 execute C=A ;
others are idle. Step 2 — processors where B≠0 execute C=A/B ; others are idle. This
two-step execution wastes cycles. (Textbook Example 2.11)

Examples of SIMD Machines

• Illiac IV, MasPar MP-1, CM-2 (historical)

• Intel MMX / SSE / AVX extensions (modern SIMD in CPUs)

• GPU Streaming Multiprocessors (SIMT = Single Instruction Multiple Threads)

1.6 MIMD Systems

In an MIMD architecture (Figure 2.3b from the textbook), each processing element has its
own control unit and can execute a completely different program on different data,
independently and asynchronously.
 Figure 1.2: Typical MIMD architecture — each processor has independent control (Grama et al., Fig.
2.3b).

SPMD — Single Program Multiple Data

In practice, a simpler variant called SPMD is widely used: multiple processors execute
the same program but on different portions of data. Control flow diverges through
conditional checks on the processor's rank/ID. SPMD has the same expressive power as
full MIMD but is much easier to program.
MIMD Sub-classifications by Memory Organisation

Type Memory Communication Example

Shared-Memory All processors Read/Write shared Multicore

MIMD share one global variables; use locks for CPUs, SGI
address space synchronisation Origin

Distributed- Each processor has Explicit message passing IBM SP,

Memory MIMD private local (MPI, PVM) Beowulf
memory clusters

Table 1.3: MIMD memory organisations.

1.7 Interconnection Networks

Interconnection networks connect processors to each other and to memory. They are
classified as static (fixed point-to-point links) or dynamic (switches configure links on
demand).

Key Evaluation Metrics

Diameter
Maximum shortest path between any two nodes. Lower is better (less latency).

Bisection Width
Minimum links to cut the network into two equal halves. Higher is better (more
bandwidth).

Arc Connectivity
Minimum links that must fail to disconnect the network. Higher means more fault
tolerance.

Cost
Total number of links/switches.
Important Static Topologies

• Linear Array: each node connects to left and right neighbours. Simple but high
diameter (p−1).

• Ring: linear array with wrap-around. Diameter = ⌊p/2⌋.

• 2-D Mesh: nodes arranged in a grid; each node connects to up to 4 neighbours.

Natural for matrix computations. Diameter = 2(√p − 1).

• 2-D Torus: mesh with wrap-around in both dimensions. Bisection width = 2√p.

• Hypercube: d-dimensional cube; p = 2d nodes. Diameter = log p, bisection width = p/

2. Very well connected.

• Tree / Fat-Tree: hierarchical; bandwidth increases towards root in fat-tree designs.

• Completely Connected: every node directly linked to every other. Diameter = 1 but
cost = p(p−1)/2.

Figure 1.3: Construction of hypercubes — 0-D through 4-D (Grama et al., Fig. 2.17).

Bisection Arc Cost

Network Diameter
Width Connectivity (links)

Completely- 1 p²/4 p−1 p(p−1)/2

Connected
 Bisection Arc Cost
Network Diameter
Width Connectivity (links)

Star 2 1 1 p−1

Complete Binary 2 log((p+1)/ 1 1 p−1

Tree 2)

Linear Array p−1 1 1 p−1

2-D Mesh (no 2(√p−1) √p 2 2(p−√p)

wrap)

2-D Wraparound 2⌊√p/2⌋ 2√p 4 2p

Mesh

Hypercube log p p/2 log p (p log p)/2

Table 1.4: Characteristics of static network topologies (Grama et al., Table 2.1).

Dynamic Networks

• Bus: simplest; all nodes share one medium. Cost O(p), distance O(1), but bandwidth
bottleneck at large p.

• Crossbar: non-blocking grid of switches. Cost Θ(p²) — expensive but maximum

performance.

• Multistage (Omega) Network: intermediate cost Θ(p log p); uses perfect-shuffle
interconnect + 2×2 switches across log p stages. Blocking network.

Omega Network

Figure 1.4: An 8-input Omega network with perfect-shuffle interconnect (Grama et al., Fig. 2.12).

1.8 Cache Coherence

When multiple processors each have their own cache but share a global memory, the same
variable may have multiple copies in different caches. If one processor updates its copy,
the other copies become stale. Cache coherence is the mechanism that keeps all copies
consistent.

Figure 1.5: Cache coherence — Invalidate (a) vs Update (b) protocols (Grama et al., Fig. 2.21).

Two Main Protocols

Invalidate Protocol
When a processor writes to a variable, all other cached copies of that variable are
marked invalid. Other processors must re-fetch on next access. Used by most modern
processors (lower steady-state bandwidth).

Update Protocol
When a processor writes, all other cached copies are updated with the new value.
Good when variables are read by many processors frequently.

Three Cache States (Simple Protocol)

• Shared: block is resident in one or more caches; all copies are clean (consistent with
memory).

• Invalid: the block in this cache is outdated and may not be used.
 • Dirty (Modified): this cache holds the only valid copy; memory is stale.

False Sharing

Two processors update different variables that happen to reside in the same cache line.
The coherence protocol treats the whole line as shared/modified, causing unnecessary
invalidations and cache-line "ping-pong" between processors. This is a major source of
hidden overhead in shared-memory parallel programs.

Implementation Mechanisms

Snoopy Caches
Used with broadcast networks (bus/ring). All processors monitor (snoop) the bus; they
update their own cache state when they detect relevant transactions. Simple but limited
scalability.

Directory-Based Coherence
A directory stores the state (presence bits) of every memory block — which caches
hold it, and whether it is dirty. Coherence messages go only to processors that hold a
copy, making it scalable for large systems. Distributed directories spread this metadata
across processors.

1.9 Shared-Memory vs. Distributed-Memory

Feature Shared-Memory Distributed-Memory

Memory access All processors access one Each processor has private local
global address space memory; no direct access to remote
memory

Communication Read/write shared variables Explicit message passing (send/

receive)

Programming Threads (Pthreads), MPI, PVM

model OpenMP
 Feature Shared-Memory Distributed-Memory

Scalability Limited (memory bus Highly scalable to thousands of nodes

becomes bottleneck)

Ease of Easier (no explicit data More complex (must manage data
programming movement) distribution)

Subtypes UMA (uniform latency), MPP (massively parallel), clusters

NUMA (non-uniform
latency)

Example SGI Origin 2000, Sun Ultra IBM SP, Beowulf cluster, Cray T3E
platforms HPC, multicore PCs

Table 1.5: Shared-Memory vs. Distributed-Memory comparison.

UMA vs. NUMA

UMA (Uniform Memory Access)

Every processor takes the same time to access any memory location. Simple bus-based
multiprocessors.

NUMA (Non-Uniform Memory Access)

Accessing "local" memory (physically near a processor) is faster than accessing
"remote" memory. Examples: SGI Origin 2000. Algorithms must exploit locality for
good performance.
 Figure 1.6: UMA shared-address-space architecture (Grama et al., Fig. 2.5a).

1.10 Coordinating Processes / Threads

Shared-Memory Coordination

• Mutex (Lock): ensures only one thread at a time enters a critical section. Operations:
lock() / unlock() .

• Semaphore: generalised lock with a counter; supports signalling between threads.

• Barrier: all threads block at this point until every thread has arrived, then they all
proceed.

• Condition Variable: allows a thread to sleep until a condition is signalled by another.

• Monitor: combines mutual exclusion and condition synchronisation.

Message-Passing Coordination

• Send / Receive: explicit transfer of data between processes.

 • Synchronous (blocking): sender waits until receiver has received the message.

• Asynchronous (non-blocking): sender continues immediately; communication

happens in background.

• Collective operations: Broadcast, Scatter, Gather, Reduce, All-to-All — involve

groups of processes.

SPMD pattern: all processes call MPI_Comm_rank() to get their unique ID, and
MPI_Comm_size() to get total number of processes. They then use if-else on

rank to specialise behaviour.

MODULE 2

GPU Programming, MIMD Performance &

Scalability
GPU programming · Programming hybrid systems · MIMD systems & GPUs · Performance
— Speedup, Efficiency, Amdahl's Law · Scalability in MIMD systems · Taking timings ·
GPU performance

2.1 Introduction to GPU Programming

A Graphics Processing Unit (GPU) was originally designed to accelerate graphics

rendering. Modern GPUs contain thousands of small, simple processing cores optimised
for throughput rather than single-thread latency. They follow the SIMT (Single
Instruction, Multiple Threads) model — groups of threads execute the same instruction
simultaneously.
CPU vs. GPU — Key Differences

Feature CPU GPU

Core count Few (2–64 high- Thousands of smaller cores (e.g., 5000+
performance cores) in modern GPUs)

Optimised for Low latency, serial High throughput, massively parallel

execution execution

Cache Large caches (L1/L2/L3) Small per-core caches; relies on high

memory bandwidth

Programming Threads / OpenMP / MPI CUDA, OpenCL, HIP

model

Memory Lower (e.g., ~50–100 GB/ Much higher (e.g., 500–1000+ GB/s
bandwidth s) HBM)

Use case Control-heavy, branchy Regular, data-parallel computations

code (ML, simulation)

Table 2.1: CPU vs. GPU design philosophy.

2.2 Programming Hybrid Systems

Modern high-performance computers combine multiple CPUs and GPUs in the same
node (e.g., NVIDIA DGX systems, supercomputer nodes with GPUs). Programming such
hybrid systems requires:

• Using MPI between nodes (distributed memory).

• Using OpenMP or Pthreads within a CPU (shared memory on a node).

• Using CUDA/OpenCL to offload compute-intensive kernels to the GPU.

 Hybrid programming pattern: outer loop over MPI ranks (distributed nodes) →
inner OpenMP threads (CPU cores per node) → CUDA kernels (GPU cores per
accelerator). Each level targets a different memory hierarchy.

2.3 Speedup and Efficiency in MIMD Systems

Let Ts = best serial execution time, Tp = parallel execution time using p processors.

Speedup S(p)

S(p) = Ts / Tp

Ideal (linear) speedup: S(p) = p. Superlinear speedup (S > p) can occur if parallelism
reveals cache effects.

Efficiency E(p)

E(p) = S(p) / p = Ts / (p × Tp)

E = 1 (100%) is ideal. In practice E < 1 due to overhead.

Cost C(p)

C(p) = p × Tp

Total processor-time spent. A parallel algorithm is cost-optimal if C(p) = O(Ts).

Overhead Functions

The parallel overhead To captures all extra work done by a parallel program that is absent
in the serial version:

To = p × Tp − Ts
Sources of overhead include: inter-process communication, idle time (load imbalance),
synchronisation barriers, and redundant computation.

2.4 Amdahl's Law

Proposed by Gene Amdahl in 1967, this law expresses the theoretical maximum speedup
when only a fraction of the code can be parallelised.

Let f = fraction of the program that is inherently serial (cannot be parallelised, 0 ≤

f ≤ 1). Then the speedup on p processors is:

S(p) = 1 / ( f + (1 − f)/p )

Implications of Amdahl's Law

• As p → ∞ (unlimited processors): Smax = 1/f

• If 10% of code is serial (f = 0.1): maximum speedup = 10, regardless of how many
processors are used.

• If 1% is serial (f = 0.01): maximum speedup = 100.

• This is often called the serial bottleneck.

Serial Fraction Max Speedup (1/ Speedup with 8 Speedup with 64

(f) f) CPUs CPUs

0.50 (50%) 2 1.78 1.97

0.25 (25%) 4 2.91 3.81

0.10 (10%) 10 4.71 7.54

0.01 (1%) 100 7.41 39.3

0.001 (0.1%) 1000 7.94 60.2

Table 2.2: Speedup values for different serial fractions (Amdahl's Law).
 Remember: Amdahl's Law explains why simply adding more processors yields
diminishing returns. Reducing the serial fraction f is more impactful than increasing p.

2.5 Scalability in MIMD Systems

A parallel system is said to be scalable if it can maintain efficiency as both problem size
and number of processors grow proportionally.

Isoefficiency Function
The rate at which problem size W must grow as p increases to maintain constant
efficiency. A smaller isoefficiency function indicates a more scalable system.

Strong Scaling (Amdahl)

Fixed problem size, increase p. Speedup is limited by the serial fraction.

Weak Scaling (Gustafson's Law)

Problem size grows proportionally with p. Efficiency can be maintained much better.

Gustafson's Law

Sscaled(p) = p − f × (p − 1)

This reformulation argues that for fixed time (not fixed problem size), increasing
processors allows solving proportionally larger problems — giving near-linear speedup.
This is more realistic for scientific computing.

2.6 Taking Timings of MIMD Programs

Measuring parallel program performance requires care:

• Use wall-clock time (not CPU time) because MIMD programs run on multiple
processors simultaneously.

• Time the slowest processor (the one that finishes last), as this determines overall
execution time.
 • Exclude initialization and I/O from timing if comparing computational kernels.

• In MPI: use MPI_Wtime() which returns wall-clock seconds. Follow with

MPI_Barrier() before and after timed region to ensure all processes are

synchronised.

• Run multiple trials and report the minimum or average to reduce noise.

/* MPI timing pattern */

MPI_Barrier(MPI_COMM_WORLD);
start = MPI_Wtime();
/* ... parallel computation ... */
MPI_Barrier(MPI_COMM_WORLD);
elapsed = MPI_Wtime() - start;

2.7 GPU Performance

GPU performance is characterised by:

FLOPS (Floating-Point Operations Per Second)

Peak compute rate. Modern GPUs reach tens of TFLOPS for FP32 and hundreds of
TFLOPS for FP16/INT8.

Memory Bandwidth
GPUs use High Bandwidth Memory (HBM) or GDDR achieving hundreds of GB/s.
Many GPU programs are memory-bandwidth bound.

Occupancy
Fraction of maximum possible warps actively running on a Streaming Multiprocessor
(SM). High occupancy hides memory latency.

Roofline Model
A visual performance model comparing compute intensity (FLOPs/byte) against peak
FLOPS and memory bandwidth to identify whether a kernel is compute-bound or
memory-bound.
Factors Affecting GPU Performance

• Coalesced memory access: threads in a warp should access consecutive memory

addresses for maximum bandwidth.

• Thread divergence: if threads in a warp follow different code paths (if-else), paths
are serialised — hurts performance.

• Shared memory usage: on-chip shared memory is ~100× faster than global memory
— use it as a software-managed cache.

• Host-Device transfers: PCIe bandwidth (~16 GB/s) is far lower than GPU memory
bandwidth; minimise data transfers between CPU and GPU.

MODULE 3

Distributed Memory Programming with MPI

MPI functions · The trapezoidal rule in MPI · Dealing with I/O · Collective communication
· MPI derived datatypes · Performance evaluation · A parallel sorting algorithm

3.1 Principles of Message-Passing Programming

The message-passing model consists of p processes, each with its own private address
space. Processes communicate only by sending and receiving explicit messages. There is
no shared memory — all data exchange must be programmed explicitly.

Basic Four Operations

send(data, destination)
Transfer data to destination process.

receive(data, source)
Receive data from source process (or any source).
whoami() → rank
Return this process's unique identifier (rank) within the communicator.

numprocs() → p
Return total number of processes in the communicator.

Any message-passing program can be written using just these four primitives.

3.2 Key MPI Functions

MPI (Message Passing Interface) is the industry-standard library for distributed-

memory parallel programming. Below are the most essential functions:

Function Purpose

MPI_Init(&argc, &argv) Initialise the MPI execution

environment. Must be the first MPI call.

MPI_Finalize() Clean up MPI state. Must be the last

MPI call.

MPI_Comm_rank(comm, &rank) Get this process's rank (0 to p−1) in

communicator comm .

MPI_Comm_size(comm, &size) Get total number of processes in

comm .

MPI_Send(buf, count, dtype, dest, tag, Blocking send: returns when buffer can
comm) be reused.

MPI_Recv(buf, count, dtype, src, tag, Blocking receive: blocks until message
comm, &status) arrives.

MPI_Isend(..., &request) Non-blocking send — returns

immediately; use MPI_Wait to check
completion.

MPI_Irecv(..., &request) Non-blocking receive.

 Function Purpose

MPI_Wtime() Return wall-clock time in seconds (for

performance measurement).

MPI_Barrier(comm) Block all processes until everyone calls

this function.

MPI_Bcast(buf, count, dtype, root, Root broadcasts data to all processes.

comm)

MPI_Reduce(sbuf, rbuf, count, dtype, Combine values from all processes

op, root, comm) using operator op into root.

MPI_Scatter(sbuf, scount, sdtype, Root distributes different data chunks to

rbuf, rcount, rdtype, root, comm) each process.

MPI_Gather(sbuf, scount, sdtype, rbuf, Each process sends data to root; root
rcount, rdtype, root, comm) collects all.

MPI_Allreduce(sbuf, rbuf, count, Reduce + broadcast: result goes to all

dtype, op, comm) processes.

Table 3.1: Commonly used MPI functions.

MPI Program Structure

#include <mpi.h>
#include <stdio.h>

int main(int argc, char* argv[]) {

int rank, size;
MPI_Init(&argc, &argv);
MPI_Comm_rank(MPI_COMM_WORLD, &rank);
MPI_Comm_size(MPI_COMM_WORLD, &size);

printf("Hello from process %d of %d\n", rank, size);

MPI_Finalize();
 return 0;
}

3.3 The Trapezoidal Rule in MPI

The trapezoidal rule is a standard numerical integration technique used in the textbook to
illustrate MPI programming. It approximates:

∫ab f(x) dx ≈ h × [f(x0)/2 + f(x1) + f(x2) + ... + f(xn-1) + f(xn)/2]

where h = (b−a)/n is the subinterval width and n is the number of trapezoids.

Parallel Strategy

1. Process 0 (root) reads a, b, n and broadcasts them to all processes.

2. Each process computes a local sum for its assigned sub-range of x.

3. All local sums are reduced (summed) to process 0 using MPI_Reduce .

4. Process 0 prints the result.

/* Parallel Trapezoidal Rule — MPI sketch */

MPI_Bcast(&a, 1, MPI_DOUBLE, 0, MPI_COMM_WORLD);
MPI_Bcast(&b, 1, MPI_DOUBLE, 0, MPI_COMM_WORLD);
MPI_Bcast(&n, 1, MPI_INT, 0, MPI_COMM_WORLD);

h = (b - a) / n;
local_n = n / comm_sz; /* trapezoids per process */
local_a = a + rank * local_n * h;
local_b = local_a + local_n * h;

local_integral = Trap(local_a, local_b, local_n, h);

MPI_Reduce(&local_integral, &total, 1,
MPI_DOUBLE, MPI_SUM, 0, MPI_COMM_WORLD);
 if (rank == 0)
printf("With n = %d trapezoids, integral = %.15e\n", n, total);

3.4 Dealing with I/O in MPI

MPI provides no special I/O primitives in its basic specification. Standard practice:

• Process 0 handles all I/O: it reads input and broadcasts to others; it receives results
and writes output. This avoids conflicts when multiple processes write to the same
terminal.

• Parallel I/O (MPI-IO): advanced feature ( MPI_File_open , MPI_File_read )

allowing all processes to read/write a shared file in parallel with coordinated access.

• Use MPI_Bcast after reading to distribute parameters to all processes.

3.5 Collective Communication

Collective operations involve all processes in a communicator. They must be called by

every process.

Operation Description MPI Function

Broadcast Root sends same data to all processes MPI_Bcast

Scatter Root sends different data to each process MPI_Scatter

Gather Each process sends data to root MPI_Gather

All-Gather Each process sends data; all receive all MPI_Allgather

Reduce Combine (sum/min/max/etc.) all values at root MPI_Reduce

All-Reduce Reduce, but result goes to all processes MPI_Allreduce

Scan (Prefix) Inclusive prefix operation across ranks MPI_Scan

Barrier Synchronise all processes MPI_Barrier

 Operation Description MPI Function

All-to-All Every process sends unique data to every other MPI_Alltoall

Table 3.2: MPI collective communication operations.

MPI Reduce Operators

Pre-defined operators for MPI_Reduce / MPI_Allreduce :

• MPI_SUM — summation

• MPI_PROD — product

• MPI_MAX / MPI_MIN — maximum / minimum

• MPI_LAND / MPI_LOR — logical AND / OR

• MPI_MAXLOC / MPI_MINLOC — max/min with location

3.6 MPI Derived Datatypes

MPI has built-in types: MPI_INT , MPI_DOUBLE , MPI_CHAR , etc. Derived datatypes
allow you to send non-contiguous or structured data in a single MPI call, avoiding manual
packing.

MPI_Type_contiguous
A block of identical elements.

MPI_Type_vector
Elements with a fixed stride between them (e.g., columns of a row-major matrix).

MPI_Type_struct
A C struct with mixed-type fields.

MPI_Type_indexed
Arbitrary list of blocks at specified offsets.

/* Sending a column of a 2D array (strided data) */

MPI_Type_vector(rows, 1, cols, MPI_DOUBLE, &column_type);
MPI_Type_commit(&column_type);
 MPI_Send(&A[0][col], 1, column_type, dest, tag, MPI_COMM_WORLD);
MPI_Type_free(&column_type);

3.7 Performance Evaluation of MPI Programs

Key cost model for point-to-point communication:

Tcomm = tstartup + n × tdata

where tstartup = message latency (startup overhead), n = message size in words, tdata =
per-word transfer time (inverse of bandwidth).

Important Observations

• For small messages: latency (tstartup) dominates — avoid many small messages.

• For large messages: bandwidth (1/tdata) dominates — try to coalesce data into fewer,
larger messages.

• Overlapping communication with computation (using non-blocking MPI_Isend /

MPI_Irecv ) can hide latency.

3.8 A Parallel Sorting Algorithm (Sample Sort)

A simple parallel sorting strategy using MPI is Sample Sort (Bucket Sort variant):

1. Each process sorts its local data.

2. Each process selects p−1 equally-spaced splitter samples from its sorted local data.

3. All samples are gathered at process 0, which selects p−1 global splitters.

4. Global splitters are broadcast to all processes.

5. Each process partitions its local data into p buckets based on the splitters.

6. Each process sends bucket i to process i ( MPI_Alltoallv ).

7. Each process locally merges its received data.

 Remember: The parallel sorting algorithm combines local computation with collective
communication. Choosing good splitters is crucial for load balance.

MODULE 4

Shared-Memory Programming with OpenMP

OpenMP pragmas and directives · Trapezoidal rule · Scope of variables · Reduction clause ·
Loop-carried dependency · Scheduling · Producers and consumers · Caches, cache
coherence & false sharing · Tasking · Thread safety

4.1 OpenMP Pragmas and Directives

OpenMP (Open Multi-Processing) is a directive-based API for shared-memory parallel

programming in C, C++, and Fortran. It uses compiler directives ( #pragma omp ... )
to annotate regions of code to be run in parallel. The programmer specifies parallelism;
the compiler and runtime manage threads.

Fork-Join Model: An OpenMP program begins with a single master thread. When
a parallel region is encountered, the master forks a team of threads. At the end of
the parallel region, threads synchronise and the master continues alone (join).

Core OpenMP Directives

Directive Purpose

#pragma omp parallel Create a parallel region; all threads execute the block.

#pragma omp parallel Parallelise a for loop — iterations divided among threads.
for
 Directive Purpose

#pragma omp critical Mutual exclusion — only one thread at a time executes this
block.

#pragma omp atomic Atomic memory update — lighter than critical for single
operations.

#pragma omp barrier All threads wait until all reach this point.

#pragma omp single Only one thread executes this block; others wait at the end.

#pragma omp master Only the master thread executes this block; no implicit
barrier.

#pragma omp sections Different sections executed by different threads.

#pragma omp task Create an independent task (may run asynchronously).

#pragma omp taskwait Wait for all child tasks to complete.

#pragma omp flush Make thread's view of memory consistent with others.

Table 4.1: Core OpenMP directives.

Runtime Library Functions

• omp_get_thread_num() — returns thread's rank within team (0 to nthreads−1).

• omp_get_num_threads() — returns total number of threads in current team.

• omp_set_num_threads(n) — set number of threads for next parallel region.

• omp_get_wtime() — wall-clock time in seconds.

Minimal OpenMP Program

#include <stdio.h>
#include <omp.h>

int main() {
#pragma omp parallel
{
 int id = omp_get_thread_num();
int total = omp_get_num_threads();
printf("Thread %d of %d\n", id, total);
}
return 0;
}

Compile: gcc -fopenmp program.c -o program

4.2 The Trapezoidal Rule with OpenMP

Using #pragma omp parallel for with a reduction clause:

double h = (b - a) / n;
double integral = (f(a) + f(b)) / 2.0;

#pragma omp parallel for reduction(+:integral)

for (int i = 1; i < n; i++) {
integral += f(a + i * h);
}
integral *= h;

The reduction(+:integral) clause creates a private copy of integral for each

thread, then sums all private copies into the shared variable at the end — correctly and
without race conditions.

4.3 Scope of Variables in OpenMP

Variables in an OpenMP parallel region are classified as shared or private:

shared(var)
All threads access the same memory location for var . This is the default for
variables declared outside the parallel region.

private(var)
 Each thread gets its own local copy, uninitialised. Changes do not affect the original
variable.

firstprivate(var)
Each thread gets a private copy initialised with the value of var at the start of the
parallel region.

lastprivate(var)
After the parallel region, the shared variable is updated with the value from the last
iteration/section.

default(none)
Forces the programmer to explicitly declare the scope of every variable —
recommended for correctness.

Remember: The loop variable in #pragma omp parallel for is automatically

private. Always use default(none) and declare all variables explicitly to avoid
accidental sharing bugs.

4.4 The Reduction Clause

reduction(op:var) is a special clause that safely performs a collective operation

across all threads. Each thread works on a private copy; at the end, all copies are
combined using op .

Supported operators: + , * , - , & , | , ^ , && , || , min , max .

int sum = 0;
#pragma omp parallel for reduction(+:sum)
for (int i = 0; i < n; i++)
sum += a[i];
/* sum now contains the correct total */
4.5 Loop-Carried Dependency

A loop-carried dependency (LCD) occurs when an iteration of a loop depends on data

produced by a previous iteration. This prevents parallel execution of that loop.

/* Loop WITH LCD — cannot parallelise directly */

for (int i = 1; i < n; i++)
a[i] = a[i-1] + b[i]; /* a[i] depends on a[i-1] */

/* Loop WITHOUT LCD — safe to parallelise */

for (int i = 0; i < n; i++)
c[i] = a[i] + b[i]; /* each iteration independent */

Remember: Reduction variables (like a running sum) look like LCDs but are handled
correctly by the reduction clause. True dependence chains (like Fibonacci
recurrences) cannot be trivially parallelised.

4.6 Scheduling in OpenMP

The schedule clause controls how loop iterations are assigned to threads:

Schedule
Behaviour Use When
Type

static Iterations divided into equal-size chunks, Equal iteration costs — good
assigned to threads in round-robin order at load balance with no
compile time. overhead.

static, Chunks of size k, assigned round-robin. Control chunk size explicitly.

k

dynamic Threads claim chunks from a work queue Unequal iteration costs (e.g.,
at runtime; default chunk = 1. variable-length inner loops).
 Schedule
Behaviour Use When
Type

dynamic, Dynamic with chunk size k (fewer queue Balance overhead vs. load
k operations). imbalance.

guided Chunk size starts large and decreases Unknown or irregular

exponentially. Good for irregular work. iteration costs.

runtime Schedule determined at runtime via Tuning without

OMP_SCHEDULE environment variable. recompilation.

auto Compiler/runtime chooses. Let the system decide.

Table 4.2: OpenMP loop scheduling options.

#pragma omp parallel for schedule(dynamic, 4)

for (int i = 0; i < n; i++)
process(tasks[i]); /* tasks have variable cost */

4.7 Producers and Consumers

In the producer-consumer pattern, some threads produce data items that other threads
consume. This requires a shared queue and synchronisation:

• Producer calls enqueue(item) ; consumer calls dequeue(item) .

• The queue must be protected by a mutex (or atomic operations).

• A semaphore or condition variable signals consumers when data is available.

• In OpenMP, this is typically implemented using #pragma omp critical sections

around queue operations.

/* Producer-Consumer with OpenMP */

omp_lock_t queue_lock;
omp_init_lock(&queue_lock);

#pragma omp parallel sections

 {
#pragma omp section /* Producer */
{
for (int i = 0; i < N; i++) {
item = produce();
omp_set_lock(&queue_lock);
enqueue(q, item);
omp_unset_lock(&queue_lock);
}
}
#pragma omp section /* Consumer */
{
for (int i = 0; i < N; i++) {
omp_set_lock(&queue_lock);
item = dequeue(q);
omp_unset_lock(&queue_lock);
consume(item);
}
}
}

4.8 Caches, Cache Coherence, and False Sharing in OpenMP

Even though OpenMP targets shared-memory machines, cache behaviour critically affects
performance.

False Sharing in OpenMP

If two threads update different elements of an array, but those elements happen to be in
the same cache line (typically 64 bytes = 8 doubles), the hardware coherence protocol
causes the cache line to be invalidated and reloaded for every write — even though the
threads are writing to different variables. This is false sharing.

/* BAD: false sharing — all elements in same cache line */

double partial_sum[num_threads];
#pragma omp parallel
{
int t = omp_get_thread_num();
 partial_sum[t] += compute(); /* different index, same cache line! */
}

/* GOOD: use padding to separate to different cache lines */

#define PAD 8 /* doubles per cache line */
double partial_sum[num_threads][PAD];
#pragma omp parallel
{
int t = omp_get_thread_num();
partial_sum[t][0] += compute(); /* each on its own cache line */
}

Remember: False sharing can reduce parallel speedup dramatically. Use private
variables, the reduction clause, or padding to avoid it. The reduction clause
avoids false sharing automatically.

The flush Directive

OpenMP threads may have local (register/cache) copies of shared variables. #pragma
omp flush(var) forces a thread to write its copy of var back to shared memory and

invalidate its local copy. Without flush, threads may see stale values.

4.9 Tasking in OpenMP

OpenMP 3.0 introduced the task construct for expressing irregular parallelism (e.g.,
tree traversal, recursive algorithms) where loop-based parallelism is insufficient.

/* Parallel recursive Fibonacci with tasks */

int fib(int n) {
if (n < 2) return n;
int x, y;
#pragma omp task shared(x)
x = fib(n - 1);
#pragma omp task shared(y)
y = fib(n - 2);
#pragma omp taskwait
 return x + y;
}

Each call to fib creates two tasks that can run independently (and in parallel).
taskwait ensures both child tasks complete before summing.

Task vs. Parallel For

• parallel for — regular parallelism, known loop bounds, equal-ish work per
iteration.

• task — irregular parallelism, dynamic workloads, recursive algorithms, while-

loops.

4.10 Thread Safety

A function or library is thread-safe if it can be called concurrently by multiple threads

without causing race conditions or incorrect results.

Sources of Thread Unsafety

• Global / static variables: if multiple threads read and write the same global variable
without synchronisation, results are undefined.

• Non-reentrant library functions: e.g., strtok() , rand() , localtime()

maintain internal state. Use thread-safe alternatives: strtok_r() , rand_r() ,
localtime_r() .

• Dynamic memory: malloc/free are generally thread-safe in modern

implementations, but use with care in tight loops.

Race Condition Example

/* Race condition — multiple threads increment counter */

int counter = 0;
#pragma omp parallel for
for (int i = 0; i < N; i++)
counter++; /* NOT SAFE: read-modify-write is not atomic */
 /* Fix 1: atomic */
#pragma omp parallel for
for (int i = 0; i < N; i++)
#pragma omp atomic
counter++;

/* Fix 2: critical */
#pragma omp parallel for
for (int i = 0; i < N; i++) {
#pragma omp critical
counter++;
}

/* Fix 3: reduction (best performance) */

#pragma omp parallel for reduction(+:counter)
for (int i = 0; i < N; i++)
counter++;

MODULE 5

GPU Programming with CUDA

GPUs and GPGPU · GPU architectures · Heterogeneous computing · Threads, blocks, and
grids · NVIDIA compute capabilities · Vector addition · Returning results from CUDA
kernels · CUDA trapezoidal rule

5.1 GPUs and GPGPU

GPGPU (General-Purpose computing on Graphics Processing Units) refers to using

the GPU for computations beyond graphics. Modern NVIDIA GPUs with CUDA provide
a programmable parallel computing platform capable of teraflops throughput on scientific
and machine-learning workloads.
Why Use GPUs for General Computing?

• Thousands of cores — suited for data-parallel problems.

• Very high memory bandwidth (HBM2/3 reaches 1–3 TB/s).

• Cost-effective FLOPS per dollar vs. CPU clusters.

• Standard CUDA ecosystem: cuBLAS, cuFFT, cuDNN, Thrust.

5.2 GPU Architectures

NVIDIA GPU Architecture Hierarchy

Hardware Unit Description

GPU Device The entire GPU chip. Contains many SMs.

Streaming A collection of CUDA cores, shared memory, registers, warp

Multiprocessor schedulers. CUDA threads run on SMs.
(SM)

CUDA Core A single arithmetic unit (1 FP32 + 1 INT32 op per cycle). Each SM
has 32–128 CUDA cores depending on generation.

Warp The basic unit of scheduling: 32 threads that execute the same
instruction simultaneously (SIMT). The GPU is most efficient when
all 32 warp threads follow the same code path.

Shared Memory / Fast, on-chip memory shared among all threads in a block (~32–96
L1 Cache KB per SM). Programmer-managed.

Global Memory Large off-chip DRAM (e.g., 16–80 GB on modern GPUs). Accessible
(VRAM) by all threads but high latency (~300–600 cycles). Use coalesced
access.

Registers Private per-thread. Very fast. Limited per SM — using too many
reduces occupancy.

Constant Memory
 Hardware Unit Description

64 KB read-only, cached. Good for parameters broadcast to all

threads.

Texture Memory Read-only, cached with 2D spatial locality — good for image
processing.

Table 5.1: NVIDIA GPU memory and compute hierarchy.

5.3 Heterogeneous Computing

In CUDA, computation is split between:

Host
The CPU and its memory. Runs serial control code, manages memory transfers, and
launches kernels.

Device
The GPU and its memory. Executes the parallel kernel code across thousands of
threads.

Typical CUDA Program Flow

1. Allocate memory on the device: cudaMalloc()

2. Copy input data from host to device: cudaMemcpy(...,

cudaMemcpyHostToDevice)

3. Launch kernel on device: kernel<<<gridDim, blockDim>>>(args)

4. Wait for kernel to finish: cudaDeviceSynchronize()

5. Copy results back to host: cudaMemcpy(..., cudaMemcpyDeviceToHost)

6. Free device memory: cudaFree()

5.4 Threads, Blocks, and Grids

CUDA organises the parallel work into a three-level hierarchy:

Thread
The smallest unit of execution. Each thread runs the kernel function independently and
has its own registers and local memory. Identified by threadIdx.x (and .y, .z for
multi-dim).

Block (Thread Block)

A group of threads (up to 1024 per block in modern GPUs) that execute on the same
SM. Threads within a block can synchronise via __syncthreads() and share on-
chip shared memory. Identified by blockIdx.x .

Grid
A collection of blocks launched by a single kernel call. All blocks run the same kernel.
Blocks in different grids cannot communicate directly or synchronise during
execution.

Unique Global Thread Index

For a 1D grid of 1D blocks:

global_idx = blockIdx.x × blockDim.x + threadIdx.x

/* Kernel execution configuration */

int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
my_kernel<<<blocksPerGrid, threadsPerBlock>>>(d_array, N);

Built-in Variable Meaning

gridDim.x Number of blocks in grid (x dimension)

blockDim.x Number of threads per block (x dimension)

blockIdx.x Index of this block within the grid

threadIdx.x Index of this thread within the block

Table 5.2: CUDA built-in dimension variables.

5.5 NVIDIA Compute Capabilities and Device Architectures

NVIDIA assigns each GPU a compute capability (e.g., 3.5, 7.0, 8.0) that indicates the
supported CUDA features. Higher versions support more features and generally have
more SMs/cores.

Compute
Architecture Key Features Example GPU
Cap.

Kepler 3.x Dynamic parallelism, Hyper-Q Tesla K40

Maxwell 5.x Unified memory improvements GTX 980

Pascal 6.x NVLink, HBM2, FP16 support Tesla P100

Volta 7.0 Tensor Cores (matrix ops for DL), Tesla V100
MIG

Turing 7.5 RT Cores (ray tracing), Tensor Cores RTX 2080

2nd gen

Ampere 8.0, 8.6 TF32, BF16 Tensor Cores, sparsity A100, RTX
support 3090

Hopper 9.0 Transformer Engine, FP8 Tensor H100

Cores

Table 5.3: NVIDIA GPU architectures and compute capabilities.

/* Query device properties */

cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, 0);
printf("Compute capability: %d.%d\n", [Link], [Link]);
printf("Max threads per block: %d\n", [Link]);
printf("Shared mem per block: %zu bytes\n", [Link]);
5.6 Vector Addition in CUDA

Vector addition is the "Hello World" of CUDA: compute c[i] = a[i] + b[i] for all i
in parallel.

/* CUDA Kernel: each thread handles one element */

__global__ void vecAdd(double* A, double* B, double* C, int n) {
int i = blockDim.x * blockIdx.x + threadIdx.x;
if (i < n)
C[i] = A[i] + B[i];
}

int main() {
int n = 1024;
size_t size = n * sizeof(double);

/* Allocate host memory */

double *h_A = malloc(size), *h_B = malloc(size), *h_C = malloc(size);
/* ... initialise h_A and h_B ... */

/* Allocate device memory */

double *d_A, *d_B, *d_C;
cudaMalloc(&d_A, size);
cudaMalloc(&d_B, size);
cudaMalloc(&d_C, size);

/* Copy host to device */

cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);

/* Launch kernel */
int threadsPerBlock = 256;
int blocks = (n + threadsPerBlock - 1) / threadsPerBlock;
vecAdd<<<blocks, threadsPerBlock>>>(d_A, d_B, d_C, n);

/* Copy result back to host */

cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

/* Free memory */
cudaFree(d_A); cudaFree(d_B); cudaFree(d_C);
free(h_A); free(h_B); free(h_C);
 return 0;
}

Key Points

• __global__ — kernel runs on GPU, called from CPU.

• __device__ — function runs on GPU, called from GPU.

• __host__ — function runs on CPU, called from CPU (default).

• The guard if (i < n) prevents out-of-bounds access when n is not a multiple of

blockDim.x.

5.7 Returning Results from CUDA Kernels

CUDA kernels are void — they cannot return values directly. Results are returned by:

1. Writing to device memory and copying back with cudaMemcpy .

2. Atomic operations: use atomicAdd , atomicMax , etc., to accumulate into a single

device variable.

3. Parallel reduction: a multi-step kernel reduces an array to a scalar on the device.

/* Returning a sum via atomicAdd */

__global__ void sum_kernel(double* a, double* result, int n) {
int i = blockDim.x * blockIdx.x + threadIdx.x;
if (i < n)
atomicAdd(result, a[i]); /* atomic global accumulation */
}

/* OR: better — use shared memory reduction per block */

__global__ void block_reduce(double* a, double* partial_sums, int n) {
extern __shared__ double s[];
int i = blockDim.x * blockIdx.x + threadIdx.x;
s[threadIdx.x] = (i < n) ? a[i] : 0.0;
__syncthreads();

/* Parallel tree reduction within the block */

for (int stride = blockDim.x / 2; stride > 0; stride /= 2) {
 if (threadIdx.x < stride)
s[threadIdx.x] += s[threadIdx.x + stride];
__syncthreads();
}
if (threadIdx.x == 0)
partial_sums[blockIdx.x] = s[0];
}

Important: __syncthreads() synchronises all threads within a block. It must

be called by all threads in the block or none — calling it conditionally causes
deadlock.

5.8 CUDA Trapezoidal Rule

Applying numerical integration using the trapezoidal rule on the GPU:

Basic Version

__device__ double f(double x) {

return x * x; /* Example: integrate x^2 */
}

__global__ void trap_kernel(double a, double h, int n, double* partial) {

extern __shared__ double s[];
int i = blockDim.x * blockIdx.x + threadIdx.x;
double x = a + i * h;
s[threadIdx.x] = (i > 0 && i < n) ? f(x) : (f(a) + f(a + n * h)) / 2.0;
__syncthreads();

/* Block-level reduction */
for (int stride = blockDim.x/2; stride > 0; stride /= 2) {
if (threadIdx.x < stride)
s[threadIdx.x] += s[threadIdx.x + stride];
__syncthreads();
}
if (threadIdx.x == 0)
partial[blockIdx.x] = s[0];
 }

/* Host: launch kernel, then sum partial results */

/* partial[0..gridDim-1] summed on host × h = total integral */

Improving Performance: Multiple Warps per Block

Use blockDim.x = 256 or 512 to keep all warp schedulers busy on each SM. The shared
memory reduction tree runs in O(log blockDim) steps within the block.

Three Versions Progression (VTU Syllabus)

Version Description Key Improvement

CUDA Basic global memory version; one Establishes correctness.

Trapezoidal Rule thread per trapezoid.
I

CUDA Add shared memory to accumulate Reduces global memory

Trapezoidal Rule within each block before writing to writes by blockDim factor.
II global memory.

CUDA Use blocks with multiple warps Maximises SM occupancy

Trapezoidal Rule (blockDim > 32); use tree reduction in and throughput.
III shared memory.

Table 5.4: Progressive improvements to CUDA Trapezoidal Rule (VTU Syllabus, Module 5).

5.9 Module 5 Summary

Concept Key Point

GPGPU Use GPU's thousands of cores for non-graphics compute tasks.

SIMT 32 threads (one warp) execute same instruction; divergent paths are
serialised.

Thread Hierarchy
Concept Key Point

Thread → Block → Grid. Threads in a block share SM and shared

memory.

Memory Registers > Shared mem > L2 cache > Global mem (increasing
Hierarchy latency).

Kernel Launch kernel<<<gridDim, blockDim>>>(args)

Coalescing Threads in a warp should access consecutive addresses for bandwidth

efficiency.

Shared Memory ~100× faster than global; use as software-managed cache for reused
data.

__syncthreads() Block-level barrier; ensures all threads finish before proceeding.

atomicAdd Thread-safe addition to a single global variable.

Reduction Parallel tree reduction in shared memory is preferred over atomic

operations.

Table 5.5: CUDA key concepts summary.

QUICK REVISION

Key Formulas and Definitions at a Glance

For all 5 modules — exam ready summary

Topic Formula / Definition

Speedup S(p) = Ts / Tp
Topic Formula / Definition

Efficiency E(p) = S(p) / p = Ts / (p × Tp)

Amdahl's Law S(p) = 1 / (f + (1−f)/p); Smax = 1/f as p→∞

Gustafson's Law Sscaled(p) = p − f(p−1)

SIMD Single Instruction, Multiple Data — one control unit, many data streams

MIMD Multiple Instruction, Multiple Data — fully independent processors

UMA Uniform Memory Access — all processors have equal memory latency

NUMA Non-Uniform Memory Access — local memory faster than remote

Cache Coherence Protocol ensuring all copies of a shared variable stay consistent

False Sharing Two threads updating different data on the same cache line — causes
unnecessary invalidations

MPI Comm cost Tcomm = tstartup + n × tdata

CUDA global i = blockIdx.x × blockDim.x + threadIdx.x

index

Warp size 32 threads execute same instruction simultaneously (SIMT)

Hypercube log p (where p = 2d nodes)

diameter

Omega network Θ(p log p) switches; log p stages

cost

OpenMP parallel #pragma omp parallel for [clauses]

for

OpenMP reduction(op:var) — private copies, combined at end

reduction

Thread safety
 Topic Formula / Definition

Function safe to call from multiple threads without synchronisation

issues

Quick Revision Table: all key formulas and definitions.

BCS702 Parallel Computing — Module-wise Notes | VTU, Semester VII, B.E. CSE
Content summarised and structured from: Introduction to Parallel Computing (2nd Ed.) by Ananth Grama,
Anshul Gupta, George Karypis & Vipin Kumar, Addison-Wesley/Pearson, 2003. Modules 4–5 (OpenMP/
CUDA) also reflect standard curriculum content as per VTU BCS702 syllabus. All figures are from the
textbook (Grama et al.) as extracted from the uploaded PDF.