Unit-5 - Gpu Architecture and Programming

Ananth Grama, Anshul Gupta, George Karypis, and Vipin

Kumar
To accompany the text “Introduction to Parallel
Computing”,
Addison Wesley, 2003
 Topic Overview

1.​ Hardware Architecture

2.​ Integrated GPUs
3.​ Multi GPUs
4.​ GPU Architecture
5.​ Memory Handling with CUDA:
a.​ Shared Memory
b.​ Global Memory
c.​ Constant Memory
d.​ Texture Memory
6.​ Introduction to CUDA C
7.​ Parallel Programming in CUDA C
8.​ Thread Cooperation – Shared Memory and Synchronization

​
 Hardware Architecture
Definition: The physical design and arrangement of a computer's
components. It defines how components are connected and how they interact.
Core Physical Components:

●​ CPU (Central Processing Unit): The "brain" that executes instructions.

●​ Memory: (e.g., RAM) Stores data and instructions for the CPU.
●​ I/O Devices: (Input/Output) How the computer interacts with the user
and other devices (e.g., keyboard, monitor, hard drive).
●​ Buses: The communication pathways (Control, Address, and Data Bus)
that transfer data between components.

Fundamental Models:

●​ Von Neumann Architecture: The classic model where instructions and

data are stored in the same memory. (Basis for most modern
computers).
●​ Harvard Architecture: Uses separate memories for instructions and
data.

The 3 Categories of Architecture

1. System Design:

●​ Defines the entire system's structure.

●​ Includes the CPU, memory, I/O, and how they are all connected.

2. Instruction Set Architecture (ISA):

●​ The "programmer's view" of the computer.

●​ Defines the set of all commands (instructions) that the CPU can
understand and execute.
●​ Examples: RISC (Reduced Instruction Set) and CISC (Complex
Instruction Set).

3. Microarchitecture:

●​ The internal design of the CPU itself.

●​ Describes how the ISA is physically implemented (e.g., pipelining,
cache design).
●​ It's the "hardware" that makes the "software" (ISA) run.
Importance of Architecture Design

●​ Performance: A well-designed architecture optimizes data flow and

processing speed, leading to a faster, more responsive system.
●​ Power Efficiency: Smart design and low-power components reduce
energy consumption, which is critical for battery life and environmental
impact.
●​ Scalability & Flexibility: Allows for easier upgrades and expansion (e.g.,
adding more RAM, a better CPU) to adapt to future needs.
●​ Customization: Enables "specialized" hardware tailored for specific
tasks, such as:
a.​ GPUs/TPUs: Designed for high-performance AI and Machine
Learning.
b.​ Embedded Systems: Designed for low-power, specific-task
devices (like in a car or microwave).

Key Trends in Modern Architecture

●​ AI/ML Acceleration: Creating specialized chips (GPUs, TPUs, FPGAs)

to handle the massive calculations required for Artificial Intelligence.
●​ Edge Computing: Processing data locally on the device (like a smart
camera or phone) instead of sending it to the cloud. This reduces
latency and improves privacy.
●​ Heterogeneous Computing: Integrating different types of processors
(e.g., a CPU, a GPU, and an AI accelerator) into one system, using the
best chip for each specific task.
●​ Memory-Centric Design: As data sets grow, new designs focus on
reducing the "bottleneck" of moving data between memory and the CPU
(e.g., In-Memory Computing).
●​ Hardware-Level Security: Building security features directly into the
processor (like Secure Enclaves) to protect sensitive data from cyber
threats
 Integrated GPUs (iGPUs)

What is an Integrated GPU (iGPU)?

●​ Definition: An Integrated Graphics Processing Unit (GPU) is a graphics

processor that is built directly into the same chip as the Central
Processing Unit (CPU).
●​ No Separate Card: Unlike a "dedicated" GPU, it is not a separate, large
card that you plug into the motherboard.
●​ Shared Resources: The iGPU shares system RAM (your computer's
main memory) with the CPU, rather than having its own dedicated video
memory (VRAM).
●​ Found In: Most standard laptops, budget desktops, tablets, and
smartphones.

●​

Who is an iGPU For?

Integrated graphics are the perfect choice for everyday users and tasks, such
as:

●​ Web browsing, email, and streaming video (e.g., YouTube, Netflix).

●​ Office applications (Word, Excel, PowerPoint).
●​ Light photo editing.
●​ Casual or older 2D/3D games.
It is NOT suitable for:

●​ Serious gamers who want to play modern, graphically intense titles (like
Cyberpunk 2077 or Starfield).
●​ Video editors working with high-resolution 4K/6K+ footage.
●​ AI/Machine Learning researchers.
●​ Professional 3D artists and animators.

