# CelerisLab Performance Analysis Report

Date: 2026-05-31
GPU: Tesla V100-SXM2-16GB (CUDA 12.4)
Test grid: 384x192 D2Q9, MRT, double_buffer, FP32
Tools: cProfile (Python), Nsight Systems (GPU timeline), Nsight Compute (kernel metrics)

## 1. Performance Tree (384x192, MRT, 1 cylinder + 1 sensor)

### Per-step breakdown (batch=100, 100 steps per launch group)

```
┌──────────────────────────────────────────────────────────────┐
│                    One LBM Step (~100 steps, ~7.8ms total)   │
├──────────────────────────────────────────────────────────────┤
│                                                              │
│  Python Layer (CPU)                   ~40 μs/step (0.5%)    │
│  ├── _launch_curved()                 ~6 μs                  │
│  ├── step_fn() (cuLaunchKernel)       ~22 μs                 │
│  ├── _launch_sensor()                 ~5 μs                  │
│  └── loop overhead (Python for _ in range) ~5 μs             │
│                                                              │
│  GPU Kernel Layer (V100)              ~13 μs/step (99.5%)   │
│  ├── CurvedBoundaryKernel             ~4.0 μs  (30%)         │
│  │   grid=(2,1,1) block=256  regs=31  288 threads total     │
│  │   → latency-bound (too few threads to fill SM)            │
│  ├── OneStep                           ~5.9 μs  (44%)        │
│  │   grid=(2,192,1) block=256 regs=39  73,728 threads       │
│  │   → mixed compute+memory (53% DRAM bw, 22% SM util)       │
│  └── SensorKernel                      ~3.5 μs  (26%)        │
│      grid=(2,1,1) block=256 regs=21   305 threads total      │
│      → latency-bound                                          │
│                                                              │
│  Memory Transfers (negligible, <1 μs/step)                   │
│  ├── set_body():   72 bytes H2D (action upload)              │
│  └── read_force(): 60 bytes D2H (obs download, async)        │
└──────────────────────────────────────────────────────────────┘
```

### Wall-clock contributions (2000 steps, batch=100)

| Layer | Total time | Per step |
|-------|-----------|----------|
| Python (cProfile) | 0.901s (warmup+measured, 70 batches) | ~12.9ms/100-step batch |
| GPU kernel (Nsight) | 0.071s (5300 kernel invocations) | 13.4 μs |
| Batch NVTX range (nsys) | 7.8ms for batch=100 | 78 μs/step (incl. Python) |

**Key insight:** The per-step GPU kernel time is only ~13 μs. The nsys batch timing (78 μs/step) includes Python overhead, kernel launch overhead, and stream synchronization.

## 2. Python Layer Details (from cProfile, T4)

| Function | Cumulative time | Calls | Per call | % of run time |
|----------|----------------|-------|----------|---------------|
| Simulation.__init__ | 1.125s | 1 | 1.125s | 62% (one-time) |
| LBMStepper.step | 0.901s | 70 | 12.9ms | 50% of step time |
| pycuda function_call | 0.769s | 21001 | 36.6 μs | 42% of step time |
| _launch_sensor (inside step) | 0.409s | 7000 | 58 μs | — |
| _launch_curved (inside step) | 0.311s | 7000 | 44 μs | — |
| pycuda _build_arg_buf | 0.357s | 21001 | 17 μs | 20% of step time |
| manager.sync_to_gpu | 0.017s | 1 | — | one-time (init) |
| set_body | <0.001s | 20 | <50 μs | negligible |
| read_force | <0.001s | 20 | <50 μs | negligible |

**Conclusion:** The Python layer is extremely thin. The dominant cost in `stepper.step()` is `pycuda.driver.function_call` (~37 μs per kernel launch) which is the `cuLaunchKernel` overhead from the Python-CUDA bridge. This is an unavoidable cost of using pycuda.

## 3. GPU Kernel Details (from Nsight Systems)

### OneStep (MRT, 384x192, 5300 invocations)

| Metric | Value |
|--------|-------|
| Grid | (2, 192, 1) |
| Block | (256, 1, 1) |
| Registers per thread | 39 |
| Min duration | 5.57 μs |
| **p50 duration** | **5.86 μs** |
| p90 duration | 6.11 μs |
| Max duration | 7.65 μs |
| Total (5300 steps) | 31.2 ms |

