CUDA Programming for NVIDIA H100s – Comprehensive Course

Learn CUDA programming for NVIDIA Hopper GPUs. You will study to build businesslike WGMMA pipelines and leverage Cutlass optimizations to execute the monolithic matrix multiplications that powerfulness modern AI. Beyond single-chip performance, the program covers multi-GPU scaling and NCCL primitives basal for training trillion-parameter models. To get the astir retired of these lessons, you should person a foundational grasp of C++ syntax and linear algebra, peculiarly really matrices are tiled and multiplied. - Course website: https://cudacourseh100.github.io - Course repo: https://github.com/cudacourseh100/H100-Course - X: https://x.com/_PrateekShukla_ - GitHub Sponsors: https://github.com/sponsors/prateekshukla1108 ✏️ Developed by @Prateek_Shukla ❤️ Support for this transmission comes from our friends astatine Scrimba – the coding level that's reinvented interactive learning: https://scrimba.com/freecodecamp Contents - 0:00:00 Course Introduction - 0:07:27 Table of Contents & Course Overview - 0:23:30 LESSON 1 — H100 Hopper GPU Architecture - 0:25:47 H100 Specifications: HBM3, Bandwidth & Power - 0:26:22 Tensor Cores Overview - 0:27:18 Tensor Memory Accelerator (TMA) - 0:34:44 Transformer Engine - 0:34:58 L2 Cache Architecture - 0:35:21 GPCs, TPCs & SM Layout - 0:37:00 Thread Block Clusters - 0:46:22 Distributed Shared Memory - 0:52:44 SM Sub-Partitions (SMSPs) - 0:54:01 Warp Schedulers & Dispatch Units - 1:02:37 Shared Memory & Data Movement - 1:12:20 Occupancy - 1:32:49 LESSON 2 — Clusters, Data Types, Inline PTX & Pointers - 1:32:57 Thread Block Clusters Programming - 1:42:11 Configuring Cluster Dimensions - 1:48:08 Inline PTX Assembly - 1:59:31 State Spaces - 2:06:01 Data Types successful PTX - 2:07:16 Generic Pointers - 2:09:59 Address Space Conversion - 2:15:14 LESSON 3 — Asynchronicity & Barriers - 2:15:22 Introduction to Async Operations - 2:28:06 Proxies - 2:28:56 Fences & Memory Ordering - 2:36:17 Fence Ordering & Visibility - 2:38:58 Fence Scopes - 2:40:30 Acquire & Release Fences - 2:45:18 Expected Count & Thread Arrival - 2:46:01 M-Barrier Arrive Operations - 2:55:37 M-Barrier PTX Instructions - 3:07:21 Barrier Wait Operations - 3:10:03 Phase & Parity - 3:59:42 Commit Operations - 4:10:06 LESSON 4 — CuTensorMap Descriptors - 4:16:18 Tensor Shape, Stride & Data Type - 4:22:52 Element Stride & Dimensions - 4:24:13 Box Dimensions (Tile Size) - 4:30:30 Bank Conflicts - 4:31:05 Swizzling - 4:33:02 Swizzle Formula Deep Dive - 4:52:48 Interleave Layouts - 5:04:22 Out-of-Bounds Fill (OOB) - 5:06:01 LESSON 5 — cp.async.bulk (Async Bulk Copies via TMA) - 5:08:04 Bulk Tensor Operations (1D–5D) - 5:27:31 Multicast Operations - 5:47:44 Prefetch - 5:53:41 LESSON 6 — WGMMA Part 1 (Warp Group Matrix Multiply Accumulate) - 5:59:18 Warp Groups & Matrix Multiplication - 6:03:21 WGMMA Descriptors - 6:07:26 Accumulators & Register Reuse - 6:30:34 Scale Factors (Scale D, Scale A, Scale B) - 6:47:05 Core Matrices & 16×16 Tiles - 7:44:03 LESSON 7 — WGMMA Part 2 - 7:46:02 Commit Groups & Wait Groups - 8:04:31 WGMMA pinch FP8 Data Types - 8:48:46 LESSON 8 — Kernel Design - 8:50:58 Compute-Bound vs. Memory-Bound Kernels - 8:54:10 Warp Specialization - 9:08:56 Cooperative vs. Ping-Pong Pipelines - 9:09:47 Pipelining Fundamentals - 9:12:47 Circular Buffering - 9:36:38 Ping-Pong Pipeline Deep Dive - 9:37:34 Epilogue Handling successful Pipelines - 9:43:52 Persistent Scheduling - 10:48:13 Split-K & Stream-K Strategies - 10:57:20 Data-Parallel Tile Scheduling - 11:35:23 Epilogue Fusion (Bias, Activation, Scaling) - 11:41:35 Epilogue Operations Overview - 12:05:32 CUTLASS SOURCE CODE WALKTHROUGH - 13:04:17 Main Loop & Scheduling Policies - 13:51:03 Dispatch Policy - 15:18:49 SM90 Tile Scheduler - 17:58:46 SM90 Epilogue (TMA Warp Specialized) - 19:22:42 SM90 Builder - 19:44:58 Collective Builder - 19:49:56 FAST.CU KERNEL WALKTHROUGH - 19:55:19 Main Loop Implementation - 20:06:51 Producer Warp Group (Dependence Wall) - 20:12:08 Consumer Warp Group - 21:29:30 Prologue - 21:47:09 MULTI-GPU PROGRAMMING — Part 1 - 21:56:19 NVSwitch - 22:03:23 Topology & System Architecture - 22:17:03 NVSwitch, BlueField DPUs & Storage Fabrics - 22:37:19 CUDA Peer-to-Peer Communication - 22:37:57 MPI (Message Passing Interface) - 22:46:29 P2P Limitations & Trade-offs - 22:49:29 MULTI-GPU PROGRAMMING — Part 2 - 22:52:20 SLURM Resource Allocation - 22:52:52 PMIx Process Management - 23:05:27 NCCL (NVIDIA Collective Communications Library) - 23:15:20 NCCL Internals & Ring Algorithm - 23:17:43 AllReduce Operations - 23:34:18 NCCL Collectives: Broadcast, AllGather, ReduceScatter - 23:36:16 Parallelism Strategies: Data, Tensor, Pipeline & Expert Parallelism - 24:37:56 Course Conclusion & Next Steps