AI Factory Glossary

GPU Multi-Process,Service MPS,sharing,virtualization

**GPU Multi-Process Service MPS** is **an NVIDIA GPU feature enabling multiple CPU processes to concurrently utilize GPU resources through time-slicing and context management — enabling higher GPU utilization by preventing GPU idleness during CPU process switching and improving throughput for workloads with many small GPU kernels**. GPU multi-process service addresses the limitation that traditional GPU execution isolates each CPU process with exclusive access to GPU, preventing concurrent execution of kernels from different processes and leaving GPU idle during context switch delays. The MPS system uses proxy connections where multiple processes communicate with single connection to GPU, with central MPS daemon managing GPU resource allocation and scheduling across connected processes. The concurrency level in MPS is limited by GPU architecture and resource constraints, with typical implementations supporting 16-32 concurrent process contexts depending on GPU generation. The performance characteristics of MPS depend on workload mixing and GPU resource availability, with processes having incompatible resource requirements potentially causing contention and reduced overall throughput. The isolation guarantees in MPS are reduced compared to exclusive process contexts, with multiple processes sharing execution resources and potentially exhibiting cache interference and other contention effects. The performance prediction with MPS is challenging due to dynamic scheduling and resource contention, requiring careful measurement and profiling to validate application performance with MPS enabled. The power efficiency improvements from MPS come from higher GPU utilization reducing idle time and associated power consumption, often resulting in significant energy savings despite slightly reduced per-application performance. **GPU multi-process service MPS enables concurrent GPU access by multiple CPU processes through resource sharing and scheduling, improving aggregate system throughput.**

gpu nvlink interconnect,nvswitch all to all,nvlink 4.0 bandwidth,nvlink c2c chip to chip,nvidia dgx h100 nvlink

**NVLink and NVSwitch GPU Interconnects: High-Bandwidth All-to-All GPU Networking — specialized interconnect technology enabling 900 GB/s per-GPU communication for tightly-coupled multi-GPU systems** **NVLink 4.0 Specifications** - **Bandwidth per Direction**: 450 GB/s (bidirectional = 900 GB/s total), 5.6× faster than PCIe Gen5 x16 (64 GB/s) - **Scalability**: up to 8 GPUs per node (NVLink links 3-4× per GPU, some shared), full bisection bandwidth between pairs - **Latency**: sub-microsecond GPU-to-GPU communication (vs 1-2 µs PCIe latency), enables fine-grain synchronization - **Power Efficiency**: 900 GB/s with modest power (~20% of GPU compute power), superior to PCIe (higher power for lower bandwidth) - **Protocol**: extends PCIe protocol (NVLink 3.0 based on PCIe 4.0, NVLink 4.0 on PCIe 5.0 electrical) **NVSwitch 3.0 Architecture** - **All-to-All Connectivity**: 8-way crossbar switch (full mesh within node), any GPU pair achieves 900 GB/s simultaneously - **Bisection Bandwidth**: 57.6 TB/s total (8 GPUs × 900 GB/s × 8 directions), non-blocking (no contention) - **Scalability**: single switch per 8 GPUs (typical), larger clusters cascade switches (rack-level switches for multi-rack) - **Switching Latency**: minimal (sub-microsecond), transparent to GPU communication - **Design**: custom switch ASIC (not Ethernet switch), optimized for GPU protocols **DGX H100 Superchip Node Architecture** - **8 H100 GPUs**: full 8-way NVSwitch 3.0 connectivity, all-to-all GPU communication at 900 GB/s - **CPU**: 12-core Intel Xeon (or AMD EPYC), connected to GPU cluster via NVLink-C2C (see below) - **Memory**: 141 GB total GPU memory (16 GB HBM3 per GPU, shared via NVSwitch), coherent memory model - **Power**: ~10.2 kW for 8 H100s + CPU (8 GPUs × 700 W + 500 W CPU), thermal challenge - **Performance**: 141 TFLOPS FP32 aggregate (8 GPUs × 17.5 TFLOPS each), 700+ TFLOPS with sparsity/quantization **NVLink-C2C (Chip-to-Chip)** - **Grace-Hopper Superchip**: Grace CPU (ARM-based, 144 cores) + Hopper GPU (132 SMs) on single module (not separate dice) - **Integration**: CPU + GPU share high-bandwidth interconnect (900 GB/s), coherent memory (CPU accesses GPU HBM, vice versa) - **Use Case**: CPU for system services (PCIe control, memory management), GPU for compute, tight coupling enables efficient communication - **Deployment**: Frontier compute nodes use Grace-Hopper (experimental, limited volume) **NVLink vs PCIe Comparison** - **Bandwidth**: NVLink 4.0 (900 GB/s) vs PCIe Gen5 x16 (64 GB/s), 14× advantage - **Latency**: NVLink <1 µs vs PCIe 1-2 µs, 2× improvement - **Power**: NVLink more power-efficient (lower power per Gbps), benefits multi-GPU workloads - **Cost**: NVLink expensive (specialized silicon), justified for HPC/AI (not consumer) - **Industry Support**: NVLink proprietary (NVIDIA only), vs PCIe open standard (AMD, Intel) **NVLink over Fiber** - **NVLink-f**: optical NVLink (fiber-based), enables long-distance GPU communication (100+ meters) - **Use Case**: disaggregated GPU clusters (GPUs in separate racks), avoids copper interconnect limitations - **Latency**: fiber adds ~10-100 ns per meter, acceptable for across-datacenter links - **Adoption**: still experimental (research deployments), future potential for flexible GPU pools **Multi-GPU Scaling in Deep Learning** - **Data Parallelism**: batch split across GPUs, each GPU gradient computed independently, allreduce synchronizes gradients - **Model Parallelism**: model split across GPUs (layers on different GPUs), forward pass traverses GPUs (serial communication) - **Pipeline Parallelism**: layers pipelined (GPU 0→1→2→3 stage-by-stage), reduces synchronization overhead - **Gradient Aggregation**: allreduce critical bottleneck (all GPUs exchange gradients), NVLink reduces latency/bandwidth penalty **Communication Efficiency** - **Gradient Bandwidth**: 8 GPUs × 40 GB gradients = 320 GB gradients per step, allreduce requires 2× (reduce + broadcast) - **NVLink Advantage**: 900 GB/s enables allreduce in ~700 ns (640 GB / 900 GB/s), negligible vs 100+ ms compute per batch - **Scalability**: 100 nodes × 8 GPUs = 800 GPUs, allreduce scales O(log 800) = 10 steps (vs 800 steps if sequential) - **Overhead**: allreduce <5% of training time (with NVLink), vs 10-20% without NVLink optimization **NVIDIA GH200 Superchip (Future)** - **Integration**: Grace CPU + Hopper GPU stacked 3D (face-to-face), higher bandwidth + lower latency than separate chips - **Memory**: 141 GB HBM shared (CPU + GPU), coherent access model - **Expected Performance**: 4-5× memory bandwidth vs separate Grace+H100 (via 3D stacking) - **Deployment**: targeted at AI (training + inference), emerging 2024-2025 **Challenges** - **Heat Dissipation**: 8 H100s in single node = 10+ kW power density (requires liquid cooling), thermal management critical - **Scalability Beyond 8**: beyond-8-GPU scaling requires multi-level NVSwitch (rack-level switches), introduces latency hierarchy - **Synchronization**: tightly-coupled GPUs require frequent synchronization (allreduce every few steps), latency-sensitive **Future Roadmap**: NVLink generation per GPU generation (+50% bandwidth typically), optical interconnect NVLink-f emerging, heterogeneous GPU clusters (mix of CPU+GPU types) requiring flexible interconnects.

gpu occupancy optimization,cuda occupancy calculator,register pressure gpu,shared memory occupancy,thread block sizing

**GPU Occupancy Optimization** is **the practice of maximizing the number of active warps per Streaming Multiprocessor (SM) relative to the hardware maximum — balancing register usage, shared memory allocation, and thread block configuration to keep the GPU's warp scheduler fully utilized and hide memory access latency**. **Occupancy Definition:** - **Theoretical Occupancy**: ratio of active warps per SM to maximum warps supported by the hardware; e.g., A100 supports 64 warps (2048 threads) per SM; if a kernel achieves 32 active warps, occupancy is 50% - **Achieved Occupancy**: actual runtime average of active warps per cycle, accounting for block launch timing and completion; typically 5-15% lower than theoretical due to partial waves and resource fragmentation - **Sufficient Occupancy**: diminishing returns above ~50-60% occupancy for compute-bound kernels; memory-bound kernels benefit from higher occupancy (more warps to hide memory latency); the exact threshold is workload-dependent **Resource Limiters:** - **Register Usage**: each SM has a fixed register file (e.g., 65536 registers on A100); kernel using 64 registers per thread limits occupancy to 1024 threads (32 warps, 50% of max); reducing to 32 registers enables full occupancy but may increase register spilling to local memory - **Shared Memory Per Block**: each SM has limited shared memory (e.g., 164 KB on A100 configurable vs L1); a kernel using 48 KB shared memory per block can fit 3 blocks per SM; increasing to 96 KB limits to 1 block per SM - **Thread Block Size**: block size must be a multiple of warp size (32); small blocks (32 threads) may not fill SM due to max-blocks-per-SM limits; large blocks (1024 threads) may underutilize SM if resource usage per block is high - **Block Count Limitation**: each SM supports maximum blocks (e.g., 32 on A100); very small blocks (32 threads each) with low resource usage may still be limited by block count **Optimization Strategies:** - **CUDA Occupancy Calculator**: NVIDIA provides API (cudaOccupancyMaxActiveBlocksPerMultiprocessor) and spreadsheet tool; input register count, shared memory, block size → output occupancy percentage and limiting factor - **Launch Bounds**: __launch_bounds__(maxThreadsPerBlock, minBlocksPerMultiprocessor) directive hints the compiler to limit register usage to achieve target occupancy; may increase instruction count due to spilling but improves parallelism - **Register Pressure Reduction**: restructuring code to reduce live variable count; using shared memory for intermediate results; compiler flag --maxrregcount limits register allocation globally - **Dynamic Shared Memory**: using extern __shared__ with dynamic allocation rather than fixed arrays allows block size flexibility; combined with occupancy API to select optimal configuration at runtime **Beyond Occupancy:** - **Latency vs Throughput**: some kernels achieve peak performance at low occupancy by maximizing per-thread register usage and instruction-level parallelism; ILP can hide latency as effectively as thread-level parallelism - **Memory Bandwidth Saturation**: memory-bound kernels may saturate bandwidth at 50-75% occupancy; higher occupancy adds warps that compete for the same bandwidth without improving throughput - **Instruction Mix**: compute-bound kernels with high arithmetic intensity need fewer warps to saturate compute pipelines; memory-bound kernels need maximum warps to generate enough outstanding memory requests GPU occupancy optimization is **a crucial but nuanced aspect of CUDA performance tuning — high occupancy is necessary for memory-bound kernels to hide latency, but blindly maximizing occupancy at the expense of per-thread efficiency can hurt compute-bound kernels — the optimal balance requires understanding the kernel's arithmetic intensity and profiling with Nsight Compute**.

gpu occupancy optimization,occupancy calculator,warp occupancy,thread block size,sm utilization

**GPU Occupancy Optimization** is the **process of maximizing the ratio of active warps to the maximum possible warps per Streaming Multiprocessor (SM)** — achieved by carefully choosing thread block sizes and managing resource usage (registers, shared memory) to ensure enough warps are resident on each SM to hide memory latency through warp switching, though maximum occupancy does not always yield maximum performance. **Understanding Occupancy** - $\text{Occupancy} = \frac{\text{Active Warps per SM}}{\text{Max Warps per SM}}$ - Example (A100): Max 64 warps per SM. If kernel runs with 32 active warps → 50% occupancy. **What Limits Occupancy?** | Resource | A100 Limit per SM | How It Limits | |----------|------------------|---------------| | Threads/block | 1024 max | Limits threads per block | | Warps per SM | 64 max | Hard cap on active warps | | Registers per SM | 65536 | If kernel uses 64 regs/thread, 256 threads max → 8 warps | | Shared memory per SM | 164 KB (configurable) | If block uses 48 KB → only 3 blocks fit | | Blocks per SM | 32 max | Even tiny blocks: max 32 | **Register Pressure Example** - Kernel uses 32 registers per thread. - 65536 registers / 32 = 2048 threads max = 64 warps → 100% occupancy. - Kernel uses 64 registers per thread. - 65536 / 64 = 1024 threads = 32 warps → 50% occupancy. - Kernel uses 128 registers per thread. - 65536 / 128 = 512 threads = 16 warps → 25% occupancy. **Shared Memory Example** - SM has 164 KB shared memory. - Block uses 48 KB → 164/48 = 3 blocks max. - If block has 256 threads = 8 warps → 24 active warps → 37.5% occupancy. **Choosing Block Size** | Block Size | Warps/Block | Pros | Cons | |-----------|-------------|------|------| | 32 (1 warp) | 1 | Minimal shared memory | Max 32 blocks = 32 warps | | 128 (4 warps) | 4 | Good balance | Common default | | 256 (8 warps) | 8 | High occupancy | Higher shared memory/block | | 512 (16 warps) | 16 | Fewer blocks needed | Limits block count per SM | | 1024 (32 warps) | 32 | Max threads/block | Only 2 blocks possible per SM | **Occupancy vs. Performance** - Higher occupancy → more warps to switch between → better latency hiding. - BUT: Higher occupancy may force fewer registers → more register spilling → slower. - **Sweet spot**: Often 50-75% occupancy. Going from 75% to 100% rarely helps. - **Profile-driven**: Use Nsight Compute to measure actual performance vs. occupancy. **Tools** - **CUDA Occupancy Calculator**: Spreadsheet/API that computes occupancy from kernel resource usage. - `cudaOccupancyMaxPotentialBlockSize()`: API to auto-select block size for max occupancy. - **Nsight Compute**: Reports achieved occupancy, register/shared memory usage, and limiting factor. GPU occupancy optimization is **a necessary but not sufficient condition for high GPU performance** — while insufficient occupancy leaves the SM unable to hide memory latency, blindly maximizing occupancy at the cost of register spilling or reduced per-thread work can actually decrease performance, requiring empirical tuning guided by profiling.

gpu occupancy optimization,register pressure gpu,occupancy limiter,latency hiding gpu,active warps per sm

**GPU Occupancy** is the **ratio of active warps on a Streaming Multiprocessor (SM) to the maximum number of warps the SM can support — a key performance metric that determines the GPU's ability to hide memory latency through warp switching, where insufficient occupancy (too few active warps) leaves the SM idle during memory stalls while excessive resource usage per thread (registers, shared memory) is the primary factor that limits occupancy**. **Why Occupancy Matters** GPU performance relies on latency hiding through massive multithreading. When one warp stalls on a memory access (~400 cycles), the SM instantly switches to another ready warp at zero cost (hardware warp scheduling). But this only works if there are enough warps ready to execute. If occupancy is too low (e.g., 25%), the SM exhausts ready warps and stalls. **Occupancy Limiters** Each SM has fixed resources. The occupancy is the MINIMUM imposed by any resource: 1. **Registers**: Each SM has a register file (e.g., 65,536 registers on Ampere). If a kernel uses 64 registers/thread and threads come in warps of 32: 64 × 32 = 2,048 registers per warp. Max warps = 65,536 / 2,048 = 32 (but SM max may be 48). Reducing register usage to 48/thread: 48 × 32 = 1,536/warp → 42 warps. Higher occupancy. 2. **Shared Memory**: If a block uses 48 KB of shared memory, and the SM has 164 KB configured as shared, max 3 blocks per SM. If block size is 256 threads (8 warps): 3 × 8 = 24 active warps out of 48 max = 50% occupancy. 3. **Thread Block Size**: If block size is 64 (2 warps) and max blocks per SM is 16: 16 × 2 = 32 warps. Larger blocks (256 threads) may allow higher occupancy if other resources permit. **The Occupancy Trap** Higher occupancy does NOT always mean higher performance: - A kernel at 50% occupancy using more registers per thread may outperform 100% occupancy with register spilling (register values stored to/loaded from slow local memory). - A kernel with extensive shared memory reuse at 25% occupancy may be compute-bound and fully utilizing the ALUs. - The goal is ENOUGH occupancy to hide latency — typically 40-60% is sufficient for many kernels. **Tuning Tools** - **CUDA Occupancy Calculator**: Given kernel resource usage, computes theoretical occupancy. Available as spreadsheet and `cudaOccupancyMaxActiveBlocksPerMultiprocessor()` API. - **Nsight Compute**: Reports achieved occupancy, active warps, and identifies the limiting resource (registers, shared memory, or block count). - **Launch Configuration**: `__launch_bounds__(maxThreadsPerBlock, minBlocksPerSM)` hints to the compiler to limit register usage for target occupancy. **GPU Occupancy is the resource-constrained balancing act of GPU programming** — trading per-thread resource richness (registers, shared memory) against parallelism (active warps), where the optimal balance depends on whether the kernel is memory-latency-bound, compute-bound, or bandwidth-bound.

gpu occupancy optimization,sm occupancy,register pressure gpu,thread block size selection,occupancy calculator

**GPU Occupancy Optimization** is the **performance tuning discipline that maximizes the number of active warps per Streaming Multiprocessor (SM) — measured as the ratio of active warps to the SM's maximum supported warps — to ensure that the GPU's warp scheduler always has warps ready to execute, hiding memory latency through context switching between warps rather than stalling on any single memory request**. **Why Occupancy Matters** GPU SMs operate by rapidly switching between active warps. When one warp stalls on a global memory access (~500 cycles), the scheduler immediately switches to another ready warp. With enough active warps (high occupancy), the SM stays busy while stalled warps wait for data. With too few warps (low occupancy), all warps may be stalled simultaneously → the SM sits idle. **Resources That Limit Occupancy** Each SM has fixed quantities of three resources shared among all active thread blocks: | Resource | H100 SM Limit | How It Limits Occupancy | |----------|---------------|------------------------| | **Registers** | 65,536 per SM | Kernel using 64 regs/thread × 256 threads/block = 16,384 regs/block → max 4 blocks/SM | | **Shared Memory** | 228 KB per SM | Kernel using 48KB/block → max 4 blocks (192KB used) | | **Thread Blocks** | 32 per SM | Hard limit regardless of resource usage | | **Warps** | 64 per SM | Maximum occupancy = 64 warps × 32 threads = 2048 threads/SM | Occupancy is limited by whichever resource is exhausted first. **Register Pressure** Registers are the most common occupancy limiter. A complex kernel with many variables may use 128 registers per thread, limiting occupancy to 2 blocks of 256 threads (25% occupancy). Reducing register usage (via `__launch_bounds__`, algorithmic simplification, or register spilling to local memory) increases occupancy — but spilling registers to memory adds latency. The optimum is usually 50-75% occupancy, not maximum occupancy. **Diminishing Returns** Occupancy beyond 50% often provides minimal additional performance because: 1. Enough warps already exist to hide memory latency. 2. Cache thrashing increases as more blocks compete for the same L1/shared memory. 3. Register spilling to achieve higher occupancy adds local memory traffic that offsets the latency-hiding benefit. The right approach: start at maximum occupancy, benchmark, then systematically trade occupancy for more registers/shared memory per thread if it improves IPC. **Tooling** - **CUDA Occupancy Calculator**: Excel spreadsheet or `cudaOccupancyMaxActiveBlocksPerMultiprocessor()` API. Takes kernel register count, shared memory, and block size → reports achievable occupancy. - **Nsight Compute**: Profiles actual vs. theoretical occupancy and identifies the limiting resource. Shows achieved occupancy (affected by workload) vs. theoretical (resource-limited). GPU Occupancy Optimization is **the art of balancing resource allocation per thread against total active parallelism** — giving each thread enough registers and shared memory to work efficiently while ensuring enough warps exist to keep the SM continuously busy.

gpu oom, out of memory, cuda error, gradient checkpointing, quantization, memory optimization, batch size

**GPU Out of Memory (OOM)** errors occur when **model weights, activations, or intermediate computations exceed available GPU VRAM** — a common issue in ML development that requires understanding memory usage patterns and applying techniques like gradient checkpointing, mixed precision, quantization, and batch size reduction to fit models within constraints. **What Is GPU OOM?** - **Error**: "CUDA out of memory" or "RuntimeError: CUDA error: out of memory." - **Cause**: GPU memory (VRAM) exhausted by model/data/activations. - **Context**: Training uses more memory than inference. - **Resolution**: Reduce memory usage or increase available VRAM. **Why OOM Happens** - **Model Weights**: Large models need gigabytes for parameters. - **Activations**: Saved for backpropagation during training. - **Optimizer States**: Adam stores 2x parameters in memory. - **Gradients**: Same size as parameters. - **KV Cache**: For inference, grows with sequence length. - **Batch Size**: More samples = more memory. **Diagnosis** **Check Current Usage**: ```bash # Current GPU memory nvidia-smi # Real-time monitoring watch -n1 nvidia-smi # Detailed per-process nvidia-smi pmon -s m ``` **Python Memory Tracking**: ```python import torch # Check memory usage print(f"Allocated: {torch.cuda.memory_allocated() / 1e9:.2f} GB") print(f"Cached: {torch.cuda.memory_reserved() / 1e9:.2f} GB") # Get detailed snapshot torch.cuda.memory_summary() ``` **Memory Estimation**: ``` Model memory (FP16) ≈ 2 × parameters Example (7B model): Parameters: 7B × 2 bytes = 14 GB Optimizer (Adam): 14 GB × 2 = 28 GB Gradients: 14 GB Activations: Variable (~10-20 GB) ───────────────────────────── Training total: ~70-80 GB (won't fit on single 80GB GPU!) ``` **Solutions** **Reduce Batch Size** (First try): ```python # If batch_size=32 OOMs: batch_size = 16 # Try smaller # Or even batch_size = 1 with gradient accumulation ``` **Gradient Accumulation** (Same effective batch): ```python accumulation_steps = 8 for i, batch in enumerate(dataloader): loss = model(batch) / accumulation_steps loss.backward() if (i + 1) % accumulation_steps == 0: optimizer.step() optimizer.zero_grad() ``` **Gradient Checkpointing** (Trade compute for memory): ```python # PyTorch native model.gradient_checkpointing_enable() # Hugging Face model = AutoModel.from_pretrained( "model-name", gradient_checkpointing=True ) # Savings: 2-3× memory reduction # Cost: ~20% slower training ``` **Mixed Precision Training**: ```python from torch.cuda.amp import autocast, GradScaler scaler = GradScaler() for batch in dataloader: with autocast(): # FP16 compute loss = model(batch) scaler.scale(loss).backward() scaler.step(optimizer) scaler.update() # Savings: ~2× memory for activations ``` **Quantization** (For inference): ```python # bitsandbytes 4-bit from transformers import BitsAndBytesConfig bnb_config = BitsAndBytesConfig( load_in_4bit=True, bnb_4bit_compute_dtype=torch.float16 ) model = AutoModelForCausalLM.from_pretrained( "model-name", quantization_config=bnb_config ) # 7B model: 14 GB → ~4 GB ``` **Clear Cache**: ```python # Clear unused cached memory torch.cuda.empty_cache() # Delete unused variables del large_tensor torch.cuda.empty_cache() # Use context manager for temporary tensors with torch.no_grad(): # Inference without saving gradients output = model(input) ``` **Memory-Efficient Techniques Summary** ``` Technique | Memory Savings | Trade-off ---------------------|----------------|------------------- Smaller batch size | Linear | More iterations Gradient accumulation| None (same effect)| Code complexity Gradient checkpointing| 2-3× | 20% slower Mixed precision (FP16)| 2× activations| Minor precision Quantization (INT4) | 4× weights | Quality varies Flash Attention | ~2× attention | None DeepSpeed ZeRO | Split across GPUs| Multi-GPU needed ``` **Inference OOM** ```python # Use vLLM for efficient inference from vllm import LLM llm = LLM( model="model-name", quantization="awq", # 4-bit quantization gpu_memory_utilization=0.9 # Use 90% of VRAM ) # Reduce context length if needed llm = LLM(model="model", max_model_len=4096) ``` GPU OOM is **the most common issue in ML development** — understanding where memory goes and systematically applying reduction techniques enables running larger models on available hardware, making memory optimization skills essential for ML engineers.

gpu operator,device plugin,nvidia

