
Explore CUDA programming with a hands-on suite of sample projects in Visual Studio, and leverage CUDA toolkit documentation and the CUDA Zone community to solve real-world problems.
Learn the fundamentals of parallel programming with CUDA, covering von Neumann architecture, processes and threads, context switching, and the difference between concurrency and true parallelism, including task- and data-level parallelism.
Explore how parallel computing powers supercomputing and why GPUs evolved as general purpose devices, comparing CPU and GPU design, throughput, and heterogeneous computing with examples Blue Gene and Sunway TaihuLight.
Install the CUDA toolkit on your system and run your first CUDA program. Verify the toolkit with nvcc and explore a simple vector add kernel in a cuda project.
Learn the basic elements of a CUDA program, including host and device code, kernel launch, and dim3 grid and block configurations.
Explore how CUDA initializes threadIdx per thread block within 1D and 2D grids, showing x, y, z coordinates for each thread.
Explore how blockIdx, blockDim, and gridDim organize threads in a CUDA program, with dim3 coordinates and 1D and 2D grids, and see how values are printed per thread block.
Modify code from the previous video to print threadIdx, blockIdx, and gridDim for a 3d grid with four per dimension and a 2x2x2 block, then submit the main and kernel.
Compute unique global indices for each CUDA thread by combining threadIdx, blockIdx, and blockDim across 1D grids. Learn to access array elements with gid and verify with multi-block setups.
Compute the global index for a 2D grid by combining row offset, block offset, and the thread id. Show 2D indexing with block and grid offsets to produce unique indices.
Explore unique global index calculation for a 2d grid using 2d thread blocks and 2d grid, derive tid and offsets to enable consecutive memory access and uniquely index an array.
Explore memory transfer between host memory and device memory in a CUDA program, using cudaMemCpy with hostToDevice and DeviceToHost, allocate with cudaMalloc, and cudaDeviceSynchronize() after computation.
Implement a cuda c++ exercise that transfers a 64-element array to the device, launches a 3d grid, computes a unique row-major index, and prints each value.
Learn to implement a CUDA parallel sum of two arrays using a 1D grid and 128 threads per block. Verify results with a CPU implementation to perform a validity check.
Learn how to handle runtime and compile-time errors in CUDA programs by using cudaError, cudaGetErrorString, and gpuErrorChk macros to check device operations and report errors, including kernel launches.
Compare CPU and GPU execution times for a sum array by measuring clock cycles and memory transfer. Experiment with block sizes 128, 256, 512, 1024 to identify the best configuration.
Query device properties with cudaGetDeviceProperties to read device name, major and minor compute capability, and total memory, then use these values to set max threads per block and grid limits.
Explore parallel programming with CUDA, explaining host and device roles, kernels, grids and blocks, memory transfers via cudaMemCpy, asynchronous kernel launches, cudaDeviceSynchronize, error handling, and performance timing.
Explore how CUDA hardware architecture drives kernel performance by examining streaming multiprocessors, memory hierarchy, and SIMT execution across Maxwell and Volta GPUs.
Explore how CUDA threads are organized into 32-thread warps within blocks, how SIMT execution and warp scheduling impact performance on SMs, and why block sizes should align to 32.
Analyze warp divergence in CUDA SIMT execution caused by if-else and its impact on parallelism. Use nvprof to measure branch efficiency and observe optimization masking divergence.
Explore resource partitioning and latency hiding on GPUs by examining how registers and shared memory limit resident blocks and warps, and how the warp scheduler achieves zero-cost context switching.
Assess how resource partitioning and latency hiding use many eligible warps to hide arithmetic and memory latencies, improving SM utilization and distinguishing bandwidth-bound from computation-bound CUDA apps.
Occupancy measures active warps per SM; register and shared memory usage and warp granularity determine limits, and you can use the CUDA occupancy calculator to optimize kernels.
Explore nvprof profiling for profile-driven optimization of CUDA programs by experimenting with 1D and 2D grids, block sizes, and memory transfers, observing performance variations in a sum_array kernel.
Explore CUDA thread synchronization using cudaDeviceSynchronize and syncthreads() through a parallel reduction example. Learn how to compute partial sums across thread blocks with neighboring pairs and ensure safe synchronization.
Address warp divergence in CUDA parallel reduction by comparing neighboring-thread summation and interleaved pair approaches, with 8- and 128-element block examples and verified GPU–CPU results.
Master loop and thread block unrolling in CUDA to cut instruction overhead and speed up parallel reduction by unrolling two, four, or eight data blocks with interleaved pairs and syncthread.
Learn how to optimize CUDA parallel reduction with warp unrolling, using interleaved pairs, thread blocks of 128, and careful handling of warp divergence to accumulate efficiently.
Completely unroll the reduction for loop for a 1024-thread block to perform offset-based sums (512, 256, 128, 64) and remove warp divergence while preserving correct results.
Compare reduction kernel implementations and assess performance, highlighting template-parameter optimizations and unrolled blocks to minimize runtime instructions. Demonstrate that the completely unrolled kernel yields the best speed so far.
Explore CUDA dynamic parallelism, launching child kernels from a parent kernel with parent and child grids, runtime grid sizing, and synchronization.
Explore dynamic parallelism to implement reduction with interleaved pair method, using recursive child kernels, in place partial sums, and synchronization concepts, then analyze kernel launch overhead and performance implications.
Explore warp execution, resource partitioning, and latency hiding in CUDA programs, covering warp divergence, shared memory, registers, occupancy, synchronization, and dynamic parallelism.
Explore the cuda memory model and memory hierarchy, examining locality, global memory efficiency, throughput, and how block size affects performance using nvprof.
Explore the memory types in CUDA, from registers and shared memory to global, texture, constant, and local memory. See how on-chip versus DRAM locations and usage influence performance and occupancy.
Explore memory management in cuda programs by comparing pageable host memory with pinned memory, using cudaMallocHost and cudaFreeHost to boost host-device transfer bandwidth.
Learn how zero copy memory lets host and device access shared mapped pinned memory via cudaHostAlloc and cudaHostGetDevicePointer, with synchronization requirements and potential PCIe performance drawbacks.
Explore unified memory in cuda programming, a managed memory pool accessible by both host and device with automatic data migration, reducing explicit transfers and simplifying code.
Explore how global memory access patterns affect throughput in CUDA programs, focusing on aligned and coalesced accesses, cache pathways (L1/L2), and bus utilization.
Understand how global memory writes bypass L1 cache and are cached in L2 in 32-byte segments, with warp-aligned 128-byte ranges delivering best efficiency, while misaligned writes incur clear penalties.
Compare array of structures and structure of arrays in CUDA programming, showing how SOA enables coalesced memory access and higher bandwidth than AOS, evidenced by nvprof profiling.
Implement a CUDA-based matrix transpose, using row-major indexing and 2D thread grids, comparing row-major and column-major reads to analyze coalesced vs uncoalesced memory access.
Demonstrates unrolled CUDA matrix transpose kernels with a fourfold unrolling factor, exploring row-major and column-major variants, using ti and to indices and boundary checks to verify same results.
Explore a diagonal coordinate system to implement matrix transpose in CUDA, reducing partition champing and improving memory coalescing by spreading data across DRAM partitions.
Explore the cuda memory model, focusing on global memory types and access patterns, coalesced transactions, AOS vs SOA, pinned and unified memory, and a diagonal coordinate approach to matrix transpose.
Learn to optimize CUDA kernels using shared memory as a program-managed cache and scratch pad to reduce global memory latency, with static and dynamic declarations.
Explore how shared memory is divided into 32 banks and how warp access patterns affect bandwidth and bank conflicts, including 32-bit and 64-bit modes on Fermi and Kepler devices.
Explore how row-major versus column-major access to shared memory affects bank conflicts and transaction counts on a 2d matrix, with 32-bit and 64-bit modes on a gpu.
Shows dynamic shared memory in cuda, converting row major to column major access on a two dimensional grid using one dimensional storage with extern and syncthread.
Explore how memory padding eliminates shared memory bank conflicts in CUDA by adding extra columns to 2D shared memory, enabling row-major and column-major access without conflicts.
Enhance parallel reduction with shared memory to outperform the previous global-memory implementation. Profile with nvprof to compare gld_transaction and gst_transaction counts at block size 1024.
Explore CUDA memory ordering in a weakly ordered model, showing cross-thread visibility and how memory fences like __threadfence_block() and __threadfence() enforce order.
Optimize matrix transpose with shared memory in CUDA by loading row-major input into shared memory, transposing in column-major order, and writing row-major output, achieving a twofold speedup and 80% fewer global memory transactions.
Master cuda constant memory: a read-only device memory with warp-synchronized access, and apply stencil computations with halo handling and shared memory.
Enhance matrix transpose performance in CUDA by adding shared memory padding to reduce bank conflicts and memory load transactions, and boost efficiency with unrolling using a 1D shared memory array.
Explore CUDA warp shuffle instructions that exchange register data within a warp with low latency, using __shfl_sync, __shfl_up_sync, __shfl_down_sync, and __shfl_xor_sync. Understand lane id, width, and full_mask.
Explore implementing parallel reduction using warp shuffle instructions, with warp unrolling and shuffle down patterns, comparing to a shared-memory kernel and analyzing shared memory transactions and performance via nvprof.
Explore the CUDA memory model, covering shared memory banks, static and dynamic allocation, strategies to avoid bank conflicts with padding, plus warp shuffle, constant memory, and reduction.
Learn how to achieve grid-level parallelism by launching multiple kernels on a CUDA device concurrently, overlapping memory transfers with kernel execution through streams and asynchronous operations.
Explore how to overlap memory transfers with kernel execution using cudaMemCpyAsync, streams, and pinned memory to reduce execution time, and learn how to launch kernels on multiple streams with synchronization.
Learn to use CUDA streams for asynchronous operations, overlap kernel execution with memory transfers, and launch multiple kernels on separate streams while profiling with nvvp.
Learn to overlap memory transfers with kernel execution using CUDA streams for accelerated array summation. Use asynchronous copies with pinned memory, divide inputs into chunks, and launch per-stream kernels.
Explore blocking and non-blocking streams relative to the null stream, and how explicit and implicit synchronization governs CUDA program execution using cudaStreamCreate and cudaStreamCreateWithFlags.
Explore explicit synchronization in CUDA, including grid, stream, and event synchronization with cudaDeviceSynchronize, cudaStreamSynchronize, and cudaEventSynchronize. Note implicit synchronization from memory operations like cudaMallocHost and cudaHostAlloc that block prior work.
Learn how to declare, create, and manage cuda events, record and synchronize them in streams, and measure execution time with cudaEventElapsedTime.
Explore configurable cuda events and inter-stream dependencies, create events with flags like default and blocking sync, disable timing, then use cudaStreamWaitEvent to synchronize streams and observe parallel kernel execution.
Explore CUDA arithmetic functions, including floating point, intrinsic, and atomic operations, and learn how precision, accuracy, and compiler optimizations affect throughput and correctness.
Explore how floating point values behave in CUDA under IEEE 754, comparing float and double precision. Understand accuracy and performance trade-offs, including memory cost and kernel speed.
Compare cuda intrinsic versus standard functions, note that intrinsic is faster but less precise than standard, and learn to control mad optimizations with nvcc and fmad flags.
Explore CUDA atomic functions for mutual exclusive access to shared memory, including atomic add and compare-and-swap, and learn about performance trade-offs in parallel kernels.
Explore CUDA GPU parallel patterns by learning prefix sum, including inclusive and exclusive scan, and compare lean, mean sequential CPU implementations with parallel approaches.
Implement a simple in-place parallel inclusive scan in CUDA using a naive single-block kernel up to 1024 elements, based on reduction and iterative offsets.
Learn a work-efficient parallel prefix-sum algorithm using a balanced-tree pattern with reduction (upsweep) and downsweep phases, to implement an exclusive (and inclusive) scan in a single-block CUDA kernel.
Develop a work-efficient, two-phase inclusive scan using reduction and downsweep, illustrated on a 16-element array and optimized with shared memory for CUDA block processing.
Extend inclusive scan for large data sets by dividing array into data blocks, performing pre scan and exclusive scan on an auxiliary array, then applying post scan to finalize results.
Learn how to perform parallel compaction on the GPU to remove zeros from sparse matrices, using a reversible CSR storage scheme and compare sequential and parallel implementations.
Explore the basics of digital image processing with CUDA, from imaging concepts and camera obscura origins to sampling, pixels, and image resolution.
Explore the image processing life cycle from acquisition and light balancing through display, printing, compression, enhancement, and information extraction, and see how CUDA enables tone mapping and parallel image processing.
Explore digital image processing through practical examples, from Mars rover image compression and contrast enhancement to panoramic and 360 image construction, facial recognition, privacy blurring, and movie post-processing.
Explore the human perception system, including the retina with cones and rods, fovea, and adaptation, and learn how Weber law and Mach bands shape image processing.
Explore how grayscale and rgb images are represented in computers through sampling and quantization, including rgb channels, mosaic sensors, and interpolation that influence resolution, saturation, and noise.
Install OpenCV, an open source computer vision library, from sourceforge, set the dll path, and configure a Visual Studio CUDA project with include and library paths, linking OpenCV_world343d.lib.
This course is an in-depth, unofficial guide to parallel programming using GPU computing techniques with C++. We'll begin by exploring foundational concepts such as the GPU programming model, execution structure, and memory hierarchy. From there, you’ll dive into hands-on development, implementing advanced parallel algorithms optimized for high-performance graphics processors.
Since performance is at the heart of GPU-based computing, this course places a strong emphasis on optimization techniques. You’ll learn how to fine-tune your code for maximum speed and efficiency, and apply industry-standard tools for profiling and debugging, including nvprof, nvvp, memcheck, and GDB-based GPU debuggers.
The course includes the following core sections:
Introduction to GPU programming concepts and execution models
Understanding execution behavior on parallel processors
Deep dive into memory systems: global, shared, and constant memory
Using streams to manage concurrent execution
Fine-tuning instruction-level behavior for performance
Implementing real-world algorithms using GPU acceleration
Profiling and debugging tools overview
To reinforce learning, this course includes programming exercises and quizzes designed to help you internalize each concept.
This is the first course in a masterclass series on GPU-based parallel computing. The knowledge you gain here will form a strong foundation for exploring more advanced topics in future courses.
As GPUs continue to drive innovation in fields like AI and scientific computing, mastering these tools and techniques will set you apart in the tech industry.
Note: This course is not affiliated with or endorsed by NVIDIA Corporation. CUDA is a registered trademark of NVIDIA Corporation, used here solely for educational reference purposes.