profile-analyzer

---
name: profile-analyzer
description: >
  Use this skill to understand why a Sirius query is slow, identify GPU bottlenecks, or detect
  performance regressions. Generates reports with kernel occupancy, memory bandwidth, operator
  attribution, and cross-run comparisons. Trigger when the user mentions profiling, nsys, GPU
  utilization, kernel analysis, performance reports, or wants to compare query timings across runs.
  This skill focuses on measurement and reporting — for mapping hotspots to source code fixes,
  use optimization-advisor instead.
---

# Sirius nsys Profile Analyzer

You are analyzing GPU performance profiles for Sirius, a GPU-accelerated SQL query engine built on DuckDB. The profiles come from NVIDIA Nsight Systems (nsys) and are stored as SQLite databases.

## Profiling Overhead Warning

**nsys profiling adds measurable overhead to query execution times.** Timings captured during a profiled run (cold/hot in `summary.json`) are inflated and should NOT be used to determine whether an optimization actually improved performance. Instead:

1. **Profiled runs** → Use for GPU analysis (kernels, operators, occupancy, memory, bottlenecks)
2. **Non-profiled runs** → Use for accurate performance timing (cold/hot comparisons, regression detection)

Always run both when comparing performance across code changes.

## Workflows

There are four workflows — choose based on what the user wants:

### Workflow A: Full Performance Analysis (recommended)

This is the complete workflow: profile for GPU analysis, then run without profiling for accurate timings. Both runs should use the same queries, scale factor, and iteration count.

**Step 1: Profiled run** (for GPU analysis data)
```bash
# Profile TPC-H queries via nsys_report.sh (delegates to performance_test.py --precmd nsys)
bash test/tpch_performance/nsys_report.sh --sf <scale_factor> [query_numbers...]
```

`nsys_report.sh` orchestrates the profiling under the hood: it calls `performance_test.py --precmd nsys` (the same runner the `benchmark` skill uses) to produce one `.nsys-rep` + `.sqlite` per query, then runs `nsys_analyze.sh` to emit `report.md` and `summary.json`.

**Step 2: Non-profiled timing run** (for accurate cold/hot times)
```bash
export SIRIUS_CONFIG_FILE=<path_to_config>

# Sirius-only timing
pixi run python test/tpch_performance/performance_test.py \
    --input <parquet_dir> --scale-factor <SF> --engine gpu --iterations <N> [--queries 1,3,6-10]

# DuckDB vs Sirius timing + result validation in one shot
pixi run python test/tpch_performance/performance_test.py \
    --input <parquet_dir> --scale-factor <SF> --engine both --iterations <N> --validation
```

The non-profiled run produces a long-format `<bench>/csv/runtimes.csv` (`engine,query,iteration,runtime_s`) and per-query `result.txt`. `--validation` additionally byte-compares the saved Sirius vs DuckDB results (with a small `abs_tol` on floats).

**Data source (parquet or duckdb).** `--input` accepts either a parquet directory (`--data-source parquet`, the default) or a single `.duckdb` file (`--data-source duckdb`, exercising the GPU-native `seq_scan`); `--pin gpu|host` works for both. Profiling works for both sources through the **same** orchestrator — `nsys_report.sh --data-source duckdb --duckdb-file <file>.duckdb` (or just `--sf <SF>`, which defaults to `test_datasets/tpch_sf<SF>.duckdb`). See `test/tpch_performance/CLAUDE.md`.

**When comparing across runs:**
- Use the non-profiled timings to determine if performance actually improved or regressed
- Use the profiled data to understand *why* performance changed (kernel times, operator attribution, occupancy, memory patterns)
- The profiled `summary.json` timings are useful for relative comparisons within the same profiled run (e.g., which query is slowest) but not across runs with different code

**Full options for the profiled run:**
```bash
bash test/tpch_performance/nsys_report.sh \
    --sf <scale_factor> \
    --parquet-dir <path_to_parquet_dir> \
    --data-source parquet|duckdb \
    --duckdb-file <path_to.duckdb> \
    --output-dir ./reports \
    --label <custom_name> \
    --iterations 4 \
    --query-timeout 120 \
    --compare <baseline_report_dir> \
    [query_numbers...]
```

`--parquet-dir` defaults to `$PROJECT_DIR/test_datasets/tpch_parquet_sf<SF>`; with `--data-source duckdb`, `--duckdb-file` defaults to `$PROJECT_DIR/test_datasets/tpch_sf<SF>.duckdb`.

