1.

With a neat diagram Explain the architecture of modern GPUs in detail and
Discuss the concepts of SIMD, SMs, SPs, SIMT.

⏷ Modern GPU Architecture

<
_1
+̇

The GPU architecture follows a hierarchical structure built around numerous simple
processing units.
1. Global Memory (Device Memory)
This is the Video RAM (VRAM), typically high-speed memory like GDDR6 or HBM (High
Bandwidth Memory). It is the largest and highest latency memory on the GPU, shared by all
processing elements.
2. L2 Cache
This unified, high-capacity cache acts as a shared pool for all Streaming Multiprocessors
(SMs), sitting between them and the slower Global Memory.
3. Streaming Multiprocessor (SM)
The Streaming Multiprocessor (SM) (or Compute Unit (CU) in AMD terminology) is the
fundamental building block of the GPU. A single GPU chip contains dozens to hundreds of
SMs. An SM is a self-contained, multi-threaded processor designed to handle multiple blocks
of parallel work concurrently.
Key Components inside an SM:
• CUDA Cores / Stream Processors (SPs): These are the individual, simple Arithmetic
Logic Units (ALUs) that perform the bulk of the floating-point and integer
calculations.
• Warp Schedulers: These units are responsible for selecting Warps (groups of threads)
that are ready to execute and issuing the same instruction to all active.

• Shared Memory / L1 Cache: A small, ultra-fast, user-managed memory that is

exclusive to the threads running on that specific SM
• Register File: Ultra-fast on-chip storage for thread-specific data and intermediate
results.
• Tensor Cores (NVIDIA): Specialized processing units optimized for matrix
multiplication and accumulation, critical for deep learning and AI operations.

□ Parallel Execution Concepts

1. Single Instruction, Multiple Data (SIMD)
• Concept: A single instruction is executed simultaneously on multiple data elements
by a single processing unit (or a set of vector lanes).
2. Single Instruction, Multiple Threads (SIMT)
• Concept: This is the core programming model used by GPUs (first coined by NVIDIA).
A single instruction is executed across multiple independent threads.
3. Streaming Multiprocessor (SM)
• As detailed above, the SM is the physical unit that houses the execution resources. It
concurrently manages and executes multiple warps/wavefronts and provides the
memory and control logic for the threads.
4. Stream Processor (SP) / CUDA Core
• The SP or CUDA Core is the execution unit (the ALU) that corresponds to a single
thread's execution capability within the SIMT model. When a warp is scheduled, each
active thread in that warp uses one SP to perform the computation mandated by the
single instruction issued by the SM's scheduler.
2. Write a CUDA program that prints greetings from the threads explain with
Compiling and running the program.

CUDA Program: hello_cuda.cu

This program defines a kernel that launches multiple threads. Each thread calculates its
unique ID and prints a greeting along with that ID.

#include <stdio.h>