### CurvedBoundaryKernel (1 cylinder at center, 5300 invocations)

| Metric | Value |
|--------|-------|
| Grid | (2, 1, 1) |
| Block | (256, 1, 1) |
| Registers per thread | 31 |
| Total threads | 288 active out of 512 launched (56% utilization) |
| **p50 duration** | **4.00 μs** |
| Total (5300 steps) | 21.2 ms |

### SensorKernel (1 sensor, 5300 invocations)

| Metric | Value |
|--------|-------|
| Grid | (2, 1, 1) |
| Block | (256, 1, 1) |
| Registers per thread | 21 |
| **p50 duration** | **3.55 μs** |
| Total (5300 steps) | 18.9 ms |

### Batch Scaling (from Nsight NVTX ranges)

| Batch size | Total time (100 steps) | Effective per step |
|------------|----------------------|-------------------|
| 1 step/launch | 8.08 ms | 80.8 μs |
| 10 steps/launch | 7.86 ms | 78.6 μs |
| 100 steps/launch | 7.78 ms | **77.8 μs** |

**Observation:** Batch size has only a 4% effect on total time between batch=1 and batch=100. This is because the GPU kernel itself is only ~13μs of the total ~78μs. The remaining ~65μs is the pycuda launch overhead + stream sync.

## 4. SRT vs MRT Comparison (from cProfile + ncu)

| Metric | SRT (T1) | MRT (T2) | Ratio |
|--------|----------|----------|-------|
| stepper.step total (70×100 steps) | 0.190s | 0.191s | **1.01x** |
| pycuda function_call count | 7001 | 7001 | 1.0x |
| pycuda function_call time | 0.146s | 0.147s | 1.01x |
| ncu Duration | 10.05 μs | 10.53 μs | 1.05x |
| Registers per thread | — | 39 | — |
| SM Throughput | 18.88% | 22.23% | 1.18x |
| Memory Throughput | 52.73% | 52.92% | 1.00x |
| FMA pipe utilization | 27.3% | 32.3% | 1.18x |
| Executed IPC (active) | 1.18 | 1.24 | 1.05x |
| L1/TEX hit rate | 6.82% | 6.82% | 1.0x |
| L2 hit rate | 54.03% | 54.06% | 1.0x |
| Simulation.__init__ | 1.113s | 1.298s | 1.17x |

**Key finding: MRT is only 5% slower than SRT at this grid size.** Both are dominated by memory access (53% DRAM bandwidth utilization), not compute. MRT's extra FMA operations increase SM throughput from 18.9% to 22.2% but don't elongate wall time because the kernel is limited by the L1TEX scoreboard stalls (49.9% of warp cycles). The top stall reason is **waiting for L1TEX data** (scoreboard dependency), not waiting for compute pipelines.

The grid (384 blocks) is too small to fill the GPU — only 0.8 full waves across 80 SMs. At production scales (3000×300), the SM throughput for MRT would approach 50-60% and the 5% gap over SRT would widen to ~15-20%.

## 5. CUDA API Breakdown (from Nsight Systems)

| CUDA API | Total time (over whole run) | Per call | Notes |
|----------|---------------------------|----------|-------|
| cuCtxCreate_v2 | 123 ms | 123 ms | One-time, CUDA context creation |
| cuModuleLoad | 1.46 ms + 1.02 ms | 1.2 ms | Kernel compilation (one-time per compile) |
| cuLaunchKernel (warmup + measured) | ~47 μs per call | 47 μs | pycuda overhead from cProfile ~37 μs + CUDA driver overhead |
| cuMemcpyHtoD (init: 14KB) | 21-45 μs | varies | Config data upload |
| cuMemcpyHtoD (action: 72 bytes) | ~1.15 μs | 1.15 μs | Body rotation upload |
| cuMemcpyDtoH (DDF: 2.6MB) | 1.02 ms + 2.26 ms | 1-2 ms | Full DDF download (get_macroscopic) — one-time in profiled run |

**DTOH overhead:** The 72-byte action H2D takes ~1.15 μs, and the 60-byte obs D2H would take ~1 μs. These are truly negligible at any scale.

