High-Performance GPU Building Blocks for Radio Astronomy (and beyond)

High Performance GPU building blocks for radio astronomy (and beyond) John W. Romein, Bram Veenboer CASPER Workshop September 8th–12th, 2025 2 Outline •“radio blocks”: –tensor-core correlator –tensor-core beam former –GPU filter •solving the I/O challenge –DPDK demo correlator 3 RADIOBLOCKS 4 RADIOBLOCKS RADIOBLOCKS noun, plural UK: /ˈreɪ.di.əʊ.bloks/ US: /ˈreɪ.di.oʊ.bla:ks/ 0.1 name of an EU-funded project 0.2 building blocks from which one can build a radio telescope 0.3 efficient (GPU) libraries for signal processing 5 RADIOBLOCKS RADIOBLOCKS noun, plural UK: /ˈreɪ.di.əʊ.bloks/ US: /ˈreɪ.di.oʊ.bla:ks/ 0.1 name of an EU-funded project 0.2 building blocks from which one can build a radio telescope 0.3 efficient (GPU) libraries for signal processing not our only funding agency! ESO, NLeSC, NWO 6 RADIOBLOCKS RADIOBLOCKS noun, plural UK: /ˈreɪ.di.əʊ.bloks/ US: /ˈreɪ.di.oʊ.bla:ks/ 0.1 name of an EU-funded project 0.2 building blocks from which one can build a radio telescope 0.3 efficient (GPU) libraries for signal processing 7 RADIOBLOCKS •1-3-'23 – 28-2-'27 •work packages: WP2 novel detectors and components WP3 digital receivers WP4 data transport and correlation WP5 data processing tool kit for advanced radio astronomy 000 8 000 RADIOBLOCKS •1-3-'23 – 28-2-'27 •work packages: WP2 novel detectors and components WP3 digital receivers WP4 data transport and correlation WP5 data processing tool kit for advanced radio astronomy 9 WP4 data transport and correlation station processing combine data (real time) calibration, imaging, etc. disk station processing station processing “correlator” WAN this talk's focus 16 GPU Tensor Cores •tensor core compute power increases rapidly source: NVIDIA tensor cores 17 GPU tensor cores •performs 16x16 * 16x16 multiplication in one go •decompose large matrix multiplications •use cuBLAS, pyTorch, …, or program directly x B CA = + 32 bitsD16/8/4 bits 16/8/4 bits 32 bits 18 Use tensor cores for signal processing? •yes, if: –algorithm translates to matrix-matrix multiplications •correlator: ✔ •beam former: ✔, ✘ •FIR filter: ? •DFT: ✔ •FFT: ✘ (radix 8: ✔) –operates on few bits ✔ 19 The Tensor-Core Correlator1 •GPU library –performs (tensor-core) computations –not full application (no I/O etc.) –highly optimized –hides nasty details •open source2 •rapidly adopted by radio telescopes worldwide 1) J.W. Romein, The Tensor-Core Correlator, A&A 656(A52), Dec 2021 2) https://git.astron.nl/RD/tensor-core-correlator 20 Implementation challenges 1) complex numbers 2) triangular output format 3) fast data fetching 4) ptx assembly hacking •all hidden from the user not supported by tensor cores 21 Complex-valued matrix multiplication c = a * conj(b) = (ar + iai) * conj(br + ibi) = (arbr + aibi) + i(aibr - arbi) ⋯⋯⋯⋯arai⋯⋯⋯⋯ crci ⋮⋮ ⋮⋮ ⋮⋮ ⋮⋮ br-bi bibr ⋮ ⋮ ⋮ ⋮ ⋮⋮ ⋮⋮ X = 22 Complex-valued matrix multiplication •copy, conjugate, permute columns in B matrix –(or rows in A matrix) r0i0r1i1r2i2⋯ ⋯ r7i7 crci r0-i0 i0r0 r1-i1 i1r1 r2-i2 i2r2 ⋮ ⋮ ⋮ ⋮ r7-i7 i7r7 X = 23 Triangles •result is Hermitian compute and store 1 side→ •tensor cores: rectangles only •entries near diagonal treated differently •minimize redundant computations •storing results requires ugly hacks 64x64 receivers 24 Runtime compilation •compile GPU code at runtime –target GPU known –many variables constants→ –improves code readability greatly better performance 25 What's new? •new GPUs: Ada, Hopper, Blackwell (sm120) •new data types: fp8 (e4m3, e5m2) & fp4 (e2m1) –fp8 disabled on Ada inaccurate← •optimizations 32 Possible future improvements •support for tcgen05 (= faster tensor cores on some Blackwell GPUs) –significant performance improvement –lot of work –not forward compatible •support for single-pol input •support for XX, YY output only •AMD GPU support –only makes sense if filter supports AMD (see later) let me know what you need→ 33 The Tensor-Core Beam Former 34 The new GPU filter 35 GPU filter library •new library: filter •connects to TCC, TCBF • receivers Filter TCBF TCC 36 GPU filter library ●filter splits frequency band into channels ●also: delay & bandpass corrections, transpose PolyPhase Filter (PPF) bank receivers Delay Compensation FFTFIR filter Bandpass Correction Transpose Correlator Beam Former 37 Filter input •i4, i8, i16, fp4, fp8, fp16, fp32 input •real complex, complex complex→ → •supports any data layout: provide your own access funtion (optional) PolyPhase Filter (PPF) bank receivers Delay Compensation FFTFIR filter Bandpass Correction Transpose Correlator Beam Former 38 PolyPhase Filter bank •reduces channel leakage •0 or 16 FIR filter taps •64, 256 channels •weights can be generated by library •mirror channels (optional); fft shift (optional) PolyPhase Filter (PPF) bank receivers Delay Compensation FFTFIR filter Bandpass Correction Transpose Correlator Beam Former : wide range 39 cuFFTDx •uses cuFFT cuFFTDx library –call FFT from CPU GPU kernel –FFT and other tasks in single GPU kernel read & write data →3x 1x –significant performance improvement –AMD GPUs PolyPhase Filter (PPF) bank receivers Delay Compensation FFTFIR filter Bandpass Correction Transpose Correlator Beam Former 40 Delay compensation •fine delays (optional) –phase correction –nth-order polynomial PolyPhase Filter (PPF) bank receivers Delay Compensation FFTFIR filter Bandpass Correction Transpose Correlator Beam Former 41 GPU filter library •amplitude correction (optional) –remove ripple from previous filter PolyPhase Filter (PPF) bank receivers Delay Compensation FFTFIR filter Bandpass Correction Transpose Correlator Beam Former 48 I/O Challenge 49 The I/O challenge ●tensor cores made GPU correlators 100x faster than a decade ago ●I/O must scale proportionally ●receiving >40 Gb/s data: prohibitive OS overhead –wait for faster processors –need new I/O techniques 50 Motivation for this work •higher data rates ↔ more powerful telescopes ↔ more astronomical science •for example: Atacama Large Millimeter Array (ALMA) GPU correlator study –63 Tb/s 51 Data flow ●digitizer network (corner turn) correlator→ → ●typically: UDP ●stream data directly FPGA GPU→ 52 Correlator processing ring buffer correlatechannel filter GPU CPU or GPU Ethernet packet processing 53 Grace Hopper •powerful GPU •7x more CPU GPU bandwidth than PCIe gen 5→ •up to 3 400GbE NICs source: NVIDIA 54 Exploring potential solutions ●Remote DMA (RoCE v2) ●DOCA GPUNetIO ●DOCA DPA ●Data Plane Development Kit ●(Holoscan ANO: on top of these) ●implemented demo correlator1 different advantages & disadvantages 1) J.W. Romein, Breaking the I/O Barrier: 1.2 Tb/s Ethernet Packet Processing on a GPU. Euro-Par'25, Dresden, Aug. 27-29, 2025 55 RDMA vs DPDK RDMA DPDK protocol complex any Ethernet type FPGA firmware complex less complex CPU load & energy use low few cores polling IP multicast no yes session setup yes no secure yes no superuser privileges needed no yes 56 Exploring potential solutions ●Remote DMA (RoCE v2) ●DOCA GPUNetIO ●DOCA DPA ●Data Plane Development Kit ●(Holoscan ANO: on top of these) ●implemented demo correlator1 different advantages & disadvantages 1) J.W. Romein, Breaking the I/O Barrier: 1.2 Tb/s Ethernet Packet Processing on a GPU. Euro-Par'25, Dresden, Aug. 27-29, 2025 57 DPDK: Data Plane Development Kit •toolkit for high-bandwidth, low-latency Ethernet communication •well supported by major NIC vendors •NIC under application control –user space no OS overhead→ •API: send/receive packets, allocate/free packet buffers •protocol stack (UDP/IP, ICMP, …): application responsibility 64 Are we happy with the results? •DPDK GPUdev approach: ☺ exceptionally high data rates  CPU + GPU overhead  0-copy difficult to integrate into existing correlator software •future work: explore alternatives 66 Conclusion •radio blocks: high-performance signal processing library –Tensor-Core Correlator –Tensor-Core Beam Former –GPU filter •demonstrated 1.2 Tb/s FPGA GPU communication→ –DPDK –fast, but with some CPU and GPU overhead 67 Acknowledgements •This work was funded by –European Southern Observatory (ALMA GPU Correlator Study) –EU HORIZON INFRA-TECH (RADIOBLOCKS, Grant Agreement nr. 101093934) –Netherlands eScience Center (PADRE, RECRUIT) –NWO Netherlands Foundation for Scientific Research (DAS-6)