**Output directory structure (profiled):**
```
reports/<label>_<YYYYMMDD_HHMMSS>/
  report.md        - Human-readable analysis with all metrics
  summary.json     - Machine-readable per-query metrics (profiled timings — use for analysis, not perf comparison)
  metadata.json    - Hardware, git commit, config, driver version
  comparison.md    - (if --compare used) Regression/improvement analysis
  profiles/        - Flat raw artifacts (one set per query, flattened from the
                     per-query subdirs produced by performance_test.py --precmd nsys)
    q1.sqlite, q1.nsys-rep, q1_timings.csv, q1_result.txt, ...
  profiles_tree/   - Original per-query layout from performance_test.py:
    sirius/q<N>/{nsys.nsys-rep, nsys.sqlite, nsys.sql, timings.csv, log_dir/}
```

**Output from non-profiled run** (`performance_test.py` benchmark layout): a timestamped `<benchmark_dir>/` with `metadata.json` (including `scale_factor`), effective `queries/q<N>.sql`, `csv/runtimes.csv` (`engine,query,iteration,runtime_s`), per-query `<engine>/q<N>/result.txt`, and per-query `sirius/q<N>/sirius.log`. See `test/tpch_performance/CLAUDE.md` for the full layout.

### Workflow B: Generate Report from Existing Profiles

```bash
bash test/tpch_performance/nsys_report.sh --profile-dir <path_to_profiles>
```

### Workflow C: Quick Analysis (no archival)

Use `nsys_analyze.sh` directly for quick, one-off analysis without creating a report directory.

```bash
bash test/tpch_performance/nsys_analyze.sh <path_to_profiles_or_file> [query_numbers...]
```

### Workflow D: Compare Two Existing Reports

```bash
bash test/tpch_performance/nsys_compare.sh <baseline_report_dir> <current_report_dir> [--threshold PCT]
```

Default threshold is 10%. Values beyond the threshold are flagged:
- **REGRESSION**: current is >threshold% slower
- **IMPROVED**: current is >threshold% faster
- **FIXED**: query failed in baseline but passes now
- **BROKEN**: query passed in baseline but fails now

**Important**: The timings in `summary.json` (used by `nsys_compare.sh`) are from profiled runs and include nsys overhead. These comparisons are useful for spotting large changes but should be validated with non-profiled timing runs before concluding a real regression or improvement exists. For definitive performance comparison, compare the `csv/runtimes.csv` files from a non-profiled `performance_test.py` run.

## Before Running

- **Ask the user** for any paths you don't know. Do NOT assume paths.
- For profiling (Workflow A without `--profile-dir`), the Sirius config must be set:
  ```bash
  export SIRIUS_CONFIG_FILE=<path_to_config>
  ```
- For full options on the profiled run (parquet dir override, custom output dir, label, comparison baseline), see `bash test/tpch_performance/nsys_report.sh --help`.

## Analysis Sections

The report contains these sections per query:

All analysis is **scoped to the query execution window** — the time span from the first Sirius operator start to the last operator end. Init overhead (CUDA context creation, cudaHostAlloc, cuFile init) and cleanup (cudaFreeHost, pool destruction) are excluded from the main metrics and shown separately.

| Section | What it Shows |
|---------|---------------|
| **Execution Time Breakdown** | Trace duration vs query execution time vs init vs cleanup |
| **GPU Hardware** | GPU model, SM count, VRAM, compute capability |
| **NVTX Domain Summary** | Time breakdown across software layers (Sirius, libkvikio, libcudf, CCCL, cuFile) |
| **Sirius Physical Operators** | Per-operator call counts, wall times, percentages |
| **Top GPU Kernels** | Hottest kernels by total GPU time |
| **Kernel Occupancy Estimation** | Theoretical SM occupancy per kernel, limiting factor (registers/shared_mem/warps) |
| **Register Spill Analysis** | Kernels using local memory (register spilling to slow memory) |
| **GPU Kernel Time Summary** | Aggregate kernel stats, stream count, device count |
| **GPU Utilization Overview** | Kernel time as % of query execution time (not total trace) |
| **Memory Transfer Breakdown** | H2D/D2H/D2D with Pageable vs Pinned src/dst, bandwidth in GB/s |
| **CUDA Runtime API Hotspots** | Slowest CUDA API calls *during query execution only* |
| **Host Memory Allocation During Query** | Only alloc calls during runtime (init allocs excluded) |
| **Device Memory Allocation — cudaMalloc Counts** | Total `cudaMalloc` calls split init / during-query / cleanup — during-query calls bypass the RMM pool |
| **Init/Cleanup Overhead** | What was excluded — cudaHostAlloc, cudaFreeHost, context creation, etc. |
| **GPU Kernel Attribution** | Maps GPU kernel time back to Sirius operators via correlation chain |
| **Top Kernels per Operator** | Which specific kernels each operator launches |
| **GPU Stream Utilization** | Per-stream busy% = kernel_time / stream_active_span |
| **Synchronization Analysis** | GPU sync wait times by type |
| **NVTX Operations by Domain** | Top operations per software layer |
| **Cross-Query Comparison** | (multi-file only) Side-by-side query overview |