## 6. Bottleneck Identification

### Current bottleneck (384x192): pycuda kernel launch overhead

At this grid size, GPU compute time (13 μs/step) is dominated by pycuda's cuLaunchKernel overhead (~37 μs per kernel × 3 kernels = ~111 μs + Python loop ~7 μs = ~~78 μs/step from nsys, ~129 μs/step from cProfile). The **GPU is idle about 80% of the time** waiting for the next kernel launch.

### What would change at larger grids

| Grid size | GPU time/step (est.) | pycuda overhead | GPU idle |
|-----------|--------------------|----------------|----------|
| 384x192 | ~13 μs | ~37 μs × 3 = 111 μs | **80%** |
| 600x600 | ~33 μs | ~37 μs × 3 = 111 μs | 55% |
| 1000x500 | ~47 μs | 111 μs | 35% |
| 3000x300 (exp_ctrl) | ~530 μs | 111 μs | **17%** |
| 5000x500 | ~1.6 ms | 111 μs | 7% |

**At your production scale (3000x300), GPU utilization would be ~83%.** The pycuda overhead becomes acceptable.

### Where optimization would matter

1. **Reducing pycuda kernel launch overhead** — This is the #1 bottleneck for small grids. The 37 μs/launch comes from pycuda's argument buffer packing (`_build_arg_buf: 17 μs`) + CUDA driver overhead. Switching to `pycuda.driver.LaunchKernel` with pre-packed arguments could reduce this, but it would require significant changes to pycuda.

2. **Combining curved boundary into OneStep** — Would eliminate 2 of 3 kernel launches per step (save ~74 μs). For your production grid (3000x300), this would save about 12% of total time.

3. **Using larger batch sizes** — Already confirmed: batch=100 vs batch=1 saves about 4% total time on 384x192. The savings increase with grid size.

### Where optimization would NOT matter

1. **`set_body()` / `read_force()` overhead** — Already <1 μs. Not worth any optimization effort.
2. **GPU memory bandwidth** — The DDF read/write (5.3 MB/step) is only 2.5% of V100 HBM2 bandwidth. Not the bottleneck.
3. **Python loop overhead** — 5 μs/step loop from `for _ in range(n)` is negligible.
4. **Registry / Body API** — dict lookups and property access are sub-microsecond. Not relevant.

## 7. Performance Model

```
Per-step time (384x192 MRT) = 
    Python_overhead(~40 μs)    ← fixed for any grid size
    + CurvedBoundaryKernel(~4 μs per link-group)
    + OneStep(~5.9 μs)         
    + SensorKernel(~3.5 μs per sensor-group)

Per-step time (general) ≈ 
    pycuda_launch_kernels_3 × 37 μs
    + grid_cells × (pull_load + collide + pull_store) / V100_FLOPs
    
Scaling:
    - 384x192:  ~78 μs/step (GPU 17%, pycuda 83%)
    - 3000x300: ~640 μs/step (GPU 83%, pycuda 17%)
    - 6000x600: ~2.5 ms/step (GPU 96%, pycuda 4%)
```

## 8. Nsight Compute Deep-Dive Metrics

### OneStep (MRT) — Nsight Compute full metrics

| Metric | Value | Interpretation |
|--------|-------|----------------|
| Duration (ncu) | 10.53 μs | (ncu includes replay passes; nsys p50=5.86 μs is the real runtime) |
| Registers per thread | 39 | Well below V100 limit of 64. **No register spill to local memory.** |
| Grid | (2, 192, 1) × 256 = 384 blocks | Only 0.8 full waves — grid too small to fill 80 SMs |
| Memory Throughput | 52.9% of HBM2 peak | Moderate bandwidth usage |
| SM Throughput | 22.2% of peak | Low — limited by L1TEX scoreboard stalls |
| FMA pipe active | 32.3% of active cycles | Most-used pipeline, but not saturated |
| L1/TEX hit rate | 6.82% | Low — DDF reads are scattered across the grid |
| L2 hit rate | 54.0% | Moderate — half of cache-line requests hit L2 |
| Top stall reason | L1TEX data wait (49.9%) | Scoreboard dependency on global memory reads |
| IPC (active) | 1.24 inst/cycle | Far below V100's theoretical 4 IPC |
| Active warps/scheduler | 6.09 | Out of 16 max — low occupancy |
| Eligible warps/scheduler | 0.97 | Only 1 warp per cycle is ready to issue |

