gpu-optimization

# 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](mdc:.cursor/skills/cmake-workflow/SKILL.md). **Build type**: Always use `Release` for benchmarking and profiling. `Debug` is too slow; `RelWithDebInfo` adds debug info overhead that skews results. ## 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. ## CLI Commands ```bash # 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.) | File | Source | What it tells you | |------|--------|-------------------| | `report/report.md` | `run` | Wall-clock time breakdown by stage — **what matters most** | | `timer_frames.json` | `run` | Per-frame timer tree as JSON — **source of truth for hierarchy** | | `<scene>_report.md` | `profile` | Kernel GPU metrics — **what's inefficient** | | `<scene>_report.json` | `profile` | Structured kernel data with `source_hint` and `optimization_hints` | | `<scene>.ncu-rep` | `profile` | Binary — do NOT read, for Nsight Compute GUI only | ### 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](mdc:.cursor/skills/cmake-workflow/SKILL.md)), 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: ```cpp #include <uipc/common/timer.h> void MySystem::do_something() { { Timer timer{"MySystem::phase_A"}; // ... GPU work ... } { Timer timer{"MySystem::phase_B"}; // ... GPU work ... } } ``` 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 [sim_engine_do_advance.cu](mdc:src/backends/cuda/engine/sim_engine_do_advance.cu), [global_linear_system.cu](mdc:src/backends/cuda/linear_system/global_linear_system.cu), and [linear_pcg.cu](mdc:src/backends/cuda/linear_system/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/`. Legacy note: before the named-kernel migration, kernels were launched via `muda::ParallelFor().apply(N, lambda)` and appeared as `parallel_for_kernel<Class::method()::lambda>` (with `#N` for the Nth lambda in a function). The `_shorten_kernel_name()` function in [nsight.py](mdc:python/src/uipc/profile/nsight.py) was written for that scheme; with named kernels it passes the readable name through unchanged. **Buffer operations** (`buffer_fill_kernel<T>`, `buffer_copy_kernel<T>` in `cuda_tool`) are memory operations, usually not optimization targets. ## 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 ```python 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 ```python 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 ```python 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**: Kernel fusion or batching via `muda::ParallelFor`. 5. **High launch overhead**: Reduce host-device sync points. 6. **High register count** (>64/thread): Simplify kernel logic, split into multiple passes, use `__launch_bounds__(blockSize, minBlocksPerSM)`. ## Scenes Scenes are loaded from [HuggingFace: MuGdxy/uipc-assets](https://huggingface.co/datasets/MuGdxy/uipc-assets). Each asset has a `scene.py` with `build_scene(scene)`. The asset module is at [assets/__init__.py](mdc:python/src/uipc/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: ```json { "name": "StacklessBVH::calcExtNodeSplitMetrics", "launches": 3, "total_duration_ns": 123456.0, "registers_per_thread": 40, "avg_sm_pct": 100.0,