**GPU Management in Kubernetes** **NVIDIA GPU Operator** Automates GPU driver installation, container toolkit, and device plugins in Kubernetes. **Installation** ```bash # Add NVIDIA Helm repo helm repo add nvidia https://helm.ngc.nvidia.com/nvidia # Install GPU Operator helm install gpu-operator nvidia/gpu-operator --namespace gpu-operator --create-namespace ``` **Components Installed** | Component | Purpose | |-----------|---------| | Driver | NVIDIA GPU drivers | | Container Toolkit | nvidia-container-runtime | | Device Plugin | Expose GPUs to K8s scheduler | | DCGM Exporter | GPU metrics for Prometheus | | MIG Manager | Multi-Instance GPU config | **Requesting GPUs in Pods** ```yaml apiVersion: v1 kind: Pod metadata: name: llm-server spec: containers: - name: inference image: llm-inference:latest resources: limits: nvidia.com/gpu: 1 # Request 1 GPU ``` **Multiple GPUs** ```yaml resources: limits: nvidia.com/gpu: 4 # Multi-GPU for large models ``` **Node Selectors for GPU Types** ```yaml spec: nodeSelector: nvidia.com/gpu.product: "NVIDIA-A100-SXM4-80GB" containers: - name: model resources: limits: nvidia.com/gpu: 1 ``` **GPU Sharing (Time-Slicing)** ```yaml # ConfigMap for time-slicing apiVersion: v1 kind: ConfigMap metadata: name: time-slicing-config data: any: |- version: v1 sharing: timeSlicing: resources: - name: nvidia.com/gpu replicas: 4 # 4 pods can share each GPU ``` **MIG (Multi-Instance GPU)** Split A100/H100 into multiple instances: ```yaml resources: limits: nvidia.com/mig-3g.20gb: 1 # 3GB compute, 20GB memory slice ``` **Monitoring GPUs** ```bash # Check GPU allocation kubectl describe nodes | grep nvidia.com/gpu # View GPU metrics kubectl logs -n gpu-operator dcgm-exporter-xxx # GPU utilization in Grafana via DCGM metrics ``` **Best Practices** - Use GPU Operator for consistent setup - Set appropriate GPU limits - Use node selectors for GPU types - Monitor GPU utilization - Consider time-slicing for dev environments - Use MIG for flexible resource allocation

gpu performance profiling nsight,nvtx annotation,roofline model gpu,achieved bandwidth occupancy,gpu bottleneck analysis

**GPU Performance Profiling** encompasses **systematic measurement and analysis of kernel execution, memory access patterns, and hardware utilization using Nsight tools, roofline models, and application-specific metrics to identify bottlenecks and guide optimization.** **Nsight Compute and Nsight Systems Overview** - **Nsight Compute**: Kernel-centric profiler. Analyzes single kernel execution: register/shared memory usage, L1/L2 cache hit rates, warp stall reasons, SM efficiency. - **Nsight Systems**: System-wide profiler. Timeline view of entire application: kernel launches, memory transfers, CPU-GPU synchronization, context switches, power consumption. - **Guided Analysis Workflow**: Nsight Compute recommends optimizations based on measured metrics (e.g., "warp occupancy 50%, increase shared memory usage to 75%"). - **Overhead**: Profiling adds ~5-50% runtime overhead depending on metric set. Light profiling (SM efficiency) minimal; heavy profiling (register spills) substantial. **NVTX Annotations for Custom Metrics** - **NVTX (NVIDIA Tools Extension)**: API to annotate application code. Marks user-defined ranges, domains, events with custom names. - **Range Annotation**: nvtxRangePush/Pop() delineate code sections. Nsight timeline shows annotated regions, enabling user-level performance tracking. - **Domain Separation**: nvtxDomainCreate() organizes related annotations. Example: separate domains for preprocessing, compute, postprocessing. - **Color and Category**: Annotations assigned colors (visual grouping) and categories (filtering). Facilitates timeline analysis of complex multi-threaded applications. **Roofline Model for GPU Analysis** - **Roofline Concept**: 2D plot of achievable GFLOP/s vs arithmetic intensity (FLOP per byte transferred). Machine peak provides "roofline" ceiling. - **Peak Compute Roofline**: GPU compute peak (theoretical FP32 FLOP/s). Ampere A100: 312 TFLOP/s peak. - **Peak Bandwidth Roofline**: GPU memory bandwidth (theoretical throughput). A100 HBM2e: 2 TB/s peak. Roofline ceiling = MIN(peak_compute, intensity × peak_bandwidth). - **Application Characterization**: Measure kernel arithmetic intensity (FLOP count / memory bytes transferred). Points below roofline indicate under-utilization. **Achieved Occupancy and Bottleneck Analysis** - **Occupancy Metric**: Percentage of SM warp slots filled. Occupancy = (resident_warps / max_warps_per_sm) × 100%. Max warp/SM: 64 (Volta), 48 (Ampere). - **Limiting Factors**: Register pressure (32k limit per SM), shared memory allocation (96KB per SM), thread blocks per SM (varies by GPU). - **Occupancy vs Performance**: Higher occupancy generally improves performance (more warps hide memory latency), but not always. Some high-register kernels benefit from lower occupancy. - **Warp Stall Reasons**: Nsight reports stall causes (memory, dependency, execution resource, synchronization). Prioritize fixing most-common stall. **Memory Bandwidth Utilization** - **Effective Bandwidth**: Measured memory bytes (profiler) vs theoretical peak. Typical ratios: 50-90% depending on access pattern. - **Coalescing Efficiency**: Consecutive threads accessing consecutive memory addresses coalesce into single transaction. Scattered access wastes bandwidth (cache-only reuse). - **Bank Conflicts**: Shared memory bank conflicts serialize accesses. All 32 threads accessing same bank → 32x slowdown. Proper access pattern avoids conflicts. - **L2 Cache Effectiveness**: L2 cache hit rate impacts bandwidth. Reuse distance (iterations between data access) determines cache utility. **Cache Utilization and Patterns** - **L1 Cache**: Per-SM cache (32-96KB depending on config). Caches load/store operations if enabled. Bank conflicts similar to shared memory. - **L2 Cache**: Shared across all SMs (4-40 MB depending on GPU). Victim cache for L1, also receives uncached loads. - **Hit Rate Interpretation**: High L1 hit rate (>80%) indicates locality; low ratio indicates poor spatial/temporal locality. - **Profiler L2 Analysis**: Misses per 1k instructions metric. Aim for <2-5 misses/1k instructions for well-optimized kernels. **SM Efficiency and Load Balancing** - **SM Efficiency**: Percentage of SM slots executing useful instructions. Idle slots due to warp stalls, divergence, or under-occupancy. - **Warp Divergence Analysis**: Branch divergence metrics show divergence frequency and impact. Serialization within warp reduces throughput. - **Grid-Level Load Balancing**: Blocks distributed unevenly → some SMs idle while others compute. Profiler shows block-per-SM histogram. - **Dynamic Parallelism Overhead**: Child kernels launched from kernel require synchronization overhead. Impacts SM efficiency if child kernels small. **Optimization Workflows** - **Memory-Bound Analysis**: If roofline point below bandwidth line, kernel memory-bound. Optimize: improve coalescing, increase data reuse, prefetching. - **Compute-Bound Analysis**: If roofline point below compute line, kernel compute-bound. Optimize: reduce instruction count, use tensor cores, improve ILP. - **Iterative Refinement**: Profile → identify bottleneck → optimize → re-profile. Typical 5-10 iteration cycle for 2-5x speedup.

GPU Persistent,Threads,pattern,kernel design

**GPU Persistent Threads Pattern** is **an advanced GPU kernel design pattern where single kernel launch creates threads that persist across multiple iterations of input data processing — enabling sophisticated state management, dynamic load balancing, and algorithmic flexibility impossible in conventional bulk-synchronous GPU programming models**. The persistent threads pattern addresses the limitation of conventional GPU programming where kernel launch overhead (microseconds) can become significant for kernels with short execution time, with persistent kernel design amortizing launch overhead across many iterations. The persistent kernel structure typically involves loop within kernel where threads iterate over input data, processing multiple items per thread rather than strictly one block per item conventional decomposition. The dynamic load balancing enabled by persistent kernels allows threads to request additional work items dynamically rather than statically predetermined work decomposition, enabling natural load balancing for irregular algorithms. The state accumulation across iterations in persistent kernels enables sophisticated state management and aggregation patterns, supporting algorithms with iterative refinement or multi-pass processing. The synchronization patterns in persistent kernels are more complex than conventional kernels, requiring careful attention to prevent deadlock or excessive synchronization overhead. The performance characteristics depend critically on loop iteration count, thread block geometry, and whether computation is memory-bound or compute-bound, requiring careful tuning. The debugging and correctness verification of persistent kernels is more challenging than conventional kernels due to complex control flow and state management. **GPU persistent threads pattern enables sophisticated kernel design with dynamic load balancing and state management through persistent loop-based kernels.**

gpu power management,dvfs gpu,gpu power limit,gpu frequency scaling,gpu thermal throttle

**GPU Power Management and DVFS** is the **dynamic adjustment of GPU clock frequency and voltage to balance performance, power consumption, and thermal limits** — where modern GPUs continuously modulate their operating point hundreds of times per second based on workload demand, power budget, and temperature, with the GPU's actual clock speed often differing significantly from its advertised "boost" frequency. **GPU Power States** | State | Frequency | Voltage | Power | Usage | |-------|----------|---------|-------|-------| | Idle | 210 MHz | 0.65V | 10-30W | Desktop/idle | | Light Load | 800-1200 MHz | 0.75V | 50-100W | Video, light compute | | Base Clock | 1200-1800 MHz | 0.85V | 150-250W | Sustained all-core | | Boost Clock | 1800-2500 MHz | 0.95-1.1V | 250-400W | Thermal/power headroom | | Max Boost | 2500-3000 MHz | 1.05-1.1V | 400-700W | Transient, single SM | **DVFS on GPUs** - **Dynamic Voltage and Frequency Scaling**: GPU firmware continuously adjusts V and F. - $P_{dynamic} \propto C \times V^2 \times F$ — reducing voltage provides quadratic power savings. - GPU firmware reads: Temperature sensors, power sensors, workload monitors. - Decision every: ~1 ms — adjusts clock speed in real time. **Power Limiting Mechanisms** 1. **TDP (Thermal Design Power)**: Maximum sustained power the cooling solution can handle. - RTX 4090: 450W TDP. A100 SXM: 400W TDP. H100 SXM: 700W TDP. 2. **Power Limit**: Software-configurable cap. If GPU hits limit → reduce frequency. - `nvidia-smi -pl 300` — set power limit to 300W. 3. **Thermal Throttling**: If junction temperature exceeds limit (83-95°C) → reduce clock. 4. **Voltage Limit**: Maximum safe voltage for the silicon → caps max boost. **Undervolting and Overclocking** - **Undervolting**: Reduce voltage at given frequency → less power, same performance. - Risk: Instability if voltage too low for the specific silicon sample. - **Overclocking**: Increase power limit + frequency offset. - Diminishing returns: 10% more power → 3-5% more performance (voltage scaling). **Data Center GPU Power Management** - **NVIDIA MIG Power Isolation**: Each MIG instance has proportional power budget. - **Power Capping for TCO**: Data centers cap GPU power at 70-80% of max → significantly reduces cooling cost with only 5-10% performance loss. - **nvidia-smi queries**: - `nvidia-smi --query-gpu=power.draw,clocks.gr,temperature.gpu --format=csv` **GPU Power Efficiency Trend** | Generation | Performance/Watt Improvement | |-----------|----------------------------| | Kepler → Maxwell | ~2x | | Maxwell → Pascal | ~1.5x | | Pascal → Turing | ~1.5x | | Turing → Ampere | ~1.5x | | Ampere → Hopper | ~2x (FP8 ops) | GPU power management is **the invisible governor of GPU performance** — understanding how DVFS, power limits, and thermal throttling interact is essential for anyone benchmarking, deploying, or optimizing GPU workloads, as the actual sustained performance can be 20-30% below peak specifications.

gpu power management,gpu energy efficiency,power capping,gpu tdp,thermal design power gpu

**GPU Energy Efficiency and Power Management** is the **set of hardware and software mechanisms that dynamically control GPU power consumption to maximize performance within thermal and electrical constraints** — balancing the competing demands of peak computational throughput, thermal dissipation limits, power supply capacity, and data center energy budgets, where modern data center GPUs consume 300-1000W each and power/cooling costs represent 40-60% of total data center operating expenses. **GPU Power Components** | Component | Typical % of Total | Scaling | |-----------|-------------------|--------| | Compute (SM/CU cores) | 50-60% | Scales with utilization and frequency | | Memory (HBM/GDDR) | 15-25% | Scales with access rate | | Interconnect (NVLink, PCIe) | 5-10% | Scales with communication volume | | Leakage | 10-20% | Always present, increases with temperature | | I/O and misc | 5-10% | Relatively fixed | **Power Management Mechanisms** | Mechanism | Level | What It Controls | |-----------|-------|------------------| | DVFS | Hardware | Voltage and frequency per SM | | Clock gating | Hardware | Disable clocks to idle units | | Power gating | Hardware | Cut power to unused blocks | | Power capping | Software | Enforce max power limit | | Boost clocks | Firmware | Raise frequency when thermal headroom exists | | MIG (Multi-Instance GPU) | Firmware | Partition GPU into isolated instances | **NVIDIA GPU Power States** ```bash # Query current power and clocks nvidia-smi -q -d POWER,CLOCK # Set power cap to 300W (from default 400W TDP) nvidia-smi -pl 300 # Lock clocks for reproducible benchmarking nvidia-smi --lock-gpu-clocks=1200,1200 # Monitor power in real-time watch -n 1 nvidia-smi --query-gpu=power.draw,temperature.gpu,clocks.sm --format=csv ``` **Power Capping Trade-offs** | Power Cap (% of TDP) | Performance Loss | Energy Savings | Use Case | |----------------------|-----------------|---------------|----------| | 100% (default) | 0% | 0% | Maximum throughput | | 80% | 5-10% | 20% | Good efficiency point | | 60% | 20-30% | 40% | Power-constrained DC | | 40% | 40-50% | 60% | Extreme power limits | - Key insight: Power-performance is NOT linear. - Reducing power by 20% often costs only 5-10% performance → excellent efficiency. - Diminishing returns at low power: 50% cap may lose 30%+ performance. **Data Center GPU Power** | GPU | TDP | Peak Perf (FP16) | Perf/Watt | |-----|-----|-------------------|----------| | A100 (80GB) | 400W | 312 TFLOPS | 780 GFLOPS/W | | H100 (80GB) | 700W | 990 TFLOPS | 1414 GFLOPS/W | | B200 | 1000W | 2250 TFLOPS | 2250 GFLOPS/W | | MI300X (AMD) | 750W | 1300 TFLOPS | 1733 GFLOPS/W | **Energy-Efficient Training Strategies** - **Lower precision**: FP16/BF16 → 2× throughput at similar power → 2× energy efficiency. - **Power-capped long runs**: Run at 80% power → 5% slower but 15% less total energy. - **Batch size tuning**: Larger batches → better GPU utilization → more FLOPS per joule. - **Dynamic scaling**: Scale down GPUs during communication phases (gradient sync). GPU power management is **the critical constraint shaping data center AI infrastructure** — with a single AI training cluster consuming megawatts of power (enough for a small town), optimizing the energy efficiency of GPU computation is both an economic imperative and an environmental responsibility, where techniques like power capping and precision reduction can reduce total training energy by 20-40% with minimal impact on model quality.

GPU Power,thermal management,cooling efficiency

**GPU Power and Thermal Management** is **a critical GPU system design and programming discipline ensuring that GPU power consumption and heat dissipation remain within system thermal and power budgets while maintaining performance — requiring cooperation between hardware design, system architecture, and software optimization**. GPU power consumption scales dramatically with operating frequency and supply voltage, with power dissipation reaching hundreds of watts in high-performance GPUs, creating thermal challenges requiring sophisticated cooling infrastructure. The power capping mechanisms limit total GPU power to specified budgets, with automatic frequency/voltage adjustment reducing performance when power approaches limits, preventing thermal runaway but potentially impacting application performance. The dynamic power management features including dynamic voltage and frequency scaling (DVFS) enable runtime adjustment of GPU clock frequency and supply voltage based on workload demands, reducing power consumption for light workloads while enabling peak performance for compute-intensive tasks. The thermal throttling automatically reduces clock frequency when GPU temperature exceeds safe limits, providing protection against hardware damage but potentially causing unpredictable performance variations. The cooling system design for GPUs involves heat sink sizing, thermal interface materials, and airflow management ensuring effective heat transfer from GPU die to ambient environment. The measurement and profiling of GPU power consumption and thermal characteristics enable identification of power-hungry kernels and opportunities for power optimization. The algorithmic optimization for power reduction including reduced precision (16-bit or 8-bit arithmetic), lower frequency execution, or memory access pattern optimization can reduce power without proportional performance loss for many applications. **GPU power and thermal management through dynamic frequency scaling and thermal monitoring ensures safe operation within system constraints while maintaining performance.**

gpu profiling debugging,nsight compute profiling,nsight systems timeline,cuda profiling tools,gpu performance analysis

**GPU Profiling and Debugging** is **the systematic analysis of GPU application performance and correctness using specialized tools that provide detailed metrics, timeline visualization, and error detection** — where NVIDIA Nsight Compute delivers kernel-level analysis with 1000+ metrics covering memory bandwidth (achieved vs peak 1.5-3 TB/s), compute throughput (achieved vs peak 20-80 TFLOPS), occupancy (50-100%), and warp efficiency (target >90%), while Nsight Systems provides system-wide timeline showing CPU-GPU interaction, kernel launches, memory transfers, and API calls, enabling developers to identify bottlenecks (memory-bound, compute-bound, latency-bound), optimize resource utilization, and achieve 2-10× performance improvement through data-driven optimization, making profiling essential for GPU development where intuition often misleads and measurement is the only path to understanding actual performance characteristics. **Nsight Compute (Kernel Profiling):** - **Purpose**: detailed single-kernel analysis; 1000+ metrics; memory, compute, occupancy, warp efficiency; identifies kernel bottlenecks - **Launch**: ncu ./app or ncu --set full ./app; GUI: ncu-ui; command-line or graphical interface - **Metrics**: Memory Throughput (GB/s), Compute Throughput (TFLOPS), SM Efficiency (%), Occupancy (%), Warp Execution Efficiency (%), Branch Efficiency (%) - **Sections**: Memory Workload Analysis, Compute Workload Analysis, Launch Statistics, Occupancy, Scheduler Statistics, Warp State Statistics **Nsight Systems (System Profiling):** - **Purpose**: system-wide timeline; CPU-GPU interaction; kernel launches, memory transfers, API calls; identifies system-level bottlenecks - **Launch**: nsys profile ./app or nsys profile --trace=cuda,nvtx ./app; generates .qdrep file; open in nsys-ui - **Timeline View**: visualizes all GPU activity; shows overlaps, gaps, synchronization points; identifies idle time - **Use Cases**: multi-GPU profiling, stream concurrency, CPU-GPU overlap, kernel launch overhead, memory transfer analysis **Memory Profiling:** - **Memory Throughput**: achieved bandwidth / peak bandwidth; target 80-100% for memory-bound kernels; A100: 1.5-2 TB/s, H100: 2-3 TB/s - **Memory Replay**: indicates uncoalesced access; replay >1.5 means poor coalescing; restructure data layout - **L1/L2 Hit Rate**: cache effectiveness; high hit rate (>80%) good for reused data; low hit rate indicates streaming access - **Global Load/Store Efficiency**: percentage of useful bytes loaded; low efficiency (<50%) indicates wasted bandwidth; improve coalescing - **Bank Conflicts**: shared memory bank conflicts; high conflicts (>10%) cause serialization; add padding or change access pattern **Compute Profiling:** - **Compute Throughput**: achieved TFLOPS / peak TFLOPS; target 50-80% for compute-bound kernels; A100: 19.5 TFLOPS FP32, 312 TFLOPS FP16 - **SM Efficiency**: percentage of time SMs are active; target 80-100%; low efficiency indicates insufficient work or poor scheduling - **Tensor Core Utilization**: percentage of time Tensor Cores active; target 50-80% for matrix operations; 312 TFLOPS on A100 - **IPC (Instructions Per Cycle)**: instructions executed per cycle; higher is better; target 2-4 for well-optimized kernels **Occupancy Analysis:** - **Achieved Occupancy**: percentage of maximum warps active; target 50-100%; higher occupancy hides latency - **Theoretical Occupancy**: maximum possible based on resource usage; limited by registers, shared memory, block size - **Occupancy Limiter**: identifies limiting factor (registers, shared memory, block size); guides optimization - **Occupancy Calculator**: CUDA Occupancy Calculator spreadsheet; predicts occupancy from resource usage; useful for tuning **Warp Efficiency:** - **Warp Execution Efficiency**: percentage of active threads in executed warps; target >90%; low efficiency indicates divergence - **Branch Efficiency**: percentage of branches without divergence; target >90%; divergent branches cause serialization - **Predication Efficiency**: percentage of instructions not predicated off; target >90%; high predication indicates divergence - **Optimization**: minimize divergence; use ballot/shuffle for divergent code; restructure algorithms **Roofline Model:** - **Concept**: plots achieved performance vs arithmetic intensity; shows whether memory-bound or compute-bound - **Memory Roof**: horizontal line at peak memory bandwidth; memory-bound kernels hit this ceiling - **Compute Roof**: diagonal line at peak compute throughput; compute-bound kernels hit this ceiling - **Optimization**: move toward upper-right (higher intensity, higher performance); tiling increases intensity **Timeline Analysis:** - **Kernel Gaps**: idle time between kernels; indicates launch overhead or synchronization; use streams to overlap - **Memory Transfer Gaps**: idle time during transfers; use async transfers and streams; overlap with compute - **CPU-GPU Sync**: cudaDeviceSynchronize() causes gaps; minimize synchronization; use events for fine-grained control - **Multi-GPU**: visualize cross-GPU communication; identify load imbalance; optimize data distribution **NVTX Markers:** - **Purpose**: annotate code regions; shows in Nsight Systems timeline; helps identify bottlenecks in application logic - **API**: nvtxRangePush("label"), nvtxRangePop(); marks code regions; nvtxMark("event") for single events - **Use Cases**: mark training iterations, data loading, preprocessing, inference; correlate with GPU activity - **Best Practice**: annotate all major code sections; hierarchical markers; color-code by category **Debugging Tools:** - **cuda-memcheck**: detects memory errors; out-of-bounds access, race conditions, uninitialized memory; run with cuda-memcheck ./app - **Compute Sanitizer**: newer tool replacing cuda-memcheck; more features; memcheck, racecheck, initcheck, synccheck modes - **CUDA_LAUNCH_BLOCKING=1**: serializes all operations; easier debugging; disables async; use only for debugging - **cuda-gdb**: command-line debugger; breakpoints, watchpoints, inspect variables; cuda-gdb ./app **Performance Metrics:** - **Achieved Bandwidth**: GB/s of memory traffic; compare to peak (1.5-3 TB/s); target 80-100% for memory-bound - **Achieved TFLOPS**: floating-point operations per second; compare to peak (20-80 TFLOPS); target 50-80% for compute-bound - **Kernel Time**: total kernel execution time; identify slow kernels; focus optimization efforts - **Launch Overhead**: time between kernel launches; target <1% of total time; use CUDA Graphs to reduce **Bottleneck Identification:** - **Memory-Bound**: <50% compute throughput, high memory throughput; optimize memory access patterns, use shared memory, reduce accesses - **Compute-Bound**: <50% memory throughput, high compute throughput; use Tensor Cores, increase ILP, reduce divergence - **Latency-Bound**: low occupancy, low throughput; increase occupancy, reduce register usage, increase block size - **Instruction-Bound**: high instruction overhead; reduce branches, use warp primitives, optimize control flow **Optimization Workflow:** - **Profile**: run Nsight Compute and Nsight Systems; identify bottleneck (memory, compute, latency) - **Analyze**: examine relevant metrics; memory throughput, compute throughput, occupancy, warp efficiency - **Optimize**: apply targeted optimizations; memory coalescing, shared memory, occupancy tuning, divergence reduction - **Measure**: re-profile; verify improvement; compare metrics before and after - **Iterate**: repeat for next bottleneck; diminishing returns after 3-5 iterations; 2-10× total speedup typical **Common Profiling Patterns:** - **Baseline**: profile unoptimized code; establish baseline metrics; identify major bottlenecks - **Incremental**: optimize one aspect at a time; measure impact; easier to attribute improvements - **Comparison**: compare against reference implementation (cuBLAS, cuDNN); identify gaps; target 80-95% of library performance - **Regression**: profile after code changes; detect performance regressions; maintain performance over time **Multi-GPU Profiling:** - **Nsight Systems**: visualizes all GPUs simultaneously; shows cross-GPU communication; identifies load imbalance - **NCCL Profiling**: NCCL_DEBUG=INFO shows communication details; bandwidth, latency, algorithm selection - **Per-GPU Metrics**: profile each GPU separately; identify stragglers; optimize slowest GPU first - **Scaling Analysis**: measure scaling efficiency; compare 1 GPU vs N GPUs; target 80-95% efficiency **Advanced Profiling:** - **Sampling**: sample-based profiling for long-running applications; lower overhead; nsys profile --sample=cpu,cuda - **Metrics Collection**: collect specific metrics; ncu --metrics sm__throughput.avg.pct_of_peak_sustained_elapsed; reduces overhead - **Kernel Replay**: replay kernel with different configurations; find optimal launch parameters; ncu --launch-count 1 --replay-mode kernel - **Source Correlation**: correlate metrics with source code; identify hot spots; ncu --source-level-analysis **Performance Targets:** - **Memory Bandwidth**: 80-100% of peak (1.5-3 TB/s); coalesced access, minimal bank conflicts - **Compute Throughput**: 50-80% of peak (20-80 TFLOPS); use Tensor Cores, high ILP, minimal divergence - **Occupancy**: 50-100%; balance register and shared memory usage; 256 threads per block typical - **Warp Efficiency**: >90%; minimize divergence; uniform control flow - **Kernel Time**: <1ms for small kernels, <100ms for large; longer kernels risk timeout; split if necessary **Best Practices:** - **Profile Early**: profile from the start; avoid premature optimization but measure early; establish baseline - **Profile Often**: profile after each optimization; verify improvement; catch regressions - **Use Both Tools**: Nsight Compute for kernel details, Nsight Systems for system view; complementary insights - **Focus on Bottlenecks**: optimize slowest kernels first; 80/20 rule applies; 20% of kernels often account for 80% of time - **Measure, Don't Guess**: intuition often wrong; always measure; data-driven optimization **Common Mistakes:** - **Optimizing Wrong Thing**: optimizing fast kernels instead of slow ones; profile to identify bottlenecks - **Ignoring Occupancy**: assuming higher occupancy always better; balance with resource usage; profile to find optimal - **Over-Optimizing**: diminishing returns after 3-5 iterations; 2-10× total speedup typical; know when to stop - **Not Profiling**: relying on intuition; guessing bottlenecks; always measure actual performance **Real-World Impact:** - **Matrix Multiplication**: profiling reveals 20% of peak; optimization achieves 80-95% of peak; 4-5× speedup - **Reduction**: profiling shows bank conflicts; optimization eliminates conflicts; 2-3× speedup; 60-80% of peak - **Convolution**: profiling reveals memory-bound; shared memory tiling achieves 70-90% of peak; 5-10× speedup - **Custom Kernels**: profiling guides optimization; 2-10× improvement typical; achieves 50-80% of peak GPU Profiling and Debugging represent **the essential tools for GPU performance optimization** — by providing detailed metrics, timeline visualization, and error detection through Nsight Compute and Nsight Systems, developers identify bottlenecks, optimize resource utilization, and achieve 2-10× performance improvement through data-driven optimization, making profiling the difference between GPU code that achieves 10% or 80% of theoretical peak performance where measurement is the only path to understanding actual performance characteristics and intuition often misleads.

