cuda-profiling

Files (2)

• LICENSE.source
• SKILL.md

SKILL.md (first lines)

# CUDA Profiling

## Purpose

Guide agents through profiling CUDA applications with Nsight Systems (timeline-level) and Nsight Compute (kernel-level metrics), using the NCU CLI for automated metric collection, interpreting roofline models, and diagnosing whether kernels are memory-bound or compute-bound.

## When to Use

- A CUDA kernel is slower than expected and you need bottleneck identification
- Comparing kernel variants (tiling strategies, block sizes)
- Building CI performance regression checks with `ncu` metrics
- Correlating CPU and GPU activity in multi-stream pipelines
- Annotating application phases with NVTX for timeline visibility
- Interpreting occupancy, memory throughput, and SM utilization metrics

## Workflow

### 1. Choose profiling tool

```
What do you need?
├── System-wide timeline (CPU+GPU+CUDA API) → Nsight Systems (nsys)
├── Per-kernel deep metrics (occupancy, memory) → Nsight Compute (ncu)
└── Quick metric from CLI in CI → ncu --metrics ...
```

### 2. Nsight Systems — timeline profiling

```bash
# Profile entire application
nsys profile --trace=cuda,nvtx,osrt --output=report ./my_cuda_app

# Open report
nsys-ui report.nsys-rep

# CLI summary
nsys stats report.nsys-rep
```

What to look for in the timeline:
- Gaps between kernel launches (CPU bottleneck or sync points)
- `cudaDeviceSynchronize` stalls
- Overlap between H2D copies and kernel execution across streams
- CUDA API call overhead

```bash
# Capture with CUDA graph info
nsys profile --capture-range=cudaProfilerApi ./my_cuda_app
```

### 3. NVTX range annotations

```cpp
#include <nvtx3/nvToolsExt.h>

void pipeline(void) {
    nvtxRangePushA("H2D copy");
    cudaMemcpyAsync(d_in, h_in, size, cudaMemcpyHostToDevice, stream);
    nvtxRangePop();

    nvtxRangePushA("kernel");
    my_kernel<<<grid, block, 0, stream>>>(d_in, d_out, n);
    nvtxRangePop();

    nvtxRangePushA("D2H copy");
    cudaMemcpyAsync(h_out, d_out, size, cudaMemcpyDeviceToHost, stream);
    nvtxRangePop();
}
```

Compile with `-lnvToolsExt` or link nvtx3 header-only. Ranges appear as colored bands in Nsight Systems.

### 4. Nsight Compute — kernel analysis

```bash
# Profile all kernels, save report
ncu -o kernel_report ./my_cuda_app

# Profile specific kernel by name
ncu --kernel-name regex:matmul_tiled ./my_cuda_app

# Launch UI
ncu-ui kernel_report.ncu-rep
```

Key sections in NCU report:
- **Speed of Light**: SM throughput vs memory throughput vs peak
- **Occupancy**: Active warps vs hardware limit
- **Memory Workload Analysis**: L1/L2 hit rates, coalescing efficiency
- **Warp State Statistics**: Stall reasons (memory, barrier, dispatch)

### 5. NCU CLI metrics

```bash
# Essential metrics set
ncu --metrics \
  sm__throughput.avg.pct_of_peak_sustained_elapsed,\
  dram__throughput.avg.pct_of_peak_sustained_elapsed,\
  sm__warps_active.avg.pct_of_peak_sustained_active,\
  l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum,\
  smsp__sass_thread_inst_executed_op_ffma_pred_on.sum \
  ./my_cuda_app

# CSV export for CI
ncu --csv --metrics dram__bytes_read.sum,dram__bytes_write.sum ./my_cuda_app

# Set kernel replay mode for accurate counters
ncu --kernel-replay-mode application ./my_cuda_app
```

### 6. Memory-bound vs compute-bound diagnosis

```
Roofline interpretation
├── dram__throughput near peak AND sm__throughput low → memory-bound
│   └── Fix: coalescing, shared mem tiling, reduce traffic
├── sm__throughput near peak AND dram low → compute-bound
│   └── Fix: tensor cores, loop unrolling, ILP
└── Both low → launch config, occupancy, or sync overhead

From OzcanOrhanDemirci/Mode_Code, Apache-2.0.

More from OzcanOrhanDemirci/Mode_Code