IMPORTANT QUESTION - List the advantages and limitations of using

integrated GPUs for CUDA applications.

Advantages

1.​Unified Memory Space: Integrated GPUs often share the main system
memory with the CPU. This architecture (like AMD's HSA) eliminates
distinct memory pools.
2.​No Communication Overhead: Because memory is shared, there's no
need to explicitly transfer data between CPU memory and GPU memory
over a bus like PCIe, which introduces latency and overhead with
discrete GPUs.
3.​Simplified Programming: Developers don't need to manage the
complexities of explicit data transfers between host and device memory.
4.​Better Resource Integration: The close coupling of CPU and GPU
cores on the same die potentially allows for better and more efficient
utilization of computing resources, assigning tasks to the most suitable
unit (CPU for latency-sensitive, GPU for throughput-sensitive) without
bus transfer penalties.

(Page Reference: Primarily pages 26, 30, 391-392.)

Limitations

1.​CUDA Incompatibility (Implied): The text explicitly states that CUDA is

an Nvidia-specific technology. The primary examples of integrated
GPUs discussed for their unified memory benefits are AMD APUs and
Intel integrated graphics (mentioned with OpenCL, not CUDA).
Therefore, based solely on this text, standard integrated GPUs from
AMD or Intel wouldn't run CUDA applications, which is a major limitation
for this specific scenario. (While Nvidia does have integrated solutions
like Tegra, the text doesn't discuss their CUDA capabilities in this
comparative context).
2.​Lower Performance: Integrated GPUs typically share the main system
RAM (like DDR3/DDR4) with the CPU, which usually has lower
bandwidth compared to the dedicated GDDR memory found on discrete
graphics cards. They also compete with the CPU for this memory
bandwidth.
3.​Compromised Design: Integrating CPU and GPU cores onto the same
silicon die involves design trade-offs. Both the CPU and GPU
components might be less powerful than their discrete counterparts due
to limitations on die space, power consumption, and heat dissipation.
4.​Target Market: Integrated solutions are often found in lower-cost or
power-constrained systems (laptops, mainstream desktops) and
generally don't offer the raw computational power needed for
high-performance computing tasks compared to high-end discrete
GPUs.

(Page Reference: Primarily pages 26, 391-392.)

 Multi GPUs
Key takeaways:

●​ Multi-GPU computing involves using two or more graphics processing

units in parallel to tackle a computational task, significantly boosting
processing power and speed for compatible workloads.
●​ This approach is common in AI model training, high-end gaming, and
simulations, where dividing the workload among multiple GPUs can
reduce execution time and handle larger datasets or more complex
computations.
●​ Effective multi-GPU computing requires specialized software or
frameworks to distribute tasks across GPUs and manage
communication between them, ensuring the workload is balanced and
results are synchronized.

What is multi-GPU computing?

Multi-GPU computing is a hardware setup that uses multiple graphics

processing units (GPUs) together to handle complex computational tasks.
This parallel processing approach distributes workloads across two or more
GPUs to simultaneously process different portions of the same task.
Multi-GPU systems split the work to achieve faster processing times and
handle larger datasets.
Types of multi-GPU configurations

Multi-GPU systems come in several configurations. Each is designed for

specific use cases and performance requirements. The choice between these
setups depends on factors like workload type, scalability, and budget.

●​ Single-system multi-GPU: Multiple GPUs installed in one workstation,

connected through the motherboard. This setup is common for
individual developers and small teams working on AI model
development or graphics rendering.
●​ Direct GPU-to-GPU connection: GPUs communicate directly through
high-speed interconnects like NVIDIA NVLink or AMD Infinity Fabric.
This configuration minimizes latency and maximizes bandwidth.
●​ Network-based GPU clusters: Multiple computers (each with one or
more GPUs) connected through high-speed networks. Often used in
data centers and cloud environments.
●​ Distributed multi-GPU systems: GPUs spread across different physical
locations but working together through specialized software.
●​ Hybrid GPU configurations: A mix of different GPU types working
together, such as combining consumer and professional-grade GPUs.

IMPORTANT QUESTION - Describe the difference between multiple GPUs in

separate PCI slots, dual GPU boards, and integrated vs. discrete GPUs with
respect to memory access and communication overhead

The difference between multiple GPUs in separate PCI slots, dual-GPU

boards, and integrated vs. discrete GPUs mainly lies in how they access
memory and the amount of communication overhead involved.

1.​Multiple GPUs in Separate PCI Slots:​

Each GPU has its own dedicated memory and connects to the CPU
through the PCI Express (PCIe) bus. Although this allows independent
operation, communication between GPUs involves data transfer across
the PCIe link, which is relatively slow compared to on-chip
communication. This results in high communication overhead and
explicit data transfer between host and device memories.​

2.​Dual-GPU Boards:​
A dual-GPU board integrates two GPUs on the same physical card
and may use technologies such as NVIDIA SLI or AMD CrossFire.
 These setups allow direct data transfer between the GPUs’
memories without always involving the CPU, leading to lower latency
and faster inter-GPU communication than separate PCI-slot GPUs.
However, kernel execution is still limited to one GPU at a time, and data
coordination requires careful management.​

3.​ Integrated vs. Discrete GPUs:​

○​ Integrated GPUs, such as those found in APUs (Accelerated

Processing Units), share the same physical memory as the
CPU. This eliminates the need for explicit data transfers, resulting
in minimal communication overhead. However, performance is
limited by shared bandwidth and slower memory subsystems.​

○​ Discrete GPUs, on the other hand, have dedicated

high-bandwidth memory (GDDR5/GDDR6) and operate
independently of system memory. They offer higher
computational performance, but data must be copied between
CPU and GPU memory, increasing latency and overhead.

In summary:

●​ Integrated GPUs → Low overhead, shared memory, but lower

performance.
●​ Discrete GPUs (single or dual) → High performance, but higher data
transfer cost.
●​ Multiple GPUs in separate PCI slots → Independent but require
high-overhead communication through PCIe.

Textbook Reference:

Chapter 6 – GPU Programming, pp. 391–393 and pp. 480–482

(Sections on “Existing Architectures for CPU-GPU Systems” and “CUDA and

MPI – Multi-GPU Systems”)
 Memory Handling with CUDA

●​ What It Is: The main, large graphics memory (VRAM) on the GPU card.
This is the primary storage for your large datasets.
●​ Scope: Accessible by all threads in the grid, as well as the host (CPU).
●​ Use Case:
○​ Initial data transfer from CPU to GPU.
○​ Storing large input datasets and final output results.
●​ Key Characteristics:
○​ Pro: Largest capacity (measured in Gigabytes).
○​ Con: Highest latency (slowest access speed). Accessing it is a
major performance bottleneck.
○​ Best Practice: Accesses should be "coalesced," meaning threads
in a block access memory in a sequential, aligned pattern to
maximize bandwidth.
Shared Memory

●​ What It Is: A very small, very fast on-chip memory.

●​ Scope: Shared by all threads within a single block. Data is not visible to
other blocks.
●​ Use Case:
○​ Inter-thread communication: Allows threads in a block to
cooperate and exchange data.
○​ Data reuse: A block can load a "tile" of data from slow Global
Memory into fast Shared Memory once, and then all threads in
the block can access it many times.
●​ Key Characteristics:
○​ Pro: Extremely fast (approaching register speed).
○​ Con: Very small (measured in Kilobytes per block).
○​ Analogy: It's like a tiny, shared "scratchpad" for each team (block)
of workers (threads).

Constant Memory

●​ What It Is: A small, read-only memory space on the GPU.

●​ Scope: Accessible by all threads in the grid.
●​ Use Case:
○​ Storing data that is the same for all threads and does not change
during the kernel's run.
○​ Examples: Filter coefficients, physical constants, algorithm
parameters.
●​ Key Characteristics:
○​ Pro: Very fast when all threads in a "warp" (group of 32 threads)
read the same address. The value is broadcast to all threads.
○​ Pro: It is cached, so subsequent reads are even faster.
○​ Con: Read-only and small (typically ~64 KB).

Texture Memory

●​ What It Is: A read-only memory space optimized for a specific access

pattern.
●​ Scope: Accessible by all threads in the grid.
●​ Use Case:
○​ Access patterns with spatial locality, common in image processing
or simulations.
○​ Hardware-accelerated filtering (e.g., linear interpolation) and
addressing modes (e.g., clamping or wrapping at edges).
 ●​ Key Characteristics:
○​ Pro: Has a dedicated on-chip cache.
○​ Pro: Specialized hardware can perform filtering and boundary
checks automatically.
○​ Con: Read-only. (Note: Modern CUDA also has "Surfaces" which
are similar but read-write).
○​ Analogy: Think of it as a "smart" memory that understands 2D/3D
space, not just a 1D array.

○​

IMPORTANT QUESTION - Describe the memory hierarchy in NVIDIA CUDA

GPU architecture, including global memory, shared memory, and caches. How
do access patterns affect performance?

Memory Hierarchy Components

1.​Registers: These are the fastest memory available, located on-chip

within each Streaming Multiprocessor (SM). Each thread has its own
private set of registers, lasting for the thread's lifetime. They are used
for frequently accessed automatic variables declared within a kernel.
Overuse can lead to spilling into slower local memory .
2.​Local Memory: Despite the name, local memory physically resides in
off-chip DRAM (the same chips as global memory). It is used when
registers spill or for automatic variables (like arrays) declared within a
kernel that are too large for registers. Access is private per thread and
lasts for the thread's lifetime. Access latency is high, similar to global
memory, although it can be cached in L1 and L2 caches.
 3.​Shared Memory: This is a fast on-chip RAM located within each SM.
It's accessible only by threads within the same block and lasts for the
block's lifetime. It offers much lower latency than global/local memory
(around 20-32 cycles vs 400-600) and is used for fast data sharing and
cooperation between threads in a block, acting as a user-managed
cache.
4.​L1 and L2 Caches: These are hardware-managed caches. L1 cache is
on-chip within each SM, often physically sharing the same RAM as
shared memory (the split is configurable). L2 cache is typically larger,
on-chip, and shared across all SMs. They cache accesses to local and
global memory, helping to reduce latency. Access is transparent to the
programmer.

5.​Global Memory: This is the main GPU memory, residing in off-chip

DRAM. It has the largest capacity but also the highest latency
(400-600 cycles). It is accessible by all threads in the grid as well as
the host (CPU) via memory copy operations. Its lifetime persists across
kernel launches until explicitly freed.
6.​Constant Memory: Resides in off-chip DRAM but is cached on-chip
(e.g., 8KB per SM). It's limited in size (e.g., 64KB total) and is read-only
from the device perspective. It's efficient when multiple threads read the
same location (broadcast). Lifetime is for the application duration.
7.​Texture/Surface Memory: Resides in off-chip DRAM and is also
cached on-chip. It's accessed via special hardware units offering
features like filtering and addressing modes useful for graphics and
some computations. Lifetime is for the application duration.

How Access Patterns Affect Performance

Access patterns significantly impact performance due to the hardware's

characteristics:

●​ Global Memory Coalescing: The GPU hardware tries to group

(coalesce) memory accesses from threads within a warp (typically 32
threads) into fewer, larger transactions (e.g., 32-byte or 128-byte
segments) . Performance is best when threads in a warp access
contiguous, aligned memory locations. Scattered or misaligned
accesses result in multiple transactions, drastically reducing effective
bandwidth. Using pitched memory allocation (cudaMallocPitch) for 2D
 arrays and preferring Structure-of-Arrays (SoA) over Array-of-Structures
(AoS) can improve coalescing .
●​ Shared Memory Bank Conflicts: Shared memory is divided into banks
(e.g., 32 banks). If multiple threads within a warp access different
locations within the same memory bank simultaneously, a bank conflict
occurs. The accesses are serialized, causing stalls and reducing
throughput. Broadcasting (one location read by multiple threads) or
accessing different banks causes no conflicts. Padding data structures
or using staggered access patterns can mitigate bank conflicts .
●​ Cache Utilization: Access patterns exhibiting good spatial locality
(accessing nearby memory locations) and temporal locality (accessing
the same location multiple times) benefit most from the L1 and L2
caches, reducing effective global memory latency. Constant and texture
memory benefit significantly from their dedicated caches when access
patterns are suitable (e.g., uniform access across a warp for constant
memory).

Efficient CUDA programming requires structuring memory accesses to

maximize coalescing for global memory, minimize bank conflicts for shared
memory, and leverage cache locality.

(Page Reference: Overall: 410-418, 424-425, 432, 446-448, 454-458)

IMPORTANT QUESTION - Explain the difference between global and

shared memory in CUDA with respect to:​
a) Access latency​
b) Scope​
c) Lifetime​
d) Typical usage