gpu profiling nsight compute,kernel performance profiling,gpu occupancy analysis,memory throughput profiling,instruction throughput analysis

**GPU Profiling with Nsight Compute** is **the systematic analysis of GPU kernel performance characteristics — including compute throughput, memory throughput, occupancy, stall reasons, and instruction mix — to identify bottlenecks and guide optimization decisions using the detailed hardware performance counters available on NVIDIA GPUs**. **Key Profiling Metrics:** - **SM Throughput**: percentage of peak compute throughput achieved — low values indicate instruction-level inefficiency (poor ILP, warp divergence, or stalls) - **Memory Throughput**: percentage of peak memory bandwidth utilized — high values indicate a memory-bound kernel; optimization should focus on reducing memory traffic or improving access patterns - **Occupancy**: ratio of active warps to maximum warps per SM — higher occupancy helps hide latency but isn't always necessary; some kernels achieve peak performance at 50% occupancy with good data reuse - **Warp Execution Efficiency**: average number of active threads per warp instruction — values below 32 indicate divergence; target >28 for well-optimized kernels **Stall Analysis:** - **Memory Dependency Stalls**: warps waiting for memory load/store completion — indicates insufficient occupancy or poor memory access patterns (uncoalesced, cache misses) - **Execution Dependency Stalls**: warps waiting for previous instruction result — indicates long instruction latency chains (transcendental functions, integer division) without sufficient parallelism to hide latency - **Synchronization Stalls**: warps waiting at __syncthreads() or atomics — indicates load imbalance within a block or excessive atomic contention - **Instruction Fetch Stalls**: instruction cache misses, typically from very large kernels or low I-cache locality — rare but occurs with complex control flow and large instruction footprints **Memory Analysis:** - **L1/L2 Cache Hit Rate**: percentage of loads served from cache vs. DRAM — low hit rates suggest poor data locality or working set larger than cache capacity - **Sector Utilization**: percentage of bytes in each cache sector (32 bytes) actually used by the requesting warp — low utilization indicates poor coalescing or wasted bandwidth from partial cache line usage - **Shared Memory Efficiency**: transactions per request — 1.0 means no bank conflicts; higher values indicate N-way conflicts reducing shared memory bandwidth - **DRAM Read/Write Ratio**: excessive writes relative to reads may indicate unnecessary store operations or write-back traffic — read-heavy workloads are more common for inference-style kernels **GPU profiling with Nsight Compute is the indispensable diagnostic tool for GPU performance engineering — without quantitative profiling data, kernel optimization is guesswork; with it, engineers can systematically identify and eliminate bottlenecks to approach the theoretical performance ceiling defined by the roofline model.**

gpu profiling optimization,nsight profiler,occupancy analysis,kernel optimization gpu,performance bottleneck gpu

**GPU Performance Profiling and Optimization** is the **systematic analysis methodology that identifies and eliminates performance bottlenecks in GPU kernels — using hardware performance counters, execution traces, and roofline analysis to determine whether a kernel is limited by compute throughput, memory bandwidth, latency, or occupancy, then applying targeted optimizations that can improve kernel performance by 2-10x compared to a naive implementation**. **Profiling Tools** - **NVIDIA Nsight Compute**: Kernel-level profiler that collects hundreds of hardware metrics per kernel launch. Reports achieved throughput vs. peak (compute utilization, memory throughput), warp execution efficiency, register usage, shared memory usage, and detailed pipeline stall reasons. - **NVIDIA Nsight Systems**: System-level profiler showing the timeline of GPU kernel launches, memory transfers, CPU activity, and API calls. Identifies gaps where the GPU is idle (kernel launch latency, synchronization waits, host-device transfer bottlenecks). - **AMD ROCprofiler / Omniperf**: Equivalent profiling tools for AMD GPUs, providing similar hardware counter access and roofline analysis. **Key Performance Metrics** - **SM Occupancy**: Ratio of active warps to maximum warps per SM. Higher occupancy helps hide memory latency (more warps to switch to while waiting for memory). Limited by register usage, shared memory usage, and block size. - **Compute Throughput**: % of peak FLOPS achieved. Low compute throughput on a compute-bound kernel indicates instruction-level inefficiencies (poor ILP, warp divergence). - **Memory Throughput**: % of peak memory bandwidth achieved. Low memory throughput on a memory-bound kernel indicates uncoalesced accesses, bank conflicts, or insufficient in-flight memory requests. - **Warp Execution Efficiency**: % of active lanes per warp instruction. Below 100% indicates branch divergence — threads within a warp taking different paths. **Common Bottlenecks and Optimizations** - **Uncoalesced Memory Access**: Adjacent threads access non-adjacent global memory addresses, causing multiple memory transactions instead of one. Fix: restructure data layout (AoS → SoA — Array of Structures to Structure of Arrays). - **Shared Memory Bank Conflicts**: Multiple threads in a warp access the same shared memory bank simultaneously. Fix: pad shared memory arrays (add one element per row) to shift access patterns across banks. - **Low Occupancy**: Kernel uses too many registers (>64 per thread), limiting the number of concurrent warps. Fix: reduce register pressure by simplifying per-thread computation or using launch_bounds to hint the compiler. - **Kernel Launch Overhead**: Many small kernels create a stream of short launches with GPU idle gaps between them. Fix: fuse kernels, use CUDA graphs to batch launches, or increase per-kernel work. - **Branch Divergence**: Conditional branches cause warp serialization. Fix: restructure computation so all threads in a warp take the same path, or reorganize data so divergent work is across warps rather than within. **The Optimization Cycle** 1. Profile → identify the bottleneck (compute? memory? latency?). 2. Optimize the identified bottleneck. 3. Re-profile → verify improvement and identify the new bottleneck. 4. Repeat until reaching the roofline ceiling. GPU Profiling is **the empirical science of GPU performance** — because intuition about where bottlenecks lie is almost always wrong in the complex, highly parallel execution environment of a modern GPU, and only measurement-driven optimization reliably delivers the performance gains that justify the GPU's hardware investment.

GPU Profiling,Nsight,performance analysis,optimization

**GPU Profiling Nsight Performance Analysis** is **a comprehensive GPU performance profiling and analysis toolkit enabling detailed measurement and visualization of kernel execution, memory access patterns, and hardware utilization — identifying performance bottlenecks and guiding optimization efforts**. NVIDIA Nsight Tools provide GPU profiling across multiple levels of abstraction, from high-level timeline visualization showing kernel execution and memory transfers, to low-level instruction-level profiling showing execution of individual GPU instructions. The kernel timeline profiling shows when each kernel executes, how long kernels run, dependencies between kernels, and overlapping execution of multiple concurrent kernels, enabling identification of under-utilized GPU and opportunities for parallelism improvement. The warp efficiency metrics show what fraction of warps are actively computing versus idle waiting for memory, cache misses, or synchronization, with low warp efficiency indicating potential optimization opportunities. The memory bandwidth profiling shows actual achieved memory bandwidth compared to theoretical maximum and identifies whether kernels are memory-bound or compute-bound, guiding optimization focus to the limiting resource. The cache statistics show cache hit rates and cache miss distributions across different cache levels, identifying potential benefits from memory hierarchy optimization like increased data reuse. The hardware counter profiling measures diverse GPU performance metrics (instructions executed, cache misses, stalls) enabling identification of specific performance bottlenecks and validation of optimization hypotheses. The source-level profiling correlates performance metrics back to specific lines of code, enabling direct correlation of performance measurements to source code enabling straightforward identification of bottlenecks. **GPU profiling with Nsight tools enables comprehensive performance analysis and identification of optimization opportunities through detailed measurement and visualization.**

gpu programming model,cuda thread block,warp execution,thread hierarchy gpu,cooperative groups

**GPU Programming Model and Thread Hierarchy** is the **software abstraction that organizes millions of GPU threads into a hierarchical structure — grids of thread blocks (each containing hundreds of threads organized into warps of 32) — where the programmer expresses parallelism at the thread block level while the hardware scheduler dynamically maps blocks to Streaming Multiprocessors (SMs), enabling a single program to scale from a 10-SM laptop GPU to a 132-SM data center accelerator without code changes**. **Thread Hierarchy** ``` Grid (Kernel Launch) ├── Block (0,0) ← Thread Block: 32-1024 threads, scheduled on one SM │ ├── Warp 0 (threads 0-31) ← 32 threads executing in SIMT lockstep │ ├── Warp 1 (threads 32-63) │ └── ... ├── Block (0,1) ├── Block (1,0) └── ... (up to 2^31 blocks) ``` - **Thread**: The finest granularity of execution. Each thread has its own registers and program counter (logically — physically, warps share a PC). - **Warp (32 threads)**: The hardware scheduling unit. All 32 threads execute the same instruction simultaneously (SIMT). Divergent branches cause warp serialization. - **Thread Block (32-1024 threads)**: The programmer-defined grouping. All threads in a block execute on the same SM, share shared memory (up to 228 KB on H100), and can synchronize with __syncthreads(). - **Grid**: All thread blocks in a kernel launch. Blocks execute independently in any order — the GPU hardware schedules them dynamically. **Why This Hierarchy Works** - **Scalability**: The programmer specifies blocks, not SM assignments. A grid of 1000 blocks runs on a 10-SM GPU with 100 blocks per SM (time-sliced) or a 100-SM GPU with 10 blocks per SM (all concurrent). The same kernel binary scales automatically. - **Synchronization Scope**: Threads within a block can synchronize (barrier) and communicate (shared memory). Threads in different blocks cannot synchronize (no global barrier within a kernel) — this independence is what enables the scheduler's flexibility. **Cooperative Groups (CUDA 9+)** Extends the programming model beyond the block level: - **Thread Block Tile**: Partition a block into fixed-size tiles (e.g., 32 threads = warp) with tile-level sync and collective operations. - **Grid Group**: All blocks in a kernel can synchronize using cooperative launch (grid-wide barrier). Requires all blocks to be resident simultaneously — limits the number of blocks. - **Multi-Grid Group**: Synchronization across multiple kernel launches. **Occupancy and Scheduling** The SM scheduler assigns as many blocks to each SM as resources allow (registers, shared memory, max threads per SM). For example, if each block uses 64 registers per thread × 256 threads = 16,384 registers per block, and the SM has 65,536 registers, then 4 blocks can be resident simultaneously. Higher occupancy (more warps in-flight) helps hide memory latency. **Thread Indexing** ``` int gid = blockIdx.x * blockDim.x + threadIdx.x; // Global thread ID int lid = threadIdx.x; // Local (block) ID ``` The global ID maps each thread to a unique data element. The local ID selects shared memory locations. Multi-dimensional indexing (3D grids and blocks) naturally maps to 2D/3D data structures. The GPU Programming Model is **the abstraction that makes massively parallel hardware programmable** — hiding the complexity of warp scheduling, SM assignment, and hardware resource management behind a clean hierarchical model that lets programmers focus on the parallel algorithm rather than the machine architecture.

gpu ray tracing compute,cuda ray tracing optimization,optix ray tracing,gpu graphics compute,rtx ray tracing performance

**GPU Ray Tracing and Graphics Compute** is **the parallel implementation of ray tracing algorithms and graphics computations on GPUs** — where dedicated RT Cores on NVIDIA RTX GPUs accelerate ray-triangle intersection (10-100× faster than shader cores) and bounding volume hierarchy (BVH) traversal achieving 10-100 billion rays/second, while OptiX framework provides optimized ray tracing pipeline with denoising (10-50× noise reduction), path tracing (physically accurate lighting), and hybrid rasterization-ray tracing that enables real-time ray tracing at 30-60 FPS for 1080p-4K resolutions, making GPU ray tracing essential for photorealistic rendering, real-time graphics, and scientific visualization where ray tracing provides accurate shadows, reflections, and global illumination that are impossible with rasterization and proper optimization through BVH construction, ray coherence, denoising, and hybrid rendering determines whether applications achieve 1 FPS or 60 FPS at production quality. **RT Cores:** - **Hardware Acceleration**: dedicated ray-triangle intersection and BVH traversal units; 10-100× faster than shader cores - **Performance**: 10-100 billion rays/second on RTX 4090; 1-10 billion on RTX 3090; 0.1-1 billion on RTX 2080 - **Throughput**: 1-10 rays per clock per SM; 80-132 SMs on modern GPUs; massive parallelism - **Use Cases**: primary rays, shadow rays, reflection rays, global illumination; any ray-geometry intersection **BVH (Bounding Volume Hierarchy):** - **Structure**: tree of bounding boxes; hierarchical spatial partitioning; log(N) traversal for N triangles - **Construction**: CPU or GPU; GPU construction 5-20× faster; 100-500 million triangles/second - **Quality**: SAH (Surface Area Heuristic) for optimal quality; 2-5× fewer traversal steps; slower construction - **Update**: dynamic BVH for animated scenes; refit (fast) or rebuild (slow); 10-100 million triangles/second **OptiX Framework:** - **Pipeline**: ray generation, intersection, any-hit, closest-hit, miss shaders; programmable pipeline - **Denoising**: AI-based denoiser; 10-50× noise reduction; enables 1 sample per pixel; real-time quality - **Performance**: 10-100 billion rays/second; 80-95% of hardware capability; highly optimized - **Integration**: works with CUDA, OpenGL, Vulkan; flexible deployment; production-ready **Ray Types:** - **Primary Rays**: camera to scene; 1 per pixel; 1-10 billion rays/second; fully coherent; optimal performance - **Shadow Rays**: surface to light; 1-10 per pixel; 1-10 billion rays/second; partially coherent; good performance - **Reflection Rays**: surface to reflected direction; 0-5 per pixel; 0.5-5 billion rays/second; less coherent; moderate performance - **Diffuse Rays**: surface to random direction; 1-100 per pixel; 0.1-10 billion rays/second; incoherent; challenging performance **Ray Coherence:** - **Coherent Rays**: similar origin and direction; optimal BVH traversal; 10-100 billion rays/second - **Incoherent Rays**: random origins and directions; poor cache locality; 0.1-10 billion rays/second; 10-100× slower - **Optimization**: sort rays by direction; batch similar rays; 2-10× speedup for incoherent rays - **Primary Rays**: fully coherent; optimal performance; shadow rays partially coherent; diffuse rays incoherent **Path Tracing:** - **Algorithm**: trace rays recursively; accumulate lighting; Monte Carlo integration; physically accurate - **Performance**: 0.1-10 billion rays/second; depends on scene complexity and ray depth; 1-10 samples per pixel for real-time - **Convergence**: 100-10000 samples for noise-free; denoising enables 1-10 samples; 10-100× speedup with denoising - **Use Cases**: photorealistic rendering, global illumination, caustics; film, architecture, product visualization **Denoising:** - **AI Denoiser**: OptiX AI denoiser; 10-50× noise reduction; enables 1 sample per pixel; real-time quality - **Performance**: 1-10ms per frame at 1080p; 5-20ms at 4K; 5-10% of frame time; acceptable overhead - **Quality**: near-converged quality with 1-10 samples; 100-1000× faster than brute force; production quality - **Integration**: post-process after ray tracing; works with any renderer; easy to integrate **Hybrid Rendering:** - **Rasterization + Ray Tracing**: rasterize primary visibility; ray trace shadows, reflections, GI; 30-60 FPS at 1080p-4K - **Performance**: 10-100× faster than pure ray tracing; 2-10× better quality than pure rasterization; best of both worlds - **Use Cases**: real-time games, interactive visualization; balance quality and performance - **Techniques**: screen-space reflections + ray traced reflections; shadow maps + ray traced shadows; hybrid GI **BVH Optimization:** - **SAH Construction**: Surface Area Heuristic; optimal quality; 2-5× fewer traversal steps; 2-10× slower construction - **Fast Construction**: spatial median split; 5-20× faster construction; 2-5× more traversal steps; good for dynamic scenes - **Refitting**: update BVH for animated geometry; 10-100× faster than rebuild; acceptable quality for small changes - **Compaction**: remove empty nodes; 10-30% memory reduction; 10-20% traversal speedup **Memory Optimization:** - **BVH Size**: 10-100 bytes per triangle; 1-10GB for 10-100 million triangles; significant memory usage - **Compression**: compress BVH nodes; 2-4× reduction; 10-20% traversal slowdown; acceptable trade-off - **Streaming**: stream geometry and BVH; enables scenes larger than GPU memory; 2-10× slower; necessary for huge scenes - **LOD**: level of detail for distant objects; reduces triangle count; 2-10× speedup; acceptable quality loss **Shading Optimization:** - **Material Complexity**: simple materials faster; complex materials (subsurface scattering, volumetrics) 10-100× slower - **Texture Sampling**: texture cache important; coherent access 10-100× faster than random; sort rays by material - **Shader Divergence**: minimize divergence; similar materials together; 2-10× speedup - **Inline Ray Tracing**: inline ray tracing in compute shaders; more control; 10-30% faster for some patterns **Real-Time Ray Tracing:** - **Target**: 30-60 FPS at 1080p-4K; 16-33ms per frame; challenging for complex scenes - **Techniques**: 1 sample per pixel + denoising; hybrid rendering; adaptive sampling; temporal accumulation - **Performance**: 10-100 billion rays/second; 1-10 rays per pixel; 1-10 billion pixels/second - **Quality**: near-photorealistic with denoising; acceptable for real-time; production quality **Offline Rendering:** - **Target**: photorealistic quality; 100-10000 samples per pixel; minutes to hours per frame - **Performance**: 0.1-10 billion rays/second; depends on scene complexity; 10-1000 rays per pixel - **Quality**: converged, noise-free; film quality; no compromises - **Use Cases**: film, architecture, product visualization; quality over speed **Multi-GPU Ray Tracing:** - **Data Parallelism**: each GPU renders subset of pixels; 70-85% scaling efficiency; simple implementation - **Scene Parallelism**: distribute scene across GPUs; 50-70% efficiency; complex implementation; necessary for huge scenes - **NVLink**: 900 GB/s between GPUs; enables efficient scene sharing; 70-85% efficiency - **Use Cases**: very large scenes, offline rendering; real-time multi-GPU challenging **Performance Profiling:** - **Nsight Graphics**: profiles ray tracing pipeline; shows ray counts, BVH traversal, shading time - **Metrics**: rays/second, traversal steps, shader time; target 10-100 billion rays/second - **Bottlenecks**: incoherent rays, complex shaders, poor BVH quality; optimize based on profiling - **Tuning**: adjust BVH quality, ray coherence, shader complexity; profile to find optimal **Best Practices:** - **Use RT Cores**: always use hardware ray tracing; 10-100× faster than software - **Optimize BVH**: use SAH for static scenes; fast construction for dynamic; 2-5× speedup - **Ray Coherence**: sort rays, batch similar rays; 2-10× speedup for incoherent rays - **Denoising**: use AI denoiser; 10-50× noise reduction; enables real-time quality - **Hybrid Rendering**: combine rasterization and ray tracing; 10-100× faster than pure ray tracing **Performance Targets:** - **Primary Rays**: 10-100 billion rays/second; fully coherent; optimal performance - **Shadow Rays**: 1-10 billion rays/second; partially coherent; good performance - **Path Tracing**: 0.1-10 billion rays/second; depends on depth and complexity - **Real-Time**: 30-60 FPS at 1080p-4K; 1-10 rays per pixel; with denoising **Real-World Applications:** - **Games**: real-time ray traced shadows, reflections, GI; 30-60 FPS at 1080p-4K; RTX 3080/4080/4090 - **Film**: offline path tracing; photorealistic quality; minutes to hours per frame; render farms - **Architecture**: interactive visualization; real-time or near-real-time; 10-30 FPS at 4K - **Scientific Visualization**: accurate lighting for data visualization; 10-60 FPS; depends on complexity GPU Ray Tracing and Graphics Compute represent **the revolution in photorealistic rendering** — by leveraging dedicated RT Cores that accelerate ray-triangle intersection (10-100× faster than shader cores) and BVH traversal achieving 10-100 billion rays/second, combined with OptiX framework providing AI denoising (10-50× noise reduction) and hybrid rendering techniques, developers enable real-time ray tracing at 30-60 FPS for 1080p-4K resolutions and photorealistic offline rendering where proper optimization through BVH construction, ray coherence, denoising, and hybrid rendering determines whether applications achieve 1 FPS or 60 FPS at production quality.');

gpu ray tracing hardware,rt core bvh traversal,hardware ray triangle intersection,real time ray tracing,bvh acceleration structure