## Architecture Context

### Sirius Execution Model
Sirius intercepts DuckDB query plans and offloads them to GPU via cuDF:
1. **DuckDB** parses SQL and creates a logical plan
2. **Sirius** converts it to a physical GPU plan with operators like `sirius_physical_table_scan`, `sirius_physical_hash_join`, etc.
3. **cuDF** (libcudf) provides GPU-accelerated DataFrame primitives
4. **CCCL** (CUB/Thrust) provides GPU algorithm primitives underneath cuDF
5. **CUDA kernels** execute on the GPU

### NVTX Domain Hierarchy

Domain numbers are not registered by Sirius — they're discovered from each nsys profile at analysis time (`nsys_analyze.sh` reads `NVTX_EVENTS` and groups by `domainId`). The mapping below is typical; what shows up depends on which libraries emitted NVTX ranges in your particular run.

- **Domain 0 (Sirius)**: Physical operator execution (`sirius_physical_*::execute/sink`)
- **Domain 1 (libkvikio)**: GPU-Direct Storage I/O operations (`posix_host_read`, `task`)
- **Domain 2 (cuFile)**: cuFile handle management
- **Domain 3 (libcudf)**: cuDF operations (`aggregate`, `binary_operation`, `materialize_all_columns`, etc.)
- **Domain 4 (CCCL)**: CUB/Thrust primitives (`DeviceFor::ForEachN`, `thrust::transform`, etc.)

### Key Relationships
- **Operator -> Kernel attribution**: Correlates via `CUPTI_ACTIVITY_KIND_RUNTIME.correlationId` (links runtime API calls to kernels; runtime call timestamps fall within NVTX operator ranges)
- **Cold vs Hot**: First iteration is "cold" (I/O, JIT). Subsequent iterations are "hot" (cached).
- **Multi-stream**: Sirius uses a dynamic CUDA stream pool sized by the `pipeline.num_threads` setting in its YAML config (constructed in `src/pipeline/gpu_pipeline_executor.cpp`). The actual count depends on your config.

### Sirius Physical Operators
| Operator | Purpose |
|----------|---------|
| `table_scan` | Read input via cuDF — parquet (`read_parquet`) or DuckDB-native tables (`seq_scan`, `src/op/scan/duckdb_native_gpu_ingestible.cpp`) |
| `projection` | Evaluate column expressions |
| `filter` | Row filtering |
| `hash_join` | Hash-based join |
| `grouped_aggregate` | Group-by aggregation |
| `grouped_aggregate_merge` | Merge partial aggregates across partitions |
| `partition` | Data partitioning |
| `order` | ORDER BY |
| `merge_sort` | Merge sorted partitions |
| `sort_partition` / `sort_sample` | Sort-based repartitioning |
| `top_n` / `top_n_merge` | LIMIT processing |
| `concat` | Concatenation |
| `materialized_collector` / `result_collector` | Final result materialization |

## Interpretation Guide

### Occupancy
- **100%**: Maximum warps per SM. Block size and register usage fit perfectly.
- **50-100%**: Generally acceptable. Check if the limiter can be relaxed.
- **Below 50%**: Potential bottleneck. Check the `limiter` column:
  - `registers`: Kernel uses too many registers per thread. Compiler flag `--maxrregcount` could help, or algorithmic changes to reduce register pressure.
  - `shared_mem`: Shared memory per block limits active blocks. The `shmem_b` column shows the *driver-allocated* amount (aligned to SM partition granularity, often 16KB/32KB/64KB/102KB). Reducing shared memory usage or using dynamic allocation may help.
  - `warps`: Block is too large relative to SM warp capacity.
  - `hw_limit`: Hit the max blocks per SM limit (e.g. 24 on Ada Lovelace). `nsys_analyze.sh` reads `maxBlocksPerSm` from the profile, so the comparison auto-adapts per GPU.
- **Note**: Low occupancy doesn't always mean poor performance — compute-bound kernels can achieve peak throughput at low occupancy if they have high ILP (instruction-level parallelism).

### Bandwidth
- **Pinned H2D**: Expect 12-14 GB/s on PCIe 4.0 x16. Below 8 GB/s suggests contention or small transfers.
- **Pageable H2D**: Much slower (~0.05 GB/s). Indicates un-pinned host allocations — a major performance issue if significant data goes through this path.
- **D2D**: Internal GPU memory shuffles. Should achieve 100-400 GB/s depending on transfer patterns.
- **D2H**: Usually small volumes (query results). Bandwidth similar to H2D.