global void hello_cuda_kernel() {

int thread_id = blockIdx.x * blockDim.x + threadIdx.x;

printf("Hello from CUDA thread ID: %d\n", thread_id);

int main() {

int threads_per_block = 256; // Launch 256 threads in each block

int num_blocks = 4; // Launch 4 blocks

int total_threads = threads_per_block * num_blocks;

printf("Launching %d total threads (%d blocks of %d threads each)...\n",

total_threads, num_blocks, threads_per_block);

hello_cuda_kernel<<<num_blocks, threads_per_block>>>();

cudaDeviceSynchronize();

cudaError_t err = cudaGetLastError();

if (err != cudaSuccess) {

fprintf(stderr, "CUDA error: %s\n", cudaGetErrorString(err));

return 1;

return 0;

Explanation of Key CUDA Concepts (Short)

⬛ Host and Device Code
1

• Host (CPU) runs the main program → sets up execution.

• Device (GPU) runs the kernel function marked with global → executed by many
parallel threads.

⬛ Kernel Launch Configuration

2
Syntax:

kernel<<<gridDim, blockDim>>>(args);

In example:

• num_blocks = 4 (grid size)

• threads_per_block = 256 (block size)

• Total threads:

 4 × 256 = 1024 threads

³ Thread Indexing
⬛
Unique thread ID:

id = blockIdx.x * blockDim.x + threadIdx.x

So threads are numbered from 0 to 1023 uniquely.

Compiling the Program (Short)

Use nvcc:
nvcc [Link] -o output ./output

3. Define Threads, blocks, and grids and write a CUDA program that
prints greetings from threads in multiple blocks.

1. Thread
• Definition: A Thread is the smallest unit of execution in CUDA. It executes the kernel
function independently.
2. Block (Thread Block)
 • Definition: A Block (or Thread Block) is a group of concurrent threads that execute
the same kernel code.
3. Grid
• Definition: A Grid is the top-level collection of all thread blocks launched by a single
kernel call.

□ CUDA Program: Multiple Block Greetings

#include <stdio.h>
global void hello_multi_block_kernel() {

int thread_id = blockIdx.x * blockDim.x + threadIdx.x;

int block_index = blockIdx.x;
printf("Block %d, Thread %d: Hello from Global ID %d\n",
block_index, threadIdx.x, thread_id);
}
int main() {
const int THREADS_PER_BLOCK = 128; // Block Dimension
const int NUM_BLOCKS = 8; // Grid Dimension
int total_threads = THREADS_PER_BLOCK * NUM_BLOCKS;
printf("--- CUDA Kernel Launch Setup ---\n");
printf("Launching %d total threads (Grid: %d blocks, Block: %d threads)\n",
total_threads, NUM_BLOCKS, THREADS_PER_BLOCK);
printf(" \n");
hello_multi_block_kernel<<<NUM_BLOCKS, THREADS_PER_BLOCK>>>();
cudaDeviceSynchronize();
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
fprintf(stderr, "CUDA error: %s\n", cudaGetErrorString(err));
return 1;
}
printf(" \n");
 printf("Kernel execution finished successfully.\n");
return 0;
}

COMPLIE: nvcc multi_block_greeting.cu -o multi_block_greeting

RUN: ./multi_block_greeting
OUTPUT:
Block 0, Thread 0: Hello from Global ID 0
Block 4, Thread 0: Hello from Global ID 512

4. Explain Nvidia GPU compute capabilities and architectures, and how

compute capability affects thread limits, block sizes, and CUDA
compatibility more information

Compute
Architecture Capability Key Features Introduced/Enhanced
(CC)

Volta 7.0 First Tensor Cores (Specialized AI/DL units)

First Ray Tracing (RT) Cores; Integer (INT32) and

Turing 7.5
Floating Point (FP32) pipelines decoupled.

Gen 3 Tensor Cores (supporting TF32, BF16); Gen 2

Ampere 8.0, 8.6, 8.7 RT Cores; Higher Streaming Multiprocessor (SM)
throughput.

Ada Gen 4 Tensor Cores; Gen 3 RT Cores; Focus on

8.9
Lovelace increased clock speed and power efficiency.

Transformer Engine; DPX Instructions; Major

Hopper 9.0
datacenter-focused architectural shift.

Gen 5 Tensor Cores; CUDA Tile programming model

Blackwell 10.x, 12.x
introduced (CC 10.x/12.x).
How Compute Capability Affects CUDA Parameters
1. Thread Limits (Block Size)
The CC dictates the maximum number of threads allowed in a single Thread Block. This
value has remained consistent for many modern architectures, but it is a hardware-enforced
limit.
2. Block Size and Scheduling
While the maximum thread block size is 1024, the CC also determines the resources per
Streaming Multiprocessor (SM), which affects how many blocks can be concurrently resident
on an SM, impacting Occupancy.
3. CUDA Compatibility

• Forward Compatibility (PTX): The CUDA compiler (nvcc) can generate intermediate
code called PTX (Parallel Thread Execution). PTX is like a virtual assembly language
• Binary Compatibility (CUBIN): The final binary code (CUBIN) compiled for a
specific CC (e.g., CC 8.6) is only guaranteed to run on GPUs with the same or higher
minor revision within the same major revision
• Feature Availability: If your code uses advanced features like Tensor Cores or FP64
(double-precision), the kernel must be compiled for a CC that supports those features

5. Write a program to show Kernel and main function of a CUDA program

that adds two vectors.
This program adds two vectors, and , element-wise to produce a result vector , where .
C[i] = A[i] + B[i].

#include <stdio.h>

#include <stdlib.h> // For malloc, exit

#include <cuda_runtime.h>
global void vectorAdd(float *C, const float *A, const float *B, int N) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < N) {
// Perform the addition
C[i] = A[i] + B[i];
}
}
int main() {
// 1. Setup Parameters
const int N = 1 << 20; // Vector size: 2^20 elements (1,048,576)
const size_t bytes = N * sizeof(float); // Total memory size in bytes

printf("Vector size (N): %d elements\n", N);

float *h_A = (float*)malloc(bytes);
float *h_B = (float*)malloc(bytes);
float *h_C = (float*)malloc(bytes);

if (h_A == NULL || h_B == NULL || h_C == NULL) {

fprintf(stderr, "Host memory allocation failed!\n");
exit(EXIT_FAILURE);
}
for (int i = 0; i < N; ++i) {
h_A[i] = 1.0f;
h_B[i] = 2.0f;

}
float *d_A, *d_B, *d_C;
// Allocate memory on the GPU (device)
cudaMalloc((void**)&d_A, bytes);
cudaMalloc((void**)&d_B, bytes);
cudaMalloc((void**)&d_C, bytes);
cudaMemcpy(d_A, h_A, bytes, cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, bytes, cudaMemcpyHostToDevice);

const int NUM_BLOCKS = (N + THREADS_PER_BLOCK - 1) /

THREADS_PER_BLOCK;
printf("Launching kernel with %d blocks and %d threads per block...\n",
vectorAdd<<<NUM_BLOCKS, THREADS_PER_BLOCK>>>(d_C, d_A, d_B, N);
 cudaMemcpy(h_C, d_C, bytes, cudaMemcpyDeviceToHost);
int errors = 0;
for (int i = 0; i < N; ++i) {

if (h_C[i] != 3.0f) { // Expected result is 1.0f + 2.0f = 3.0f

errors++;
if (errors < 5) { // Print only a few errors
printf("Error at index %d: Expected 3.0f, Got %f\n", i, h_C[i]);
}

}
}
if (errors == 0) {
printf("\nVerification successful! Vector addition completed correctly.\n");
} else {
printf("\nVerification failed! Found %d errors.\n", errors);
}
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_C);
free(h_A);
free(h_B);
free(h_C);
return 0;
}