**GPU Hardware Ray Tracing** is the **dedicated fixed-function hardware (NVIDIA RT Cores, AMD Ray Accelerators, Intel Ray Tracing Units) that accelerates the computationally intensive ray-scene intersection tests required for photorealistic rendering — traversing bounding volume hierarchies (BVH) and computing ray-triangle intersections at hundreds of billions of tests per second, enabling real-time ray tracing for reflections, shadows, ambient occlusion, and global illumination in games, film production, and scientific visualization**. **Why Hardware Acceleration** Ray tracing requires testing each ray against potentially millions of triangles. Software BVH traversal on shader cores achieves ~1-5 billion ray-box tests/second. RT Cores achieve 50-300 billion tests/second — a 10-100× speedup. This hardware acceleration transforms ray tracing from offline rendering (hours per frame) to real-time (16-33 ms per frame at 30-60 FPS). **BVH (Bounding Volume Hierarchy)** The acceleration structure that makes ray tracing tractable: - **Construction**: Scene triangles are recursively partitioned into groups, each enclosed by an axis-aligned bounding box (AABB). The root AABB contains the entire scene; leaf nodes contain 1-8 triangles. - **Traversal**: A ray tests against the root AABB. If it intersects, test both children. Recursively descend into intersected nodes, skip non-intersected subtrees. Average complexity: O(log N) triangle tests instead of O(N). - **Quality vs. Build Time**: SAH (Surface Area Heuristic) BVH construction produces optimal traversal trees but is expensive to build. LBVH (Linear BVH) uses Morton codes for fast construction (suitable for dynamic scenes). TLAS/BLAS split: Bottom-Level AS (BLAS) per object (rebuilt rarely), Top-Level AS (TLAS) for scene arrangement (rebuilt every frame). **RT Core Architecture (NVIDIA)** - **BVH Traversal Unit**: Dedicated hardware that traverses the BVH tree, testing ray-AABB intersections at each node. One traversal step per clock — 2-3 AABB tests per cycle. Operates concurrently with shader execution on SM cores. - **Ray-Triangle Intersection Unit**: Computes Möller-Trumbore ray-triangle intersection for leaf nodes. Reports hit distance, barycentric coordinates, and triangle ID. - **Opacity Micro-Map (Hopper+)**: Hardware-accelerated alpha-test evaluation. Encodes per-micro-triangle opacity, allowing RT Cores to skip fully transparent triangles and classify semi-transparent regions — 2× speedup for foliage and particle effects. **Ray Tracing Pipeline (DXR/Vulkan RT/OptiX)** 1. **Ray Generation Shader**: Launches rays (one per pixel for primary rays, additional for reflections/shadows). 2. **BVH Traversal** (hardware): RT Core traverses TLAS → BLAS hierarchy. 3. **Intersection Shader** (optional): Custom intersection test for non-triangle primitives (spheres, curves, SDF). 4. **Any-Hit Shader**: Called for each potential hit — used for alpha-test transparency. Can accept or reject the hit. 5. **Closest-Hit Shader**: Called for the nearest intersection. Computes shading (material, lighting, launches secondary rays). 6. **Miss Shader**: Called when no intersection found — returns environment/sky color. **Performance Metrics** NVIDIA RTX 4090: 191 billion RT Core TFLOPS equivalent, ~30-60 FPS in fully ray-traced scenes at 4K with DLSS. AMD RDNA 3 (RX 7900 XTX): significant improvement over RDNA 2 but still trails NVIDIA in pure RT throughput. Intel Arc provides competitive RT performance in its class. GPU Hardware Ray Tracing is **the fixed-function acceleration that transformed photorealistic rendering from an offline computation to a real-time capability** — dedicated silicon that makes the physically-based lighting, reflections, and shadows of ray tracing achievable within the millisecond-per-frame budgets of interactive applications.

gpu reduction patterns,parallel reduction cuda,warp reduction optimization,cuda reduce performance,hierarchical reduction gpu

**GPU Reduction Patterns** are **the parallel algorithms for combining array elements into single value through associative operations** — where hierarchical reduction using warp primitives (__shfl_down_sync) for intra-warp (500-1000 GB/s), shared memory for inter-warp (300-600 GB/s), and atomic operations for inter-block (200-400 GB/s) achieves 60-80% of peak memory bandwidth and 2-10× speedup over naive implementations, making reduction optimization critical for applications like sum, max, min, dot product that appear in 40-80% of GPU kernels and proper implementation using warp-level primitives instead of shared memory, minimizing synchronization, and hierarchical patterns determines whether reductions achieve 100 GB/s or 1000 GB/s throughput. **Reduction Fundamentals:** - **Associative Operations**: sum, max, min, product, AND, OR, XOR; order doesn't affect result; enables parallelization - **Commutative**: most reductions also commutative; further optimization opportunities; non-commutative requires careful ordering - **Tree Pattern**: log2(N) steps to reduce N elements; each step halves active threads; optimal work complexity O(N) - **Memory Bound**: reductions are memory-bound; limited by bandwidth (1.5-3 TB/s); not compute; optimize memory access **Warp-Level Reduction:** - **Shuffle Down**: use __shfl_down_sync() in loop; 5 iterations for 32 threads; no shared memory; 2-10× faster than shared memory - **Code Pattern**: for (int offset = 16; offset > 0; offset /= 2) { val += __shfl_down_sync(0xffffffff, val, offset); } - **Performance**: 500-1000 GB/s; 60-80% of peak bandwidth; 2-5× faster than shared memory; no synchronization overhead - **Use Cases**: reduce within warp; building block for block and grid reductions; critical optimization **Block-Level Reduction:** - **Two-Stage**: warp reduction → shared memory → warp reduction; optimal at each level; 300-600 GB/s - **Pattern**: each warp reduces to single value; write to shared memory; first warp reduces shared memory; single thread has result - **Shared Memory**: 32-64 elements in shared memory (one per warp); minimal usage; enables high occupancy - **Performance**: 300-600 GB/s; 40-60% of peak bandwidth; 2-5× faster than pure shared memory **Grid-Level Reduction:** - **Three-Stage**: warp reduction → block reduction → atomic or multi-kernel; 200-400 GB/s - **Atomic Approach**: each block atomics final result; simple but contention; 200-400 GB/s for low contention - **Multi-Kernel**: first kernel reduces to per-block results; second kernel reduces blocks; no contention; 300-600 GB/s - **Cooperative Groups**: grid.sync() enables single-kernel multi-block; 200-400 GB/s; simpler than multi-kernel **Optimization Techniques:** - **Warp Primitives**: always use __shfl for intra-warp; 2-10× faster than shared memory; 500-1000 GB/s - **Minimize Sync**: reduce synchronization points; use warp primitives (no sync needed); 20-40% improvement - **Coalesced Access**: ensure coalesced memory reads; 128-byte aligned; achieves 100% bandwidth - **Multiple Elements Per Thread**: each thread processes multiple elements; reduces overhead; 20-50% improvement **Unrolling and Specialization:** - **Loop Unrolling**: unroll reduction loops; reduces overhead; 10-20% speedup; #pragma unroll - **Template Specialization**: specialize for block size; enables compile-time optimization; 10-30% speedup - **Warp Unrolling**: fully unroll warp reduction (5 iterations); eliminates loop overhead; 10-20% speedup - **Compile-Time Constants**: use template parameters for sizes; enables aggressive optimization; 20-40% improvement **Multiple Accumulators:** - **Pattern**: use multiple accumulators to increase ILP; reduces dependency chains; 30-60% speedup - **Code**: float sum1 = 0, sum2 = 0; for (int i = tid; i < N; i += stride) { sum1 += data[i]; sum2 += data[i+offset]; } - **Merge**: combine accumulators at end; single reduction; amortizes overhead - **Performance**: 30-60% improvement; exposes instruction-level parallelism; hides latency **Reduction with Transformation:** - **Pattern**: transform elements during reduction; fused operation; eliminates temporary storage - **Example**: sum of squares: reduce(x * x); dot product: reduce(a * b); L2 norm: sqrt(reduce(x * x)) - **Performance**: 2-5× faster than separate transform and reduce; eliminates memory traffic; 500-1000 GB/s - **Implementation**: thrust::transform_reduce() or custom kernel with inline transformation **Segmented Reduction:** - **Concept**: reduce multiple independent segments; each segment reduced separately; useful for batched operations - **Implementation**: CUB DeviceSegmentedReduce; segment offsets specify boundaries; 200-400 GB/s - **Use Cases**: per-row sum in matrix, per-group aggregation; graph algorithms; database operations - **Performance**: 200-400 GB/s; 40-60% of peak; depends on segment sizes; small segments have overhead **Thrust Reduce:** - **API**: float sum = thrust::reduce(d_vec.begin(), d_vec.end(), 0.0f, thrust::plus()); - **Performance**: 500-1000 GB/s; 70-90% of hand-tuned; 1 line vs 50-100 for custom implementation - **Customization**: custom operators supported; thrust::maximum(), custom functors - **Use Cases**: rapid development, general-purpose reduction; acceptable 10-30% performance gap **CUB Reduce:** - **API**: cub::DeviceReduce::Sum(d_temp_storage, temp_storage_bytes, d_in, d_out, N); - **Performance**: 500-1000 GB/s; 80-95% of hand-tuned; 10-20% faster than Thrust; lower-level API - **Features**: sum, min, max, custom operators; segmented reduce; flexible - **Use Cases**: performance-critical reductions; fine-grained control; production systems **Atomic Reduction:** - **Pattern**: each thread/warp/block atomics to global result; simple but contention-prone - **Warp Aggregation**: reduce within warp first; lane 0 atomics; 32× fewer atomics; 5-20× speedup - **Performance**: 200-400 GB/s with warp aggregation; 10-50 GB/s without; contention limits performance - **Use Cases**: simple reductions, low contention; when multi-kernel overhead unacceptable **Hierarchical Patterns:** - **Three-Level**: warp → block → grid; optimal at each level; 300-600 GB/s - **Warp Level**: __shfl_down_sync(); 500-1000 GB/s; no shared memory; no sync - **Block Level**: shared memory for inter-warp; 300-600 GB/s; minimal shared memory usage - **Grid Level**: atomic or multi-kernel; 200-400 GB/s; depends on contention **Performance Profiling:** - **Nsight Compute**: shows memory bandwidth, warp efficiency, occupancy; identifies bottlenecks - **Metrics**: achieved bandwidth / peak bandwidth; target 60-80% for reductions; memory-bound - **Bottlenecks**: uncoalesced access, excessive synchronization, low occupancy; optimize patterns - **Tuning**: adjust block size, elements per thread, unrolling; profile to find optimal **Common Pitfalls:** - **Shared Memory Only**: not using warp primitives; 2-10× slower; always use __shfl for intra-warp - **Excessive Sync**: too many __syncthreads(); 10-30% overhead each; minimize synchronization - **Uncoalesced Access**: stride access patterns; 10-100× slowdown; ensure coalesced reads - **Single Accumulator**: not using multiple accumulators; limits ILP; 30-60% slower **Best Practices:** - **Warp Primitives**: always use __shfl for intra-warp reduction; 2-10× faster than shared memory - **Hierarchical**: warp → block → grid; optimal at each level; 300-600 GB/s - **Multiple Accumulators**: use 2-4 accumulators; increases ILP; 30-60% improvement - **Coalesced Access**: ensure coalesced memory reads; 128-byte aligned; achieves 100% bandwidth - **Profile**: measure actual bandwidth; compare with peak; optimize only if bottleneck **Performance Targets:** - **Warp Reduction**: 500-1000 GB/s; 60-80% of peak; 2-5× faster than shared memory - **Block Reduction**: 300-600 GB/s; 40-60% of peak; optimal for 256-512 threads - **Grid Reduction**: 200-400 GB/s; 30-50% of peak; limited by atomic contention or multi-kernel overhead - **Overall**: 500-1000 GB/s for large arrays (>1M elements); 60-80% of peak bandwidth **Real-World Applications:** - **Sum**: array sum, vector norm; 500-1000 GB/s; 60-80% of peak; critical building block - **Dot Product**: vector dot product; 500-1000 GB/s; fused multiply-add and reduce; 60-80% of peak - **Max/Min**: find maximum/minimum; 500-1000 GB/s; 60-80% of peak; used in normalization - **Statistics**: mean, variance, standard deviation; 400-800 GB/s; multiple reductions; 50-70% of peak GPU Reduction Patterns represent **the fundamental parallel primitive** — by using hierarchical reduction with warp primitives for intra-warp (500-1000 GB/s), shared memory for inter-warp (300-600 GB/s), and atomic operations for inter-block (200-400 GB/s), developers achieve 60-80% of peak memory bandwidth and 2-10× speedup over naive implementations, making reduction optimization critical for GPU applications where reductions appear in 40-80% of kernels and proper implementation using warp-level primitives instead of shared memory determines whether reductions achieve 100 GB/s or 1000 GB/s throughput.

gpu register file optimization,register spilling local memory,occupancy register tradeoff,ptx register allocation,kernel register count

