CelerisLab/docs/performance_analysis.md

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.