Skip to content

perf: overhaul GPU hot path (7.7x at HERA scale) + revive parity tests - #130

Open
steven-murray wants to merge 27 commits into
mainfrom
speed
Open

perf: overhaul GPU hot path (7.7x at HERA scale) + revive parity tests#130
steven-murray wants to merge 27 commits into
mainfrom
speed

Conversation

@steven-murray

@steven-murray steven-murray commented Jul 17, 2026

Copy link
Copy Markdown
Contributor

Description

This PR overhauls the GPU hot path for large simulations. At the production-relevant scale (350 antennas, 350 unique beams, polarized, single precision, gridded beams), per-chunk GPU time drops 505 ms → 65 ms (7.7×) and steady-state wall time per integration drops 15.2 s → 2.06 s, with GPU utilization rising from ~35% to ~95%. It also revives the CPU/GPU parity test suite (silently skipped for years), fixes single-precision GPU support with gridded beams, adds an fp32-vs-fp64 validation test, adds a reusable benchmarking/profiling kit, and adds user-facing performance documentation.

Benchmarks below were measured on an RTX A2000 laptop GPU (4 GB, Ampere) with per-chunk CUDA events at the "production-slice" config: nant=nbeam=350, polarized, gridded beams, 10⁶ sources in 30 chunks, fp32, ERFA coordinates. Reproduce with profiling/run-canonical.sh.

Headline numbers (per-chunk CUDA-event timing)

Stage Before (ms) After (ms) Speedup
beam interp 224 9.9 22×
tau 23 2.1 11×
Z 13 5.7 2.3×
matprod 124 46.8 2.7×
chunk total 505 65.3 7.7×

After the change, ~70% of GPU time is the cuBLAS cherk kernel itself, i.e. the pipeline sits at the library roofline (verified against a bare-cuBLAS benchmark of the same shape). Larger relative gains are expected on data-centre GPUs (V100/A100), since the removed host-side overheads don't shrink with faster hardware.

How the baseline was diagnosed