**GPU Register Optimization** addresses the **critical trade-off between register availability for instruction-level parallelism and kernel occupancy, directly impacting throughput and latency hiding in GPU applications.** **Register File Architecture and Limits** - **Register File Size**: Per-SM registers (Ampere: 256 KB, Volta: 128 KB). Shared across all resident warps. Maximum per thread varies by GPU model (255 registers Ampere, 255 registers Volta). - **Register Banking**: 32 banks (one per thread in warp). Concurrent register access for all 32 threads requires bank conflict-free address patterns. Same-bank concurrent accesses serialize. - **Register Allocation**: Compiler allocates registers to variables. Scalar variables (float) 1 register; double 2 registers; arrays need consecutive registers. - **Allocation Pressure**: More live variables → more registers. Compiler optimizes to minimize register count (without harming ILP). **Register Spilling to Local Memory** - **Spilling Mechanism**: When register count exceeds budget (--maxrregcount), excess data spilled to local memory (on-chip cache hierarchy: L1 → L2 → HBM). - **Performance Impact**: Local memory ~100-500 cycles latency (vs ~10 cycles for register). Spilled values effectively become cache misses. Throughput drops 10-100x. - **Spill Detection**: Profiler reports spill rate (registers spilled per thread). Nonzero spill rate indicates register pressure. Target: 0 spills for performance-critical kernels. - **Reduce Spilling**: Decrease --maxrregcount (fewer blocks resident, less parallelism), rewrite code to reduce register pressure (reorganize loops, fuse operations). **Occupancy-Register Tradeoff** - **Occupancy Definition**: Percentage of SM warp slots filled. More registers per thread → fewer resident warps → lower occupancy. - **Occupancy Curve**: Register count vs occupancy (inverse relationship). Register count = N → occupancy = (SM_reg_size / (N × threads_per_warp × warps_per_block)). - **Latency Hiding**: High occupancy (many resident warps) hides memory latency. Low occupancy (few register, few warps) relies on few warps to hide latency. - **Optimal Point**: Often exists between extremes. Too low register (occupancy 100%) = stalls on memory. Too high register (occupancy 25%) = stalls on latency. **PTX ISA Register Model** - **PTX Register Classes**: %r (32-bit register), %rd (64-bit), %p (predicate), %f (float), %d (double). Abstract model (not tied to specific GPU architecture). - **Virtual Registers**: PTX compilation assigns unlimited virtual registers; target-specific compiler (NVCC, PTXAS) maps to physical registers. - **Physical Constraints**: Target GPU (SM 7.0, 8.0, 9.0) determines physical register count per warp, occupancy implications. Same PTX code → different occupancy on different GPUs. - **ISA Compatibility**: PTX forward/backward compatible within reason. Code compiled to PTX can target multiple GPU architectures (with occupancy variation). **Compiler Register Allocation Strategies** - **Register Pressure Analysis**: Compiler builds interference graph (variables live simultaneously). Graph coloring assigns registers; chromatic number = min registers needed. - **Spilling Decision**: When variables exceed registers, spill to local memory. Decisions impact performance; algorithm heuristic-based (not optimal). - **Loop Unrolling Effect**: Unrolling increases register count (multiple loop iterations's variables live simultaneously). Trade-off: faster loop (fewer branches) vs higher register pressure. - **Optimization Passes**: Multiple passes refine allocation. LICM (loop-invariant code motion), CSE (common subexpression elimination), dead code elimination reduce register pressure. **Kernel Register Count Reduction Techniques** - **Refactor Loops**: Break long loops into smaller loops (reduce simultaneous live variables). Example: Process array in 256-element chunks instead of full array. - **Array Privatization**: Private arrays (private to thread) expensive (registers). Replace with scalars, iterate instead of bulk allocation. - **Use Functions**: Inline functions increase register pressure; non-inlined functions transfer data via memory (cheaper than spilling). Trade-off: function call overhead vs register savings. - **Reduce Precision**: float (1 register) vs double (2 registers). Use float where possible; promote to double only when necessary. **Warp-Level Register Sharing and Limits** - **Warp Register Pool**: All threads in warp share 32 registers (Ampere) conceptually. Thread i gets registers r_{i*N}, r_{i*N+1}, ..., r_{i*N+N-1} (N = registers per thread). - **Cross-Warp Sharing**: Register file shared among multiple warps (SM occupancy). Warp 0 occupies registers 0-4095, Warp 1 occupies 4096-8191, etc. - **Bank Conflict Minimization**: Register accesses within warp sequential (thread i accesses bank i). Careful allocation avoids conflicts. **Profiling and Optimization Workflow** - **Nsight Metrics**: "Register per Thread" metric shows allocation. "Registers per Inst Executed" indicates spilling (>4 typical, >8 severe). - **Occupancy Analysis**: Nsight reports occupancy-limiting factor (registers, shared memory, threads-per-block). - **Optimization Priority**: Eliminate spilling first (highest impact). Then reduce registers if occupancy < 50% (may improve performance).

GPU Register File,optimization,allocation,spilling

**GPU Register File Optimization** is **a critical low-level GPU optimization technique managing allocation and utilization of per-thread register storage — preventing register spilling to shared/global memory that would cause dramatic performance degradation from cache misses and increased instruction latency**. Modern GPU register files are shared among all active threads on a streaming multiprocessor (SM), with total register file capacity (typically 256KB per SM for current NVIDIA GPUs) divided among active threads, determining maximum number of simultaneous threads. The register allocation is performed automatically by GPU compiler based on kernel requirements, with compiler attempting to minimize registers while providing sufficient storage for all kernel variables. The register spilling occurs when kernel requires more registers than available per thread, causing compiler to spill excess values to local memory (typically in global memory with cache), causing dramatic performance degradation from memory latency and pressure on memory hierarchy. The register pressure reduction techniques including instruction scheduling optimization, register reuse through careful variable management, and algorithmic changes to reduce intermediate values can minimize register requirements and prevent spilling. The awareness of register usage limitations during algorithm design enables selection of algorithms with lower register requirements even if they require slightly more total operations, often resulting in better overall performance. The compiler flags controlling register usage (e.g., maxrregcount) enable explicit limitation of register usage to force lower occupancy with better per-warp performance if that proves beneficial for specific kernels. The measurement of actual register spilling through profiling tools enables identification of problematic kernels and validation that optimization efforts successfully eliminate spilling. **GPU register file optimization through careful algorithm design and compiler-directed register pressure management prevents memory spilling and maintains performance.**

gpu register pressure,register allocation gpu,register spill,occupancy register,gpu register file

**GPU Register Pressure** is the **conflict between a kernel's per-thread register demand and the GPU's fixed register file capacity** — where each additional register per thread reduces the number of concurrent threads (occupancy), potentially hiding less memory latency, while reducing registers may cause spills to slow local memory, creating a critical optimization tradeoff for GPU kernel performance. **GPU Register File Architecture** - Each NVIDIA SM (Streaming Multiprocessor) has a **fixed register file**: 65,536 32-bit registers (typical). - Registers are **partitioned** among all active threads on the SM. - More registers per thread = fewer threads per SM = lower occupancy. **Occupancy Example (NVIDIA A100)** | Registers/Thread | Max Threads/SM | Occupancy (of 2048 max) | |-----------------|---------------|------------------------| | 32 | 2048 | 100% | | 64 | 1024 | 50% | | 128 | 512 | 25% | | 255 (max) | 256 | 12.5% | - At 255 registers/thread: Only 256 threads active (8 warps) — very little latency hiding. - At 32 registers/thread: Full 2048 threads (64 warps) — maximum latency hiding potential. **Register Spilling** - When kernel needs more registers than allocated → compiler **spills** excess to local memory. - Local memory is actually device DRAM (L1/L2 cached) — 100x slower than register access. - Spilling causes significant performance degradation: 2-10x slowdown for spill-heavy kernels. **Optimization Strategies** - **Limit register count**: `__launch_bounds__(maxThreadsPerBlock, minBlocksPerMultiprocessor)` or `--maxrregcount=N` compiler flag. - **Reduce live variables**: Recompute values instead of storing them. Reorder operations to reduce simultaneous live registers. - **Use shared memory**: Move some per-thread data to shared memory (explicitly managed cache). - **Loop unrolling control**: Aggressive unrolling increases register usage — `#pragma unroll` factor tuning. **Profiling Register Usage** - `nvcc --ptxas-options=-v` reports register count per kernel. - NVIDIA Nsight Compute shows register usage, spills, and occupancy. - CUDA occupancy calculator: Interactive tool to find optimal register/thread configuration. **Register Pressure vs. ILP** - Some kernels benefit from low occupancy + high ILP (instruction-level parallelism per thread). - Heavy compute kernels (matrix math): Fewer threads with more registers can outperform many threads with spilling. - **Principle**: Occupancy is not the only metric — achieved throughput is what matters. GPU register pressure is **one of the most impactful performance-limiting factors in GPU programming** — understanding and managing the register-occupancy-spill tradeoff is essential for extracting peak performance from GPU hardware.

gpu scan prefix sum,parallel scan cuda,cuda prefix sum optimization,inclusive exclusive scan,scan algorithm gpu

**GPU Scan (Prefix Sum)** is **the parallel algorithm that computes cumulative sums or other associative operations across array elements** — where inclusive scan produces [a0, a0+a1, a0+a1+a2, ...] and exclusive scan produces [0, a0, a0+a1, ...], achieving 400-800 GB/s throughput (50-70% of peak bandwidth) through hierarchical implementation using warp primitives (__shfl_up_sync) for intra-warp (500-1000 GB/s), shared memory for inter-warp (300-600 GB/s), and multi-pass algorithms for large arrays, making scan essential for algorithms like stream compaction (removing elements), radix sort (computing output positions), and sparse matrix operations where scan appears in 30-60% of advanced GPU algorithms and proper implementation using warp-level primitives and minimizing global memory accesses determines whether applications achieve 100 GB/s or 800 GB/s throughput. **Scan Fundamentals:** - **Inclusive Scan**: output[i] = input[0] + input[1] + ... + input[i]; includes current element; natural for cumulative sums - **Exclusive Scan**: output[i] = input[0] + input[1] + ... + input[i-1]; excludes current element; useful for computing positions - **Associative Operations**: sum, max, min, product, AND, OR, XOR; order matters for scan; enables parallelization - **Applications**: stream compaction, radix sort, sparse matrix, work distribution; fundamental building block **Warp-Level Scan:** - **Shuffle Up**: use __shfl_up_sync() in loop; 5 iterations for 32 threads; no shared memory; 2-5× faster than shared memory - **Code Pattern**: for (int offset = 1; offset < 32; offset *= 2) { int temp = __shfl_up_sync(0xffffffff, val, offset); if (lane >= offset) val += temp; } - **Performance**: 500-1000 GB/s; 60-80% of peak bandwidth; 2-5× faster than shared memory; no synchronization overhead - **Use Cases**: scan within warp; building block for block and grid scans; critical optimization **Block-Level Scan:** - **Two-Stage**: warp scan → inter-warp scan → add offsets; optimal at each level; 300-600 GB/s - **Pattern**: each warp scans independently; scan warp sums; add warp offsets to elements; requires two syncs - **Shared Memory**: store warp sums (32-64 elements); minimal usage; enables high occupancy - **Performance**: 300-600 GB/s; 40-60% of peak bandwidth; 2-5× faster than pure shared memory **Large Array Scan:** - **Multi-Pass**: first pass scans blocks independently; second pass scans block sums; third pass adds offsets; 400-800 GB/s - **Three-Kernel**: kernel 1 scans blocks; kernel 2 scans block sums; kernel 3 adds offsets; no global sync needed - **Single-Kernel**: use cooperative groups grid.sync(); simpler but requires all blocks fit on GPU; 400-800 GB/s - **Performance**: 400-800 GB/s for large arrays (>1M elements); 50-70% of peak bandwidth; memory-bound **Optimization Techniques:** - **Warp Primitives**: always use __shfl_up for intra-warp; 2-5× faster than shared memory; 500-1000 GB/s - **Bank Conflict Avoidance**: pad shared memory arrays; prevents bank conflicts; 10-30% improvement - **Coalesced Access**: ensure coalesced memory reads/writes; 128-byte aligned; achieves 100% bandwidth - **Multiple Elements Per Thread**: each thread processes multiple elements; reduces overhead; 20-50% improvement **Inclusive vs Exclusive:** - **Inclusive**: simpler implementation; natural for cumulative sums; output[i] includes input[i] - **Exclusive**: useful for positions; output[i] = sum of elements before i; radix sort, compaction - **Conversion**: exclusive = shift(inclusive, 1) with 0 at start; inclusive = exclusive + input; trivial conversion - **Performance**: same performance; choice based on application needs; exclusive more common in algorithms **Segmented Scan:** - **Concept**: scan multiple independent segments; each segment scanned separately; useful for batched operations - **Implementation**: CUB DeviceSegmentedScan; segment flags or offsets specify boundaries; 300-600 GB/s - **Use Cases**: per-row scan in matrix, per-group aggregation; graph algorithms; sparse matrix operations - **Performance**: 300-600 GB/s; 40-60% of peak; depends on segment sizes; small segments have overhead **Thrust Scan:** - **API**: thrust::inclusive_scan(d_in.begin(), d_in.end(), d_out.begin()); or thrust::exclusive_scan() - **Performance**: 400-800 GB/s; 60-80% of hand-tuned; 1 line vs 100-200 for custom implementation - **Customization**: custom operators supported; thrust::plus(), thrust::maximum(), custom functors - **Use Cases**: rapid development, general-purpose scan; acceptable 20-40% performance gap **CUB Scan:** - **API**: cub::DeviceScan::InclusiveSum(d_temp_storage, temp_storage_bytes, d_in, d_out, N); - **Performance**: 400-800 GB/s; 70-90% of hand-tuned; 10-30% faster than Thrust; lower-level API - **Features**: inclusive, exclusive, custom operators; segmented scan; flexible - **Use Cases**: performance-critical scans; fine-grained control; production systems **Stream Compaction:** - **Pattern**: scan predicate flags; use scan result as output positions; write elements to compacted array - **Code**: flags = predicate(input); positions = exclusive_scan(flags); if (flags[i]) output[positions[i]] = input[i]; - **Performance**: 300-600 GB/s; 40-60% of peak; scan is bottleneck; 2-3× faster than CPU - **Use Cases**: removing elements, filtering; sparse matrix, graph algorithms; data preprocessing **Radix Sort Integration:** - **Histogram**: count elements per bin; 300-600 GB/s - **Scan**: exclusive scan of histogram; computes output positions; 400-800 GB/s - **Scatter**: write elements to sorted positions; 200-300 GB/s - **Performance**: scan is 20-40% of radix sort time; critical for overall performance **Work Distribution:** - **Pattern**: scan work counts; use scan result to assign work to threads; load balancing - **Code**: work_counts = compute_work(input); positions = exclusive_scan(work_counts); assign work based on positions - **Performance**: 300-600 GB/s for scan; enables balanced work distribution; 20-50% improvement in irregular algorithms - **Use Cases**: irregular parallelism, dynamic work assignment; graph algorithms; sparse operations **Hierarchical Patterns:** - **Three-Level**: warp → block → grid; optimal at each level; 400-800 GB/s - **Warp Level**: __shfl_up_sync(); 500-1000 GB/s; no shared memory; no sync - **Block Level**: shared memory for inter-warp; 300-600 GB/s; minimal shared memory usage - **Grid Level**: multi-pass or cooperative groups; 400-800 GB/s; depends on array size **Performance Profiling:** - **Nsight Compute**: shows memory bandwidth, warp efficiency, occupancy; identifies bottlenecks - **Metrics**: achieved bandwidth / peak bandwidth; target 50-70% for scans; memory-bound - **Bottlenecks**: uncoalesced access, bank conflicts, excessive synchronization; optimize patterns - **Tuning**: adjust block size, elements per thread, padding; profile to find optimal **Common Pitfalls:** - **Shared Memory Only**: not using warp primitives; 2-5× slower; always use __shfl_up for intra-warp - **Bank Conflicts**: unpadded shared memory; 2-10× slowdown; add padding to avoid conflicts - **Uncoalesced Access**: stride access patterns; 10-100× slowdown; ensure coalesced reads/writes - **Too Many Syncs**: excessive __syncthreads(); 10-30% overhead each; minimize synchronization **Best Practices:** - **Warp Primitives**: always use __shfl_up for intra-warp scan; 2-5× faster than shared memory - **Hierarchical**: warp → block → grid; optimal at each level; 400-800 GB/s - **Pad Shared Memory**: avoid bank conflicts; 1-2 elements padding; 10-30% improvement - **Coalesced Access**: ensure coalesced memory reads/writes; 128-byte aligned; achieves 100% bandwidth - **Profile**: measure actual bandwidth; compare with peak; optimize only if bottleneck **Performance Targets:** - **Warp Scan**: 500-1000 GB/s; 60-80% of peak; 2-5× faster than shared memory - **Block Scan**: 300-600 GB/s; 40-60% of peak; optimal for 256-512 threads - **Grid Scan**: 400-800 GB/s; 50-70% of peak; for large arrays (>1M elements) - **Overall**: 400-800 GB/s for large arrays; 50-70% of peak bandwidth; memory-bound **Real-World Applications:** - **Stream Compaction**: removing invalid elements; 300-600 GB/s; 40-60% of peak; used in rendering, physics - **Radix Sort**: computing output positions; 400-800 GB/s; 20-40% of sort time; critical for performance - **Sparse Matrix**: CSR format construction; 300-600 GB/s; 30-50% of construction time - **Work Distribution**: load balancing irregular work; 300-600 GB/s; 20-50% improvement in irregular algorithms GPU Scan (Prefix Sum) represents **the essential parallel primitive for position computation** — by using hierarchical implementation with warp primitives for intra-warp (500-1000 GB/s), shared memory for inter-warp (300-600 GB/s), and multi-pass algorithms for large arrays, developers achieve 400-800 GB/s throughput (50-70% of peak bandwidth) and enable algorithms like stream compaction, radix sort, and sparse matrix operations where scan is fundamental building block and proper implementation using warp-level primitives determines whether applications achieve 100 GB/s or 800 GB/s throughput.');

gpu scheduling heterogeneous, heterogeneous task scheduling, cpu gpu co-scheduling, device affinity

**Heterogeneous Task Scheduling** is the **algorithmic and runtime discipline of assigning computational tasks to the most appropriate processing element — CPU cores, GPU compute units, FPGAs, or specialized accelerators — based on task characteristics, device capabilities, and system-wide optimization objectives** such as throughput, latency, energy efficiency, and fairness. Modern computing platforms are fundamentally heterogeneous: a single node may contain multi-core CPUs, discrete GPUs, integrated GPUs, NPUs, and FPGAs. Efficiently utilizing all resources simultaneously requires scheduling algorithms far more sophisticated than traditional homogeneous CPU scheduling. **Scheduling Dimensions**: | Dimension | Options | Impact | |-----------|---------|--------| | Device selection | CPU vs GPU vs accelerator | Throughput, energy | | Task granularity | Kernel, sub-task, pipeline stage | Overhead vs utilization | | Data placement | Host RAM, GPU VRAM, unified | Transfer cost | | Preemption | Cooperative vs preemptive | Latency, fairness | | Priority | Deadline, throughput, fairness | QoS guarantees | **CPU-GPU Co-Scheduling Strategies**: - **Static partitioning**: Assign task types to devices at compile time or configuration time. Simple but cannot adapt to runtime workload variation. - **Dynamic work-stealing**: Idle devices steal work from busy devices' queues. Requires portable task representations (e.g., OpenCL kernels that run on both CPU and GPU). - **Predictive scheduling**: Profile task execution time on each device, use performance models to assign tasks to minimize total completion time. Accounts for data transfer overhead and device contention. - **Feedback-driven**: Monitor actual execution times and adjust device allocation ratios online. EMA (exponential moving average) smoothing handles variability. **GPU Scheduling Specifics**: GPU scheduling operates at multiple levels: **application-level** (which kernels to launch and when), **driver-level** (ordering kernel submissions in hardware queues), **hardware-level** (SM/CU allocation among concurrent kernels via MPS or hardware partitioning). GPU preemption granularity varies: NVIDIA supports context-level preemption and instruction-level preemption (since Pascal), enabling real-time GPU sharing. **Frameworks and Runtimes**: **CUDA MPS** (Multi-Process Service) enables spatial GPU sharing among processes; **NVIDIA MIG** (Multi-Instance GPU) provides hardware-isolated GPU partitions; **AMD ROCm** supports similar multi-tenancy; **OpenCL** provides device-agnostic task dispatch; **StarPU** and **Legion** offer task-based heterogeneous runtimes with automatic data management; and **Kubernetes device plugins** handle cluster-level GPU scheduling. **Challenges**: **Performance portability** — the same algorithm may have 10-100x different performance on CPU vs GPU; **data gravity** — moving data between devices costs time and energy (PCIe ~32 GB/s vs GPU memory ~3 TB/s); **tail latency** — heterogeneous execution creates variable completion times that complicate deadline guarantees; and **resource fragmentation** — partially utilizing multiple devices may be worse than fully utilizing one. **Heterogeneous scheduling is the key to unlocking the full computational potential of modern hardware — systems that intelligently match workloads to devices can achieve 2-5x higher throughput and energy efficiency compared to naive CPU-only or GPU-only execution.**

gpu shared memory bank conflict,shared memory optimization cuda,bank conflict avoidance,shared memory padding,warp shared memory access

**GPU Shared Memory Bank Conflicts** represent **the performance hazard that occurs when multiple threads within a warp simultaneously access different addresses mapped to the same shared memory bank — serializing what should be parallel memory accesses and degrading shared memory bandwidth by factors proportional to the conflict degree**. **Bank Architecture:** - **Bank Organization**: shared memory is divided into 32 banks (matching warp width), each 4 bytes wide; consecutive 4-byte words map to consecutive banks (bank = (address/4) mod 32) - **Conflict-Free Access**: when all 32 threads access addresses in 32 different banks, or when all threads access the exact same address (broadcast), the access completes in a single cycle - **N-Way Conflict**: when N threads access different addresses in the same bank, the hardware serializes into N sequential accesses — a 32-way conflict (all threads hit bank 0) is 32× slower than conflict-free - **Broadcast Mechanism**: when multiple threads read the identical address, the hardware broadcasts the single read to all requesting threads in one cycle — this is NOT a conflict **Common Conflict Patterns:** - **Stride-Based Access**: accessing shared memory with stride 32 (or any multiple of 32) causes all threads to hit the same bank; stride 1 is conflict-free; stride 2 produces 2-way conflicts - **Matrix Column Access**: storing a 32×32 matrix in shared memory row-major, then reading columns produces 32-way bank conflicts — the classic transpose problem - **Reduction Operations**: naive tree-based reduction where stride doubles each step encounters bank conflicts at specific reduction levels - **Histogram Binning**: multiple threads atomically updating the same histogram bin in shared memory creates serialized atomic conflicts **Conflict Avoidance Techniques:** - **Padding**: adding one extra element per row of a 2D shared memory array shifts column addresses across banks — declaring float smem[32][33] instead of float smem[32][32] eliminates column-access conflicts with minimal memory overhead - **Index Permutation**: XOR-based index remapping (bank = threadIdx XOR some_value) distributes accesses across banks for specific access patterns like matrix transpose - **Access Reordering**: restructuring algorithms so each warp accesses shared memory with stride-1 pattern wherever possible; converting AoS to SoA layout in shared memory - **Warp-Level Primitives**: using __shfl_sync for register-to-register communication eliminates shared memory bank conflicts entirely for warp-local data exchange **Profiling and Diagnosis:** - **Nsight Compute Metrics**: l1tex__data_pipe_lsu_wavefronts_mem_shared reports actual wavefront count; comparing to ideal (1 per instruction) reveals conflict ratio - **Bank Conflict Ratio**: (actual_wavefronts / issued_instructions) - 1 gives the average number of additional serialized accesses per instruction; values above 0.2 warrant optimization - **Occupancy Impact**: severe bank conflicts do not reduce occupancy but extend instruction latency, stalling dependent operations and reducing instruction-level parallelism within each warp GPU shared memory bank conflicts are **a subtle but significant performance bottleneck that can reduce shared memory throughput by up to 32× — understanding bank mapping, applying padding or index permutation, and profiling with Nsight Compute are essential skills for achieving peak shared memory performance in CUDA kernels**.

gpu shared memory bank conflict,shared memory optimization,bank conflict resolution,shared memory access pattern

**GPU Shared Memory Bank Conflicts** are the **performance penalties that occur when multiple threads in a warp simultaneously access different addresses that map to the same shared memory bank** — forcing the accesses to be serialized rather than served simultaneously, reducing effective bandwidth by a factor equal to the degree of the conflict, and representing one of the most common and impactful GPU optimization targets. **Shared Memory Bank Architecture** - Shared memory is divided into **32 banks** (matching warp size). - Banks are interleaved: Address 0 → Bank 0, Address 4 → Bank 1, ..., Address 124 → Bank 31, Address 128 → Bank 0. - (For 4-byte words: bank = (address / 4) % 32.) - Each bank can service one address per cycle. - **No conflict**: All 32 threads access 32 different banks → 1 cycle (full bandwidth). - **N-way conflict**: N threads access same bank, different addresses → N cycles (serialized). - **Broadcast**: Multiple threads access SAME address in same bank → 1 cycle (broadcast, no conflict). **Conflict Examples** ``` // No conflict — stride 1 (consecutive access) shared[threadIdx.x] // thread 0→bank 0, thread 1→bank 1, ... // 2-way conflict — stride 2 shared[threadIdx.x * 2] // thread 0→bank 0, thread 16→bank 0 (conflict!) // 32-way conflict — stride 32 (worst case) shared[threadIdx.x * 32] // ALL threads hit bank 0 → fully serialized // No conflict — stride that is odd shared[threadIdx.x * 3] // Odd stride → all banks hit uniquely ``` **Bank Conflict Rule** - **Conflict occurs when**: stride is a multiple of any power of 2 that divides 32. - **No conflict when**: stride is odd (coprime with 32). - Stride 1: No conflict. Stride 2: 2-way. Stride 4: 4-way. Stride 32: 32-way. - Stride 3: No conflict. Stride 5: No conflict. Stride 7: No conflict. **Common Conflict Scenarios and Fixes** | Scenario | Problem | Fix | |----------|---------|-----| | Matrix column access | Stride = matrix width (power of 2) | Pad shared array: `shared[N][N+1]` | | Struct array | Struct size = power of 2 bytes | Pad struct or use SoA layout | | Reduction tree | Half-warp accesses same bank | Use sequential addressing, not interleaved | | Histogram | Multiple threads update same bin | Use privatization, then merge | **Padding Technique (Most Common Fix)** ```cuda // Problem: 32x32 matrix, column access = stride 32 = 32-way conflict __shared__ float tile[32][32]; // column access: 32-way conflict // Fix: Pad each row by 1 element __shared__ float tile[32][32 + 1]; // column access: stride 33 (odd) → no conflict! ``` **Diagnosing Bank Conflicts** - **NVIDIA Nsight Compute**: Reports shared memory bank conflicts per kernel. - **Metric**: `l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_{load,store}`. - Target: 0 conflicts. Acceptable: < 1 conflict per instruction on average. GPU shared memory bank conflicts are **one of the most frequent micro-architectural performance pitfalls** — a single line of code using a power-of-2 stride can reduce shared memory throughput by 32x, making bank conflict analysis and padding/layout optimization essential skills for GPU performance engineers.

gpu shared memory optimization,shared memory bank conflict,shared memory tiling,scratchpad memory gpu,shared memory programming

**GPU Shared Memory Optimization** is the **performance-critical programming technique that uses the GPU's fast, software-managed on-chip memory (shared memory / scratchpad) to cache frequently-accessed data, enabling data reuse across threads within a thread block while avoiding repeated expensive global memory accesses — where proper use can improve kernel performance by 5-20x but improper use (bank conflicts, insufficient occupancy) can negate the benefits entirely**. **Shared Memory Architecture** Shared memory is a low-latency (~5 cycles), high-bandwidth on-chip SRAM organized into 32 banks (each 4 bytes wide). All threads in a thread block share the same shared memory instance (configurable 48-164 KB per SM on modern GPUs). Access latency is ~100x lower than global memory (HBM: ~500 cycles). **Bank Conflicts** The 32 banks can each serve one 4-byte access per cycle simultaneously. If two or more threads in the same warp access different addresses in the same bank, the accesses are serialized (N-way bank conflict → N cycles). Conflict-free access patterns: - **Linear stride-1**: Thread i accesses word i → each thread hits a different bank. No conflict. - **Linear stride-2**: Threads 0,16 both hit bank 0; threads 1,17 both hit bank 1 → 2-way conflict everywhere. - **Stride-32**: All threads hit the same bank → 32-way conflict (worst case). Solution: Pad the shared memory array to offset the stride. For a 32 × 32 float array, declaring it as float tile[32][33] shifts each row by one bank, eliminating conflicts for column access. **Common Optimization Patterns** - **Tiling (Matrix Multiply)**: Load a tile of matrix A and matrix B into shared memory. Each thread in the block reuses these tiles for multiple multiply-accumulate operations. For a 32×32 tile, each global memory load is reused 32 times, reducing global memory traffic by 32x. - **Stencil Computation**: Load a tile plus halo (boundary elements needed by border threads) into shared memory. Compute the stencil entirely from shared memory. Avoids redundant global memory reads of overlapping halo regions. - **Histogram / Reduction**: Accumulate partial results in shared memory across the thread block, then write a single consolidated result to global memory. **Configuration Tradeoffs** - **Shared Memory vs. L1 Cache**: Modern GPUs allow configuring the partition between shared memory and L1 cache (e.g., prefer 48KB shared + 112KB L1, or 164KB shared + 0KB L1 on H100). Kernels with explicit tiling benefit from more shared memory; kernels with irregular access patterns benefit from more L1. - **Occupancy Impact**: More shared memory per block means fewer blocks can run concurrently per SM. If a kernel uses 48KB shared/block and the SM has 164KB total, only 3 blocks run simultaneously. Reducing shared memory usage to 32KB allows 5 blocks → higher occupancy and better latency hiding. GPU Shared Memory is **the parallel programmer's most powerful tool for bridging the bandwidth gap between compute and memory** — a manually-managed cache that, when used correctly, transforms memory-bound kernels into compute-bound kernels.

gpu shared memory optimization,shared memory bank conflict,tiling gpu kernel,shared memory usage cuda,local data share

**GPU Shared Memory Optimization** is the **critical CUDA/GPU programming technique of using on-chip shared memory (32-228 KB per Streaming Multiprocessor) as a programmer-managed cache to reduce global memory accesses — where properly tiled algorithms using shared memory achieve 5-50x speedup over naive global memory implementations because shared memory provides ~20 cycle latency and ~100 TB/s aggregate bandwidth compared to global memory's ~400 cycle latency and ~2-8 TB/s bandwidth**. **Shared Memory Architecture** - **Location**: On-chip SRAM within each SM/CU, shared among all threads in a thread block/workgroup. - **Size**: 48-228 KB per SM (configurable split with L1 cache on NVIDIA GPUs). Ampere: up to 164 KB. Hopper: up to 228 KB. - **Bandwidth**: 128 bytes per clock per bank. With 32 banks operating at ~1.5 GHz: ~6 TB/s per SM. - **Latency**: ~20-30 cycles. Comparable to L1 cache, 10-20x faster than global memory. **Bank Conflicts** Shared memory is organized into 32 banks (NVIDIA). Consecutive 4-byte words map to consecutive banks. If multiple threads in a warp access different addresses in the same bank in the same cycle, the accesses serialize (bank conflict): - **No conflict**: Each thread accesses a different bank. Full bandwidth. - **2-way conflict**: Two threads hit the same bank. Half bandwidth. - **32-way conflict**: All threads hit the same bank. 1/32 bandwidth (serial access). **Common conflict patterns**: - Stride-32 access: threads access every 32nd word — all map to the same bank. Worst case. - Fix: Pad the shared memory array by one element per row: `__shared__ float tile[32][33];` — the extra column shifts each row's bank mapping, eliminating conflicts. **Tiling Pattern** The canonical optimization pattern for matrix operations: 1. **Load tile**: Threads cooperatively load a tile of input data from global memory into shared memory (coalesced global reads). 2. **__syncthreads()**: Barrier ensures all threads have completed loading. 3. **Compute**: Threads read from shared memory (fast, reusable) to compute their outputs. Each element loaded once from global memory but read multiple times from shared memory. 4. **__syncthreads()**: Barrier before the next tile load (prevent overwriting data still in use). 5. **Repeat**: Iterate over tiles until the full input is processed. **GEMM Example** Naive GEMM: each element of C reads an entire row of A and column of B from global memory — N³ global reads for an N×N matrix multiply. Tiled GEMM with shared memory: load a TILE_SIZE × TILE_SIZE block of A and B into shared memory, compute partial products, iterate over tiles. Global memory reads drop from N³ to N³/TILE_SIZE — a 16-32x reduction for typical tile sizes. **GPU Shared Memory is the key lever that transforms memory-bound GPU kernels into compute-bound ones** — enabling the data reuse patterns that are essential to achieve a significant fraction of the GPU's peak computational throughput.

gpu shared memory,shared memory optimization,smem,scratchpad

**GPU Shared Memory Optimization** — using the fast, programmer-managed on-chip memory (shared memory / SMEM) within each GPU Streaming Multiprocessor (SM) to drastically reduce global memory accesses. **Shared Memory Properties** - Location: On-chip SRAM within each SM - Size: 48–228 KB per SM (configurable vs. L1 cache) - Latency: ~20-30 cycles (vs. ~400 cycles for global memory) - Bandwidth: ~10 TB/s aggregate (vs. ~2 TB/s for HBM) - Scope: Shared among all threads in a thread block **Classic Pattern: Tiled Matrix Multiply** ``` 1. Load tile of A from global → shared memory 2. Load tile of B from global → shared memory 3. __syncthreads() // All threads in block sync 4. Compute partial result using fast shared memory reads 5. Repeat for next tile ``` - Without shared memory: Each element read from slow global memory multiple times - With shared memory: Each element loaded once from global, reused many times from SMEM - Speedup: 10-20x for matrix multiply **Bank Conflicts** - Shared memory divided into 32 banks - Threads in a warp accessing different banks → simultaneous (fast) - Multiple threads accessing same bank → serialized (bank conflict, slow) - Solution: Pad shared memory arrays to avoid conflict patterns **Best Practices** - Use shared memory for data reused across threads in a block - Always `__syncthreads()` between write and read phases - Avoid bank conflicts by careful indexing **Shared memory** is the #1 optimization technique in CUDA programming — mastering it is what separates a 10x kernel from a 100x kernel.

gpu sm architecture,streaming multiprocessor,cuda core,gpu compute unit,sm design

**GPU Streaming Multiprocessor (SM) Architecture** is the **fundamental compute building block of NVIDIA GPUs, where each SM contains a set of CUDA cores, warp schedulers, register files, shared memory, and cache** — with the entire GPU composed of tens to hundreds of SMs that independently execute thread blocks, and understanding SM architecture is essential for optimizing kernel occupancy, register usage, shared memory allocation, and achieving peak throughput on any CUDA workload. **SM Components (H100 Example)** ``` ┌─────────────────── Streaming Multiprocessor (SM) ───────────────────┐ │ [Warp Scheduler 0] [Warp Scheduler 1] [Warp Scheduler 2] [WS 3] │ │ ↓ ↓ ↓ ↓ │ │ ┌─────────┐ ┌─────────┐ ┌─────────┐ ┌─────────┐ │ │ │ 32 FP32 │ │ 32 FP32 │ │ 32 FP32 │ │ 32 FP32 │ = 128 FP32 │ │ │ cores │ │ cores │ │ cores │ │ cores │ cores/SM │ │ └─────────┘ └─────────┘ └─────────┘ └─────────┘ │ │ ┌─────────┐ ┌─────────┐ │ │ │ 16 FP64 │ │ 16 FP64 │ = 32 FP64 cores/SM (if enabled) │ │ └─────────┘ └─────────┘ │ │ ┌────────────────────────┐ │ │ │ 4 Tensor Cores (4th gen)│ = Matrix multiply acceleration │ │ └────────────────────────┘ │ │ ┌────────────────────────┐ │ │ │ 16 SFU (Special Func) │ = sin, cos, rsqrt, etc. │ │ └────────────────────────┘ │ │ ┌────────────────────────┐ │ │ │ 16 LD/ST Units │ = Memory load/store │ │ └────────────────────────┘ │ │ Register File: 256 KB (65536 × 32-bit registers) │ │ L1 Cache / Shared Memory: 256 KB (configurable split) │ │ Max threads: 2048 Max warps: 64 Max blocks: 32 │ └────────────────────────────────────────────────────────────────────┘ ``` **SM Evolution Across Generations** | Architecture | Year | SMs | FP32/SM | Shared Mem/SM | Registers/SM | |-------------|------|-----|---------|--------------|-------------| | Pascal (P100) | 2016 | 56 | 64 | 64 KB | 256 KB | | Volta (V100) | 2017 | 80 | 64 | 96 KB | 256 KB | | Ampere (A100) | 2020 | 108 | 64 | 164 KB | 256 KB | | Hopper (H100) | 2022 | 132 | 128 | 256 KB | 256 KB | | Blackwell (B200) | 2024 | 160+ | 128 | 256 KB | 256 KB | **Warp Scheduling** - Each SM has 4 warp schedulers (Volta+). - Each scheduler selects one warp per cycle and issues instruction. - 4 schedulers × 1 instruction/cycle = 4 instructions/cycle per SM. - When warp stalls (memory): Scheduler instantly switches to another ready warp. - This is why occupancy matters: More warps → more scheduling options → better latency hiding. **Resource Partitioning per Thread Block** ``` Thread block requests: - 256 threads (8 warps) - 32 registers per thread = 8192 registers - 4 KB shared memory SM capacity: 65536 registers, 256 KB shared mem, 64 warps → Can fit: min(65536/8192, 256K/4K, 64/8, 32 blocks) = 8 blocks → 64 warps active → 100% occupancy ``` **Performance Optimization Based on SM** | Bottleneck | Symptom | Solution | |-----------|---------|----------| | Low occupancy | Few active warps | Reduce registers or shared mem per block | | Register spill | Slow local memory access | Reduce variables, use __launch_bounds__ | | Shared mem limited | Can't fit all data | Tile the computation | | Compute bound | All cores busy | Algorithmic optimization | | Memory bound | Cores waiting | Improve coalescing, caching | GPU SM architecture is **the hardware foundation that every CUDA optimization decision ultimately targets** — understanding how warps are scheduled, how registers and shared memory are partitioned across thread blocks, and how many SMs compose a given GPU determines whether a kernel achieves 20% or 90% of theoretical peak throughput, making SM architecture knowledge the essential bridge between writing correct GPU code and writing fast GPU code.

gpu sm occupancy optimization, streaming multiprocessor, warp occupancy, gpu occupancy tuning

**GPU SM Occupancy Optimization** is the **tuning of GPU kernel resource usage (registers, shared memory, block size) to maximize the number of concurrent warps executing on each Streaming Multiprocessor (SM)**, enabling the hardware's latency-hiding mechanism — where the SM switches to a ready warp when the current warp stalls on a memory access — to maintain high throughput despite individual memory latencies of hundreds of cycles. GPU architecture depends on massive thread-level parallelism to hide latency. Unlike CPUs (which use large caches and out-of-order execution), GPUs use thousands of concurrent threads — when one warp waits for data, the SM instantly switches to another ready warp, keeping ALUs busy. Low occupancy means insufficient warps to hide latency, leaving ALUs idle. **Occupancy Limiters**: | Resource | SM Limit (A100 example) | Impact on Occupancy | |----------|----------------------|---------------------| | **Registers per thread** | 65536 per SM | More regs → fewer concurrent threads | | **Shared memory per block** | 164 KB per SM | More shmem → fewer concurrent blocks | | **Threads per block** | 1024 max | Must be multiple of 32 (warp size) | | **Blocks per SM** | 32 max | Even if resources allow more warps | | **Warps per SM** | 64 max (2048 threads) | Hard ceiling | **Occupancy Calculation Example**: SM supports 64 warps max. Kernel uses 128 registers/thread → each thread uses 128 regs × 32 threads/warp = 4096 regs/warp. With 65536 regs/SM: 65536/4096 = 16 warps → occupancy = 16/64 = 25%. Reducing to 64 regs/thread: 2048 regs/warp → 32 warps → 50% occupancy. The trade-off: fewer registers may cause spilling to slow local memory. **When High Occupancy Matters**: Occupancy is most impactful for **memory-bound kernels** where latency hiding is critical. For a kernel that spends 90% of time waiting for global memory loads, increasing occupancy from 25% to 50% can halve the stall time, improving performance by ~40%. For **compute-bound kernels** (ALUs fully utilized at low occupancy), increasing occupancy provides minimal benefit and may even hurt performance (more register spilling, more cache pressure). **Optimization Strategies**: 1. **Reduce register usage**: Use `-maxrregcount` compiler flag, simplify per-thread computation, or manually optimize register-heavy code sections. Launch bounds (`__launch_bounds__(maxThreads, minBlocks)`) give the compiler optimization hints. 2. **Reduce shared memory**: Use shared memory only for data with true reuse; replace single-use shared memory with register-to-register warp shuffles (`__shfl_sync`). 3. **Block size tuning**: Try block sizes of 128, 256, 512 — different sizes interact differently with register/shared memory limits. Non-obvious sweet spots are common. 4. **Dynamic shared memory**: Allocate shared memory dynamically (third kernel launch parameter) instead of statically — allows runtime tuning without recompilation. **Diminishing Returns**: The relationship between occupancy and performance is not linear. Going from 25% to 50% occupancy often yields significant improvement. Going from 50% to 100% typically yields diminishing returns — beyond a threshold, the SM has enough warps to keep the pipeline full. The CUDA Occupancy Calculator and Nsight Compute's occupancy analysis help identify the sweet spot. **GPU SM occupancy optimization is the art of balancing the per-thread resource budget against the need for massive parallelism — the right balance enables the GPU's latency-hiding architecture to function effectively, translating raw hardware capability into actual application throughput.**

gpu sorting algorithms,cuda radix sort,parallel sorting gpu,gpu sort performance,cuda sort optimization

**GPU Sorting Algorithms** are **the parallel implementations of sorting that leverage thousands of GPU threads to achieve 100-300 GB/s throughput** — where radix sort (optimal for integers and fixed-point) achieves 200-300 GB/s by processing multiple bits per pass and exploiting warp-level primitives, merge sort (optimal for general comparisons) achieves 100-200 GB/s through hierarchical merging, and bitonic sort (optimal for power-of-2 sizes) achieves 150-250 GB/s with fixed communication patterns, making GPU sorting 10-50× faster than CPU sorting (5-20 GB/s) and essential for applications like database operations, graph algorithms, and data preprocessing where sorting is bottleneck (20-60% of runtime) and proper algorithm selection based on data characteristics (integer vs float, key-only vs key-value, size) determines whether applications achieve 40% or 90% of theoretical peak bandwidth. **Radix Sort:** - **Algorithm**: sorts by processing k bits per pass; typically k=4-8 bits; requires ceil(32/k) passes for 32-bit integers; stable sort - **Performance**: 200-300 GB/s on A100; 60-80% of peak memory bandwidth; optimal for integers, fixed-point; 10-50× faster than CPU - **Implementation**: histogram per block → prefix sum → scatter; uses shared memory for local histogram; warp primitives for reduction - **Use Cases**: integer keys, fixed-point values; uniform distribution; large arrays (>1M elements); 80-95% of peak bandwidth **Merge Sort:** - **Algorithm**: hierarchical merging; bottom-up or top-down; log2(N) passes; stable sort; comparison-based - **Performance**: 100-200 GB/s on A100; 40-60% of peak bandwidth; optimal for general comparisons, small arrays - **Implementation**: warp-level merge → block-level merge → global merge; uses shared memory for local merging - **Use Cases**: general comparisons, custom comparators; small-medium arrays (10K-1M elements); stable sort required **Bitonic Sort:** - **Algorithm**: comparison network; fixed communication pattern; log2(N) × (log2(N)+1) / 2 comparisons; not stable - **Performance**: 150-250 GB/s on A100; 50-70% of peak bandwidth; optimal for power-of-2 sizes; predictable performance - **Implementation**: warp-level bitonic → block-level bitonic → global bitonic; uses shuffle for warp-level, shared memory for block-level - **Use Cases**: power-of-2 sizes, predictable latency; small-medium arrays; GPU-friendly communication pattern **Thrust Sort:** - **API**: thrust::sort(d_vec.begin(), d_vec.end()); automatic algorithm selection; radix sort for integers, merge sort for general - **Performance**: 100-300 GB/s; 60-80% of hand-tuned; 1 line of code vs 100-200 for custom implementation - **Customization**: thrust::sort(d_vec.begin(), d_vec.end(), thrust::greater()); custom comparators supported - **Use Cases**: rapid development, general-purpose sorting; acceptable 10-30% performance gap vs hand-tuned **CUB Sort:** - **API**: cub::DeviceRadixSort::SortKeys(d_temp_storage, temp_storage_bytes, d_keys, d_sorted, N); explicit control - **Performance**: 200-300 GB/s; 70-90% of peak bandwidth; 10-30% faster than Thrust; lower-level API - **Features**: key-only or key-value pairs; ascending or descending; segmented sort; double-buffer for in-place - **Use Cases**: performance-critical sorting; fine-grained control; production systems **Key-Value Sorting:** - **Radix Sort**: cub::DeviceRadixSort::SortPairs(); sorts keys, reorders values; 150-250 GB/s; 2× slower than key-only - **Merge Sort**: stable sort preserves value order; 80-150 GB/s; 40-60% of peak bandwidth - **Performance**: 2-3× slower than key-only; memory bandwidth limited; 2× data movement - **Use Cases**: sorting with associated data; database operations; graph algorithms **Segmented Sort:** - **Concept**: sort multiple independent segments; each segment sorted separately; useful for batched operations - **Implementation**: cub::DeviceSegmentedRadixSort::SortKeys(); segment offsets specify boundaries - **Performance**: 150-250 GB/s; 50-70% of peak; depends on segment sizes; small segments have overhead - **Use Cases**: batched sorting, per-group sorting; graph algorithms; database operations **Optimization Techniques:** - **Warp-Level Primitives**: use __shfl for warp-level sorting; 2-5× faster than shared memory; 500-1000 GB/s for small arrays - **Shared Memory**: use for block-level sorting; 100× faster than global memory; 164KB per SM on A100 - **Coalesced Access**: ensure coalesced memory access; 128-byte aligned; achieves 100% bandwidth; stride-1 access optimal - **Occupancy**: balance shared memory usage and occupancy; 50-100% occupancy typical; 256 threads per block optimal **Radix Sort Optimization:** - **Bits Per Pass**: 4-8 bits typical; more bits = fewer passes but larger histogram; 4 bits optimal for most cases - **Histogram**: per-block histogram in shared memory; warp primitives for reduction; 300-600 GB/s - **Prefix Sum**: exclusive scan of histogram; 400-800 GB/s; CUB provides optimized implementation - **Scatter**: write sorted elements to output; coalesced writes; 200-300 GB/s; double-buffering eliminates copies **Merge Sort Optimization:** - **Warp-Level Merge**: use __shfl for merging within warp; 2-5× faster than shared memory; 500-1000 GB/s - **Block-Level Merge**: shared memory for merging within block; 100-200 GB/s; minimize global memory accesses - **Global Merge**: hierarchical merging; multiple passes; 80-150 GB/s; memory-bound - **Path Decomposition**: parallel merge path algorithm; better load balancing; 20-40% speedup **Bitonic Sort Optimization:** - **Warp-Level**: use __shfl_xor for butterfly exchanges; 2-5× faster than shared memory; 500-1000 GB/s - **Block-Level**: shared memory for larger bitonic networks; 150-250 GB/s; minimize bank conflicts - **Padding**: add padding to avoid bank conflicts; 1-2 elements typical; 10-30% improvement - **Unrolling**: unroll inner loops; reduces overhead; 10-20% speedup; compiler often does automatically **Performance Comparison:** - **Radix Sort**: 200-300 GB/s; best for integers; 60-80% of peak; 10-50× faster than CPU - **Merge Sort**: 100-200 GB/s; best for general comparisons; 40-60% of peak; 5-20× faster than CPU - **Bitonic Sort**: 150-250 GB/s; best for power-of-2 sizes; 50-70% of peak; 10-30× faster than CPU - **Thrust Sort**: 100-300 GB/s; automatic selection; 60-80% of hand-tuned; easiest to use **Size Considerations:** - **Small Arrays (<10K)**: bitonic sort or warp-level sort; 150-250 GB/s; low overhead; predictable latency - **Medium Arrays (10K-1M)**: merge sort or radix sort; 150-250 GB/s; good balance; algorithm depends on data type - **Large Arrays (>1M)**: radix sort for integers, merge sort for general; 200-300 GB/s; amortizes overhead; optimal performance - **Very Large (>100M)**: out-of-core sorting; multiple passes; 100-200 GB/s; memory-limited **Stability:** - **Stable**: radix sort, merge sort; preserves relative order of equal elements; required for some applications - **Unstable**: bitonic sort, quicksort; may reorder equal elements; faster but less predictable - **Use Cases**: stable required for multi-key sorting, database operations; unstable acceptable for unique keys **Custom Comparators:** - **Merge Sort**: supports custom comparators; thrust::sort(d_vec.begin(), d_vec.end(), my_comparator()); flexible - **Radix Sort**: limited to integer keys; can use custom bit extraction; less flexible but faster - **Performance**: custom comparators 10-30% slower; function call overhead; inline when possible **Profiling and Tuning:** - **Nsight Compute**: shows memory bandwidth, occupancy, warp efficiency; identifies bottlenecks - **Metrics**: achieved bandwidth / peak bandwidth; target 60-80% for sorting; memory-bound operation - **Bottlenecks**: uncoalesced access, bank conflicts, low occupancy; optimize access patterns - **Tuning**: adjust block size, bits per pass, shared memory usage; profile to find optimal **Best Practices:** - **Use Libraries**: Thrust or CUB for most cases; 60-90% of hand-tuned; 10-100× less code - **Algorithm Selection**: radix sort for integers, merge sort for general; bitonic for power-of-2; profile to verify - **Batch Operations**: sort multiple arrays together; amortizes overhead; 20-40% improvement - **Profile**: measure actual bandwidth; compare with peak; optimize only if bottleneck - **Pre-Allocate**: allocate temporary storage once; reuse across sorts; eliminates allocation overhead **Performance Targets:** - **Radix Sort**: 200-300 GB/s; 60-80% of peak (1.5-3 TB/s); optimal for integers - **Merge Sort**: 100-200 GB/s; 40-60% of peak; optimal for general comparisons - **Bitonic Sort**: 150-250 GB/s; 50-70% of peak; optimal for power-of-2 sizes - **Key-Value**: 150-250 GB/s; 50-70% of peak; 2× slower than key-only **Real-World Applications:** - **Database Operations**: sorting query results; 200-300 GB/s with radix sort; 10-50× faster than CPU - **Graph Algorithms**: sorting edges by source/destination; 150-250 GB/s; 20-40% of graph processing time - **Data Preprocessing**: sorting features for ML; 200-300 GB/s; 10-30% of preprocessing time - **Rendering**: sorting primitives by depth; 150-250 GB/s; 5-20% of rendering time GPU Sorting Algorithms represent **the essential building block for data-intensive applications** — by leveraging thousands of parallel threads and optimized memory access patterns, GPU sorting achieves 100-300 GB/s throughput (10-50× faster than CPU) through algorithms like radix sort for integers (200-300 GB/s), merge sort for general comparisons (100-200 GB/s), and bitonic sort for power-of-2 sizes (150-250 GB/s), making GPU sorting critical for applications where sorting is bottleneck and proper algorithm selection based on data characteristics determines whether applications achieve 40% or 90% of theoretical peak bandwidth.

gpu sparse matrix operations,cuda sparse linear algebra,cusparse optimization,sparse matrix gpu performance,csr coo format gpu

**GPU Sparse Matrix Operations** are **the specialized algorithms for matrices where most elements are zero, exploiting sparsity to reduce memory and computation** — where Compressed Sparse Row (CSR) format stores only non-zero elements achieving 10-100× memory reduction and SpMV (Sparse Matrix-Vector multiplication) achieves 100-500 GB/s (20-60% of peak bandwidth) through irregular memory access patterns, while cuSPARSE library provides optimized implementations of SpMV, SpMM (Sparse Matrix-Matrix), and sparse solvers that are 5-50× faster than naive implementations, making sparse operations essential for scientific computing, graph algorithms, and machine learning where 90-99% of matrix elements are zero and proper format selection (CSR for SpMV, COO for construction, CSC for column access) and optimization techniques (vectorization, load balancing, format conversion) determine whether applications achieve 50 GB/s or 500 GB/s throughput. **Sparse Matrix Formats:** - **CSR (Compressed Sparse Row)**: stores row pointers, column indices, values; optimal for SpMV; 10-100× memory reduction; most common format - **COO (Coordinate)**: stores row indices, column indices, values; simple construction; optimal for building; easy to parallelize - **CSC (Compressed Sparse Column)**: column-major version of CSR; optimal for column access; used in some solvers - **ELL (ELLPACK)**: fixed number of non-zeros per row; regular memory access; good for uniform sparsity; wastes memory for irregular **SpMV (Sparse Matrix-Vector Multiplication):** - **Algorithm**: y = A * x where A is sparse; each row computes dot product with x; irregular memory access to x - **Performance**: 100-500 GB/s on A100; 20-60% of peak bandwidth; limited by irregular access; 5-20× faster than CPU - **CSR Implementation**: each thread/warp processes one row; loads x elements based on column indices; accumulates result - **Optimization**: warp-per-row for long rows, thread-per-row for short rows; vectorization for regular patterns; 2-5× speedup **cuSPARSE Library:** - **SpMV**: cusparseSpMV() for CSR, COO, CSC formats; automatic algorithm selection; 100-500 GB/s; 80-95% of hand-tuned - **SpMM**: cusparseSpMM() for sparse-dense matrix multiplication; 200-800 GB/s; uses Tensor Cores when possible - **Sparse Solvers**: cusparseSpSV() for triangular solve; cusparseSpSM() for multiple right-hand sides; 100-400 GB/s - **Format Conversion**: cusparseCsr2coo(), cusparseCoo2csr(); efficient conversion; 200-400 GB/s **Load Balancing:** - **Thread-Per-Row**: simple but imbalanced; short rows waste threads; long rows serialize; 50-200 GB/s - **Warp-Per-Row**: better for long rows; uses warp reduction; 100-400 GB/s; good for uniform row lengths - **Dynamic Scheduling**: work queue for rows; load balancing; 150-500 GB/s; optimal for irregular sparsity - **Hybrid**: thread-per-row for short, warp-per-row for long; 200-500 GB/s; best overall performance **Vectorization:** - **Vector Loads**: use float4, int4 for consecutive elements; 2-4× fewer transactions; 20-50% speedup - **Alignment**: align data to 128 bytes; enables vectorization; 10-30% improvement - **Padding**: pad rows to multiples of 4/8; enables vectorization; 20-40% speedup; wastes some memory - **Use Cases**: regular sparsity patterns; structured matrices; 20-50% improvement **Memory Access Optimization:** - **Coalescing**: difficult for sparse matrices; irregular column indices; use shared memory for x vector - **Shared Memory**: cache frequently accessed x elements; reduces global memory traffic; 20-50% speedup - **Texture Memory**: use texture cache for x vector; benefits from spatial locality; 10-30% speedup for some patterns - **Prefetching**: prefetch next row's data; hides latency; 10-20% improvement **Format Selection:** - **CSR**: best for SpMV; row-major access; 100-500 GB/s; most common; use for general sparse operations - **COO**: best for construction; easy parallelization; 200-400 GB/s for building; convert to CSR for SpMV - **CSC**: best for column access; transpose operations; 100-500 GB/s; use when column access dominates - **ELL**: best for uniform sparsity; regular access; 200-600 GB/s; wastes memory for irregular **Sparse Matrix Construction:** - **COO Building**: parallel insertion of non-zeros; 200-400 GB/s; sort by row then column; convert to CSR - **Atomic Operations**: use atomics for concurrent insertion; 50-200 GB/s; high contention; use warp aggregation - **Sorting**: sort COO entries; 100-300 GB/s with GPU sort; required for CSR conversion - **CSR Conversion**: scan row counts; compute row pointers; copy values and columns; 200-400 GB/s **SpMM (Sparse-Dense Matrix Multiplication):** - **Algorithm**: C = A * B where A is sparse, B is dense; multiple SpMV operations; can use Tensor Cores - **Performance**: 200-800 GB/s on A100; 30-70% of peak; benefits from dense B; Tensor Cores for large B - **Optimization**: process multiple columns of B together; use Tensor Cores when possible; 2-5× speedup - **Use Cases**: sparse neural network layers; graph neural networks; scientific computing **Sparse Solvers:** - **Triangular Solve**: cusparseSpSV(); forward/backward substitution; 100-400 GB/s; level scheduling for parallelism - **Iterative Solvers**: CG, BiCGSTAB, GMRES; SpMV is bottleneck; 100-500 GB/s; 80-95% time in SpMV - **Preconditioners**: ILU, Jacobi; improve convergence; 100-400 GB/s; critical for performance - **Multi-GPU**: distribute matrix across GPUs; NCCL for communication; 70-85% scaling efficiency **Graph Algorithms:** - **BFS/DFS**: sparse adjacency matrix; SpMV-like operations; 100-400 GB/s; irregular access patterns - **PageRank**: iterative SpMV; 100-500 GB/s; 80-95% time in SpMV; benefits from optimization - **Connected Components**: sparse matrix operations; 100-400 GB/s; irregular parallelism - **Shortest Path**: sparse matrix operations; 100-400 GB/s; dynamic parallelism helps **Performance Profiling:** - **Nsight Compute**: shows memory bandwidth, warp efficiency, occupancy; identifies bottlenecks - **Metrics**: achieved bandwidth / peak bandwidth; target 20-60% for sparse (irregular access); memory-bound - **Bottlenecks**: irregular access, load imbalance, low occupancy; optimize based on sparsity pattern - **Tuning**: adjust algorithm (thread/warp per row), vectorization, shared memory; profile to find optimal **Sparsity Patterns:** - **Uniform**: similar non-zeros per row; ELL format good; 200-600 GB/s; regular access patterns - **Power-Law**: few rows with many non-zeros; hybrid approach; 150-500 GB/s; load balancing critical - **Block-Sparse**: non-zeros in blocks; block-CSR format; 300-800 GB/s; exploits structure - **Random**: irregular sparsity; CSR format; 100-400 GB/s; difficult to optimize **Best Practices:** - **Use cuSPARSE**: highly optimized; 80-95% of hand-tuned; 10-100× less code - **Format Selection**: CSR for SpMV, COO for construction, CSC for column access; convert as needed - **Load Balancing**: use hybrid approach (thread/warp per row); 2-5× speedup over naive - **Profile**: measure actual bandwidth; compare with dense operations; optimize only if bottleneck - **Vectorization**: use when possible; 20-50% improvement for regular patterns **Performance Targets:** - **SpMV**: 100-500 GB/s; 20-60% of peak (1.5-3 TB/s); irregular access limits performance - **SpMM**: 200-800 GB/s; 30-70% of peak; benefits from dense matrix; Tensor Cores help - **Construction**: 200-400 GB/s; 30-50% of peak; sorting and conversion overhead - **Sparse Solvers**: 100-400 GB/s; 20-50% of peak; SpMV dominates; iterative methods **Real-World Applications:** - **Scientific Computing**: finite element, computational fluid dynamics; 100-500 GB/s SpMV; 80-95% of solver time - **Graph Algorithms**: social networks, web graphs; 100-400 GB/s; irregular access patterns - **Machine Learning**: sparse neural networks, embeddings; 200-800 GB/s SpMM; Tensor Cores help - **Recommendation Systems**: sparse user-item matrices; 100-500 GB/s; large-scale sparse operations GPU Sparse Matrix Operations represent **the challenge of irregular parallelism** — by exploiting sparsity through specialized formats like CSR (10-100× memory reduction) and optimized algorithms that achieve 100-500 GB/s (20-60% of peak bandwidth) despite irregular memory access, developers enable scientific computing, graph algorithms, and machine learning on matrices where 90-99% of elements are zero, making sparse operations essential where proper format selection and optimization techniques like load balancing, vectorization, and cuSPARSE library usage determine whether applications achieve 50 GB/s or 500 GB/s throughput.');