Kernel Function (vectorAdd)

This formula maps the thread's location within its block (threadIdx.x) and the block's location
within the grid (blockIdx.x) to a single, unique index corresponding to an element in the
vector.
• Element-wise Operation: The line C[i] = A[i] + B[i]; performs the addition for that
single element, utilizing the massive parallelism of the GPU.
2. Main Function (main)
• Execution Location: The main function runs entirely on the CPU (Host).
• Steps:
o Memory Management: It uses malloc for CPU memory (h_A, h_B, h_C) and
cudaMalloc for GPU memory (d_A, d_B, d_C).
o Data Transfer: cudaMemcpy is used to explicitly transfer data:
o Kernel Launch: The triple angle bracket syntax
o Cleanup: free and cudaFree are used to release the memory on their
respective systems.

6. Explain the CUDA vector addition program, describing how threads

compute individual elements and how the kernel is launched from the host.

Core Components of the CUDA Vector Addition Program

1. The Kernel Function (vectorAdd)
The kernel is the function that runs in parallel on the GPU.
• Qualifier: The global qualifier marks this function for execution on the GPU.
• Input/Output: It takes pointers to the vectors in GPU memory (d_A, d_B, d_C) and
the vector size (N).
How Threads Compute Individual Elements:
The essence of the parallel computation lies in mapping a thread's unique ID to a specific
element index in the vector.

• Global Thread Index Calculation: Each thread calculates its unique, one-
dimensional index, , across the entire Grid using the built-in CUDA variables:
• Boundary Check: The if (i < N) check is critical. This check prevents threads with
IDs from trying to access invalid memory locations.
• Parallel Work: If the index is valid, the thread performs the single operation: C[i] =
A[i] + B[i];. Since every thread calculates a different unique , the entire vector
addition is completed simultaneously.

2. The Main Function (Host)

