We take content rights seriously. If you suspect this is your content, claim it here.
Available Formats
Download as PDF or read online on Scribd
CUDA Programming in Python: Complete Handbook
Table of Contents
Introduction to CUDA and GPU Computing
‘Setting Up CUDA Environment
CUDA Basics. Kemels and Thread Hierarchy
emory Management
Thnzad ‘Synchronization and Communication
‘Optimization Techniques
Memory Patterns
and Concurrency
Debugging and
Sewn aneene
Chapter 1: Introduction to CUDA and GPU Computing {#chapter-1}
1.1 What is CUDA?
CUDA (Compute Unified Device Architecture) is NVIDIA's parallel computing platform and programming model. It
enables dramatic increases in computing performance by harnessing the power of GPUs
Key Concepts:
+ Host: The CPU and its memory
+ Device: The GPU and its memory
» Kernel: A function that runs on the GPU
* Thread: Basi ecution unit on GPU
1.2 GPU Architecture
GPUs contain thousands of smaller cores designed for parallel processing
+ Streaming Multiprocessors (SMs): Groups of CUDA cores
+ CUDA Cores: Individual processing units
* Memory Hierarchy: Registers, shared memory, global memory
1.3 When to Use CUDA
Use CUDA when:
Processing large datasets with independent operations,
Performing matrix operations
Running scientific simulations
Processing images/videos
Training neural networks Chapter 2: Setting Up CUDA Environment {#chapter-2}
2.1 Installation
[]
bash.
#t Install CUDA Toolkit (varies by OS)
# For Ubuntu
sudo apt-get install
# Install Python CUDA libraries
install
install # Replace with your CUDA version
install
2.2 Verifying Installation
[7]
python
# Using Numba
fro import
import as
# Check if CUDA is available
print("CUDA Available:"
it Get GPU information
if
print(f"GPU Name:
print(f'Compute Capability:
print(f"Total Memory: 169:.2f} GB"
2.3 Choosing a Python CUDA Library
Numba: JIT compiler, easiest to leam, Pythonic syntax CuPy: NumPy-like interface, great for array operations PYCUDA:
Low-level control, requires CUDA C knowledge Chapter 3: CUDA Basics: Kernels and Thread Hierarchy {#chapter-3}
3.1 Your First CUDA Kernel
[F]
python
from numba import cuda
import numpy as np
@eudajit
def add_kernel(s, y, out):
"Simple element-wise addition kernel"
idx =[Link](1) # Get global thread 1D
if idx < outsize:
outfidx] = xfidx] + ylidx]
# Host code
n= 100000
X= [Link](n, dtype=np.float32)
y=[Link](n, dtype=np.float32)
out = [Link](n, dtype=np.float32)
# Copy data to device
x_device = cudato_device(x)
y_device = cuda.to_device(y)
out_device = cuda.device_array(n, dtype=np.float32)
# Configure and launch kernel
threads_per_block = 256
= (n+ threads_per_block - 1) // threads_per_block
add_kernel[blocks_per_grid, threads_per_block](x_device, y_device, out_device)
blocks_per
# Copy result back to hast
out = out_device.copy_to_host()
print(f"Result; {out[:10]}") # Should be all 2.0%
3.2 Thread Hierarchy
CUDA organizes threads into a three-level hierarchy: a
v
python
@euda,it
def thread_hierarchy_demo(outpu
Demonstrates thread indexing"
# Thread indices within block
tx = [Link].x.
[Link].y
tz = [Link].z
# Block indices within grid
bx = [Link].x
by=[Link]
bz=cudablockldx.z
# Block dimensions
bw = cudablockDimx
bh = [Link]
bd = [Link].z
# Calculate global thread index
x= tx + bx * bw
yaty+by *bh
z= 12+ ba" bd
# Example: 2D grid calculation
idx = x+y * [Link].x * [Link].x
# Launch with 2D grid and blocks
threads_per_block = (16, 16)
blocks_per_grid = (4, 4)
output = cuda.device_array((64, 64))
thread_hierarchy_demo[blocks_per_grid,threads_per_block](output)
3.3 Grid Stride Loops
For processing arrays larger than the grid:
a
v python
def gi
d_stride_kernel
"Handles arbitrary array sizes’
#f Works with any array s
1900000
1024, 256)
Chapter 4: Memory Management {#chapter-4}
4.1 Memory Types
Global Memory: Large, slow, accessible by all threads Shared Memory: Fast, limited, shared within a block Local
Memory: Private to cach thread Constant Memory: Read-only, cached Registers: Fastest, very limited
4.2 Memory Allocation and Transfer
8
yj
python import numpy as np
from numba import cuda
# Host arrays
h_array = [Link](1000).astype(np.float32)
# Device allocation methods
d_arrayl = cuda.to_device(h_array) # Copy from host
d_array2 = cuda.device_array(1000, dtype=np.float32) # Allocate empty
d_array3 = cudadevice_array_like(h_atray) # Same shape/dtype
# Copy device to host
result = d_array!.copy_to_host()
# In-place copy
cuda.to_device(h_array, to=d_array2)
1 Free device memory explicitly
del d_array!, d_array2, d_array3
4.3 Shared Memory from numba import cuda
import numpy as np
@oudajit
def shared_memory_example(data, result):
"Using shared memory for faster access
#4 Allocate shared memory
shared = [Link](shape=(256,), dtype=float32)
tid = [Link].x.
bid = [Link].x
idx = tid + bid * [Link].x
# Load data into shared memory
if idx 0 else 0
right = shared[tid + 1] if tid < 255 else 0
resullfidx] = (left + shared{tid] + right) /3.0
[Link]()
# Usage
n= 10000
data = cuda.to_device([Link](n).astype(np.float32))
result = cuda.device_array(n, dtype=np.float32)
shared_memory_example[n // 256 + 1, 256](data, result)
4.4 Memory Coalescing
8
v
python @euda,jit
def coalesced_access(matrix, result)
"Good: Coalesced memory ace
pattern’
row = cudablockldx.x
col = [Link].x
# Threads in a warp access consecutive memory
if row <[Link][0] and col < [Link][1]
result[row, col] = matrix[row, col] * 2
@oudajit
def uncoalesced_access(matrix, result)
"Bad; Uncoalesced memory access pattern"""
row = [Link].x
col = [Link].x
# Threads in a warp access strided memory (slower!)
if row < matrix shape{0] and col < [Link]{1}
result{row, col] = matrix{row, col] * 2
Chapter 5: Thread Synchronization and Communication {#chapter-5}
5.1 Barrier Synchronization
F]
python from numba import cuda
import numpy as mp
@euda,jit
def parallel_prefix_sum(data, output):
"Compute prefix sum using synchronization"*"
shared = [Link](256, dtype=float32)
tid = [Link].x
# Load data into shared memory
shared{tid] = data[tid] if tid < [Link] else 0
[Link]()
# Up-sweep phase
offset = 1
d= 128
while d>0;
iftid= bins:
bbin_idx = bins = 1
# Atomic add to prevent race conditions
[Link](hist, bin_idx, 1)
# Usage
n= 1000000
data = cuda.to_device([Link](n).astype(np.float32))
hist = cuda.to_device([Link](10, dtype=np.int32))
atomic_histogram[n //256 + 1, 256](data, 10, hist)
rosult = hist.copy_to_host()
print(f"Histogram: {result}") 5.3 Warp-Level Primitives
fs)
python from numba import cuda
@euda,jit
dof warp_reduce_sum(data, result);
"Efficient reduction using warp shuffles"
tid = [Link].x
lane = tid % 32 # Lane within warp
warp_id = tid //32
# Shared memory for partial sums
shared = [Link](32, dtype=float32)
# Load data
val = data[tid] if tid < [Link] else 0.0
# Warp-level reduction using shuffle
offset = 16
while offset > 0:
val += cuda.shfl_down_syne(OxftHft, val, offset)
offset = 2
4 First thread in warp writes to shared memory
if lane == 0;
shared{warp_id] = val
[Link]()
4 First warp reduces partial sums
if warp,
val = sharedflane] if lane < 32 else 0.0
offset = 16
while offset > 0:
val += cuda.shfl_down_syne(OxfffTff, val, offset)
offset //=
if lane == 0;
result[0] = val Chapter 6: Optimization Techniques {#chapter-6}
6.1 Occupancy Optimization
[3]
v
python
from numba import cuda
# Check occupancy for different configurations
def analyze_occupaney(threads_per_block)
"Analyze kernel occupanc
@euda jit
def sample_kernel(data)
idx = [Link](1)
if idx < datasize
datalidx] *= 2
# Get occupancy information
device = cuda.get_curret
max_threads = device MAX_THREADS_PER_BLOCK
t_device()
print(f"Threads per block: {threads_per_block}")
print(f"Max threads per block: {max_threads}")
print(f"Occupancy: {threads_per_block / max_threads * 100:.1f}%6")
Test different configurations
for tpb in [64, 128, 256, 512, 1024];
analyze_occupancy(tpb)
6.2 Bank Conflict Avoidance from numba import cuda
import numpy as np
@euda,jit
def no_bank_conflict(data, result):
"Padding to avoid bank conflicts
4 Shared memory with padding (33 instead of 32)
shared = [Link]((32, 33), dtype=float32)
tx = [Link].x.
ty=[Link].y
#Load with padding
sharedtx, ty] = data[tx, ty]
cuda,syncthreads()
# Transpose (now conflict-free)
result[ty, tx] = shared, ty]
6.3 Loop Unrolling
a
v
python from numba import cuda
@euda,jit
def unrolled_dot_product(a, b, result);
"Manually unrolled loop for better performance’
idx = cudagrid(1)
if idx 0;
iftid