gpu stream event synchronization,cuda stream concurrent kernels,cuda event timing synchronization,multi stream overlap gpu,default stream synchronization cuda

**GPU Stream and Event Synchronization** is **the CUDA programming model for managing concurrent operations on the GPU by organizing kernels and memory transfers into streams (ordered sequences of operations) and using events to synchronize between them** — effective stream usage enables overlapping computation with data transfer, concurrent kernel execution, and precise timing measurements. **CUDA Stream Fundamentals:** - **Stream Definition**: a stream is a sequence of GPU operations that execute in order — operations in different streams may execute concurrently if the GPU has available resources - **Default Stream**: operations without an explicit stream use stream 0 (the default/legacy stream) — the default stream implicitly synchronizes with all other streams unless compiled with --default-stream per-thread - **Stream Creation**: cudaStreamCreate(&stream) creates a non-blocking stream — cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking) creates a stream that doesn't synchronize with the default stream - **Stream Destruction**: cudaStreamDestroy(stream) releases stream resources — any pending operations in the stream complete before destruction **Overlapping Computation and Transfer:** - **Concurrent Copy and Execute**: launch memory copies (cudaMemcpyAsync) on one stream and kernels on another — the GPU's copy engines (DMA) and compute engines (SMs) operate independently - **Double Buffering Pattern**: alternate between two buffers — while the GPU computes on buffer A, transfer results from the previous iteration and load input for the next iteration using buffer B - **Triple Buffering**: three buffers provide maximum overlap — one being computed, one being uploaded, one being downloaded — hides both upload and download latency simultaneously - **Pinned Memory Requirement**: cudaMemcpyAsync requires pinned (page-locked) host memory allocated with cudaMallocHost — unpinned memory forces synchronous copies regardless of stream assignment **Multi-Stream Concurrency:** - **Concurrent Kernels**: multiple small kernels on different streams can execute simultaneously if the GPU has enough SMs — Ampere A100 supports up to 128 concurrent kernels - **Resource Partitioning**: concurrent kernels share SM resources — each kernel gets a portion of SMs proportional to its grid size, enabling fine-grained GPU sharing - **Stream Priority**: cudaStreamCreateWithPriority(&stream, flags, priority) assigns a scheduling priority — high-priority streams preempt low-priority ones at thread block boundaries - **Breadth-First Launch**: launch all streams' operations in interleaved order (kernel1_stream1, kernel1_stream2, ...) rather than depth-first (all of stream1, then all of stream2) — ensures the GPU sees concurrent work from all streams **CUDA Events:** - **Event Creation**: cudaEventCreate(&event) creates an event marker — events are recorded into streams and can be queried or waited on from the host or other streams - **Recording**: cudaEventRecord(event, stream) inserts the event into the stream's operation queue — the event is "completed" when all preceding operations in that stream finish - **Host Synchronization**: cudaEventSynchronize(event) blocks the host thread until the event completes — more fine-grained than cudaStreamSynchronize which waits for all stream operations - **Query**: cudaEventQuery(event) returns cudaSuccess if the event is complete or cudaErrorNotReady if still pending — enables non-blocking polling from the host **Inter-Stream Synchronization:** - **cudaStreamWaitEvent**: cudaStreamWaitEvent(stream, event, 0) makes all subsequent operations in stream wait until event is complete — establishes a dependency between two streams without blocking the host - **Use Case**: launch data preparation on stream 1, record an event, then make stream 2 (which runs the compute kernel) wait on that event — ensures data is ready before computation begins - **Fork-Join Pattern**: launch multiple independent kernels on separate streams, create events for each, and have a final stream wait on all events before launching the reduction kernel — maximizes concurrency while maintaining correctness - **Graph Capture**: cudaStreamBeginCapture captures a sequence of stream operations into a CUDA Graph — replaying the graph eliminates per-launch overhead (5-10 µs per kernel launch saved) **Timing with Events:** - **Elapsed Time**: cudaEventElapsedTime(&ms, start, stop) returns the time in milliseconds between two recorded events — GPU-side timing with ~0.5 µs resolution, unaffected by host-side scheduling - **Pattern**: cudaEventRecord(start, stream) → launch kernel → cudaEventRecord(stop, stream) → cudaEventSynchronize(stop) → cudaEventElapsedTime(&ms, start, stop) — the standard GPU timing pattern - **Profiling**: events provide accurate per-kernel timing without external profiler overhead — essential for performance optimization and regression detection in production code **Common Pitfalls:** - **Default Stream Serialization**: accidentally using the default stream for one operation serializes the entire GPU — always use explicit streams for concurrent work - **Insufficient Concurrency**: launching many small kernels sequentially on one stream wastes GPU resources — distribute across multiple streams to utilize idle SMs - **Missing Synchronization**: reading GPU results on the host without synchronization leads to undefined behavior — always synchronize (event, stream, or device) before accessing results - **Pinned Memory Exhaustion**: allocating too much pinned memory with cudaMallocHost can degrade system performance — pinned memory can't be swapped, limiting available system memory **Stream and event management is the programmer's primary tool for maximizing GPU utilization — well-structured multi-stream applications achieve 90-95% GPU utilization by overlapping transfers with computation, hiding latency behind concurrent operations, and minimizing synchronization barriers.**

