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)

# Profile TPC-H queries via nsys_report.sh (delegates to performance_test.py --mode nsys-profile)
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 --mode nsys-profile (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)

export SIRIUS_CONFIG_FILE=<path_to_config>

# Sirius-only timing
pixi run python test/tpch_performance/performance_test.py \
    --input <parquet_dir> --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> --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 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 --mode nsys-profile)
    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, 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 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 test/tpch_performance/nsys_analyze.sh <path_to_profiles_or_file> [query_numbers...]

Workflow D: Compare Two Existing Reports

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:

  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)