nsys tracing at the production-slice shape (not toy scale — at small sizes ~90% of wall time is fixed setup, which is why earlier optimization attempts found nothing) showed:

  • Beam interpolation was 90% host overhead: nbeam·nfeed·nax = 1400 separate map_coordinates launches per chunk, plus per-beam coordinate-array allocation.
  • The GEMM ran on the default stream (the direct cupy_backends cublas calls never set the handle's stream), serializing against the compute stream.
  • A hidden complex128 GEMM per chunk: antpos * 1j promoted the phase-factor matmul to double precision even at precision=1, adding an fp64 GEMM + cast per chunk and a large temporary that caused OOMs.
  • The multi-stream design was a no-op: one stream per chunk, but every stage shares one set of buffers across chunks, so real overlap would be a data race — a per-chunk Device().synchronize() in the Z stage was what kept it correct, and it stalled the pipeline.

Changes

  1. matprod via cherk/zherk (gpu/_cublas.py): V = Z·Z^H is a Hermitian rank-k update — half the FLOPs of a general GEMM; a tiny kernel mirrors the computed triangle. General products (GPUVectorDot) use cublasCgemm3m (measured 2.1× over cgemm at matvis shapes). Both bound from libcublas via ctypes (cupy doesn't expose them), running on the current cupy stream. Falls back to cgemm/zgemm if the library can't be bound.
  2. One fused bilinear beam-interpolation kernel (gpu/beams.py): a single launch over all (beam, feed, axis) planes and sources; map_coordinates remains the fallback for order > 1.
  3. Fused Z kernel (gpu/getz.py): Z = A·√I·exptau in one elementwise pass (with the beam_idx gather), replacing 4 broadcast copies + a Python loop over all antennas + a full-device sync per chunk.
  4. Single compute stream (gpu/gpu.py): correct with zero device synchronization in the loop; the host queues many chunks ahead. Stages carry NVTX ranges; wall/event stats are exposed via LAST_RUN_STATS for the profiler CLI.
  5. tau precision fix (core/tau.py): keep the phase matmul at the requested precision.

Bugs fixed

  • tests/test_cpu_vs_gpu.py still guarded on importorskip("pycuda") (removed in v1.3.0), so the main CPU/GPU parity suite had been silently skipped for years. Re-enabled with cupy and extended to fp32 — which immediately exposed the next three bugs:
  • fp32 + gridded beams crashed on GPU (beam_data.set() dtype mismatch).
  • gpu.py computed its own nsrc_alloc, disagreeing with CoordinateRotation (which ignores source_buffer for chunks ≤ 1000 sources) → shape-mismatch crashes in small chunked runs.
  • Stale kx/ky spline options in the parity test (pyuvdata's map_coordinates interpolator takes order).

Validation

  • Full test suite passes (193 passed), including the revived parity suite across polarized/unpolarized × analytic/gridded × fp32/fp64 × chunking × source-buffer configurations.
  • New tests/test_precision.py: fp32 agrees with fp64 to 1e-5 of the peak visibility on both backends — the accuracy gate for running production in single precision.
  • tests/test_cublas.py extended to rectangular shapes, both dtypes, and out=/beta= accumulation (exercises the herk + mirror path).
  • Interpolation kernel verified against map_coordinates at fp32 tolerance.

Follow-up work

Filed as issues: #131 (validation on V100/ilifu), #132 (sum_chunks accumulation/pinned-memory transfer), #133 (horizon-cut compaction without a host sync), #134 (hoisting frequency-independent work out of the per-frequency loop), #135 (per-baseline source coarsening).

Checklist

I have

  • Added a test covering your new feature adequately?
  • Added a docstring (or note to a docstring) describing your feature?
  • (Optional): Added a tutorial / section to a tutorial showing usage of your new feature? (new "Performance" docs page)
  • Important: this feature does not break API compatibility.

🤖 Generated with Claude Code

Steven Murray and others added 5 commits May 29, 2026 16:49
At the production-relevant scale (350 antennas, 350 unique beams, polarized,
single precision), per-chunk GPU time drops from 505 ms to 65 ms (7.7x) and
GPU utilization rises from ~35% to ~95%. Changes:

- matprod: V = Z Z^H is a Hermitian rank-k update, so use cublasCherk/Zherk
  (half the FLOPs of a general GEMM; one triangle computed, mirrored by a
  small kernel). General products use cublasCgemm3m (measured 2.1x over
  cgemm at matvis shapes). Both are bound from libcublas with ctypes (cupy
  does not expose them) and run on the current cupy stream via
  cublasSetStream — previously the GEMM ran on the default stream,
  serializing against the compute stream. (124 -> 47 ms/chunk, at the
  cuBLAS roofline.)
- beams: one fused bilinear kernel interpolates every (beam, feed, axis)
  plane for all sources in a single launch, replacing nbeam*nfeed*nax
  map_coordinates launches (1400/chunk at 350 beams; ~90% host launch
  overhead). map_coordinates remains the fallback for order > 1.
  (224 -> 10 ms/chunk.)
- getz: new GPUZMatrixCalc computes Z = A*sqrtI*exptau in one elementwise
  pass with the beam_idx gather, replacing nfeed*nax broadcast copies, a
  Python loop over antennas, and a full-device synchronize per chunk.
- gpu loop: single in-order compute stream instead of one stream per chunk.
  The multi-stream design could never overlap (stages share buffers across
  chunks — overlap would race) and forced per-chunk synchronization; one
  stream is correct with no syncs and lets the host queue ahead. Wall-time
  and CUDA-event run stats are exposed via LAST_RUN_STATS, and stages carry
  NVTX ranges for nsys.
- tau: antpos * 1j silently promoted the phase matmul to complex128 at
  precision=1, costing a hidden fp64 GEMM plus cast per chunk and a large
  temporary that could OOM the chunk estimator's budget.

Fixes uncovered by re-enabling the parity suite:

- tests/test_cpu_vs_gpu.py had been silently skipped since the pycuda
  removal (importorskip("pycuda")); now guards on cupy and covers single
  precision.
- Single-precision GPU runs with gridded beams crashed on a dtype mismatch
  when uploading beam data.
- gpu.py computed its own nsrc_alloc, disagreeing with CoordinateRotation
  (which ignores source_buffer for chunks <= 1000 sources) and crashing
  small chunked runs; it now uses coords.nsrc_alloc like the CPU path.
- The parity test passed stale kx/ky spline options (pyuvdata's
  map_coordinates interpolator takes "order").

New tests: fp32-vs-fp64 end-to-end gate (tests/test_precision.py), herk
paths in tests/test_cublas.py (rectangular shapes, out=/beta= accumulation).

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
- matvis profile writes summary-stats-*.json (config, total/setup/loop wall
  time, per-stage line-profiler and CUDA-event timings) for before/after
  comparisons.
- profiling/: canonical benchmark configs (run-canonical.sh), speed-of-light
  micro-benchmarks (roofline.py), and cuBLAS strategy comparison
  (gemm_experiments.py), with a README covering the nsys recipe.
- Ignore benchmark outputs (profiling/results/, *.nsys-rep, *.sqlite,
  summary-stats-*.json).

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…angelog

Covers per-stage cost scaling, measured rule-of-thumb throughput, a
GEMM-bound estimation formula (including the source_buffer caveat),
precision guidance, memory/chunking, how to benchmark a configuration, and
a changelog of performance-relevant changes back to the v1.3.0 cupy rewrite.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
@codecov

codecov Bot commented Jul 17, 2026

Copy link
Copy Markdown

Codecov Report

✅ All modified and coverable lines are covered by tests.
✅ Project coverage is 99.39%. Comparing base (e931e83) to head (e781f13).
⚠️ Report is 17 commits behind head on main.

Additional details and impacted files
@@            Coverage Diff             @@
##             main     #130      +/-   ##
==========================================
+ Coverage   98.88%   99.39%   +0.50%     
==========================================
  Files          22       23       +1     
  Lines         989     1153     +164     
  Branches      103      145      +42     
==========================================
+ Hits          978     1146     +168     
+ Misses          6        4       -2     
+ Partials        5        3       -2     
Flag Coverage Δ
unittests 88.03% <99.58%> (+9.56%) ⬆️

Flags with carried forward coverage won't be shown. Click here to find out more.

☔ View full report in Codecov by Harness.
📢 Have feedback on the report? Share it here.

🚀 New features to boost your workflow:
  • ❄️ Test Analytics: Detect flaky tests, report on failures, and find test suite problems.

Steven Murray and others added 3 commits July 17, 2026 11:33
Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Codecov patch coverage flagged the GPU debug-memory-logging and
gpu_event_timing branches in gpu/gpu.py (only exercised when the logger is
at DEBUG or gpu_event_timing=True, neither of which any existing test set),
plus a handful of defensive/fallback branches in the new _cublas.py and
beams.py code that only trigger on cuBLAS errors, missing shared libraries,
higher-order spline interpolation, or dtype mismatches.

- tests/test_matvis_gpu.py: run simulate_vis with gpu_event_timing=True and
  the matvis.gpu.gpu logger at DEBUG, covering both the zero-active-chunk
  case (source always below horizon) and the normal active-chunk case
  (analytic and gridded beams, so bmfunc.use_interp is exercised both ways).
- tests/test_cublas.py: invalid-dtype errors, the _LIB=None fallback to
  cgemm/zgemm, non-zero cuBLAS status from herk/gemm3m raising RuntimeError,
  and the soname-retry loop in _load_cublas_ext.
- tests/test_beam_interp_gpu.py: the order!=1 map_coordinates fallback (for
  both real and complex beams) and interpolating a complex beam into a
  differently-dtyped output buffer (the scratch-buffer/no-sqrt path).
- tests/test_getz.py (new): direct unit tests of GPUZMatrixCalc against a
  numpy reference, including calling the same instance twice with a given
  beam_idx to exercise the cached-device-array branch, plus the two
  beam_idx=None cases (shared beam, implicit per-antenna beam).

All files touched by this PR are now at 100% statement and branch coverage;
full suite: 211 passed, 3 skipped.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
@steven-murray steven-murray self-assigned this Jul 17, 2026
Steven Murray and others added 3 commits July 17, 2026 12:26
The self-hosted GPU CI job runs `pytest -k "gpu"` to limit what executes on
that runner, matching against test node IDs (which include the module name).
test_cublas.py, test_getz.py, and test_precision.py are exclusively
cupy-dependent (the latter's GPU parametrization in particular is the
accuracy gate for running production in single precision) but didn't contain
"gpu" in their filename, so none of their tests were ever selected on the
GPU runner. Combined with the CPU-matrix jobs skipping them entirely (no
cupy installed there), this meant they got no CI coverage at all -- which is
what caused codecov's patch-coverage check to flag lines in _cublas.py and
getz.py as untested despite 100% local coverage, and, more importantly,
meant the fp32-vs-fp64 GPU validation never actually ran on real GPU
hardware in CI.

Renamed to match the existing convention (test_beam_interp_gpu.py,
test_cpu_vs_gpu.py, test_matvis_gpu.py): test_cublas_gpu.py,
test_getz_gpu.py, test_precision_gpu.py. Verified locally against the exact
CI GPU-job command (`pytest -k "gpu" --cov=matvis ...`): all files touched
by this PR now reach 100% coverage under that filter alone.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
The self-hosted GPU runner's cuBLAS build exceeded the previous rtol=1e-4 at
the largest test shape (K=5000): 2/4096 elements differed by up to 2.1e-4.
This is expected behaviour of the Gauss 3M algorithm (cublasCgemm3m), which
trades some rounding accuracy for fewer real multiplies -- documented cuBLAS
behaviour, not a correctness regression (the herk-based test_zdotz, which
doesn't use 3M, kept its tight tolerance and passed). Loosened to rtol=1e-3
for complex64 in this test only.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Measured on a GeForce GTX Titan X (Maxwell, 2015) cluster node via
profiling/run-canonical.sh, roofline.py, and gemm_experiments.py:

- Rules-of-thumb table: ~1.8s/integration at the same 350-ant/350-beam
  production-slice config as the existing A2000 entry.
- New "GEMM strategy: hardware dependence" section: on this card cgemm3m is
  slower than plain cgemm (baseline cgemm is already near-roofline, leaving
  no headroom, and the 3M decomposition's overhead becomes a net loss), and
  cherk gives no measurable gain -- both quite different from the ~2x/~2.8x
  wins measured on the dev A2000. cherk is never worse than cgemm on either
  card, so it stays a safe default; cgemm3m's benefit is not guaranteed on a
  new architecture without checking.
- Added a warning to the benchmarking section: the JSON's line-profiler
  "stages" table (in particular "Coordinate Rotation", which shares its
  bucket with the horizon-cut's blocking GPU sync) is not a reliable
  per-stage breakdown under the async GPU pipeline -- use
  run_stats.event_timing_ms / time_per_integration instead. Confirmed via
  the raw numbers Steven measured: line-profiler attributed 36.6% to
  "Coordinate Rotation" on this run, while the CUDA-event chunk_total and
  time_per_integration numbers show GPU compute is not the bottleneck there.

Note: this is a different card than the V100 (ilifu) requested in issue
#131, so it's an additional data point, not a resolution of that issue.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Steven Murray and others added 4 commits July 19, 2026 10:59
…imal

Measured on a Quadro RTX 5000 (Turing, compute capability 7.5) cluster node:

- Rules-of-thumb table: ~3.0s/integration at the 350-ant production-slice
  config. Added a note that these end-to-end numbers mix host and GPU time
  and vary by cluster CPU as well as card -- GPU-only time is only ~80% of
  the total here vs ~95% on the dev A2000, which explains why this number
  looks worse despite the card having a competitive (better, even) GEMM
  roofline.
- GEMM strategy table: on this card cgemm3m is ~1.6x faster than both plain
  cgemm and cherk, while cherk gives no measurable gain at all -- the
  opposite pattern from the Titan X (neither helps) and a stronger, more
  actionable version of the Ampere case (both help, cherk best). Since
  GPUMatMul always uses cherk, this is a real ~20-25% unrealized reduction
  in per-chunk GPU time on this hardware, not just a micro-benchmark
  curiosity. Filed issue #136 to track auto-selecting between strategies;
  linked from the docs.
- Added a tip in the benchmarking section: small/single-chunk runs (e.g. the
  dev canonical config) can be skewed by one-time cupy RawModule/RawKernel
  JIT-compilation landing inside the CUDA-event average when there are few
  samples to dilute it -- observed directly in this run's dev-config beam
  interpolation figure (182.6ms, ~12x the prodslice figure for a similar
  element count), which is why it wasn't added to the rules-of-thumb table.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
The three GPUs benchmarked so far revealed that the harness's headline
number (loop_time/ntimes) was contaminated by one-time costs: on the Quadro
RTX 5000 the first integration took 6.7s vs 1.8s steady-state (cupy kernel
compilation, cuBLAS workspace allocation, ERFA/IERS cache loads), inflating
the reported per-integration time by ~65%. The per-stage CUDA-event means
had the same problem (the "182ms beam interpolation" artifact was a compile
stall averaged over only 8 samples), and wall time conflated GPU speed with
cluster CPU speed, making cross-machine comparisons misleading.

Changes:

- matvis profile runs a small untimed warmup simulation first (same
  precision/beam-type/backends so the same kernel variants compile;
  --no-warmup to disable). After warmup, the first integration is within
  ~1.5x of steady state instead of ~80x on the dev config.
- gpu.simulate records per-integration wall times individually and reports
  steady_time_per_integration = median excluding the first integration.
- CUDA-event stage timings keep all per-chunk samples and report
  median/mean/count per stage instead of a running mean.
- The JSON gains a `derived` block with the three numbers worth quoting:
  steady_wall_per_integration, gpu_time_per_integration (median chunk total
  x nchunks; transfers across machines with the same card), and
  host_overhead_per_integration (their difference). The CLI summary prints
  these first.

Docs: the rules-of-thumb table now has separate GPU-time and wall-time
columns (replacing two prose caveats with structure), the JIT-warmup tip is
gone (the harness handles it), and the remaining line-profiler warning is
shortened. Re-measured the RTX A2000 row with the new harness (GPU 2.0s /
wall 2.1s per integration, host overhead 0.04s); the Titan X and RTX 5000
rows are marked for re-measurement since their old values predate the
warmup pass (the RTX 5000 wall entry corrects from 3.0s to ~1.8s based on
its steady-state per-integration log).

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Re-measured with the new harness: GPU 1.6 s / wall 1.7 s per integration
(host overhead 2%). Consistent with the previous run once its
warmup-contaminated first integration is accounted for. Sanity checks:
stage medians sum to chunk_total within ~1 ms, and matprod (41.0 ms/chunk)
reproduces the earlier measurement exactly; it runs ~16% above the
K=100k roofline scaling, which is normal cuBLAS kernel-selection
variation with K. Only the Quadro RTX 5000 row still awaits re-measurement.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
GPU 1.8 s / wall 1.8 s per integration (host overhead 0.8%), removing the
last pending-re-measurement footnote. The re-run validates the new harness
twice over: the steady wall time matches the estimate extracted from the old
run's per-integration log to three digits, and the earlier "~0.6 s host
overhead on this node" reading is shown to have been an artifact of
warmup-contaminated inputs (both machines are actually <=2% host overhead).

Also corrected the GEMM-strategy section with the robust medians: on the
RTX 5000, matprod is ~89% of per-chunk GPU time (not 66% as the contaminated
numbers suggested), so cgemm3m selection (issue #136) would cut GPU time
~30% -- enough to make that card the fastest of the three measured rather
than the slowest.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Comment thread docs/performance.rst
Comment on lines +65 to +79
.. list-table::
:header-rows: 1

* - Hardware
- GPU time / integration
- Wall time / integration
* - RTX A2000 laptop (Ampere, 95 W class)
- 2.0 s
- 2.1 s
* - GeForce GTX Titan X (Maxwell, 2015 workstation card)
- 1.6 s
- 1.7 s
* - Quadro RTX 5000 (Turing, 16 GB workstation card)
- 1.8 s
- 1.8 s

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Since this was generated with Claude, do you think it's easy to also report the theoretical minimum GPU time/integration? I'm guessing the "effective R" indirectly reports this for the A2000, but it might be a useful number to have around if someone down the line wants a sense of how much could potentially be gained through further optimizations.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The t_gemm formula a few lines down already is that calculation once you plug in your own measured R. I added a sentence making the connection explicit instead of leaving it for the reader to infer (72014fc).

Comment thread docs/performance.rst
Comment on lines +175 to +179
Device memory is dominated by the per-chunk :math:`Z` matrix and interpolated
beam array, each of size
:math:`N_{\rm ant/beam} N_{\rm feed} N_{\rm ax} N_{\rm src}^{\rm alloc}`
complex values, plus the raw beam grids
(:math:`N_{\rm beam} N_{\rm feed} N_{\rm ax} N_{\rm pix}`). Sources are

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Is it worthwhile to acknowledge a bit of subtlety related to the raw beam grids here? For production runs, the raw beam grids typically have $N_{\rm pix} \sim 65000$ for degree sampling, so for smaller chunk sizes this can be a dominant term, but for larger chunks this contribution might not be a very large part. (This is a genuine question—I'm not sure whether it's worthwhile to discuss this here.)

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yes, worth it — added a paragraph to "Memory and chunking": the raw beam-grid term doesn't scale with chunk size, so it can dominate for small chunks and becomes negligible once chunk-scaled terms take over (72014fc).

Comment thread docs/performance.rst
Comment on lines +190 to +192
matvis profile -a 350 -b 350 -s 1000000 -t 4 --gpu \
--interpolated-beam --single-precision --gpu-event-timing \
--coord-method CoordinateRotationERFA -o outdir

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I'm not seeing any discussion of the profiling script on the docs except for this new addition. At the very least, it's probably worthwhile to add some reference about it that explains the parameters. Even better would be a reference page on the CLI utilities.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Went with the "even better" option — added docs/cli.rst, a sphinx-click-generated reference for every matvis profile/hera-profile parameter (self-maintaining, since it reads the actual --help text), plus a real annotated JSON output example. Linked from here, the toctree, and profiling/README.md (0d6e9b7).

Comment thread docs/performance.rst
already-queued GPU work (especially "Coordinate Rotation", which shares
its bucket with the horizon-cut's blocking sync — see issue #133). Use it
only as a rough indicator for the CPU backend; for the GPU backend use
the ``derived`` and ``run_stats.event_timing_ms`` values.

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This feels very AI-coded and can probably be improved. I'm assuming the derived object (?) that is referenced several times is an entry in the output json, but I'm not about to run the profiler myself just to figure out what it's talking about. A bare minimum improvement to this section would be an example showing the contents of the json file produced via matvis profile ....

In line 206 ("``derived.gpu_time_per_integration`` ..."), I would replace "total times the" with "total multiplied by the". Using "times" to mean multiplication in a section talking about time is confusing.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Fixed the "times"/multiplication ambiguity (7048c43), and added an annotated summary-stats-*.json example on the new docs/cli.rst page rather than duplicating it inline here (0d6e9b7). I also added some extra text to the top of the performance.rst to make it clear where the CLI docs are.

Comment thread docs/performance.rst Outdated
ones to quote and compare — they are what the Rules of Thumb table reports.

The ``profiling/`` directory in the repository contains canonical benchmark
configurations, GEMM/interpolation "speed of light" micro-benchmarks, and an

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

"speed of light" micro-benchmarks? I see this phrase used again in the profiling README, and it feels weird. Does the "speed of light" modifier really help here? It feels like fluff that can potentially be confusing to readers (unless this is actually meaningful to the benchmarking community).

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Apparently Claude agreed with you and suggested "roofline micro-benchmarks" instead... which sounds just as jargony to me, but it says is the standard terminology. I changed it to roofline but also added a parenthesis to explain what that means for us non-gpu-experts.

Comment thread docs/performance.rst Outdated
timing individual Python lines, but the GPU loop is asynchronous: a line
can appear expensive simply because it's where the host next blocks on
already-queued GPU work (especially "Coordinate Rotation", which shares
its bucket with the horizon-cut's blocking sync — see issue #133). Use it

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Update the issue reference to a link to the issue if possible.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

fixed (72014fc)

Comment thread src/matvis/gpu/beams.py Outdated
# for a 350-antenna array with per-antenna beams), which left the GPU idle
# most of the time waiting on the host to issue work.
#
# Grid: x covers sources, y covers planes. Out-of-range coordinates clamp to

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

"y covers planes" means what? It looks like y is just parallelizing over beam components, while x is parallelizing over sources.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yup, reworded to "x indexes sources, y indexes (beam, feed, axis) combinations" (1ed75e6).

Comment thread src/matvis/gpu/beams.py Outdated
Comment on lines +23 to +97
_BILINEAR_MODULE = cp.RawModule(
code=r"""
#include <cupy/complex.cuh>

template<typename R, typename T>
__device__ void bilinear_all_planes(
const T* __restrict__ beam,
const R* __restrict__ az,
const R* __restrict__ za,
const R* __restrict__ daz,
const R* __restrict__ dza,
const R* __restrict__ azmin,
const int nfeed,
const int nax,
const long nza,
const long naz,
const long nsrc,
T* __restrict__ out)
{
const long s = blockIdx.x * (long)blockDim.x + threadIdx.x;
if (s >= nsrc) return;
const int p = blockIdx.y;
const int bm = p / (nax * nfeed);
const int r = p % (nax * nfeed);
const int ax = r / nfeed;
const int fd = r % nfeed;

R x = za[s] / dza[bm];
R y = (az[s] - azmin[bm]) / daz[bm];

long x0 = (long)floor(x);
long y0 = (long)floor(y);
R fx = x - x0;
R fy = y - y0;
if (x0 < 0) { x0 = 0; fx = 0; }
if (x0 > nza - 2) { x0 = nza - 2; fx = 1; }
if (y0 < 0) { y0 = 0; fy = 0; }
if (y0 > naz - 2) { y0 = naz - 2; fy = 1; }

const T* b = beam + ((((long)bm * nax + ax) * nfeed + fd) * nza + x0) * naz + y0;
T v = b[0] * ((1 - fx) * (1 - fy))
+ b[1] * ((1 - fx) * fy)
+ b[naz] * (fx * (1 - fy))
+ b[naz + 1] * (fx * fy);

out[(((long)bm * nfeed + fd) * nax + ax) * nsrc + s] = v;
}

extern "C" {
__global__ void bilinear_c64(
const complex<float>* beam, const float* az, const float* za,
const float* daz, const float* dza, const float* azmin,
int nfeed, int nax, long nza, long naz, long nsrc, complex<float>* out)
{ bilinear_all_planes<float, complex<float> >(beam, az, za, daz, dza, azmin, nfeed, nax, nza, naz, nsrc, out); }

__global__ void bilinear_c128(
const complex<double>* beam, const double* az, const double* za,
const double* daz, const double* dza, const double* azmin,
int nfeed, int nax, long nza, long naz, long nsrc, complex<double>* out)
{ bilinear_all_planes<double, complex<double> >(beam, az, za, daz, dza, azmin, nfeed, nax, nza, naz, nsrc, out); }

__global__ void bilinear_f32(
const float* beam, const float* az, const float* za,
const float* daz, const float* dza, const float* azmin,
int nfeed, int nax, long nza, long naz, long nsrc, float* out)
{ bilinear_all_planes<float, float>(beam, az, za, daz, dza, azmin, nfeed, nax, nza, naz, nsrc, out); }

__global__ void bilinear_f64(
const double* beam, const double* az, const double* za,
const double* daz, const double* dza, const double* azmin,
int nfeed, int nax, long nza, long naz, long nsrc, double* out)
{ bilinear_all_planes<double, double>(beam, az, za, daz, dza, azmin, nfeed, nax, nza, naz, nsrc, out); }
}
"""
)

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Same comment as before about C code as raw strings living in a python module.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

And same fix :-)

Comment thread src/matvis/gpu/beams.py Outdated
Comment on lines +57 to +60
if (x0 < 0) { x0 = 0; fx = 0; }
if (x0 > nza - 2) { x0 = nza - 2; fx = 1; }
if (y0 < 0) { y0 = 0; fy = 0; }
if (y0 > naz - 2) { y0 = naz - 2; fy = 1; }

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Might be helpful to human readers to add a comment above this block indicating that this is where out-of-bounds points are assigned the beam values at the input boundaries.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Added, close to your wording (1ed75e6).

Comment thread src/matvis/gpu/beams.py Outdated
Comment on lines +279 to +283
# Single fused launch over all (beam, feed, axis) planes and sources.
# The kernel writes values of the beam's own dtype; when the caller
# supplies a complex output buffer for a real (power) beam, go through
# a real-valued scratch array and cast on copy (as map_coordinates
# would have done element-wise).

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Is this comment all that helpful? It's more-or-less repeating stuff said earlier on when the bilinear interpolation routine was declared. Probably more helpful to say something like "Use the custom beam interpolation kernel. If provided a power beam and a complex output buffer, cast interpolated beam to complex on copy.".

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Used your suggested wording exactly (fc5b8da).

Comment thread src/matvis/gpu/beams.py
dza = cp.asarray(dza, dtype=rdtype)
azmin = cp.asarray(azmin, dtype=rdtype)
assert beam._c_contiguous and target._c_contiguous
block = 128

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Another comment about a seemingly arbitrary choice for block size. Was there some optimization done that resulted in this choice?

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Same answer: conventional warp-multiple default (128 here), not tuned. Commented accordingly (fc5b8da).

Comment thread src/matvis/gpu/getz.py Outdated
Comment on lines +14 to +58
_FUSED_Z_MODULE = cp.RawModule(
code=r"""
#include <cupy/complex.cuh>

template<typename R, typename T>
__device__ void fused_z(
const T* __restrict__ beam, // (nbeam, nfeed, nax, nsrc)
const T* __restrict__ exptau, // (nant, nsrc)
const R* __restrict__ sqrt_flux, // (nsrc,)
const long* __restrict__ beam_idx, // (nant,) or NULL
const long bmul, // if beam_idx NULL: 0 -> one shared beam, 1 -> beam per ant
const int nfeed,
const int nax,
const long nsrc,
const long ntot,
T* __restrict__ out) // (nant, nfeed, nax, nsrc)
{
const long i = blockIdx.x * (long)blockDim.x + threadIdx.x;
if (i >= ntot) return;
const long s = i % nsrc;
long rest = i / nsrc;
const int ax = rest % nax;
rest /= nax;
const int fd = rest % nfeed;
const long ant = rest / nfeed;
const long bm = beam_idx == NULL ? ant * bmul : beam_idx[ant];

const T a = beam[((bm * nfeed + fd) * nax + ax) * nsrc + s];
out[i] = a * exptau[ant * nsrc + s] * sqrt_flux[s];
}

extern "C" {
__global__ void fused_z_c64(
const complex<float>* beam, const complex<float>* exptau,
const float* sqrt_flux, const long* beam_idx, long bmul,
int nfeed, int nax, long nsrc, long ntot, complex<float>* out)
{ fused_z<float, complex<float> >(beam, exptau, sqrt_flux, beam_idx, bmul, nfeed, nax, nsrc, ntot, out); }

__global__ void fused_z_c128(
const complex<double>* beam, const complex<double>* exptau,
const double* sqrt_flux, const long* beam_idx, long bmul,
int nfeed, int nax, long nsrc, long ntot, complex<double>* out)
{ fused_z<double, complex<double> >(beam, exptau, sqrt_flux, beam_idx, bmul, nfeed, nax, nsrc, ntot, out); }
}
"""

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Marking another section of C code declared in a python module.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Same action taken.

Comment thread src/matvis/gpu/getz.py
)

ntot = self.nant * self.nfeed * self.nax * self.nsrc
block = 256

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Block size setting again

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Same comment again :-)

Steven Murray and others added 2 commits July 28, 2026 09:56
Measured on a Tesla V100-SXM2-32GB, the HERA production reference
architecture (issue #131) -- substantially addresses that issue's ask, though
this was not necessarily run on the ilifu box specifically, and the test
suite hasn't been run there yet.

Both JSON outputs verified self-consistent (derived values reproduce exactly
from run_stats). Production matprod (21.66 ms/chunk) is in the expected
range scaling the roofline (36.51 ms at K=1e5) down to the actual per-chunk
K, with the same ~10% super-linear efficiency gain at smaller K seen on
other cards. One incidental finding: the dev-config's tau stage showed
median=0.44ms vs mean=1.61ms (a single-outlier artifact in an 8-sample
average) -- a good real-world case for why the harness reports medians.

Rules-of-thumb table: 0.8s/integration (GPU and wall both), more than 2x
faster than any workstation/laptop card measured so far, as expected for a
data-centre part. Also updated the R-value example in the GEMM-estimation
formula (~10.7 TFLOPS vs the A2000's ~5.2).

The GEMM-strategy table's cgemm3m/cherk comparison is NOT yet added for
V100: the supplied gemm_experiments.py output only had the cgemm baseline
line (~10.9 TFLOPS), pending confirmation from a colleague on whether the
rest of the script's output was cut off or the script exited early -- noted
inline in the docs as pending re-measurement.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Full gemm_experiments.py output confirms: cgemm3m is 1.72x faster than the
cherk-based baseline (21.06ms vs 36.55ms at M=700/K=1e5), and cherk itself
gives no measurable gain (the tri-only computation is ~5% faster than
plain cgemm, but the mirror kernel's overhead erases it). This matches the
Turing pattern rather than being a one-off: two of the four architectures
sampled -- both the two most modern, and one of them the actual HERA
production reference -- get zero benefit from cherk while cgemm3m offers a
30-35% cut in GPU time per integration. Elevated this in the issue #136
comment thread accordingly.

Closed issue #131 (V100 validation) per explicit instruction: measured
0.8s/integration (GPU and wall), 2x+ faster than any workstation/laptop
card sampled, as expected for the production reference architecture.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Comment thread src/matvis/_utils.py Outdated
gpusize = {
"antpos": nant * 3 * rsize,
"flux": nsrc * rsize,
# nbeampix is already summed over beams by the caller

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Is this actually a helpful comment?

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

As written, not really — it assumed you already knew why. Expanded it to be self-contained: nbeampix is pre-summed across all beams by get_desired_chunks, so multiplying by nbeam here would double-count beam memory. (This was a real bug earlier in this PR, which is why the comment exists at all) (1ed75e6)

Comment thread src/matvis/cli.py Outdated
Comment on lines +220 to +223
# Derived headline numbers, robust to warmup and host noise. These are
# the values to quote/compare (see the docs Performance page); the
# line-profiler stage table below is indicative only, since the GPU loop
# is asynchronous.

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think the phrase "is indicative only" is throwing me off, perhaps also because I'm not quite sure what the "line-profiler stage table" is. Is it referring to the stats that are dumped to the json file? Is the correct interpretation of this something along the lines of "Because the GPU loop is asynchronous, the reported GPU times do not correspond to actual wall times"?

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yeah, this was not well worded. Reworded to say that the stages table times Python lines, but GPU work is queued asynchronously, so a line's measured time is often how long the host waited for already-queued GPU work to finish, not the cost of that line itself (1ed75e6).

Comment thread src/matvis/cli.py
Comment on lines +237 to +239
derived["host_overhead_per_integration"] = max(
derived["steady_wall_per_integration"] - gpu_time, 0.0
)

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Having this max(...) catch here feels weird to me? I suppose it's possible that there can be a scenario where the median wall time is less than the median gpu time, but because of that it feels like reporting the median isn't the right thing to do? Why are we not just reporting mean and standard deviation? Since the profiling already includes a dry run to remove one-off setup costs, do we actually expect there to be outliers that would contaminate the means?

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

About the max(...): gpu_time_per_integration is median(per-chunk-total) × nchunks, a biased estimator of the true per-integration GPU total when chunks vary in cost (e.g. the horizon cut leaves different source counts active in different chunks), while wall time sums the actual per-integration total — they can cross without either being wrong, hence the max(). Added a comment explaining that.

On mean vs. median: yes, we do see outliers even after warmup — the V100 dev-run's tau stage showed median=0.44ms vs mean=1.61ms from a single-sample outlier in only 8 warmed-up samples (see the Performance docs page). Added std alongside median/mean in the JSON so we're not limited to picking one (1ed75e6).

Comment thread src/matvis/wrapper.py Outdated
The method to use for the final matrix multiplication. Default is 'MatMul',
which simply uses matrix multiplication over the two full matrices. Currently,
the other option is `VectorLoop`, which uses a loop over the antenna pairs,
one other option is `VectorDot`, which uses a loop over the antenna pairs,

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Changing "the" to "one" is weird because right now it is still the only other alternative? Is the name change from VectorLoop to VectorDot that helpful? Isn't this also an API change?

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The change from VectorLoop to VectorDot happened long ago -- this was just an incorrect vestige that Claude found and fixed. You're right about "the"/"one" — reverted that back to "the other option" since there is only one (1ed75e6).

Comment thread src/matvis/gpu/beams.py
Comment on lines 333 to 335
if not complex_beam: # power beam
cp.sqrt(beam_at_src, out=beam_at_src)
beam_at_src = beam_at_src.astype(ctype)

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I didn't catch this earlier, but this will silently do the wrong thing when interpolating a cross-polarized power beam (I believe also for the co-pols in the case where the power beam is constructed from an e-field beam with calc_cross_pols=True). In the mindset of "explicit is better than implicit", shouldn't the e-field or power behavior be controlled with an argument that is passed to the function instead of an inference of the beam type based on the beam dtype?

@steven-murray steven-murray Aug 11, 2026

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Traced this through carefully. It can't currently happen via the real pipeline: core/beams.py's _wrangle_beams requires beam_type == "efield" whenever polarized=True (power beams never reach the interpolator in that case), and prepare_beam_unpolarized calls as_power_beam(include_cross_pols=False, ...) for the unpolarized path — so cross-pol data is excluded before this function ever sees it. That said, your underlying point stands: gpu_beam_interpolation is a public function called directly by tests and roofline.py, so this is a fragile API regardless of whether our calling functions happen to be safe. Added an explicit power_beam: bool | None parameter (default None keeps the current inference for existing callers); GPUBeamInterpolator now passes it explicitly as not self.polarized. The reworked test suite (see the test_beam_interp_gpu.py thread) exercises the new parameter directly (fc5b8da).

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

These tests should instead use pytest fixtures for the input beam used for the tests. Maybe something like an E-field beam and its corresponding power beam sampled sparsely but over the full sphere. It may be a good idea to have a real-valued E-field beam and a complex-valued E-field beam to fully cover the bases.

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Reworked the whole file around this. New fixtures build on conftest.py's existing uvbeam/uvbeam_unpol (the full test beam already used elsewhere in the suite) rather than inventing synthetic arrays. See the other three threads on this file for what replaced the old tests (d8b2d05).

Comment thread tests/test_beam_interp_gpu.py Outdated
Comment on lines +91 to +113
def test_order_gt_1_uses_map_coordinates_fallback():
"""Order != 1 must fall back to (and correctly use) cupyx map_coordinates.

The fused bilinear kernel only handles order=1; higher orders go through
the original per-plane ``map_coordinates`` loop. At grid nodes, a spline
of any order should exactly reproduce the input values, same as order=1.
"""
za = np.linspace(0, 1, 5)
az = np.linspace(0, 2 * np.pi, 8)

AZ, ZA = np.meshgrid(az, za, indexing="xy")
beam = np.array([1 - ZA**2])
beam = beam[:, np.newaxis, np.newaxis] # nax=1, nfeed=1
dza = za[1] - za[0]
daz = az[1] - az[0]

new_beam = gpu_beam_interpolation(
beam, [daz], [dza], [0.0], AZ.flatten(), ZA.flatten(), order=2
).get()

np.testing.assert_allclose(
np.sqrt(beam[0, 0, 0].flatten()), new_beam[0, 0, 0], atol=1e-6
)

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

A better test here is probably to check that the interpolated beam with ndimage.map_coordinates gives the same thing as interpolating the UVBeam object directly with the az_za_map_coordinates setting for the spatial interpolation (with the appropriate corresponding set of spline_opts)?

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Done — test_order_gt_1_matches_uvbeam_interp now cross-checks gpu_beam_interpolation(..., order=2) against UVBeam.interp(interpolation_function="az_za_map_coordinates", spline_opts={"order": 2}), evaluated at non-node points (offset from the native grid) rather than at grid nodes — the old test's node-only check was a tautology any interpolator passes regardless of order, since it just reproduces its own input (d8b2d05).

Comment thread tests/test_beam_interp_gpu.py Outdated
Comment on lines +116 to +132
def test_order_gt_1_with_complex_beam():
"""Order != 1 with an already-complex beam must skip the power-beam sqrt."""
za = np.linspace(0, 1, 5)
az = np.linspace(0, 2 * np.pi, 8)

AZ, ZA = np.meshgrid(az, za, indexing="xy")
beam_real = 1 - ZA**2
beam = np.array([beam_real + 1j * beam_real]).astype(np.complex128)
beam = beam[:, np.newaxis, np.newaxis] # (nbeam=1, nax=1, nfeed=1, nza, naz)
dza = za[1] - za[0]
daz = az[1] - az[0]

new_beam = gpu_beam_interpolation(
beam, [daz], [dza], [0.0], AZ.flatten(), ZA.flatten(), order=2
).get()

np.testing.assert_allclose(beam[0, 0, 0].flatten(), new_beam[0, 0, 0], atol=1e-6)

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Also not a fan of this test. I think for something that's testing "Does this take the square root when it's supposed to?", a better idea would be along the following lines:

  • Provide the test function with a real-valued e-field beam, its associated power beam, and a parametrized efield_or_power argument (or something similar) that says whether to do the interpolation on the e-field beam or the power beam
  • For both cases, do a no-op interpolation (i.e., evaluate the beam at the coordinates it already samples) and check that the returned beam is the same as the e-field beam

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Implemented essentially as you described — test_noop_interpolation_matches_input, parametrized over efield_or_power, evaluates each beam at exactly its own grid nodes and checks the output reproduces the input (with sqrt applied for the power case) (d8b2d05).

Comment thread tests/test_beam_interp_gpu.py Outdated
Comment on lines +135 to +162
def test_complex_beam_with_mismatched_output_dtype():
"""A complex beam interpolated into a differently-typed output buffer.

This forces the fused kernel to write into a scratch buffer of the beam's
own dtype and cast on copy, without applying sqrt (the beam is already
complex, e.g. an E-field beam, not a power beam).
"""
za = np.linspace(0, 1, 3)
az = np.linspace(0, 2 * np.pi, 6)
AZ, ZA = np.meshgrid(az, za, indexing="xy")

beam_real = 1 - ZA**2
beam = np.array([beam_real + 1j * beam_real]).astype(np.complex128)
beam = beam[:, np.newaxis, np.newaxis] # (nbeam=1, nax=1, nfeed=1, nza, naz)
dza = za[1] - za[0]
daz = az[1] - az[0]

az_flat, za_flat = AZ.flatten(), ZA.flatten()
out = cp.zeros((1, 1, 1, az_flat.size), dtype=np.complex64)

new_beam = gpu_beam_interpolation(
beam, [daz], [dza], [0.0], az_flat, za_flat, beam_at_src=out
).get()

assert new_beam.dtype == np.complex64
np.testing.assert_allclose(
new_beam[0, 0, 0], beam[0, 0, 0].flatten().astype(np.complex64), atol=1e-6
)

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This should probably be extended to test all pairs of input beam dtype and output beam dtype. Same idea as with earlier tests where the output should be compared to the provided e-field beam (but now cast to the same dtype as the output).

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The same no-op test above is also parametrized over out_dtype (complex64/complex128/None), covering the input × output dtype matrix for both the power and e-field cases (d8b2d05).

Comment thread tests/test_cublas.py

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Kind of weird deleting the old test_cublas.py file and making a new test_cublas_gpu.py file when the cublas tests are specifically for the GPU (i.e., the _gpu extension doesn't tell us anything new).

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

You were right that the filename itself carries no information — the actual reason was CI plumbing, not a naming preference: the self-hosted GPU job selected tests via pytest -k "gpu", a substring match against the filename, so a file without "gpu" in its name silently got zero GPU-runner coverage (which is exactly how this file lost coverage earlier in this PR). Fixed the root cause instead of just explaining it: registered a real pytest.mark.gpu marker, marked every GPU-only test explicitly, switched CI to pytest -m gpu, and renamed this back to test_cublas.py now that the filename doesn't need to encode "gpu" for CI selection. Verified the marker-based selection reproduces the old filename-based selection exactly (diffed both collection outputs, zero difference) (d8b2d05).

Steven Murray and others added 10 commits August 10, 2026 20:25
… beam-interp tests

Root-cause fix for a fragility discovered earlier in this PR: the self-hosted
GPU CI job selected tests via `pytest -k "gpu"`, a substring match against
test node IDs. That's why test_cublas.py/test_getz.py/test_precision.py had
to be renamed with a `_gpu` suffix to get GPU coverage at all -- a mechanism
that already caused one real, hard-to-spot coverage gap this session, and
that a reviewer flagged as needlessly coupling test *names* to CI plumbing.

- Register `gpu` as a real pytest marker (pyproject.toml).
- Mark every GPU-only test: whole-file `pytestmark` for files that are
  100% GPU (test_beam_interp_gpu.py, test_cpu_vs_gpu.py, test_cublas.py,
  test_getz.py, test_matvis_gpu.py, test_precision.py), and per-case
  `pytest.mark.gpu` on parametrize entries in the mixed files
  (test_matprod.py's GPU matprod methods, test_coordrot.py's
  GPUCoordinateRotationERFA cases via its existing `requires_gpu`
  attribute, test_wrapper.py's use_gpu=True case, test_beams.py's
  GPU-specific class/function).
- Switch the CI workflow from `pytest -k "gpu"` to `pytest -m gpu`.
- Rename test_cublas_gpu.py/test_getz_gpu.py/test_precision_gpu.py back to
  their natural names, now that the filename doesn't need to encode "gpu"
  for CI selection (a reviewer separately flagged this exact redundancy).

Verified exact parity: collected `pytest -m gpu` node IDs match the old
`pytest -k "gpu"` set 1:1 (after accounting for the three renames) via a
diff of both collection outputs -- zero difference.

Also reworks test_beam_interp_gpu.py per review feedback: replaces seven ad
hoc synthetic-array tests with fixtures built on the real e-field/power test
beam (conftest.py's uvbeam/uvbeam_unpol), a proper no-op-interpolation check
parametrized across the full input x output dtype matrix, and an order>1
cross-check against UVBeam.interp() at non-node points (the previous
grid-node check was a tautology any interpolator passes regardless of
order). Verified 100% coverage maintained on gpu/beams.py via the full
`-m gpu` suite.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…ixes

Per review feedback ("is it best practice to define C code as raw strings
within a python module?"): the three CUDA kernels (Hermitian mirror, fused
bilinear beam interpolation, fused Z matrix) move out of Python into
src/matvis/gpu/kernels/*.cu, loaded via Path(__file__).parent / "kernels" /
"<name>.cu" (matching the existing convention for src/matvis/data/*.fits --
a plain directory, no __init__.py, plain Path read rather than
importlib.resources).

Verified packaging end-to-end, not just assumed: built sdist and wheel both
contain the new kernels/ directory (confirmed via tar/unzip inspection), and
a --no-deps install into a clean venv resolves the installed package's
KERNELS_PATH correctly (i.e. this isn't just working by accident in an
editable install).

Also deletes src/matvis/gpu_src/{beam_interpolation,measurement_equation}.cu
-- found while doing this: dead code left over from the pre-cupy PyCUDA era,
untracked by anything, unrelated to any review comment, removed per
discussion rather than left to sit alongside the new kernels/ directory
under a confusingly similar name.

Collateral in the same files (both flagged in review):
- gpu/beams.py: gpu_beam_interpolation gains an explicit `power_beam: bool
  | None` parameter rather than inferring sqrt-or-not from beam.dtype alone.
  The dtype inference can't currently produce a wrong answer through the
  real pipeline (core.beams._wrangle_beams requires efield beams whenever
  polarized=True, and prepare_beam_unpolarized excludes cross-pols via
  include_cross_pols=False for the unpolarized path), but the function is
  public and called directly by tests/profiling scripts, so an explicit
  override closes the gap for future/direct callers rather than relying on
  an invariant enforced two call-frames away.
- gpu/_cublas.py, gpu/beams.py, gpu/getz.py: fixed block sizes (256, 128)
  now have a comment stating plainly they're conventional warp-multiple
  defaults, not empirically tuned -- a reviewer asked whether there was
  hidden precision behind the choice; there wasn't.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Reply-only items (no bug, but the diff alone didn't explain why): added a
comment at each of the four `Device().synchronize()` removal sites
(core/coords.py, cpu/coords.py x2, core/getz.py, gpu/matprod.py) stating why
it's safe -- gpu.simulate() now runs the whole per-time, per-chunk loop on a
single persistent stream, so kernel launches are already FIFO-ordered
without a device-wide sync; the old per-chunk-stream design needed the sync
for cross-stream ordering, this one doesn't.

wrapper.py: reverted an accidental "the other option" -> "one other option"
wording change (there's still exactly one other matprod_method option, so
"the" was correct). Also confirmed via git diff that CPUVectorDot/
GPUVectorDot were already the real class names on main before this PR --
wrapper.py's stale `VectorLoop` type hints/docstring were themselves the
bug being fixed, not a new rename.

cli.py: (a) reworded the vague "is indicative only" comment on the `stages`
table to state plainly why it's unreliable for the GPU backend (host waits
on already-queued async work); (b) added a comment explaining the
host_overhead `max(..., 0.0)` clamp -- gpu_time is a median-times-count
estimate that can cross the wall-time sum when per-chunk costs vary (e.g.
horizon-cut source counts differing between chunks), so the clamp isn't
arbitrary defensiveness.

_utils.py: expanded a terse comment ("nbeampix is already summed over beams
by the caller") to be self-contained rather than requiring PR archaeology to
understand why nbeampix isn't multiplied by nbeam.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
- roofline.py: fixed non-uniform sphere sampling in the beam-interp
  micro-benchmark (za was sampled uniformly in [0, pi/2]; correct sphere
  sampling needs cos(za) ~ Uniform, i.e. za = arccos(uniform(0,1))). Also
  wording: dropped "speed-of-light"/"micro-" (contradicted the
  "production-scale" framing a few lines later), "at the exact Z-matrix
  shape" -> "at the specified simulation size" (shape depends on the run,
  there's no single exact one), and "at nbeam=nant" -> "one beam per
  antenna by default (--nbeam to override)" since it's a default, not a
  hardcoded fact. Added a one-line reminder of what Z is for readers
  landing here without prior context.

- gemm_experiments.py: added an RMS relative-error column alongside the
  existing max-relative-error one -- max-abs is a fast worst-case eyeball
  check, RMS is what actually distinguishes "different but valid rounding"
  (e.g. cgemm3m's Gauss decomposition) from "wrong"; noted in a comment that
  neither is the real correctness gate (that's test_cublas.py's
  assert_allclose). Moved the cuBLAS function-signature doc comments down
  to sit directly above the .argtypes assignments that use them, instead of
  living separately at the top of the file. Renamed local ctypes aliases
  `ptr, i32` to `_PTR, _INT` to match _cublas.py's naming.

- run-canonical.sh: reworded the config-description comment so it can't be
  read as claiming the small `dev` config is production *scale* -- only the
  *settings* (polarized, gridded beams, one beam per antenna) match
  production; size is what differs between dev and prodslice.

- README.md: named both canonical configs explicitly instead of only
  describing one; expanded the `derived` field descriptions (what "measures
  the card" actually means, what host_overhead covers) rather than mostly
  restating the field names; replaced "back-pressure point" with the
  reviewer's own correct plain-English paraphrase; reworded roofline.py's
  description away from ad-copy phrasing into a plain statement of what it
  isolates and how to read a gap between it and a real run; added a link to
  the Performance docs page and a forward-reference to the new CLI
  reference page.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Multiple review comments asked for a reference page for `matvis profile`'s
parameters (there was none) plus a concrete example of the JSON output
(field descriptions alone don't tell you what a real number looks like).

- New docs/cli.rst, auto-generated from the actual click commands via
  sphinx-click (`.. click:: matvis.cli:main`) -- self-maintaining, since it
  reads the real --help text rather than a hand-written copy that drifts.
  Added to the docs/index.rst toctree.
- Added `help=` text to every click.option in cli.py that was missing it
  (most of them), and real docstrings for the `profile`/`hera-profile`
  commands and the top-level group (previously just "Run the script."),
  since sphinx-click renders these directly -- a reference page is only as
  good as the strings it's built from.
- docs/cli.rst also includes a real (trimmed), annotated example of a
  summary-stats-*.json output, walking through what `derived`, `run_stats`,
  and `stages` each contain and when to trust which.

Two RST gotchas fixed along the way (caught by an actual sphinx build, not
assumed): unescaped `*` in "12*nside**2" and "summary-stats-*.json" inside
click help strings were parsed as emphasis markup by the generated docs,
producing "Inline emphasis start-string without end-string" warnings;
reworded both to avoid stray asterisks.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
- Made the "theoretical minimum GPU time" connection explicit: the existing
  t_gemm formula already computes this once you plug in a measured R: said
  so directly instead of leaving the reader to infer it.
- Added a paragraph to "Memory and chunking" on the raw beam-grid
  subtlety a reviewer raised: that term doesn't scale with chunk size, so it
  can dominate total memory for small chunks and become negligible for
  large ones -- worth knowing when tuning min_chunks/memory_buffer on a
  beam-heavy, memory-constrained GPU.
- Split the PR #130 changelog entry's semicolon-chained sentence into an
  actual bulleted list; replaced "single fused bilinear beam-interpolation
  kernel" with a plainer description and defined "fused" on first use;
  disambiguated "no device syncs in the loop" to name which loop.
- Dropped "speed of light" (now "roofline micro-benchmarks", the standard
  term for this kind of peak-achievable-rate benchmarking).
- Fixed a bare `issue #133` reference to actually render as a link, like
  the other issue references on the page.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…nce.rst

Missed this instance of the "times" == multiplication ambiguity in the
previous performance.rst wording pass; profiling/README.md and docs/cli.rst
already used "×" for the same sentence.

Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
@steven-murray

Copy link
Copy Markdown
Contributor Author

@r-pascua thanks VERY much for the detailed review. I think all the questions have been answered. To be honest, I set Claude on your review, and advised it to refactor the inline CUDA code into external files, as well as to fix the way that GPU tests were collected, and to focus on humanizing the documentation.

It did OK at humanizing but not perfectly. I performed a last gloss of the performance.rst myself and removed some stuff that seemed obnoxious.

My reply comments to you were generated by Claude but I read each one individually and amended them if they were too annoying, and some I replaced wholesale with my own comments.

Thanks again -- your comments have definitely made this PR much better already.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

enhancement New feature or request

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants