GPU Optimization Pipeline

Workflow for profiling and optimizing CUDA kernels in libuipc using uipc.profile + uipc.profile.nsight + Nsight Compute CLI (ncu).

For building and installing, see cmake-workflow.

Build type: Use Release for the canonical benchmarks and profiling. Keep build type and flags identical for comparisons; a different configuration is a separate baseline.

Read the owner working agreements and performance evidence policy first.

Agent Rules

• Do NOT read .ncu-rep files — they are binary, for Nsight Compute GUI only.
• Do NOT benchmark or profile with Debug builds — results are meaningless.
• Do NOT optimize kernels in stages that take <5% of total frame time.
• Always read timer_frames.json for the actual timer hierarchy — it is dynamic and varies per scene/run.

Canonical Cross-Domain Benchmarks

Use python scripts/run_benchmark.py list and run <name> for the four versioned scenes in benchmarks/manifest.json. Throughput runs keep UIPC_BENCHMARK_TIMERS=0; collect synchronized Timer diagnostics separately. Use repeated fresh processes and inspect iteration counts as well as wall time. The runner records pre/post runtime and input hashes outside frame timing; both provenance flags must be true for a controlled comparison. Select an isolated in-repository samples worktree with --samples-directory when needed. See the benchmark guide.

CLI Commands

# List available benchmark scenes
python -m uipc.cli.benchmark list

# Run a benchmark (headless, timer-based); --scene accepts multiple names
python -m uipc.cli.benchmark run --scene <scene> [<scene2> ...] --frames <N> --output <dir>

# Profile with Nsight Compute (kernel-level GPU metrics)
python -m uipc.cli.benchmark profile --scene <scene> --frames <N> --output <dir> --count <K> --skip <S> --ncu-set <set>

# Analyze bottlenecks from a benchmark result directory
python -m uipc.cli.benchmark analyze <result_dir> --ncu-csv <csv_path>

# Compare two benchmark result directories
python -m uipc.cli.benchmark compare <before_dir> <after_dir> --output <out_dir>

profile flags:

• --ncu-set default — basic metrics (SM%, occupancy, registers). Use full for duration/memory but requires admin/root.
• --count <K> — profile first K kernel launches.
• --skip <S> — skip initial warmup/init kernels.
• --ncu-path <path> — explicit path to ncu executable. Auto-detected via NCU_PATH env var, PATH, or default NVIDIA install locations.

Note: The CLI profile defaults to --ncu-set default. The Python API nsight.run() defaults to ncu_set='full'. Be explicit to avoid confusion.

Agent Workflow

Step 1: Run Benchmark + Profile

Run both run (for SimulationStats timer data) and profile (for Nsight Compute kernel metrics). Both are needed to make good optimization decisions.

Step 2: Read Both Reports

Benchmark report (report/report.md from run):

• Shows wall-clock time per simulation stage as a hierarchical tree (e.g., Newton Iteration: 45%, Line Search: 30%)
• Tells you which stages dominate total frame time — this is where optimization has the most impact

Nsight Compute report (<scene>_report.md / <scene>_report.json from profile):

• Shows per-kernel GPU metrics (duration, SM%, occupancy, registers)
• Tells you which kernels are inefficient and why (high register usage, low occupancy, memory-bound, etc.)

Step 3: Cross-Reference to Prioritize

A kernel is worth optimizing only if it's both inefficient AND in a hot stage.

1. Read timer_frames.json or report/report.md to find which stages take the most wall-clock time (e.g., Detect DCD Candidates: 40%).
2. Read the ncu report to find which kernels in those stages have bad metrics.
3. Match kernel names to stages using the class/function names (e.g., StacklessBVH::* -> collision detection stage). Use the source_hint field in the JSON report.
4. Ignore kernels with bad metrics in stages that take <5% of frame time — not worth optimizing.
5. Prioritize kernels that are both in a hot stage AND have clear inefficiencies (high registers, low occupancy).

Step 4: Optimize

1. Follow source_hint in the JSON report to find the CUDA source.
2. Search for the kernel's class/method name in that directory.
3. Read the .cu file and optimize the kernel.
4. Rebuild with Release (see cmake-workflow), re-benchmark, and compare.

Adding Finer-Grained Timers in C++

If the benchmark report shows a hot stage but you need more detail to locate the bottleneck within it, add Timer scopes in the C++ source:

#include <uipc/common/timer.h>

void MySystem::do_something()
{
    {
        Timer timer{"MySystem::phase_A"};
        // ... GPU work ...
    }
    {
        Timer timer{"MySystem::phase_B"};
        // ... GPU work ...
    }
}

Keep Timer scopes outside CUDA-graph capture; synchronized diagnostics change wall time and belong in separate runs.

The Timer is scoped — it starts on construction and stops on destruction. Nested timers form a tree. The names appear in timer_frames.json and report/report.md, letting you drill down into which sub-phase of a hot stage is the actual bottleneck before running the more expensive ncu profiling.

Simulation Timer Hierarchy

The timer tree is dynamic — it varies per run depending on which simulation features are active (contact, friction, animation, etc.). The actual tree for any run is stored in timer_frames.json (written by run). Always read that file for the real hierarchy.

Below is a representative example with all features enabled, showing the typical nesting from advance_ipc.cu, global_linear_system.cu, and linear_pcg.cu:

Pipeline                                    engine/
├── Rebuild Scene                           engine/
│   └── Update Diff Parm                    diff_sim/               (conditional)
└── Simulation                              engine/
    ├── Clear External Forces               external_force/         (conditional)
    ├── Step Animation                      animator/               (conditional)
    ├── Compute External Force Accel.       external_force/         (conditional)
    ├── Detect DCD Candidates               collision_detection/    (conditional)
    ├── Newton Iteration                    engine/                 (LOOP)
    │   ├── Detect DCD Candidates           collision_detection/    (iter > 0)
    │   ├── Compute DyTopo Effect           dytopo_effect_system/   (conditional)
    │   │   ├── Assemble Dytopo Effect
    │   │   ├── Convert Dytopo Matrix
    │   │   └── Distribute Dytopo Effect
    │   ├── Solve Global Linear System      linear_system/
    │   │   ├── Build Linear System
    │   │   └── Solve Linear System
    │   │       └── PCG
    │   │           ├── Apply Preconditioner
    │   │           └── SpMV               (per PCG iteration)
    │   └── Line Search                     line_search/
    │       ├── Detect Trajectory Cand.     collision_detection/
    │       ├── Compute Energy              line_search/            (initial E0)
    │       ├── Filter CCD TOI              collision_detection/
    │       ├── Compute CFL Condition       contact_system/
    │       └── Line Search Iteration       line_search/            (LOOP)
    │           ├── Filter Contact Cand.    collision_detection/
    │           └── Compute Energy          line_search/
    └── Update Velocity                     time_integrator/

Parent timer durations include their children. E.g., if Newton Iteration is 80% of frame time, PCG and Line Search durations are already counted inside that 80%.

Identifying Kernels in ncu Output

All backend kernels are named __global__ functions (defined in anonymous namespaces) launched with raw <<<>>> — kernel symbols are directly readable in ncu reports, e.g.:

InfoStacklessBVHSimplexTrajectoryFilter_detect_k1_kernel
FEMLineSearchReporter_step_forward_kernel
StacklessBVH::... (older helpers may still appear as <Class>::<method>)

Naming convention: <OwningClass>_<original_function>[_kN]_kernel (_kN = Nth launch point in the same function). To find the source of a kernel: search the kernel name (or its suffix) in src/backends/cuda/.

Historical lambda-kernel reports are covered in handoff.md. Use current named-kernel symbols for new reports; do not restore removed launch wrappers.

Buffer operations (buffer_fill_kernel<T> and device-memory copies through cuda_tool) are memory operations; prioritize them only when enclosing-stage measurements identify material cost.

Python API

Package layout

uipc.profile            — benchmark runner (timer-based, in-process)
uipc.profile.nsight     — Nsight Compute profiler (kernel-level, ncu subprocess)

The primary input for both is a World. A Scene is also accepted as a convenience shortcut (a temporary Engine + World is created internally).

Setup

from uipc import Scene, Engine, World
from uipc.assets import load

scene = Scene(Scene.default_config())
load('cube_ground', scene)
engine = Engine('cuda', 'my_workspace')
world = World(engine)
world.init(scene)

uipc.profile — benchmark runner

from uipc import profile

# Simple: benchmark 10 frames from the world's current frame
result = profile.run(world, num_frames=10, name='baseline', output_dir='bench')
print(result['summary'])

# Flexible: mix warmup and benchmarking
with profile.session(world, name='baseline', output_dir='bench') as s:
    s.advance(50)    # warmup 50 frames (no stats)
    s.profile(10)    # benchmark 10 frames (collect stats)
print(s.result['summary'])

