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).
| 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 |
There was a problem hiding this comment.
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
There was a problem hiding this comment.
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).
| matvis profile -a 350 -b 350 -s 1000000 -t 4 --gpu \ | ||
| --interpolated-beam --single-precision --gpu-event-timing \ | ||
| --coord-method CoordinateRotationERFA -o outdir |
There was a problem hiding this comment.
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.
There was a problem hiding this comment.
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).
| 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. |
There was a problem hiding this comment.
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.
There was a problem hiding this comment.
| 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 |
There was a problem hiding this comment.
"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).
There was a problem hiding this comment.
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.
| 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 |
There was a problem hiding this comment.
Update the issue reference to a link to the issue if possible.
| # 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 |
There was a problem hiding this comment.
"y covers planes" means what? It looks like y is just parallelizing over beam components, while x is parallelizing over sources.
There was a problem hiding this comment.
Yup, reworded to "x indexes sources, y indexes (beam, feed, axis) combinations" (1ed75e6).
| _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); } | ||
| } | ||
| """ | ||
| ) |
There was a problem hiding this comment.
Same comment as before about C code as raw strings living in a python module.
There was a problem hiding this comment.
And same fix :-)
| 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; } |
There was a problem hiding this comment.
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.
| # 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). |
There was a problem hiding this comment.
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.".
There was a problem hiding this comment.
Used your suggested wording exactly (fc5b8da).
| dza = cp.asarray(dza, dtype=rdtype) | ||
| azmin = cp.asarray(azmin, dtype=rdtype) | ||
| assert beam._c_contiguous and target._c_contiguous | ||
| block = 128 |
There was a problem hiding this comment.
Another comment about a seemingly arbitrary choice for block size. Was there some optimization done that resulted in this choice?
There was a problem hiding this comment.
Same answer: conventional warp-multiple default (128 here), not tuned. Commented accordingly (fc5b8da).
| _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); } | ||
| } | ||
| """ |
There was a problem hiding this comment.
Marking another section of C code declared in a python module.
There was a problem hiding this comment.
Same action taken.
| ) | ||
|
|
||
| ntot = self.nant * self.nfeed * self.nax * self.nsrc | ||
| block = 256 |
There was a problem hiding this comment.
Block size setting again
There was a problem hiding this comment.
Same comment again :-)
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>
| gpusize = { | ||
| "antpos": nant * 3 * rsize, | ||
| "flux": nsrc * rsize, | ||
| # nbeampix is already summed over beams by the caller |
There was a problem hiding this comment.
Is this actually a helpful comment?
There was a problem hiding this comment.
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)
| # 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. |
There was a problem hiding this comment.
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"?
There was a problem hiding this comment.
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).
| derived["host_overhead_per_integration"] = max( | ||
| derived["steady_wall_per_integration"] - gpu_time, 0.0 | ||
| ) |
There was a problem hiding this comment.
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?
There was a problem hiding this comment.
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).
| 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, |
There was a problem hiding this comment.
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?
There was a problem hiding this comment.
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).
| if not complex_beam: # power beam | ||
| cp.sqrt(beam_at_src, out=beam_at_src) | ||
| beam_at_src = beam_at_src.astype(ctype) |
There was a problem hiding this comment.
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?
There was a problem hiding this comment.
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).
There was a problem hiding this comment.
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.
There was a problem hiding this comment.
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).
| 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 | ||
| ) |
There was a problem hiding this comment.
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)?
There was a problem hiding this comment.
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).
| 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) |
There was a problem hiding this comment.
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_powerargument (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
There was a problem hiding this comment.
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).
| 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 | ||
| ) |
There was a problem hiding this comment.
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).
There was a problem hiding this comment.
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).
There was a problem hiding this comment.
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).
There was a problem hiding this comment.
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).
… 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>
|
@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. |
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