### GPU Utilization
- `kernel_pct_of_query`: Kernel time / query execution span (excludes init/cleanup). Values of 40-60% are typical — the remainder is CPU orchestration, sync waits, and memcpy. Below 30% suggests the GPU is starved.
- `kernel_pct_of_ops`: Kernel time / Sirius operator time. Values >100% indicate GPU kernels overlap with CPU operator orchestration (normal with async execution). Low values suggest CPU-side bottlenecks within operators.

### Host Memory Overhead
- The analysis separates init-time allocations from query-runtime allocations.
- `cudaHostAlloc` during init (10+ seconds) is a one-time cost — it's shown in the Init/Cleanup Overhead section.
- If `cudaHostAlloc` appears in the "During Query Execution" section, that's a performance issue — synchronous allocation during active queries stalls the pipeline.
- `cudaStreamSynchronize` is typically the dominant cost during query execution — it represents time the CPU waits for GPU work to complete.

### Device Allocation (cudaMalloc)
- The **Device Memory Allocation — cudaMalloc Counts** section reports the total number of `cudaMalloc` calls in the trace, split into `init` / `during_query` / `cleanup` (`cudaMallocHost` is excluded — it's a host pinned alloc, covered by the host section).
- **`during_query > 0` is the red flag.** Sirius allocates GPU memory through the RMM pool; a raw `cudaMalloc` on the hot path means something bypassed the pool, and each one forces a device-wide sync that stalls the pipeline. This is a frequent source of unintended performance degradation when a new code path allocates outside the pool.
- On a warm run, expect `during_query = 0`. A nonzero count points at the operator that introduced it — cross-reference the **GPU Kernel Attribution** section to find whose operator window the call falls in. `nsys_compare.sh` flags any per-query increase (0 → N included), so it catches an allocation newly introduced between two runs.
- `init` allocations (one-time pool priming at startup) are normal and excluded from the query window.

### Register Spill
- `local_bytes_per_thread > 0` means the kernel exceeded the register file and is spilling to local memory (which actually resides in global/L2 memory — much slower).
- This is a red flag for performance-critical kernels. Solutions: reduce register usage, simplify kernel logic, or use `__launch_bounds__` to guide the compiler.

## What to Look For

1. **Where is query time going?** Compare operator time vs query execution span. Large gaps = sync waits, memcpy, or CPU orchestration overhead.
2. **GPU utilization**: Low kernel_pct_of_query (<30%) = GPU underutilized during query execution (CPU-bound, sync-bound, or I/O-bound).
3. **Occupancy hotspots**: Kernels with <50% occupancy that consume significant GPU time.
4. **Memory bandwidth**: Pageable transfers are orders of magnitude slower than pinned.
5. **Host allocation cost**: cudaHostAlloc often dominates CUDA API time.
6. **Cold vs Hot delta**: Large gap = I/O/JIT overhead. Small gap = compute-dominated. Use non-profiled timings for accurate cold/hot comparison.
7. **Stream utilization**: Low busy% with many streams = fine-grained parallelism. High busy% on few streams = load imbalance.
8. **Register spill**: Any kernel with local_bytes_per_thread > 0 is spilling.
9. **Operator attribution**: Which operators consume the most GPU time? Focus optimization here.
10. **OOM failures**: Queries failing with `std::bad_alloc` need memory optimization.
11. **Profiling vs actual performance**: Always validate profiled timing changes with non-profiled runs. nsys overhead can mask or exaggerate real performance differences.
12. **cudaMalloc during query**: the cudaMalloc count's `during_query` value should be 0 on warm runs. Any raw `cudaMalloc` inside the query window bypasses the RMM pool and forces a device-wide sync — a common source of unintended degradation. Use `nsys_compare.sh` to catch a count that newly appears (0 → N) between runs.

## Output Format

Always present findings in a structured way:
- Start with a high-level summary (pass/fail, total times, report location)
- Identify the top 3-5 bottlenecks
- For each bottleneck, explain what it means and potential causes
- Compare cold vs hot when relevant — clearly label whether timings are from profiled or non-profiled runs
- When comparing reports, highlight regressions and improvements but note that profiled timings include nsys overhead — recommend validating with non-profiled runs if not already done
- When analyzing a single query deeply, walk through the execution timeline: I/O -> operators -> kernels -> output
- When presenting performance conclusions, always distinguish between profiled timings (for analysis) and non-profiled timings (for actual performance measurement)

More agent context in sirius-db/sirius

17 other files this repository gives its agents.