Profiling & Debugging Like a Pro

Profiling & Debugging Like a Pro

Optimization without profiling is speculation. Top kernel engineers diagnose bottlenecks systematically using the NVIDIA toolchain:

The Professional Diagnostic Stack

1. compute-sanitizer: Traps out-of-bounds memory accesses, misaligned memory addresses, uninitialized shared memory, and warp synchronization race conditions.
2. NVIDIA Nsight Systems (NSYS): Timeline visualization across CPU threads, CUDA streams, memory transfers, and kernel durations.
3. NVIDIA Nsight Compute (NCU): Deep architectural profiling inside individual kernels—analyzing warp stall reasons, SM issue pipe utilization, memory chart traffic, and register spills.

Debugging Sabotaged Kernels: 3 Classic Incident Scenarios

• Incident 1: Silent Race Condition in Shared Memory: Threads write to s_data[threadIdx.x] and read from s_data[threadIdx.x + 1] without an intervening __syncthreads(). Detection: compute-sanitizer --tool racecheck ./my_kernel
• Incident 2: Shared Memory Bank Conflicts: All 32 threads in a warp access different rows in the same shared memory bank simultaneously. Detection: NCU metric l1tex__data_bank_conflicts_pipe_lsu_mem_shared.sum. Resolution: Stride padding (__shared__ float smem[32][33]).
• Incident 3: Catastrophic Register Spilling: Kernel uses 80 registers per thread when the hardware limit per thread at target occupancy is 64. Excessive variables spill to local DRAM memory, increasing latency 100x. Detection: Compiler flag --ptxas-options=-v reporting spill stores / spill loads. Resolution: Kernel launch bound annotation __launch_bounds__(256, 2).