**Why the kernel does not saturate the GPU:**
- The grid (384 blocks = 73,728 threads) is too small. V100 has 80 SMs with 64 warps each = 5120 warps capacity. This kernel only uses 384 blocks × 2 warps/block = 768 warps (15% occupancy).
- Scoreboard stalls dominate (49.9%). Each warp waits on L1TEX for nearly half its cycles. This is the DDF pull phase: reading from neighbor cells at scattered memory addresses.
- The memory access pattern is suboptimal: only 26.9 of 32 bytes per sector are utilized for global loads that miss L2.

**Optimization potential:** The ncu reports "Est. Speedup: 6.45%" for memory access pattern improvements and "47.27%" for reducing L1TEX stalls. The 47% figure is misleading — it comes from the low grid occupancy, not from actual compute inefficiency. At production grid sizes (3000×300), active warps per scheduler would increase to ~12+ and the stall ratio would drop significantly.

### CurvedBoundaryKernel — Nsight Compute full metrics

| Metric | Value | Interpretation |
|--------|-------|----------------|
| Duration (ncu) | 5.60 μs | Tiny kernel, completely latency-bound |
| Grid | (2, 1, 1) × 256 = 512 threads | Only 2 blocks — **0.0 full waves** |
| Registers per thread | (from nsys) 31 | Low — no spill |
| Memory Throughput | 0.63% of HBM2 | Essentially zero bandwidth utilization |
| SM Throughput | 0.20% of peak | The GPU is doing nothing for this kernel |
| Active warps/scheduler | 1.87 | Out of 16 max — extremely low occupancy |
| Eligible warps/scheduler | 0.06 | Nearly 0 — almost never ready to issue |
| No Eligible | 94.45% | **94% of cycles have no warp ready to issue** |
| Est. Local Speedup | 97.57% | Nsight says there's 97% headroom |

**This is a latency-bound kernel by design.** Only 288 threads for the curved links. The 5.6 μs runtime is dominated by launch overhead (context switching, instruction cache misses). There is no meaningful optimization for this kernel because it cannot use more threads — there are only 288 cut links. 

**Impact:** At 4 μs per step, this kernel adds negligible runtime. Even with 100 cut links, the total curved time would still be <10 μs.

### SensorKernel — Nsight Compute full metrics

| Metric | Value | Interpretation |
|--------|-------|----------------|
| Duration (ncu) | 4.54 μs | Even smaller than curved |
| Memory Throughput | 0.63% of HBM2 | Near-zero |
| SM Throughput | 0.14% of peak | GPU doing nothing |
| Active warps/scheduler | 1.84 | Extremely low |
| Eligible warps/scheduler | 0.04 | Almost never ready |
| No Eligible | 96.83% | 97% idle cycles |
| L1/TEX hit rate | 48.53% | Much better than OneStep (localized access pattern) |

Same conclusion: negligible overhead, no meaningful optimization possible.

## 9. Recommendations

1. **For production runs (3000x300)** — The architecture is well-optimized. Use `--batch 10` to amortize pycuda overhead. GPU utilization is ~83%.

2. **For small-grid sweeps (384x192)** — The pycuda launch overhead dominates. If you need to run many small-grid parameter sweeps, consider:
   - Grouping multiple independent simulations into a single process
   - Or accepting that small-grid runs are launch-overhead limited

3. **For kernel tuning** — At production scale (3000x300), the main kernel (OneStep MRT) uses only 39 registers (no spill), achieves 53% DRAM throughput and 22% SM throughput on 384x192. At 3000x300, SM throughput would approach 50-60% and GPU utilization would be ~83%. No further kernel optimization is needed for current use cases.

4. **Three-kernel architecture is not a problem** — Despite being split into 3 kernels (CurvedBoundary + OneStep + Sensor), the total GPU kernel time is only ~13.4 μs/step. The curved and sensor kernels are latency-bound with negligible impact. Merging them into OneStep would save ~7 μs/step but add branch divergence in the OneStep grid kernel.