The main function executes on the CPU (Host) and orchestrates the entire operation in four
key steps:
A. Data and Memory Management
• Host Allocation: Memory for the input vectors and result vector () is allocated on the
CPU using standard malloc.
• Device Allocation: Memory for the vectors is allocated on the GPU B. Data
Transfer (Host Device)
• The input vectors are copied from the slower CPU memory to the faster GPU memory
C. Kernel Launch
The kernel is launched from the host using the triple angle bracket syntax:
• (Grid Dimension): Defines the total number of blocks needed to cover all elements.
• (Block Dimension): Defines the number of threads per block (typically 128, 256, or
512 for optimization).
• Execution Flow: When the kernel is launched, the CPU does not wait for the GPU to
finish (it's an asynchronous call). The CPU continues executing the next line of code
immediately.
D. Data Transfer (Device Host)
• A synchronization point is implicitly created by thecall, which pauses the CPU until
the GPU finishes the vector addition and the result is safely copied from the GPU's
memory () back to the CPU's memory
.

7. Explain why CUDA kernels cannot return values directly to the host.
Describe the different methods used to return results from a CUDA kernel.

❌ Why Kernels Can't Return Directly

1. Asynchronous Execution: When the Host launches a kernel, the CPU thread
immediately continues executing the next instruction without waiting for the GPU to
finish.
2. Separate Memory Spaces: The Host (CPU) operates on Host Memory (RAM), and
the Device (GPU) operates on Device Memory (VRAM). A function's return value is
typically placed in a register or a stack location accessible to the caller. Methods to
Return Results from a CUDA Kernel
1. The Primary Method: Device Pointers (Global Memory)

This is the standard and most common method for returning large result sets (like vectors,
arrays, or matrices).
• Use Case: Returning large arrays or vectors (as seen in the vector addition example).
2. Returning Single/Small Values (Result Aggregation)
For returning a single summary result (like a count, a minimum value, or an error code), you
use the device pointer method, but the kernel performs an aggregation across all threads.
3. Unified Memory (Managed Memory)
Unified Memory simplifies the programming model by creating a single, managed pointer
that is visible to both the Host and the Device.
4. Asynchronous Error Checking
For checking basic kernel launch success, the Host relies on functions that check the CUDA
execution status rather than receiving a return value:
• cudaGetLastError():
• cudaDeviceSynchronize():

8. Explain the concept of CUDA trapezoidal rule I and write a serial

function implementing the trapezoidal rule for a single CPU.

The Trapezoidal Rule approximates the definite integral by dividing the area under the
curve into small trapezoids of equal width . The formula for the integral is:

Parallel Strategy (Trapezoidal Rule I)

1. Work Distribution: The total number of sub-intervals, , is divided among the
launched threads.

2. Thread's Responsibility: Each thread is assigned a subset of the trapezoids

(intervals) to calculate the function value and contribute to the central sum.
3. Local Summation: Each thread computes a local sum of its assigned terms.
4. Global Reduction: After all threads complete their local sums, a reduction operation
is performed to combine these local sums into a single global sum on the GPU.

5. Final Calculation: The Host (CPU) retrieves the global sum, adds the boundary
terms and , and multiplies the total by to get the final integral value.
 Serial CPU Function for Trapezoidal Rule

#include <stdio.h>
#include <math.h>
double function_to_integrate(double x) {
return x * x;
}
double serial_trapezoidal_rule(double a, double b, int n) {
if (n <= 0) {
return 0.0;
}
double delta_x = (b - a) / n;
double total_sum = function_to_integrate(a) + function_to_integrate(b);
for (int i = 1; i < n; i++) {
// Calculate the current x-coordinate
double x_i = a + i * delta_x;
total_sum += 2.0 * function_to_integrate(x_i);
}
double integral = (delta_x / 2.0) * total_sum;
return integral;
}
int main() {
double lower_limit = 0.0;
double upper_limit = 2.0;
int num_intervals = 100000;
double result = serial_trapezoidal_rule(lower_limit, upper_limit, num_intervals);
printf("--- Serial Trapezoidal Rule --- \n");
printf("Function: f(x) = x^2\n");
printf("Intervals (n): %d\n", num_intervals);
printf("Approximate Integral from %.2f to %.2f: %.8f\n",
lower_limit, upper_limit, result);
return 0;
}
8. Explain the concept of CUDA trapezoidal rule I and write a serial
function implementing the trapezoidal rule for a single CPU.
A CUDA Trapezoidal Rule I Concept
The Trapezoidal Rule approximates the integral by summing the area of trapezoids.
• Formula Focus: The core computation is the sum of for all interior points ( to ).
• Parallel Strategy (Method I): This is the simplest parallelization where the work of
calculating the interior terms is divided among the threads.
o Work Distribution: Each thread is assigned one or more intervals (using a
grid-stride loop).
o Local Sum: Each thread computes a small local sum of its assigned terms.
o Global Sum: The thread writes its local sum to a unique position in a large
Global Memory array.
o Drawback: The Host must then read back this entire large array and perform
the final summation (reduction) and scaling. This creates a severe Global
Memory and Host-Device transfer bottleneck.
B. Serial Trapezoidal Rule Function
#include <math.h>
double f(double x) {
return x * x;
}
double serial_trapezoidal_rule(double a, double b, int n) {
if (n <= 0) return 0.0;
double delta_x = (b - a) / n;
double total_sum = f(a) + f(b);
for (int i = 1; i < n; i++) {
double x_i = a + i * delta_x;
total_sum += 2.0 * f(x_i);
}
double integral = (delta_x / 2.0) * total_sum;
return integral;
}
9. Write a Program to Initialization, return value, and final update for
CUDA kernel and
wrapper implementing trapezoidal rule and explain problems in Cuda
implementation.
#include <stdio.h>
#include <cuda.h>
global void trapKernel(double a, double h, int n, double *d_sum)
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i > 0 && i < n) {
double x = a + i * h;
d_sum[i] = x * x; // f(x) = x^2 stored in Global Memory
}
}
int main()
{
double a = 0.0, b = 1.0;
int n = 1024;
double h = (b - a) / n;
double *d_sum, *h_sum;
h_sum = (double*)malloc(n * sizeof(double));
cudaMalloc(&d_sum, n * sizeof(double));

int blocks = 4;
int threads = 256;
trapKernel<<<blocks, threads>>>(a, h, n, d_sum);
cudaMemcpy(h_sum, d_sum, n * sizeof(double),
cudaMemcpyDeviceToHost);
double final_sum = (a*a + b*b) / 2.0;
 for(int i=1;i<n;i++)
final_sum += h_sum[i];
double result = final_sum * h;
printf("Result = %f\n", result);
cudaFree(d_sum);
free(h_sum);
return 0;
}
Step Where? Meaning
Initialization CPU Compute a, b, h and allocate memory
Kernel
GPU Each thread computes f(xi) = xi²
Execution
Copy back to
Return Value cudaMemcpy() returns computed array
host
Reduce sum + add endpoints + multiply by
Final Update CPU
h