# Compare two saved benchmark directories
md = profile.compare('bench/baseline', 'bench/optimized', output_dir='comparison')
print(md)

# Load a saved result for later analysis
data = profile.load_result('bench/baseline')

uipc.profile.nsight — Nsight Compute profiler

from uipc.profile import nsight

# Simple: profile 2 frames under ncu
result = nsight.run(world, num_frames=2, name='cube_ground',
                    output_dir='ncu_results', ncu_set='default')

# Flexible: mix warmup and profiling
with nsight.session(world, name='cube_ground') as s:
    s.profile(10)    # profiles from world.frame()
print(s.result)

When a World is passed, nsight automatically calls world.dump() and uses world.recover() in the subprocess for instant state restoration (no replay cost). The engine workspace is obtained via WorldVisitor(world).engine().workspace().

SimulationStats visualization tools

Available on the stats object in benchmark results (result['stats']):

• stats.profiler_heatmap() — sunburst chart of timer breakdown
• stats.system_dependency_graph(workspace) — directed graph of backend system dependencies
• stats.plot(keys, metric='duration') — per-frame line/bar chart
• stats.to_markdown(keys) — Markdown table of per-frame values

Key CUDA Kernel Directories

All relative to src/backends/cuda/:

• finite_element/ — FEM constitutions, gradient/Hessian assembly
• affine_body/ — Affine body dynamics
• linear_system/ — PCG solver, SpMV, preconditioners
• collision_detection/ — BVH, trajectory filtering, DCD
• contact_system/ — Contact forces (IPC barrier)
• global_geometry/ — Vertex management, bounding boxes
• time_integrator/ — Time integration, velocity update
• dytopo_effect_system/ — Dynamic topology effects
• line_search/ — Line search energy evaluation
• external_force/ — External force computation
• animator/ — Animation stepping
• diff_sim/ — Differentiable simulation
• coupling_system/ — Multi-body coupling
• implicit_geometry/ — Implicit geometry representations
• inter_primitive_effect_system/ — Inter-primitive effects
• newton_tolerance/ — Newton convergence tolerance
• engine/ — Core pipeline orchestration
• utils/ — Shared utilities

Optimization Checklist

When the bottleneck report identifies a hot kernel:

1. Memory-bound (high Mem%, low SM%): Coalesced access, shared memory, reduce data movement.
2. Compute-bound (high SM%, low Mem%): Algorithmic improvements, intrinsics.
3. Low occupancy: Reduce register pressure (__launch_bounds__), adjust block size.
4. Many small launches: Measure fusion of named raw-CUDA kernels or CUDA-graph replay; retain explicit stream dependencies.
5. High launch overhead: Reduce host-device sync points.
6. High register count (>64/thread): Measure simpler intermediates or launch bounds. Split heterogeneous work only when kernel and enclosing-stage timings both improve; keep IPC PT/EE/PE/PP Hessian assembly fused under rule 14.

Scenes

Scenes are loaded from HuggingFace: MuGdxy/uipc-assets. Each asset has a scene.py with build_scene(scene). The asset module is at assets/init.py.

Output Structure

Benchmark (run)

<output_dir>/<scene_name>/
  benchmark.json         # metadata (name, wall_time, num_frames, summary)
  timer_frames.json      # per-frame timer tree (JSON) — read this for hierarchy
  workspace_<name>/      # engine workspace (contains systems.json)
  report/
    report.md            # summary with charts — AGENT: read this
    profiler_heatmap.svg  # sunburst chart
    system_deps.svg       # system dependency graph (if workspace exists)
    *.svg                 # per-timer charts

Nsight Compute (profile)

<output_dir>/
  <scene>_report.md      # AGENT: read this — kernel hotspot table
  <scene>_report.json    # AGENT: read this — structured metrics + source_hints
  <scene>.ncu-rep        # binary — do NOT read
  <scene>.csv            # raw CSV (already parsed into reports)

JSON report per kernel:

{
  "name": "InfoStacklessBVHSimplexTrajectoryFilter_detect_k1_kernel",
  "launches": 3,
  "total_duration_ns": 123456.0,
  "registers_per_thread": 40,
  "avg_sm_pct": 100.0,
  "avg_occupancy_pct": 48.0,
  "source_hint": "src/backends/cuda/collision_detection/",
  "optimization_hints": ["Low occupancy - consider reducing registers"],
  "full_names": ["void InfoStacklessBVHSimplexTrajectoryFilter_detect_k1_kernel(...)"]
}