a) Access Latency

●​ Global Memory: Has high access latency because it resides in

off-chip DRAM. Accessing it can take 400-600 clock cycles.
●​ Shared Memory: Has low access latency because it is fast on-chip
RAM located within each Streaming Multiprocessor (SM). Accessing it
typically takes around 20-32 clock cycles.

(Page Reference: Primarily pages 414-415, discussing the memory hierarchy

and Figure 6.8.)
b) Scope

●​ Global Memory: Accessible by all threads within the entire grid

launched by a kernel. It is also accessible by the host (CPU) through
CUDA API calls for data transfers.
●​ Shared Memory: Accessible only by threads within the same thread
block. Data in one block's shared memory cannot be accessed by
threads in another block.

(Page Reference: Primarily Table 6.5 on page 415 and Section 6.6.2 on page
417.)

c) Lifetime

●​ Global Memory: Persists for the lifetime of the application or until it is

explicitly deallocated (freed) via API calls from the host.
●​ Shared Memory: Persists only for the lifetime of the thread block.
Once a block finishes execution, its shared memory contents are lost.

(Page Reference: Primarily Table 6.5 on page 415.)

d) Typical Usage

●​ Global Memory: Used for holding the main input datasets transferred
from the host, storing the final results to be transferred back to the
host, and for data that needs to be shared between different thread
blocks.
●​ Shared Memory: Used as a user-managed cache for frequently
accessed data within a block to reduce slow global memory accesses.
It's crucial for facilitating fast data sharing and cooperation between
threads within the same block
 Introduction to CUDA C

The CUDA Programming Model

●​ Heterogeneous Model: CUDA uses both the CPU and the GPU.
○​ Host: The CPU and its memory (e.g., system RAM).
○​ Device: The GPU and its memory (e.g., VRAM).
●​ Host vs. Device Roles:
○​ The Host (CPU) manages the overall application, allocates
memory on both host and device, and launches tasks.
○​ The Device (GPU) executes the parallel tasks.
●​ Kernel: A C function, marked with __global__, that runs on the Device
and is executed by many threads in parallel.

The Standard CUDA Program Flow

A typical CUDA application follows these 5 steps:

1.​ Allocate Memory: Allocate memory on both the Host (using malloc) and
the Device (using cudaMalloc).
2.​ Initialize Data: Create and initialize your input data on the Host (CPU).
3.​ Transfer Data (Host ➔ Device): Copy input data from Host RAM to
Device VRAM using cudaMemcpy.
4.​ Launch Kernel: The Host calls the __global__ kernel function to execute
on the Device (GPU).
5.​ Transfer Results (Device ➔ Host): Copy results from Device VRAM
back to Host RAM using cudaMemcpy.

(...followed by freeing memory on both Host (free) and Device (cudaFree))

Host Code: Managing the Operation (main)

The Host (main function) is the "controller."

●​ Memory Allocation: It manages two separate memory spaces.

○​ float *x, *y; // Host pointers
○​ float *d_x, *d_y; // Device pointers
○​ x = (float*)malloc(N*sizeof(float)); // Allocates on CPU
○​ cudaMalloc(&d_x, N*sizeof(float)); // Allocates on GPU
 ●​ Memory Copy: It controls data movement.
○​ cudaMemcpy(..., cudaMemcpyHostToDevice); // To GPU
○​ cudaMemcpy(..., cudaMemcpyDeviceToHost); // From GPU
●​ Cleaning Up:
○​ cudaFree(d_x); // Frees GPU memory
○​ free(x); // Frees CPU memory

Launching a Kernel

●​ Execution Configuration: The <<< ... >>> syntax tells the GPU how to
run the kernel. saxpy <<< numBlocks, threadsPerBlock >>> (N, 2.0f,
d_x, d_y);
●​ Thread Hierarchy: CUDA organizes threads into a grid of thread blocks.
●​ Parameters:
○​ numBlocks: The first argument. It sets the total number of thread
blocks in the grid (e.g., (N+255)/256).
○​ threadsPerBlock: The second argument. It sets the number of
threads inside each block (e.g., 256).

Device Code: The Parallel Kernel (__global__)

The Kernel is the code that runs on thousands of GPU threads at once.

●​ Kernel Definition:
○​ __global__ void saxpy(int n, float a, float *x, float *y)
○​ __global__ identifies it as code to be run on the Device, called
from the Host

●​ The "Magic": Finding a Unique ID

○​ How does each thread know which element i to process?
○​ int i = blockIdx.x * blockDim.x + threadIdx.x;
●​ Built-in Variables:
○​ threadIdx.x: The thread's ID within its block (e.g., 0 to 255).
○​ blockIdx.x: The block's ID within the whole grid.
○​ blockDim.x: The total number of threads in one block (e.g., 256).
○​ This formula gives every thread a unique, global index i.
 ●​ Bounds Check (Safety First!):
○​ if (i < n) y[i] = a*x[i] + y[i];
○​ This is essential! It prevents threads from trying to access
memory that doesn't exist (e.g., if N = 1000 but you launched
1024 threads).

●​

IMPORTANT QUESTION - Describe the steps involved in writing a simple

CUDA program (e.g., Vector Addition):​
a) Starting from the host​
b) Copying data​
c) Executing on the device​
d) Getting the result back

a) Starting from the host (CPU)

1.​Include Headers: Include cuda.h for CUDA functions and relevant

C/C++ headers like stdio.h and stdlib.h.
2.​Allocate Host Memory: Use standard C++ (new) or C (malloc) to
allocate memory on the host system for the input vectors (e.g., ha, hb)
and the result vector (hc).
 3.​Initialize Input Data: Populate the host input vectors with the
necessary data.

(Page Reference: Primarily pages 411-412, Listing 6.5, lines 2-4, 32-34,
42-46)