10. Explain CUDA trapezoidal rule II for improving performance by

showing Basic sum in Tree-structured communication.

CUDA Trapezoidal Rule – Method II (Improved Performance Using Shared Memory

Reduction)
In Method-I, each thread stored results in Global Memory and final summation was done on
CPU.
This was slow due to high-latency memory access.
➡ Method-II improves performance using Tree-Structured Parallel Reduction inside Shared
Memory.

★ Key Idea
✔ Each thread computes 𝑓(𝑥𝑖)
✔ Stores value into Shared Memory (fast on-chip memory)
✔ Threads reduce the values in a tree-like parallel fashion
✔ Final partial sums returned to CPU

* Tree-Structured Communication (Parallel Reduction)

✔ Number of operations reduces every step

✔ Much faster than sequential summing

11. What is Warps and warp shuffles, and explain the Tree-structured sum
using warp shuffle and Dissemination sum using shared memory.
Warps
• In NVIDIA GPUs, threads are grouped into fixed-size execution units called Warps.
• A Warp = 32 threads (execute in lockstep — same instruction at the same time).
• All threads in a warp follow the SIMT (Single Instruction Multiple Thread) model.
• If threads take different branches → branch divergence → performance loss.
➜
⬛ Warps enable fast parallel execution of thousands of threads.