gpu tensor core programming,tensor core matrix multiply,wmma api cuda,tensor core mixed precision,tensor core performance optimization

**GPU Tensor Core Programming** is **the technique of leveraging specialized matrix-multiply-and-accumulate hardware units in modern GPUs to achieve dramatic speedups for linear algebra operations — performing 4×4 or larger matrix operations in a single clock cycle with throughput exceeding 1 PFLOPS on high-end GPUs**. **Tensor Core Architecture:** - **Matrix Operation**: each tensor core performs D = A × B + C where A and B are typically FP16/BF16/TF32/INT8 and C/D are FP32 — a single SM contains 4-16 tensor cores depending on GPU generation - **Throughput Progression**: Volta: 125 TFLOPS (FP16); Ampere: 312 TFLOPS (FP16/BF16); Hopper: 989 TFLOPS (FP16) — tensor cores provide 8-16× throughput improvement over standard CUDA cores for supported operations - **Data Types Supported**: FP16×FP16→FP16/FP32 (training), BF16×BF16→FP32 (training), TF32×TF32→FP32 (easy migration from FP32), INT8×INT8→INT32 (inference), FP8×FP8→FP16/FP32 (Hopper inference/training) - **Warp-Level Operation**: tensor core operations execute at warp granularity — all 32 threads cooperatively provide matrix fragments and receive results **Programming Interfaces:** - **WMMA API (Warp Matrix Multiply-Accumulate)**: C++ API with fragment types for A, B, C matrices — load_matrix_sync, mma_sync, and store_matrix_sync operations manage fragment data; supports 16×16×16 and other tile sizes - **MMA PTX Instructions**: lower-level PTX assembly providing finer control over tensor core operations — mma.sync.aligned instruction specifies exact matrix dimensions and data types; used by library developers for maximum performance - **cuBLAS/cuDNN**: high-level libraries automatically use tensor cores when input dimensions and data types are compatible — cuBLAS GEMM with FP16 inputs automatically dispatches to tensor cores; easiest adoption path - **Cutlass**: NVIDIA template library for custom GEMM implementations using tensor cores — provides building blocks (tile iterators, warp-level MMA, epilogue fusion) for researchers needing custom matrix operation variants **Optimization Techniques:** - **Tile Size Selection**: matrix dimensions should be multiples of tensor core tile size (16 for WMMA) — padding to multiples achieves full tensor core utilization; odd dimensions waste partial tiles - **Memory Layout**: column-major or row-major layout must match the fragment loading pattern — mismatched layout requires transpose operations that reduce effective throughput - **Epilogue Fusion**: combining matrix multiply with subsequent element-wise operations (bias add, activation, scaling) in the same kernel avoids writing/reading intermediate results — improves memory efficiency by 2-3× - **Occupancy vs. Tile Size**: larger tiles improve computation efficiency but reduce SM occupancy — optimal tile size balances tensor core utilization with memory latency hiding **Tensor cores represent the primary performance driver for modern AI workloads — understanding how to structure computations to leverage tensor cores is essential for achieving the published TFLOPS ratings of modern GPUs, as standard CUDA cores provide only a fraction of this throughput.**

gpu tensor core programming,wmma api matrix multiply,mixed precision tensor core,mma ptx instruction,tensor core accumulator fp32

**GPU Tensor Core Programming** is **the practice of utilizing specialized matrix multiply-accumulate (MMA) hardware units in NVIDIA GPUs that perform small matrix operations (e.g., 16×16×16) in a single clock cycle with mixed-precision arithmetic** — Tensor Cores deliver 5-10× higher throughput than standard CUDA cores for matrix-heavy workloads like deep learning and scientific computing. **Tensor Core Hardware Architecture:** - **Matrix Operation**: each Tensor Core performs D = A × B + C where A and B are small matrices (typically 4×4 in hardware, exposed as 16×16×16 at the warp level) — inputs A, B can be FP16/BF16/TF32/INT8 while accumulator C/D is FP32 or FP16 - **Throughput per SM**: Ampere (A100) has 4 Tensor Cores per SM, each performing 256 FP16 FMA operations per cycle — total 1024 FMA ops/cycle/SM vs. 64 FMA ops/cycle/SM for CUDA cores (16× speedup) - **Supported Precisions**: FP16×FP16→FP32, BF16×BF16→FP32, TF32×TF32→FP32, FP64×FP64→FP64, INT8×INT8→INT32, INT4×INT4→INT32, FP8×FP8→FP32 (Hopper) - **Warp-Level Operation**: Tensor Core instructions are warp-cooperative — all 32 threads in a warp collectively provide the input matrix fragments and receive the output fragments **WMMA API (Warp Matrix Multiply-Accumulate):** - **Fragment Declaration**: wmma::fragment a_frag — declares a fragment for the A matrix in 16×16×16 configuration with FP16 row-major layout - **Load Operation**: wmma::load_matrix_sync(a_frag, ptr, stride) — cooperatively loads a 16×16 matrix tile from global or shared memory across all threads in the warp - **MMA Operation**: wmma::mma_sync(d_frag, a_frag, b_frag, c_frag) — performs the matrix multiply-accumulate D = A × B + C using Tensor Cores in a single warp-synchronous call - **Store Operation**: wmma::store_matrix_sync(ptr, d_frag, stride, wmma::mem_row_major) — cooperatively stores the result fragment back to memory **MMA PTX Instructions (Lower-Level):** - **mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32**: PTX instruction for 16×8×16 matrix multiply — finer granularity than WMMA, allows more flexible tiling strategies - **Register Mapping**: each thread holds specific elements of the matrix fragments in its registers — understanding the thread-to-element mapping is critical for efficient data loading - **CUTLASS Library**: NVIDIA's templated C++ library abstracts MMA instructions with compile-time tile size selection — provides optimized epilogue fusion, software pipelining, and warp specialization - **Warp Specialization**: in Hopper's programming model, warps are specialized for either data loading (producer) or computation (consumer) — decouples memory access from Tensor Core execution **Performance Optimization:** - **Shared Memory Staging**: load global memory tiles into shared memory, then load WMMA fragments from shared memory — eliminates redundant global memory accesses across warps computing adjacent output tiles - **Software Pipelining**: overlap global-to-shared memory loads for the next tile with Tensor Core computation on the current tile — maintains Tensor Core utilization at >90% for large matrices - **Register Pressure**: WMMA fragments consume significant register space (a 16×16×16 operation uses ~64 registers per thread) — balance fragment count against occupancy to maximize throughput - **Memory Layout**: Tensor Cores achieve peak performance with specific memory alignment (256-byte aligned, contiguous in the fast-changing dimension) — column-major A and row-major B avoid bank conflicts in shared memory **Mixed-Precision Training Pattern:** - **Forward Pass**: store master weights in FP32, cast to FP16/BF16 for Tensor Core GEMM operations — Tensor Cores compute in reduced precision but accumulate in FP32 - **Loss Scaling**: multiply loss by a scale factor (typically 1024-65536) before backward pass to prevent FP16 gradient underflow — dynamic loss scaling adjusts the factor based on overflow detection - **Gradient Accumulation**: accumulate gradients in FP32 even when individual gradient computations use FP16 — prevents precision loss during summation across micro-batches - **Weight Update**: apply FP32 gradients to FP32 master weights, then cast back to FP16 for next iteration — maintains model accuracy while achieving 2-3× training speedup from Tensor Cores **Tensor Cores have transformed GPU computing from a throughput-oriented architecture to a matrix-computation engine — modern AI training and inference workloads spend 90%+ of their compute time in Tensor Core GEMM operations, making their efficient utilization the single most important optimization for GPU performance.**

gpu tensor core,matrix multiply accelerator,mixed precision hardware,tensor operation unit,wmma instruction

**GPU Tensor Cores** are the **specialized matrix multiplication accelerator units embedded in modern GPU architectures (NVIDIA Volta and later, AMD Matrix Cores) that perform small matrix multiply-accumulate operations (e.g., 4×4×4 or 16×16×16) in a single clock cycle at throughput rates 8-16x higher than standard floating-point units — enabling the massive FLOPS numbers (hundreds of TFLOPS) required for deep learning training and inference**. **Architecture** Each Tensor Core performs a D = A × B + C operation on small matrix tiles in one cycle: - **Input matrices A, B**: FP16, BF16, TF32, FP8, or INT8 depending on generation. - **Accumulator matrix C/D**: FP32 or FP16 for higher precision accumulation. - **Throughput**: NVIDIA H100 delivers 989 TFLOPS at FP16 Tensor and 1,979 TFLOPS at FP8. Standard FP32 ALUs deliver 67 TFLOPS — a 15-30x gap. **Hardware Generations** | Generation | GPU | Tile Size | Precisions | Peak TFLOPS | |-----------|-----|-----------|-----------|-------------| | Volta (1st gen) | V100 | 4×4×4 | FP16→FP32 | 125 | | Turing (2nd gen) | T4/RTX 2080 | 4×4×4 | FP16,INT8,INT4 | 130 | | Ampere (3rd gen) | A100 | 8×4×8+ | FP16,BF16,TF32,FP64 | 312 | | Hopper (4th gen) | H100 | 16×16×16 | FP16,BF16,FP8,INT8 | 989 | | Blackwell (5th gen) | B200 | larger | FP4,FP6,FP8 | 4,500 | **Programming Model** - **WMMA (Warp Matrix Multiply-Accumulate)**: CUDA API where a warp cooperatively loads matrix fragments from shared memory, performs the MMA operation, and stores the result. Each thread in the warp holds a portion of the matrix fragments. - **MMA PTX Instructions**: Lower-level interface giving finer control over tile sizes and data layouts. - **cuBLAS/cuDNN**: High-level libraries that automatically use Tensor Cores for GEMM and convolution operations. The recommended interface for most users — library kernels are highly tuned for each GPU generation. **Mixed Precision Training** Tensor Cores enable mixed-precision training: forward and backward passes compute in FP16/BF16 (fast Tensor Core operations), while master weights are maintained in FP32 for gradient accumulation accuracy. Loss scaling prevents gradient underflow in FP16. Result: 2-3x training speedup with negligible accuracy loss. **Feeding Tensor Cores** Tensor Cores can execute faster than data arrives from memory. Efficient utilization requires: - **Shared Memory Tiling**: Load input matrices from global memory into shared memory tiles, then feed tiles to Tensor Cores. The software pipeline of load→compute→store must overlap to hide latency. - **Large Tile Sizes**: Larger GEMM dimensions improve Tensor Core utilization. Small matrix operations (batch size 1 inference) under-utilize Tensor Cores and are better served by standard ALUs. - **Data Layout**: Tensor Cores expect specific data layouts (column-major fragments). Memory access patterns must align with these requirements. **GPU Tensor Cores are the silicon embodiment of the observation that neural network computation is dominated by matrix multiplication** — purpose-built hardware that delivers an order of magnitude more throughput for the single operation class that matters most for AI workloads.

gpu texture memory,constant memory,read only cache,cuda texture,cuda cache

**GPU Texture and Constant Memory** are **specialized GPU memory spaces with dedicated caches** — optimized for specific access patterns that offer higher effective bandwidth than global memory for appropriate workloads. **CUDA Memory Hierarchy Summary** | Memory | Scope | Cached? | Bandwidth | |--------|-------|---------|----------| | Global | All threads | L2 | ~900 GB/s | | Shared | Block | On-chip | ~19 TB/s | | Texture | All threads | L1 tex | ~900 GB/s + spatial cache | | Constant | All threads | Const cache | Broadcast | | Local | Thread | L2 | Slow | **Texture Memory** - Cached in a separate L1 texture cache (separate from L1 data cache). - **Spatial locality caching**: Optimized for 2D access patterns — if a thread accesses (x,y), neighbors (x+1,y), (x,y+1) likely cached. - **Hardware interpolation**: GPU hardware performs bilinear/trilinear interpolation for free. - **Address modes**: Wrap, clamp, mirror — hardware boundary handling. - **Usage**: Image processing (sampling at non-integer coordinates), simulation stencils. **Texture Example** ```cuda texture tex; cudaBindTexture2D(0, tex, d_data, channelDesc, width, height, pitch); __global__ void sample_kernel() { float val = tex2D(tex, u, v); // Bilinear interpolation included } ``` **Constant Memory** - 64KB total, cached in dedicated constant cache per SM. - **Broadcast**: If all threads in a warp access the same address → single cache transaction (vs. 32 separate loads from global memory). - **Best use**: Read-only data accessed uniformly by all threads (filter coefficients, LUT, camera parameters). - **Performance**: Matching all-uniform access → as fast as registers. Divergent access → serialized (slow). **Read-Only Cache (__ldg)** - Modern alternative to texture for read-only global data. - `__ldg(&ptr[i])`: Use L1 read-only cache (separate from L1 data cache). - No setup required — simpler than texture objects. - Good for gather patterns with spatial locality. Texture and constant memory are **specialized caches that provide free speedups for specific access patterns** — image processing kernels using texture memory can achieve 2-4x better cache hit rates than equivalent global memory accesses on spatially correlated data.

gpu thermal throttling,gpu boost clock,thermal design power,gpu temperature,tdp throttle,gpu power limit

**GPU Thermal Throttling and Power Management** is the **hardware and firmware mechanism that dynamically reduces GPU clock frequency and voltage when the chip temperature or power consumption approaches or exceeds design limits** — balancing the fundamental tradeoff between maximum performance (achieved at high frequency and voltage) and reliable long-term operation within thermal and electrical safety boundaries. Understanding throttling behavior is essential for ML engineers who need sustained high-throughput training runs and hardware engineers designing GPU-based systems. **GPU Power and Thermal Limits** | GPU | TDP | Max Boost Clock | Throttle Temperature | Typical AI Workload Power | |-----|-----|----------------|---------------------|---------------------------| | NVIDIA A100 SXM | 400 W | 1410 MHz | 83°C | 350–400 W | | NVIDIA H100 SXM | 700 W | 1980 MHz | 83°C | 650–700 W | | AMD MI300X | 750 W | — | 110°C (junction) | 600–750 W | | NVIDIA RTX 4090 | 450 W | 2520 MHz | 90°C | 350–450 W | **GPU Boost Clock Algorithm (NVIDIA)** - Base clock: Guaranteed minimum frequency at TDP. - Boost clock: Maximum frequency achieved when power and thermal headroom available. - **Dynamic boost**: GPU continuously monitors: Temperature, Power consumption, Current limits, Reliability voltage guardbands. - Clock algorithm: If all metrics within limits → increase frequency; if any limit approached → reduce frequency. - Boost states: Hundreds of P-state levels → 13–26 MHz steps between states → continuous adjustment every millisecond. **Thermal Throttling Chain** ``` Normal operation → approaching TjMax → slowdown throttle Temps still rising ↓ Heavy throttle (−100 to −500 MHz) Temps still rising ↓ Critical throttle → minimum guaranteed frequency Temps still rising ↓ Emergency shutdown (hardware protection) ``` **Power Throttling** - Power limit (TDP): Set by NVIDIA at factory or adjustable by user (`nvidia-smi -pl `). - Power Brake Slowdown: When actual power > TDP → GPU throttles frequency → power reduces → temperature stabilizes. - AI training: If batch size or sequence length too large → very high memory bandwidth → power spikes → throttle → lower throughput. **Thermal Management Strategies** **1. Cooling System Design** - Data center GPU (A100, H100): Direct liquid cooling mandatory at 700 W TDP. - Cold plate: Copper liquid cold plate bonded to GPU package → water-glycol coolant → 95°C inlet water acceptable. - Air cooling: Limited to ~300 W (dual fan system) → consumer GPUs only. - Immersion cooling: Server submerged in dielectric fluid → highest density, lowest cost at scale. **2. Thermal Paste and TIM (Thermal Interface Material)** - Indium solder: Highest thermal conductivity (80 W/m·K) → used in HPC GPUs (H100). - Liquid metal: 30–70 W/m·K → high performance. - Standard TIM: 5–10 W/m·K → sufficient for lower power GPUs. **3. Power Limit Tuning** - Reduce power limit: `-pl 350` on H100 → reduces peak power by 50 W → reduces thermal load → prevents throttle. - Trade-off: Slightly lower throughput but sustained non-throttling throughput may exceed higher-power throttling throughput. - Optimal point: Usually 80–90% of TDP for sustained ML training. **Monitoring GPU Thermal State** ```bash nvidia-smi -q -d PERFORMANCE # check throttle reasons nvidia-smi dmon -s p # live power monitoring nvidia-smi -q | grep -A 4 Throttle # throttle reason flags ``` - **Throttle reasons**: `HW_SLOWDOWN`, `SW_POWER_CAP`, `THERMAL`, `RELIABILITY`. - `HW_SLOWDOWN` → GPU-detected thermal throttle → increase cooling or reduce load. - `SW_POWER_CAP` → software power limit hit → increase `-pl` or reduce batch size. **Tensor Core Efficiency Under Throttling** - Tensor Core throughput scales linearly with GPU frequency → throttling from 1980 MHz to 1600 MHz = 19% throughput loss. - Memory bandwidth is less affected (HBM frequency independent of GPU core clock in some cases). - For memory-bound workloads (LLM decode): Throttling impact is smaller than for compute-bound training. GPU thermal throttling and power management is **the physical constraint that governs maximum sustained AI computing throughput** — understanding the dynamic interplay between temperature, power, and clock frequency is essential for data center operators who must design cooling systems, for ML engineers who must size batch sizes and sequence lengths to avoid throttle, and for hardware architects who must balance peak performance claims with the practical, sustained throughput that applications actually achieve in production environments.

gpu utilization,optimization