b) Copying data (Host -> Device)

1.​Allocate Device Memory: Use cudaMalloc() to allocate memory in the

GPU's global memory for the device-side versions of the input and
result vectors (e.g., da, db, dc). Separate pointers are required for
device memory.
2.​Copy Input Data: Use cudaMemcpy() with the
cudaMemcpyHostToDevice kind specified to transfer the initialized data
from the host vectors to the corresponding device vectors.

(Page Reference: Primarily pages 410-412, Listing 6.5, lines 36-39, 49-50)

c) Executing on the device (GPU)

1.​Write the Kernel Function: Define the GPU function (kernel) using the
__global__ specifier. This function contains the code executed by each
GPU thread. Kernels called from the host must be declared void.
2.​Determine Thread Index: Within the kernel, each thread computes its
unique global index using built-in CUDA variables such as threadIdx.x,
blockIdx.x, and blockDim.x.
3.​Perform Computation: Each thread uses its index to read elements
from the input device arrays, perform the vector addition (c[myID] =
a[myID] + b[myID];), and store the result in the output device array. It's
advisable to include boundary checks to prevent out-of-bounds access
if the total number of threads exceeds the data size.
4.​Configure and Launch Kernel: On the host, calculate the necessary
grid and block dimensions (e.g., using dim3) to ensure all data elements
are processed. Launch the kernel from the host using the <<<grid,
block>>> syntax, passing device pointers and data size as arguments.
5.​Synchronize (if needed): Because kernel launches are asynchronous,
call cudaDeviceSynchronize() on the host if subsequent host operations
depend on the kernel's completion.
(Page Reference: Primarily pages 394-399, 411-413, Listing 6.5, lines 19-24,
52-53, 55)

d) Getting the result back (Device -> Host)

1.​Copy Result Data: Use cudaMemcpy() with the

cudaMemcpyDeviceToHost kind specified to transfer the computed
results from the device result vector back to the host result vector.
2.​Use/Verify Result on Host: The host can now access the computed
results in its memory for display, further processing, or verification
against a CPU-computed result.
3.​Clean Up Memory: Deallocate the device memory using cudaFree()
and the host memory using delete[] or free(). Optionally, call
cudaDeviceReset() before exiting.

(Page Reference: Primarily pages 410-414, Listing 6.5, lines 60, 63-67, 70-76)

IMPORTANT QUESTION - What is the role of the NVCC compiler in CUDA?

NVCC (Nvidia CUDA Compiler) acts as a compiler driver. Its main role is to
manage the complex process of compiling CUDA source code files (.cu files),
which contain a mix of standard C/C++ (host code) and CUDA-specific C/C++
extensions (device code).

Here's a breakdown of its functions:

