perf: overhaul GPU hot path (7.7x at HERA scale) + revive parity tests - #130
perf: overhaul GPU hot path (7.7x at HERA scale) + revive parity tests#130steven-murray wants to merge 27 commits into
Conversation
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 Report✅ All modified and coverable lines are covered by tests. 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
Flags with carried forward coverage won't be shown. Click here to find out more. ☔ View full report in Codecov by Harness. 🚀 New features to boost your workflow:
|
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>
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>
…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>
| .. 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 |
There was a problem hiding this comment.
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.
There was a problem hiding this comment.
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).
There was a problem hiding this comment.
I'm not quite sure this answers what I was asking. By "theoretical minimum," I meant "If the GPU was used 100% efficiently (i.e., the realized FLOPs are the same as the GPU's maximum FLOPs according to the device specifications, and the kernels are compute limited instead of bandwidth limited), then how long would the 'canonical run' take?".
| finally: | ||
| cublas.setPointerMode(handle, orig_mode) |
There was a problem hiding this comment.
Can we add a simple comment here explaining why we always run cublas.setPointerMode? (And in other places where this try: finally structure appears)
| Whether ``beam`` holds power (non-negative, needs a ``sqrt`` before | ||
| use as a voltage) rather than E-field values. Callers that know this | ||
| (e.g. ``GPUBeamInterpolator``, which knows ``polarized``) should | ||
| pass it explicitly. If not given, it's inferred from ``beam``'s | ||
| dtype (real => power, complex => E-field) — a reasonable default | ||
| for a real beam, but note that a complex *power* beam (e.g. one | ||
| holding cross-polarization terms) would be mis-detected as E-field | ||
| by that inference and silently skip the ``sqrt``. |
There was a problem hiding this comment.
I think this can be phrased a bit more clearly. I'd suggest something like
Whether the provided ``beam`` is in power units or E-field units. If not provided,
then it is inferred based on whether the provided ``beam`` is real- or complex-valued.
Failing to set ``power_beam=True`` and providing a power beam with cross-polarized
components will result in the interpolation routine treating the beam as if it were an
E-field beam instead of a power beam (i.e., no square root will be taken after interpolation).
| block = 256 | ||
| rdtype = np.float32 if self.ctype == np.complex64 else np.float64 | ||
| sqrt_flux = cp.ascontiguousarray(sqrt_flux, dtype=rdtype) | ||
| assert beam._c_contiguous and exptau._c_contiguous |
| See :meth:`matvis.core.getz.ZMatrixCalc.__call__` for parameters. | ||
| Unlike the base implementation, ``exptau`` is not modified in place. |
There was a problem hiding this comment.
Rather than "See ...", can we just use the combine_docstring decorator to copy the parameters over?
| "GPUCoordinateRotationERFA", | ||
| ] = "CoordinateRotationAstropy", | ||
| matprod_method: Literal["GPUMatMul", "GPUVectorLoop"] = "GPUMatMul", | ||
| matprod_method: Literal["GPUMatMul", "GPUVectorDot"] = "GPUMatMul", |
There was a problem hiding this comment.
I know that I brought this up somewhere else, but isn't this an API change? I believe Claude made some argument about this being changed in something that preceded this PR, but... it's changed here?
| simulate.__doc__ = ( | ||
| (simulate.__doc__ or "") | ||
| + f"\n{simcpu.__doc__ or ''}" | ||
| + """ | ||
| gpu_event_timing : bool, optional | ||
| If True, collect per-chunk GPU event timings for beam interpolation, | ||
| tau, Z construction, and matprod stages; log stage medians at INFO | ||
| level at the end of the run, and expose median/mean/std/count per stage | ||
| (plus per-integration wall times and a warmup-robust | ||
| ``steady_time_per_integration``) via ``LAST_RUN_STATS``. Default is | ||
| False. | ||
| """ | ||
| ) |
There was a problem hiding this comment.
Is there an equivalent way to do this with combine_docstrings?
| Neither class here calls ``cp.cuda.Device().synchronize()`` (earlier | ||
| versions did, after every ``compute``/``sum_chunks``). See the comment in | ||
| ``core.coords.CoordinateRotation.select_chunk`` for why: gpu.simulate() runs | ||
| the whole loop on a single persistent stream, so kernel launches are already | ||
| ordered without a device-wide sync. |
There was a problem hiding this comment.
See my comment about the comment in core.coords.CoordinateRotation.select_chunk.
| # compilation, cuBLAS handle+workspace creation, ERFA/IERS caches) | ||
| # are paid before any timing starts. | ||
| nwarm = min(nsource, 10_000) | ||
| logger.info("Running warmup simulation (%d sources, 1 time)...", nwarm) |
There was a problem hiding this comment.
Why are we not using an f-string here?
| # Derived headline numbers, robust to warmup and host noise. These are | ||
| # the values to quote/compare (see the docs Performance page). | ||
| # | ||
| # The `stages` table further down times Python lines, but GPU work is | ||
| # queued asynchronously -- 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. Use these derived numbers for the GPU backend | ||
| # instead. |
There was a problem hiding this comment.
Does this comment need to live here? It's a bit long, and it's a copy-paste of what's already in the rtd. Can we replace it with something like # Initialize a dictionary to track timing summary statistics.?
| # gpu_time is median(per-chunk total) x nchunks -- a biased | ||
| # estimator of the true per-integration GPU total when chunks | ||
| # vary in cost (e.g. the horizon cut leaves different numbers of | ||
| # sources active in different chunks), whereas steady wall time | ||
| # sums the *actual* per-integration total. The two can cross | ||
| # without either being wrong, hence the clamp below. |
There was a problem hiding this comment.
Similar comment to above. I'd recommend adding a comment above line 238 that is along the lines of # Estimate the per-integration GPU time based on the median per-chunk GPU time.
| print(f"Beam {i}", b2) | ||
| assert np.allclose(np.sqrt(b1.flatten()), b2) | ||
| expected = np.sqrt(d0) if is_power else d0 | ||
| expected = expected.transpose(1, 0, 2, 3).reshape(nfeed, nax, nza, naz) |
There was a problem hiding this comment.
A comment somewhere explaining the expected shape of d0 would help in understanding whether this operation actually makes sense. Maybe in prepare_for_map_coords?
| [azmin], | ||
| cp.asarray(AZ.flatten()), | ||
| cp.asarray(ZA.flatten()), | ||
| order=2, |
There was a problem hiding this comment.
Might be better to assign the spline interpolation order to a variable since it's used in two places? Especially considering that we're planning on implementing a higher order interpolation kernel in the future.
There was a problem hiding this comment.
Potential tests to add:
- Consistency check between GPU interpolation kernel results and linear order
map_coordinatesinterpolation. - Check that out-of-bounds interpolation points are assigned the boundary beam values.
- Explicit check that power beams are cast to complex beams of equivalent precision by interpolation routine.
May need to think more carefully about whether this provides complete meaningful coverage.
| cb.zdotz(cp.asarray(a), out=out) | ||
| cb.zdotz(cp.asarray(a), out=out, beta=1.0) | ||
| np.testing.assert_allclose(out.get(), 2 * expected, rtol=1e-4) |
There was a problem hiding this comment.
Why doesn't this just do
cb.zdotz(cp.asarray(a), out=out, beta=1.0)
np.testing.assert_allclose(out.get(), expected, rtol=1e-4)
|
|
||
| monkeypatch.setattr(cb, "_LIB", None) | ||
| c = cb.zdotz(cp.asarray(a)) | ||
| np.testing.assert_allclose(c.get(), np.dot(a.conj(), a.T), rtol=1e-4) |
There was a problem hiding this comment.
Wouldn't it be cleaner to instead do
expected = cb.complex_matmul(cp.asarray(a), cp.asarray(a))
np.testing.assert_equal(c.get(), expected.get())
| class FakeCDLL: | ||
| def __init__(self, name): | ||
| attempted.append(name) | ||
| if len(attempted) < 3: |
There was a problem hiding this comment.
With the current version of things, this will be fine, but there is potential for this test to fail if, for some reason, _load_cublas_ext searches for 2 or fewer libraries. Probably not an actual issue, but could be a confusing test failure in the future if it ends up breaking.
One option is to update the _cublas module to define a _libcublas_sonames list, adjust the conditional on this line to if len(attempted) < max_attempts: (with max_attempts defined in the preceding line), then check that len(attempted) == min(max_attempts, len(cb._libcublas_sonames)) at the end of this test function.
| def test_load_cublas_ext_returns_none_if_all_sonames_fail(monkeypatch): | ||
| """_load_cublas_ext should return None (not raise) if no soname loads.""" | ||
|
|
||
| class FakeCDLL: | ||
| def __init__(self, name): | ||
| raise OSError(f"cannot load {name}") | ||
|
|
||
| class FakeCtypes: | ||
| CDLL = FakeCDLL | ||
|
|
||
| monkeypatch.setattr(cb, "ctypes", FakeCtypes()) | ||
| assert cb._load_cublas_ext() is None |
There was a problem hiding this comment.
Isn't this checking one of the lines that has the # pragma: no cover flag?
| def _random_complex(rng, shape, dtype): | ||
| r = rng.standard_normal(shape) | ||
| i = rng.standard_normal(shape) | ||
| return (r + 1j * i).astype(dtype) |
There was a problem hiding this comment.
This seems like it would be a helpful function to expose to other test modules. Can it be easily accessed by other modules if it is placed in the conftest module?
| def test_fused_z_shared_beam(): | ||
| """beam_idx=None with a single beam shared by all antennas.""" | ||
| rng = np.random.default_rng(1) | ||
| nant, nfeed, nax, nsrc = 5, 2, 2, 8 | ||
| ctype = np.complex64 | ||
|
|
||
| beam = _random_complex(rng, (1, nfeed, nax, nsrc), ctype) | ||
| exptau = _random_complex(rng, (nant, nsrc), ctype) | ||
| sqrt_flux = rng.standard_normal(nsrc).astype(np.float32) | ||
|
|
||
| calc = GPUZMatrixCalc(nant=nant, nfeed=nfeed, nax=nax, nsrc=nsrc, ctype=ctype) | ||
| calc.setup() | ||
| z = calc(cp.asarray(sqrt_flux), cp.asarray(beam), cp.asarray(exptau), None) | ||
|
|
||
| expected = _reference_z(beam, exptau, sqrt_flux, None, nant, nfeed, nax, nsrc) | ||
| np.testing.assert_allclose(z.get(), expected, rtol=1e-5) |
There was a problem hiding this comment.
Can't this be combined with the previous test through an appropriate parametrization on nbeam? (Although, if updated, it should be done in a more intuitive way via a parameter called all_beams_unique or something.)
|
|
||
|
|
||
| def test_fused_z_beam_per_antenna_implicit(): | ||
| """beam_idx=None with one beam per antenna, aligned to antenna order.""" |
There was a problem hiding this comment.
This explanation doesn't feel very helpful. Also, it looks like this is just testing that nothing is shuffled when no beam indexing is provided? Surely that can be absorbed into the earlier test with a parameter.
There was a problem hiding this comment.
I think one test case that is missing from this is running a check when 1 < nbeam < nant. While this is probably an unusual use case, I think it is still an important one to cover.
| assert stats["matprod"]["mean"] > 0 | ||
|
|
||
| # Per-integration wall times and the warmup-robust steady-state metric. | ||
| assert len(LAST_RUN_STATS["integration_times"]) == 2 |
There was a problem hiding this comment.
Flagging another hardcoded variable that is referenced twice.
| # Runs both use_gpu=True and use_gpu=False cases; whole-file marking keeps | ||
| # both riding along together on the self-hosted GPU CI job, matching prior | ||
| # (filename-based) selection exactly. |
| # The error budget is fp32 round-off accumulated over the coherent source | ||
| # sum: relative to the total flux scale (~|V| at zero spacing), not to each | ||
| # individual (possibly near-zero) visibility. |
There was a problem hiding this comment.
Some parts of this comment feel a little oddly worded. Can't the (~|V| ...) parenthetical just be restated to say "the autocorrelation amplitude"? Also doesn't this set a tolerance that is quite a bit larger than the expected accumulated roundoff error? I'm just guessing here, but it feels like something along the lines of scale = max_source_flux * n_src would be closer to the right thing.
There was a problem hiding this comment.
Are these instructions to AI agents? Just curious
| over all (beam, feed, axis) planes instead of one ``map_coordinates`` | ||
| launch per plane. |
There was a problem hiding this comment.
Maybe replace "plane" with "combination"
| - The GPU loop runs on a single stream with no device synchronization, | ||
| keeping the GPU ~95% utilized. |
There was a problem hiding this comment.
I think this utilization number is one of the things I was looking to see reported in the new "Performance" section of the docs. Right now it is currently hidden in the changelog, and I think it is important enough to be highlighted somewhere earlier in the page.
| single precision is requested (this also removes a large hidden | ||
| temporary that could cause out-of-memory errors). |
There was a problem hiding this comment.
"large hidden temporary that" seems like it's missing a word
| - The profiling harness is robust to one-time costs and host noise: an | ||
| untimed warmup simulation runs first (``--no-warmup`` to disable), | ||
| per-integration wall times are recorded individually, CUDA-event stage | ||
| timings report medians as well as means, and a ``derived`` block in the |
r-pascua
left a comment
There was a problem hiding this comment.
OK, I think I have gotten around to all of the responses and completed another look through all of the changes. I think this is very close to being done, just a few more tests I would like to see and some tidying up of comments/docs.
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 withprofiling/run-canonical.sh.Headline numbers (per-chunk CUDA-event timing)
After the change, ~70% of GPU time is the cuBLAS
cherkkernel 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:
nbeam·nfeed·nax= 1400 separatemap_coordinateslaunches per chunk, plus per-beam coordinate-array allocation.cupy_backendscublas calls never set the handle's stream), serializing against the compute stream.antpos * 1jpromoted the phase-factor matmul to double precision even atprecision=1, adding an fp64 GEMM + cast per chunk and a large temporary that caused OOMs.Device().synchronize()in the Z stage was what kept it correct, and it stalled the pipeline.Changes
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) usecublasCgemm3m(measured 2.1× overcgemmat matvis shapes). Both bound fromlibcublasvia ctypes (cupy doesn't expose them), running on the current cupy stream. Falls back tocgemm/zgemmif the library can't be bound.gpu/beams.py): a single launch over all (beam, feed, axis) planes and sources;map_coordinatesremains the fallback fororder > 1.gpu/getz.py): Z = A·√I·exptau in one elementwise pass (with thebeam_idxgather), replacing 4 broadcast copies + a Python loop over all antennas + a full-device sync per chunk.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 viaLAST_RUN_STATSfor the profiler CLI.core/tau.py): keep the phase matmul at the requested precision.Bugs fixed
tests/test_cpu_vs_gpu.pystill guarded onimportorskip("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:beam_data.set()dtype mismatch).gpu.pycomputed its ownnsrc_alloc, disagreeing withCoordinateRotation(which ignoressource_bufferfor chunks ≤ 1000 sources) → shape-mismatch crashes in small chunked runs.kx/kyspline options in the parity test (pyuvdata'smap_coordinatesinterpolator takesorder).Validation
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.pyextended to rectangular shapes, both dtypes, andout=/beta=accumulation (exercises the herk + mirror path).map_coordinatesat fp32 tolerance.Follow-up work
Filed as issues: #131 (validation on V100/ilifu), #132 (
sum_chunksaccumulation/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
🤖 Generated with Claude Code