**GPU utilization** measures the percentage of a GPU's computational resources that are **actively being used** at any given moment. In the context of AI and LLM workloads, achieving high GPU utilization is critical because GPUs are extremely expensive resources — every idle cycle is wasted money. **Understanding GPU Utilization Metrics** - **SM Occupancy**: The percentage of **Streaming Multiprocessor** warps that are active. Higher occupancy generally means better utilization of compute cores. - **Compute Utilization**: How much of the GPU's raw **FLOPS** capability is being consumed — measured via tools like `nvidia-smi` or **NVIDIA Nsight**. - **Memory Bandwidth Utilization**: The fraction of available **HBM bandwidth** being used. LLM inference (especially decode) is often **memory-bandwidth bound**, meaning compute utilization may be low even when the GPU is effectively "busy." - **GPU Memory Usage**: The amount of **VRAM** occupied by model weights, KV cache, activations, and framework overhead. **Typical Utilization Patterns** - **Training**: Usually achieves **high utilization** (70–90%+) due to large batch sizes and continuous computation. - **Inference (Prefill)**: Moderate to high utilization — processing many input tokens in parallel is compute-intensive. - **Inference (Decode)**: Often **low compute utilization** (10–30%) because generating one token at a time doesn't provide enough arithmetic to saturate the GPU. This is the main bottleneck. **Improving Utilization** - **Continuous Batching**: Dynamically group multiple inference requests together to increase the effective batch size. - **Quantization**: Reduce precision to process more tokens per memory read. - **Speculative Decoding**: Generate multiple candidate tokens per step to increase arithmetic intensity. - **Right-Sizing**: Match the **GPU type and count** to the model size and expected load — over-provisioning wastes resources, under-provisioning causes queuing. Monitoring GPU utilization in production is essential for **cost optimization** and **capacity planning** in AI infrastructure.

gpu virtualization,mig multi instance,gpu sharing,vgpu,gpu partitioning

**GPU Virtualization and Multi-Instance GPU (MIG)** is the **technology enabling a single physical GPU to be partitioned into multiple isolated instances** — each with dedicated compute resources, memory, and memory bandwidth, allowing multiple users or workloads to share one GPU safely without interference, maximizing GPU utilization in cloud and enterprise environments. **Why GPU Virtualization?** - Many workloads don't need a full GPU: Inference serving, Jupyter notebooks, small training jobs. - Without sharing: A user occupying an A100 at 10% utilization wastes 90% of a $10,000+ GPU. - With MIG: Split one A100 into 7 isolated instances → 7 users, each with guaranteed resources. **NVIDIA MIG (Multi-Instance GPU)** - Available on: A100, A30, H100, H200 GPUs. - Partitions GPU into up to **7 instances** (on A100/H100). - Each instance gets: - Dedicated SM (streaming multiprocessor) slices. - Dedicated memory and memory bandwidth. - Dedicated L2 cache partition. - Fault isolation (one instance's error doesn't crash others). **MIG Partition Profiles (A100 80GB)** | Profile | GPU Memory | SMs | Use Case | |---------|-----------|-----|----------| | 1g.10gb | 10 GB | 14 SMs | Small inference | | 2g.20gb | 20 GB | 28 SMs | Medium inference/training | | 3g.40gb | 40 GB | 42 SMs | Large inference | | 4g.40gb | 40 GB | 56 SMs | Medium training | | 7g.80gb | 80 GB | 98 SMs | Full GPU (no partition) | **Other GPU Sharing Approaches** | Approach | Isolation | Overhead | Flexibility | |----------|----------|---------|------------| | MIG | Hardware-enforced | Near zero | Fixed profiles | | vGPU (NVIDIA GRID) | Driver-level | 5-15% | Time-slicing | | MPS (Multi-Process Service) | Software | Low | Concurrent kernels | | Time-Slicing | Context switching | 10-30% | Any workload | | Kubernetes GPU Sharing | Orchestration | Varies | Pod-level | **vGPU (Virtual GPU)** - NVIDIA GRID/vGPU: Hypervisor-based GPU virtualization. - GPU time-sliced between VMs — each VM sees a virtual GPU. - Used in: VDI (virtual desktops), cloud gaming, VMware/Citrix environments. - Overhead: 5-15% per VM due to context switching. **MPS (Multi-Process Service)** - Allows multiple CUDA processes to share a single GPU simultaneously. - Processes run concurrently (not time-sliced) — better utilization than context switching. - No memory isolation — one process can potentially access another's memory. - Used when: trusted workloads need to share GPU without MIG overhead. **Cloud GPU Sharing** - AWS: `p4d.24xlarge` with 8 A100s, or MIG-backed instances. - GCP: Multi-instance GPU support for A100/H100. - Azure: MIG available on ND-series VMs. GPU virtualization is **essential for economic GPU utilization in data centers** — without partitioning and sharing, the high cost of modern GPU accelerators would be wasted on workloads that use only a fraction of available compute and memory resources.

gpu warp architecture design,sm streaming multiprocessor,cuda core execution,register file gpu,warp scheduler hardware design

**GPU Streaming Multiprocessor (SM) Architecture** is the **fundamental execution unit of GPU chips, containing dozens of CUDA cores, tensor cores, warp schedulers, and hierarchical cache/memory subsystems orchestrated to achieve massive thread parallelism and memory bandwidth.** **CUDA Core and Tensor Core Organization** - **CUDA Cores**: Scalar processing elements executing FP32 (single-precision) or integer operations. Typical SM: 32-128 CUDA cores. Each core contains FP unit, integer ALU, and special function unit (SFU). - **Tensor Cores**: Specialized units performing matrix multiplication (4×4 or 8×8 matrix ops in few cycles). Recent GPUs (Volta+) dedicate substantial area to tensor cores (10-20 cores per SM). - **Special Function Units (SFU)**: Execute transcendental functions (sin, cos, reciprocal), integer operations. Typically 1 SFU per warp (32 threads) limiting throughput for special functions. **Warp Scheduling Hardware** - **Warp Concept**: Group of 32 threads executing in lockstep (SIMD). Modern GPUs issue 2-4 warps per cycle, each to different execution units. - **Warp Scheduler**: Selects ready warps (no stalls) for execution from resident warps (typically 32-64 per SM). Scheduling policies: round-robin, priority-based, or two-level hierarchical. - **Ready Warp Identification**: Tracks register availability, operand readiness, instruction fetch completion. Warp marked "stalled" when waiting for memory, synchronization, or resources. - **Dual-Issue Architecture**: Modern designs issue two independent instructions from same warp or different warps. Enables pipelining and hiding latencies. **Register File Banking and Architecture** - **Register File Size**: 64-256 KB per SM (Ampere: 256 KB). Distributed as 32 banks, one read port per bank per cycle. - **Bank Conflict**: Simultaneous accesses to same register bank by different threads. Causes serialization (pipeline stall) limiting throughput. - **Banking Layout**: Registers allocated sequentially to threads. Thread i's registers in bank (i mod 32). Stride-1 accesses have no conflicts; stride-32 accesses fully serialize. - **Register Optimization**: Compiler allocates registers to minimize bank conflicts. Unroll loops to increase register pressure but improve ILP. Register spilling to local memory expensive (~10x slower). **L1 Cache and Shared Memory Integration** - **L1 Cache**: 32-64 KB per SM. Caches all memory accesses (if enabled). Separate banks from shared memory in Ampere (flexible partitioning). - **Shared Memory**: 48-96 KB fast on-chip memory, explicitly managed by programmer. Bank-conflict free access with properly aligned patterns (sequential access best). - **Write-Through Behavior**: L1 write-through to L2 (no write-back buffering in early GPU architectures). Recent designs: write-back option for reduced memory traffic. **Load-Store Unit and Memory Subsystem** - **Load-Store Capability**: SM can issue multiple load/store instructions per cycle. Coalesced accesses (consecutive threads accessing consecutive memory addresses) merge into single bus transaction. - **Coalescing Efficiency**: 32 consecutive loads (4-byte words) coalesce into one 128-byte transaction. Scattered patterns waste bandwidth. - **Memory Latency Hiding**: 100-500 cycle memory latency hidden by scheduling other ready warps. Occupancy (resident warp count) determines latency hiding capability. **Occupancy and Latency Hiding** - **Occupancy Metric**: Percentage of maximum resident warps actually resident. Higher occupancy better hides memory latency (more warps available to schedule while others wait). - **Limiting Factors**: Register pressure, shared memory allocation per thread, block size constraints determine max occupancy (typically 50-100%). - **Ampere/Hopper Evolution**: Larger register files (256 KB), flexible shared memory partitioning, tensor float 32 (TF32) precision enable higher occupancy while maintaining performance.

gpu warp divergence,branch divergence simt,thread divergence penalty,predication gpu,warp execution efficiency

**GPU Warp Divergence** is the **performance penalty that occurs when threads within the same warp (NVIDIA, 32 threads) or wavefront (AMD, 64 threads) take different execution paths at a branch instruction — forcing the SIMT processor to serialize the divergent paths by executing each branch sequentially while masking inactive threads, potentially halving or worse the effective throughput of divergent code sections**. **How SIMT Execution Creates Divergence** GPU hardware executes one instruction across all threads in a warp simultaneously. When a conditional branch is encountered: - If ALL threads take the same path: no penalty, full throughput. - If SOME threads take the if-path and others the else-path: the hardware first executes the if-path with else-threads masked (inactive), then executes the else-path with if-threads masked. Both paths execute sequentially — the cost is the SUM of both paths, not the MAX. **Divergence Impact** ``` // High divergence — every other thread takes a different path if (threadIdx.x % 2 == 0) { path_A(); // 16 threads active, 16 masked } else { path_B(); // 16 threads active, 16 masked } // Effective utilization: 50% (both paths execute sequentially) ``` ``` // No divergence — all threads in a warp take the same path if (threadIdx.x / 32 == some_condition) { path_A(); // entire warp goes one way } else { path_B(); // different warp goes other way } // Effective utilization: 100% ``` **Mitigation Strategies** - **Data Reorganization**: Sort or bin data so that threads within a warp process similar work (e.g., particles of the same type, pixels in the same region). Coherent data produces coherent branches. - **Thread Reassignment**: Instead of assigning thread-to-data statically, use a work queue where each warp pulls homogeneous work items. - **Predication**: For short divergent code (a few instructions), compilers replace branches with predicated execution — both paths compute, and a select instruction picks the correct result. Eliminates the branch entirely at the cost of executing redundant instructions. - **Warp Specialization**: Assign different warps to different code paths rather than letting a single warp encounter the branch. More warps but each runs at full efficiency. **Nested Divergence** Nested branches compound the problem: a two-level nested if-else can reduce utilization to 25% (4 serial paths with 8 active threads each in a 32-thread warp). Deeply branching code (recursive tree traversal, interpreters) causes severe divergence and should be restructured or moved to the CPU. **Measurement** NVIDIA Nsight Compute reports "warp execution efficiency" — the ratio of active threads to total threads across all executed instructions. Values below 80% indicate significant divergence worth optimizing. **GPU Warp Divergence is the fundamental tension between the GPU's SIMT execution model and data-dependent control flow** — the performance cliff that programmers must understand and design around to achieve the throughput that makes GPU computing worthwhile.

gpu warp divergence,thread divergence cuda,branch divergence penalty,predicated execution gpu,control flow efficiency

**GPU Warp Divergence** is **the performance degradation that occurs when threads within a warp take different execution paths at a branch — forcing the hardware to serialize both paths by masking inactive threads, effectively halving or worse the warp's throughput for each divergent branch**. **Divergence Mechanics:** - **SIMT Execution Model**: all 32 threads in a warp execute the same instruction simultaneously; when a conditional branch evaluates differently across threads, the warp must execute both taken and not-taken paths sequentially - **Active Mask**: hardware maintains a bitmask indicating which threads are active for the current instruction; inactive threads execute the instruction but their results are discarded (no register writeback, no memory store) - **Reconvergence Point**: after both paths complete, the warp reconverges and resumes full-width execution; the compiler inserts synchronization stack entries to track reconvergence points - **Nested Divergence**: divergence within an already-divergent path creates further serialization; worst case is 32 unique paths executed sequentially — reducing warp throughput to 1/32 **Common Divergence Patterns:** - **Thread-ID Conditional**: if(threadIdx.x < N) creates divergence within warps where some threads satisfy the condition and others don't; only the boundary warp(s) actually diverge — warps entirely within or outside the range execute without penalty - **Data-Dependent Branching**: if(data[tid] > threshold) evaluates differently based on input data; highly irregular data causes severe divergence; sorted or clustered data reduces divergence within warps - **Loop Divergence**: while(data[tid]) where each thread iterates a different number of times; the warp continues until the last thread finishes — threads that exit early waste cycles waiting - **Switch Statements**: multi-way branches where different threads take different cases; N unique paths selected requires N serial executions of the warp **Mitigation Strategies:** - **Data Reorganization**: sorting data so adjacent threads process similar values reduces data-dependent divergence; worth the sorting overhead for kernels with many divergent branches - **Predication**: the compiler converts short branches (few instructions) into predicated execution — both paths execute but results are conditionally committed; eliminates branch divergence overhead for branches shorter than the predication threshold (~7 instructions on modern architectures) - **Warp-Level Voting**: __any_sync/__all_sync allow warps to collectively evaluate conditions before branching — if all threads agree, no divergence occurs; the fast path avoids the branch entirely - **Thread Coarsening**: assigning multiple work items per thread and processing them in a loop can convert inter-thread divergence into intra-thread sequential execution — trades parallelism for reduced divergence - **Algorithm Redesign**: replacing conditional logic with arithmetic (branchless code) eliminates divergence entirely; example: min/max using conditional assignment instead of if-else branches **Measurement and Analysis:** - **Branch Efficiency Metric**: Nsight Compute reports branch efficiency as (executed_instructions / (executed_instructions + replay_instructions)) — values below 90% indicate significant divergence - **Active Thread Occupancy**: profilers show average active threads per warp per instruction — ideal is 32; divergent code shows averages below the warp width - **Instruction Replay**: divergent warps replay instructions for each path; profiled as instruction replay overhead — high replay ratios indicate divergence as the primary performance bottleneck GPU warp divergence is **a fundamental SIMT execution constraint that requires parallel programmers to think in terms of warp-uniform control flow — in well-optimized GPU code, divergent branches are either eliminated through branchless techniques, minimized through data reorganization, or confined to boundary warps where their impact is negligible**.

gpu warp divergence,thread divergence,simt divergence,branch divergence gpu,warp efficiency

**GPU Warp Divergence** is the **performance penalty that occurs when threads within the same warp (typically 32 threads executing in lockstep) take different paths at a branch instruction** — forcing the GPU to serialize the divergent paths by executing each branch sequentially and masking inactive threads, wasting execution slots and reducing the effective parallelism that is the GPU's fundamental performance advantage. **How SIMT Execution Works** - GPU executes threads in groups called **warps** (NVIDIA, 32 threads) or **wavefronts** (AMD, 32/64 threads). - All threads in a warp execute the SAME instruction at the SAME time (Single Instruction, Multiple Threads). - No divergence: All 32 threads active → 100% utilization. - With divergence: Only a subset active per branch → utilization drops. **Divergence Example** ```cuda if (threadIdx.x < 16) { // Path A — threads 0-15 execute, 16-31 idle a[threadIdx.x] = compute_A(); } else { // Path B — threads 16-31 execute, 0-15 idle a[threadIdx.x] = compute_B(); } // Both paths reconverge here → all 32 threads active again ``` - Without divergence: 1 pass. With divergence: 2 passes → 50% efficiency. **Cost of Divergence** | Scenario | Active Threads/Warp | Efficiency | |----------|---------------------|------------| | No divergence | 32/32 | 100% | | 2-way branch (50/50) | 16/32 per pass | 50% | | 4-way branch (equal) | 8/32 per pass | 25% | | Worst case (32-way) | 1/32 per pass | 3.1% | **Sources of Divergence** - **Data-dependent branches**: `if (data[tid] > threshold)` — diverges if data varies within warp. - **Thread ID branches**: `if (tid % 4 == 0)` — predictable divergence pattern. - **Loop iteration counts**: `while (data[tid])` — threads exit loop at different times. - **Switch statements**: Multiple paths from single branch → multi-way divergence. **Minimizing Divergence** 1. **Reorganize data**: Sort/partition data so threads in same warp take same path. - Compact: Move "yes" elements together, "no" elements together → separate warps. 2. **Predication over branching**: For short branches, compute both paths and select result. - `result = (condition) ? path_A : path_B;` — no divergence, both computed. 3. **Warp-level primitives**: `__ballot_sync()`, `__shfl_sync()` — collective operations avoid branches. 4. **Algorithm redesign**: Replace branching with arithmetic (branchless min/max, bitwise selection). **Reconvergence** - After divergent section, threads must **reconverge** to resume lockstep execution. - **Stack-based reconvergence** (traditional): Hardware push/pop divergence stack. - **Independent Thread Scheduling** (Volta+): Each thread has own PC → more flexible but reconvergence still matters for performance. GPU warp divergence is **the single most common source of GPU underutilization** — understanding and minimizing divergence through data reorganization, predication, and algorithm design is essential for writing high-performance GPU kernels that achieve the theoretical throughput of the hardware.

gpu warp scheduling divergence,warp execution model cuda,thread divergence penalty,warp scheduler hardware,simt divergence handling

**GPU Warp Scheduling and Divergence** is **the hardware mechanism by which a GPU streaming multiprocessor (SM) selects warps of 32 threads for execution each cycle and handles control-flow divergence when threads within a warp take different branch paths** — understanding warp scheduling is essential for writing high-performance CUDA and GPU compute code because divergence directly reduces throughput by serializing execution paths. **Warp Execution Model:** - **Warp Definition**: a warp is the fundamental scheduling unit on NVIDIA GPUs, consisting of 32 threads that execute in lockstep under the Single Instruction Multiple Thread (SIMT) model - **Instruction Issue**: each cycle the warp scheduler selects an eligible warp and issues one instruction to all 32 threads simultaneously — a single SM typically has 2-4 warp schedulers operating in parallel - **Occupancy**: the ratio of active warps to maximum supported warps per SM — higher occupancy helps hide memory latency by allowing the scheduler to switch between warps while others wait for data - **Eligible Warps**: a warp becomes eligible for scheduling when its next instruction's operands are ready and execution resources are available — stalls occur when no warp is eligible **Thread Divergence Mechanics:** - **Branch Divergence**: when threads in a warp encounter a conditional branch (if/else) and take different paths, the warp must serialize execution — first executing the taken path while masking inactive threads, then executing the not-taken path - **Active Mask**: a 32-bit mask tracks which threads are active for each instruction — masked-off threads don't write results but still consume a scheduling slot - **Divergence Penalty**: in the worst case a warp with 32-way divergence executes at 1/32 throughput — each unique path executes sequentially while 31 threads sit idle - **Reconvergence Point**: after divergent branches complete, threads reconverge at the immediate post-dominator of the branch — the hardware stack tracks reconvergence points automatically **Warp Scheduling Policies:** - **Greedy-Then-Oldest (GTO)**: favors issuing from the same warp until it stalls, then switches to the oldest ready warp — reduces instruction cache pressure and improves data locality - **Loose Round-Robin (LRR)**: cycles through warps in a roughly round-robin fashion — provides fairness but may increase cache thrashing compared to GTO - **Two-Level Scheduling**: partitions warps into fetch groups and applies round-robin between groups while using GTO within each group — balances latency hiding with cache locality - **Criticality-Aware**: prioritizes warps on the critical path of barrier synchronization to reduce overall execution time — prevents stragglers from delaying __syncthreads() barriers **Minimizing Divergence in Practice:** - **Data-Dependent Branching**: reorganize data so that threads within a warp follow the same path — sorting input data by branch condition or using warp-level voting (__ballot_sync) to detect uniform branches - **Predication**: for short branches (few instructions), the compiler replaces branches with predicated instructions that execute both paths but conditionally write results — eliminates serialization overhead - **Warp-Level Primitives**: __shfl_sync, __ballot_sync, and __match_any_sync enable threads to communicate without shared memory, often eliminating branches entirely - **Branch-Free Algorithms**: replace conditional logic with arithmetic (e.g., using min/max instead of if/else) to maintain full warp utilization **Performance Impact and Profiling:** - **Branch Efficiency**: NVIDIA Nsight Compute reports branch efficiency as the ratio of non-divergent branches to total branches — target >90% for compute-bound kernels - **Warp Stall Reasons**: profilers categorize stalls as memory dependency, execution dependency, synchronization, or instruction fetch — guides optimization priority - **Thread Utilization**: average active threads per warp instruction indicates divergence severity — ideal is 32.0, values below 24 suggest significant divergence - **Occupancy vs. Performance**: higher occupancy doesn't always improve performance — sometimes fewer warps with better cache utilization outperform high-occupancy configurations **Modern architectures (Volta and later) introduce independent thread scheduling where each thread has its own program counter, enabling fine-grained interleaving of divergent paths and supporting thread-level synchronization primitives that weren't possible under the older lockstep model.**

gpu warp scheduling execution, simt warp divergence, warp occupancy optimization, gpu thread scheduling, streaming multiprocessor warps

**GPU Warp Scheduling and Execution Model** — GPU architectures organize threads into warps (typically 32 threads) that execute instructions in lockstep using the Single Instruction Multiple Thread (SIMT) model, where warp scheduling directly determines computational throughput. **Warp Fundamentals** — The basic execution unit in GPU computing operates as follows: - **Warp Formation** — thread blocks are divided into warps of 32 consecutive threads, each sharing a single program counter and executing the same instruction simultaneously - **SIMT Execution** — all threads in a warp fetch and execute identical instructions but operate on different data elements, achieving data-level parallelism efficiently - **Warp Context** — each warp maintains its own register state and program counter, enabling rapid context switching between warps without saving or restoring state - **Active Mask** — a per-warp bitmask tracks which threads are currently active, allowing the hardware to manage divergent execution paths transparently **Warp Scheduling Strategies** — The scheduler selects eligible warps for execution each cycle: - **Round-Robin Scheduling** — warps are selected in circular order, providing fair execution time distribution but potentially suboptimal for latency hiding - **Greedy-Then-Oldest (GTO)** — the scheduler continues executing the same warp until it stalls, then switches to the oldest ready warp, improving cache locality - **Two-Level Scheduling** — warps are divided into fetch and pending groups, with only fetch-group warps competing for execution slots to reduce cache thrashing - **Criticality-Aware Scheduling** — warps approaching barrier synchronization points receive priority to minimize idle time at synchronization boundaries **Warp Divergence and Its Impact** — Branch divergence creates significant performance challenges: - **Divergent Branches** — when threads within a warp take different branch paths, both paths must be serialized, with inactive threads masked off during each path's execution - **Reconvergence Points** — hardware identifies the earliest point where divergent paths merge, using a reconvergence stack to restore full warp utilization - **Nested Divergence** — multiple levels of divergent branches compound serialization overhead, potentially reducing effective parallelism to a single thread - **Independent Thread Scheduling** — modern architectures like NVIDIA Volta introduce per-thread program counters, enabling partial warp execution and improved divergence handling **Occupancy and Latency Hiding** — Maximizing warp-level parallelism is essential: - **Occupancy Calculation** — the ratio of active warps to maximum supported warps per streaming multiprocessor determines the potential for latency hiding - **Register Pressure** — excessive per-thread register usage reduces the number of concurrent warps, limiting the scheduler's ability to hide memory latency - **Shared Memory Allocation** — large shared memory allocations per block reduce the number of concurrent blocks and thus active warps on each multiprocessor - **Instruction-Level Parallelism** — even with low occupancy, sufficient ILP within each warp can sustain throughput by keeping functional units busy **Understanding warp scheduling and divergence behavior is essential for writing high-performance GPU kernels, as these mechanisms fundamentally determine how effectively hardware resources are utilized.**

gpu warp scheduling,simt execution,warp divergence

**GPU Warp Scheduling** — the mechanism by which a GPU's streaming multiprocessor (SM) manages and interleaves execution of warps (groups of 32 threads) to hide memory latency. **SIMT Execution** - **SIMT (Single Instruction Multiple Threads)**: All 32 threads in a warp execute the same instruction simultaneously on different data - If threads take different branches → **warp divergence** — some threads are masked off, executed serially - Divergence can halve (or worse) performance **Latency Hiding** - GPU hides memory latency (hundreds of cycles) by switching to another warp - While warp A waits for data, warp B, C, D execute - Need enough active warps to keep the SM busy → **occupancy** **Occupancy** - $Occupancy = \frac{\text{active warps}}{\text{maximum warps per SM}}$ - Limited by: registers per thread, shared memory per block, threads per block - Higher occupancy = better latency hiding (usually) - But: Sometimes lower occupancy with more registers per thread is faster **Warp Scheduling Policies** - **Round-Robin**: Each ready warp gets a turn - **Greedy-Then-Oldest (GTO)**: Execute same warp until it stalls, then switch - **Two-Level**: Group warps into fetch groups **Understanding warp behavior** is essential for writing efficient GPU code — the difference between naive and optimized kernels can be 10-100x.