1.​Code Separation: NVCC first processes the .cu file and separates the
code intended for the host (CPU) from the code intended for the device
(GPU).
2.​Host Code Compilation: It invokes a standard C/C++ host compiler
(like GCC on Linux/Mac or [Link]/Visual Studio's compiler on Windows)
to compile the host code into standard object files.
3.​Device Code Compilation: It uses internal CUDA compiler
components (cudafe, nvopencc, ptxas) to compile the device code. This
involves:
○​ Generating intermediate code (like PTX - Parallel Thread
Execution assembly).
○​ Generating binary code (CUBIN) specific to one or more target
GPU architectures.
4.​Stub Generation: NVCC generates necessary "stub" functions on the
host side that handle the setup and launching of the device kernels.
 5.​Linking: Finally, NVCC calls the host linker (like g++ or [Link]) to link
the compiled host object files, the device code (often embedded within
host object files as "fat binaries"), and necessary CUDA runtime
libraries into a single executable file.

In essence, NVCC orchestrates the entire build process, using both standard
host compilers and specialized CUDA compilation tools to produce an
executable that can run code on both the CPU and the GPU. It also handles
options like specifying target GPU architectures (-arch, -code) and enabling
separate compilation (-dc, -rdc) .

(Page Reference: Primarily pages 403-409.)

Parallel Programming in CUDA C

CUDA: A Platform for Massive Parallelism

●​ Goal: To achieve massive speedup on data-parallel problems by moving

computation from the Host (CPU) to the Device (GPU).
●​ Why a GPU? A CPU has a few powerful cores (optimized for latency). A
GPU has thousands of simpler cores (optimized for throughput).
●​ Heterogeneous Computing: CUDA is the model that allows the CPU
and GPU to work together as co-processors to solve a single problem.
●​ Kernel: The C/C++ function (__global__) that represents the parallel
part of your program.

Work Decomposition (Grid, Block, Thread)

In parallel computing, you must decompose a large problem into smaller

tasks. CUDA uses a two-level hierarchy for this:

●​ Thread: The smallest unit of parallel work. It executes one instance of

the Kernel.
●​ Thread Block: A group of threads. This is the key unit for cooperation.
○​ Threads within a block can share data via fast Shared Memory.
○​ They can coordinate using a barrier: __syncthreads().
●​ Grid: The collection of all Thread Blocks for a single Kernel launch.
 Blocks must be independent and can be scheduled in any order,
which enables massive scalability.

The Parallel Memory Model (Minimizing Latency)

Data locality is the biggest challenge in parallel performance. CUDA uses an

explicit memory hierarchy to manage this.

●​ Host (CPU) Memory: System RAM.

●​ Device (GPU) Global Memory: Large (GBs), but high latency (slow).
This is the main performance bottleneck.
●​ Device (GPU) Shared Memory: Small (KBs), but low latency (fast).
○​ Scope: Shared by all threads in one block.
●​ Parallel Strategy: Minimize slow global access. Cooperatively load a
"tile" of data from Global Memory into Shared Memory once, then have
all threads in the block access the fast Shared Memory many times.

Implementing Parallelism in CUDA C

●​ Defining Work (Kernel): The __global__ keyword defines the parallel

function that all threads will run.
○​ __global__ void myKernel(...)
●​ Distributing Work (Thread ID): Each thread finds its unique task ID.
○​ int idx = blockIdx.x * blockDim.x + threadIdx.x;
 ●​ Launching Work (Host Code): The CPU launches the entire grid of
parallel threads.
○​ myKernel<<< numBlocks, threadsPerBlock >>>( ... );
●​ Execution Model (SIMT): Threads are executed in groups of 32 (a
"warp"). All threads in a warp execute the same instruction at the same
time.

The Host-Device Interaction Pattern

This model requires explicit memory management from the programmer.

1.​ Allocate: Allocate memory on both the Host (malloc) and the Device
(cudaMalloc).
2.​ Transfer (Host ➔ Device): Manually copy input data from CPU RAM to
GPU VRAM.
○​ cudaMemcpy(..., cudaMemcpyHostToDevice);
3.​ Launch (Execute Parallel Kernel): The Host launches the Kernel on the
Device. This call is asynchronous—the Host continues immediately.
4.​ Synchronize (Wait for Parallel Work): The Host must explicitly wait for
the Device to finish.
○​ cudaThreadSynchronize();
5.​ Transfer (Device ➔ Host): Manually copy results from GPU VRAM back
to CPU RAM.
○​ cudaMemcpy(..., cudaMemcpyDeviceToHost);
 Thread Cooperation – Shared Memory and
Synchronization
Here are the key concepts from your material, distilled into five concise slides
for your presentation.

What is Shared Memory?

●​ Definition: A high-speed, on-chip memory space that is shared by all

threads within a single thread block.
●​ Keyword: Declared in CUDA C using the __shared__ qualifier.
●​ Scope: Data is visible to all threads in its block, but not visible to other
blocks.
●​ Analogy: Think of it as a small, fast "whiteboard" that a single team (a
block) can use to share notes and intermediate results.
●​ Warning: Because it's shared, it creates the risk of race conditions.

The "Race Condition" Problem

A race condition is the most common bug in parallel programming.

●​ What it is: An error that occurs when multiple threads try to read, modify,
and write to the same memory location at the same time.
●​ Example (Naive Histogram):
○​ Thread A reads output[5] (value is 10).
○​ Thread B reads output[5] (value is 10).
○​ Thread A calculates 10 + 1 = 11.
○​ Thread B calculates 10 + 1 = 11.
○​ Thread A writes 11 to output[5].
○​ Thread B writes 11 to output[5].
●​ Result: The final value is 11, but it should be 12. The program is
functionally wrong.

Solution 1: atomicAdd()

●​ What it is: A special function that performs a read-modify-write

operation as a single, uninterruptible ("atomic") step.
●​ How it works: It "locks" the memory address, performs the addition, and
"unlocks" it. This prevents race conditions.
●​ CUDA Code: atomicAdd(&(output[input[item]]), 1);
●​ Performance:
○​ Correct? Yes.
 ○​ Fast? No. When used on Global Memory, it's very slow.
○​ Why? Accessing global memory has high latency, and if many
threads target the same address, they are forced to line up and
execute one by one (serialization).

Solution 2: Shared Memory (The Better Way)

This is the optimal strategy, as it minimizes slow global memory access.

●​ The Strategy:
○​ Each block creates a private, temporary histogram in fast
__shared__ memory.
○​ All threads in the block use atomicAdd on this fast shared
memory histogram. (Conflicts are fast to resolve here).
○​ Once the block is finished, each thread copies one bin from the
shared histogram to the Global Memory histogram.
●​ Benefits:
○​ Turns thousands of slow, high-conflict global atomics...
○​ ...into thousands of fast shared atomics and one fast, conflict-free
(coalesced) write to global memory per block.
●​ Result: This is often 10x faster (or more) than the global atomicAdd
solution.
The Final Piece: __syncthreads()

●​ Problem: The shared memory strategy still has a race condition! How
do you know all threads are finished with step 2 before starting step 3?
●​ Solution: __syncthreads();
●​ What it is: A barrier. No thread in a block can execute past this line until
all threads in that block have reached it.
●​ What it guarantees:
1.​ All threads have "caught up" to the same point.
2.​ All previous writes to shared memory are now visible to all
threads in the block.
●​ Final Correct Flow:
1.​ Declare __shared__ int temp_hist[256];
2.​ Initialize shared memory (temp_hist[threadIdx.x] = 0;).
3.​ __syncthreads(); (Wait for initialization to finish)
4.​ Compute shared histogram (atomicAdd(&(temp_hist[input[item]]),
1);).
5.​ __syncthreads(); (Wait for computation to finish)
6.​ Write to global memory (atomicAdd(&(output[threadIdx.x]),
temp_hist[threadIdx.x];)

IMPORTANT QUESTION - What are grids, blocks, and threads in CUDA C,

and how are they related to expressing parallelism?

Definitions and Relationship

●​ Threads: A thread is the basic unit of execution in CUDA. Each

thread executes an instance of the kernel function independently.
Threads within the same block can cooperate using shared memory
and synchronization primitives .
●​ Blocks: Threads are grouped into blocks. A block is a collection of
threads that can cooperate efficiently via fast shared memory and can
synchronize their execution . Blocks can be organized in one, two, or
three dimensions. All threads within a block execute on the same
Streaming Multiprocessor (SM) on the GPU.
●​ Grids: Blocks are grouped into a grid. A grid represents the entirety of
the threads launched for a specific kernel invocation. Grids can also be
organized in one, two, or three dimensions. All blocks within a grid
execute the same kernel function.
The hierarchy is Grid -> Block -> Thread. A kernel is launched by specifying
the dimensions of the grid (number of blocks) and the dimensions of each
block (number of threads per block).

Expressing Parallelism

This grid/block/thread hierarchy allows CUDA to express massive parallelism

in several ways:

1.​Scalability: By launching a large grid of blocks, each containing many

threads (potentially thousands or millions in total), a computation can be
broken down into many small, independent (or semi-independent within
a block) tasks suitable for parallel execution on the GPU's numerous
cores.
2.​Mapping to Hardware: The hierarchy maps naturally to the GPU
hardware. Blocks are scheduled to run on the available SMs, and
threads within a block are executed by the cores within that SM.
3.​Thread Identification: Each thread has unique built-in index variables
(threadIdx, blockIdx) that specify its position within its block and its
block's position within the grid. Threads use these indices to determine
which portion of the overall workload or data they are responsible for
processing, enabling data parallelism where the same operation is
applied to different data elements concurrently .
4.​Cooperation Model: The block structure allows for a limited but
efficient form of cooperation between threads (synchronization and
shared memory), enabling more complex parallel algorithms than purely
independent thread execution would allow .

In essence, the grid/block/thread model provides a structured way to define

and manage a vast number of parallel threads, map them onto the GPU's
parallel architecture, and assign specific computational tasks to each thread,
thereby achieving massive parallelism.

(Page Reference: Primarily pages 394-398, 400, 417, 420-421.)

IMPORTANT QUESTION - (Case Study) — Assume we want to run a CUDA
program where multiple threads access a shared variable. Why is
Synchronization
important here?​
Name one CUDA function that helps with synchronization, and write one
drawback if synchronization is not used.

Synchronization Function

One CUDA function that helps with synchronization within a thread block is:

●​ __syncthreads()

This function acts as a barrier, ensuring that all threads within the same block
reach this point before any thread in that block is allowed to proceed further.
This guarantees that memory operations (like reads from or writes to shared
memory) initiated before the barrier are completed and visible to all threads in
the block after the barrier.

Drawback Without Synchronization

If synchronization is not used appropriately, a major drawback is obtaining

incorrect final results due to race conditions. For example, if multiple
threads increment a shared counter in shared memory without
synchronization (like using atomicAdd() or ensuring exclusive access via logic
combined with __syncthreads()), some increments might be lost because
threads could read the counter's value, increment their local copy, and write it
back, potentially overwriting increments made by other threads between their
read and write operations .

(Page Reference: Primarily pages 65-68, 72, 78-79, 417-418, 420-421,

452-453)