Warp Shuffle Operations

• Warp shuffle allows threads within the same warp to directly exchange data via
registers.
• No shared memory required → low latency and no bank conflicts.
• Common operation: shfl_down_sync() to do reductions efficiently.
➜
⬛ Best choice for intra-warp communication and reductions.
Tree-Structured Sum Using Warp Shuffle
Used for parallel reduction (e.g., summing values).
Algorithm:
1. Each thread computes partial sum.
2. Warp threads exchange values through shuffle instructions.
3. Pairwise addition happens like a binary tree.
Working:
• At first step, 16 threads add values from next 16 threads.
• Then 8 add from next 8 → 4 → 2 → 1.
• Final result stored in thread 0 of the warp.
Dissemination Sum Using Shared Memory
Used when communication is needed across multiple warps.
Steps:
1. Each warp performs its own warp-shuffle reduction.
2. Results from warp-leaders (e.g., lane 0) are written to shared memory.
3. Synchronization using syncthreads().
4. A final reduction (tree-style) is performed in shared memory.

12. Write a CUDA kernel implementing program for the trapezoidal rule
and using Warp_sum.
#include <stdio.h>

#define N 1024
#define THREADS 256

device float func(float x) {

return x * x; // Example f(x) = x²
}
inline device float warp_sum(float val) {
 for (int offset = 16; offset > 0; offset /= 2)
val += shfl_down_sync(0xFFFFFFFF, val, offset);
return val;
}
global void trapezoidalKernel(float a, float h, float *arr) {
int main() {
float a = 0.0f, b = 10.0f;
float h = (b - a) / N;
float *d_arr, h_arr[N/32];

cudaMalloc(&d_arr, (N/32) * sizeof(float));

trapezoidalKernel<<<N / THREADS, THREADS>>>(a, h, d_arr);

cudaMemcpy(h_arr, d_arr, (N/32) * sizeof(float), cudaMemcpyDeviceToHost);

float sum = 0.0f;

for (int i = 0; i < N/32; i++)
sum += h_arr[i];

// Apply trapezoidal formula

sum = h * (sum - (func(a) + func(b)) / 2.0f);
printf("Integral approx = %f\n", sum);

cudaFree(d_arr);
return 0;
}

13. Write a program for CUDA kernel implementing the trapezoidal rule
and using shared memory.
#include <stdio.h>
#define N 1024
#define THREADS 256
device float func(float x) {
 return x * x; // Example: f(x) = x²
}

global void trapKernel(float a, float h, float *result) {

shared float sdata[THREADS];
for (int stride = THREADS / 2; stride > 0; stride /= 2) {
if (threadIdx.x < stride)
sdata[threadIdx.x] += sdata[threadIdx.x + stride];
syncthreads();
}

if (threadIdx.x == 0)
result[blockIdx.x] = sdata[0]; // Block partial sum
}

int main() {
float a = 0.0f, b = 1.0f;
float h = (b - a) / N;
int blocks = N / THREADS;

float *d_res, h_res[blocks];

cudaMalloc(&d_res, blocks * sizeof(float));

trapKernel<<<blocks, THREADS>>>(a, h, d_res);

cudaMemcpy(h_res, d_res, blocks * sizeof(float), cudaMemcpyDeviceToHost);

float sum = (func(a) + func(b)) / 2.0f;

for (int i = 0; i < blocks; i++)
sum += h_res[i];

float result = sum * h;

printf("Trapezoidal Integral = %f\n", result);
 cudaFree(d_res);
return 0;
}

14. Explain how CUDA handles the trapezoidal rule in multi-warp blocks
and why Syncthreads() is necessary to avoid race conditions during partial-
sum reduction.

CUDA Trapezoidal Rule in Multi-Warp Blocks & Need for syncthreads()

When implementing the Trapezoidal Rule on a GPU, the entire numeric integration range is
divided among many threads. Each thread computes a value 𝑓(𝑥𝑖), and these values must be
summed together to get the final integral.
In a CUDA block, threads are grouped into warps (usually 32 threads/warp).
If a block has 256 threads → 8 warps operate in parallel.
❖ Multi-Warp Handling in Reduction
• Each warp independently computes a partial sum of its subset of function values.
• These partial sums are written into shared memory inside the block.
• After all warps store their results, a second reduction stage combines the values across
warps to produce one final block sum.
Thus, trapezoidal computation is performed in two levels:
1. Per-thread computation of f(x)
2. Per-warp and then inter-warp shared-memory reduction
❖ Why syncthreads() is Necessary?
syncthreads() acts as a barrier ensuring all threads in a block:
✔ Finish writing their partial sum into shared memory
✔ Do not start reading or reducing data early
✔ Prevent race conditions (one thread reading old/unwritten data)

⚠ What Happens Without syncthreads()?

• Some threads may start reducing before others have updated shared memory
• Leads to undefined / incorrect results
• Causes race conditions where threads interfere with each other’s data

15. Write use shared memory and synchronization to implement the

trapezoidal rule with large thread blocks in CUDA?
#include <stdio.h>
#define N 2048
#define TPB 256 // Threads Per Block
device float func(float x) {
return x * x; // Example function: f(x) = x²
}

global void trapShared(float a, float h, float *d_part) {

shared float sdata[TPB];
int tid = blockIdx.x * blockDim.x + threadIdx.x;
float x = a + tid * h;
float temp = 0.0f;
if (tid > 0 && tid < N)
temp = func(x);
sdata[threadIdx.x] = temp;
syncthreads();
for (int s = TPB / 2; s > 0; s >>= 1) {
if (threadIdx.x < s)
sdata[threadIdx.x] += sdata[threadIdx.x + s];
syncthreads(); // Prevent read/write conflicts
}
if (threadIdx.x == 0)
d_part[blockIdx.x] = sdata[0];
}

int main() {
float a = 0.0f, b = 1.0f;
float h = (b - a) / N;
int blocks = N / TPB;

float *d_part, h_part[blocks];

cudaMalloc(&d_part, blocks * sizeof(float));
 trapShared<<<blocks, TPB>>>(a, h, d_part);
cudaMemcpy(h_part, d_part, blocks * sizeof(float), cudaMemcpyDeviceToHost);

float sum = (func(a) + func(b)) / 2.0f;

for (int i = 0; i < blocks; i++)
sum += h_part[i];

printf("Integral = %f\n", sum * h);

cudaFree(d_part);
return 0;
}

[Link] the steps required to convert a sequential n-body solver into a parallel
version using OpenMP.
Steps to Convert a Sequential N-Body Solver to an OpenMP Parallel Version*
*1. Identify the Parallel Regions*
• Find the parts of the n-body code that can run independently.
• In an n-body solver, the main parallelizable part is:
• Computing forces between bodies
• Updating positions and velocities
• These operations are *independent per body*, so they can be parallelized.

*2. Include the OpenMP Header*

Add:
c
#include <omp.h>

*3. Add a Parallel Region*

Wrap the main computation loop with:
c
#pragma omp parallel for
This allows each iteration (each body) to be computed by a different thread.

## *4. Ensure Correct Variable Scoping*

* Arrays for positions, velocities, masses → *shared*
* Loop index and temporary force variables → *private*

Example:
c
#pragma omp parallel for private(i, fx, fy, fz) shared(pos, vel, mass)
## *5. Use Reduction if Needed*
If accumulating total energy or momentum, use:
c
#pragma omp parallel for reduction(+:total_energy)

## *6. Avoid Data Races*

• Make sure updates to shared variables happen safely.
• Typically, each thread updates *its own body*, so no critical section is needed.
• Only global accumulators need reduction.

## *7. Test Correctness*

Compare:
* Sequential results
* Parallel results
to ensure numerical accuracy is preserved.

## *8. Measure Performance*

Use omp_get_wtime():
c
double start = omp_get_wtime();
// compute forces
double end = omp_get_wtime();
printf("Parallel Time = %f\n", end - start);

Try different thread counts using:

export OMP_NUM_THREADS=4

# *Short Summary*

1. Find parallel loops

2. Add #include <omp.h>
3. Add #pragma omp parallel for
4. Set shared/private variables
5. Use reduction if needed
6. Remove data races
7. Test correctness
8. Measure speedup