Alice Weekly Meeting: Software for Hardware Accelerators
-
-
10:00
→
10:20
Discussion 20mSpeaker: David Rohr (CERN)
Color code: (critical, news from this week: blue, news from last week: purple, no news: black)
Sync reconstruction
Async reconstruction
- Need to investigate short GPU stall problem.
- Limiting factor for pp workflow is now the TPC time series, which is to slow and creates backpressure (costs ~20% performance on EPNs). Enabled multi-threading as recommended by Matthias - need to check if it works.
- Will tune existing 16-core settings, add a SITEARCH for 16core CPU, and 16coreCPU + generic NVIDIA / AMD GPU, like for 8 core.
- Will retune EPN async workflow for TPC + ITS on GPU on 2025 data.
GPU ROCm / compiler topics:
- Problem with building ONNXRuntime with MigraphX support.
- Need to find a way to build ONNXRuntime with support for CUDA and for ROCm.
- Try to find a better solution for the problem with __device__ inline functions leaking symbols in the host code.
- All gpu-ci related issues fixed, it is now compiling for the correct architectures and it is green.
- Added OpenCL build to gpu-system.sh in build mode.
- ROCm 10.1 seems as good as ROCm 10 on MI50.
- SLC10-FullCI added by Giulio/Vit: currently failing since gpu-system's build option does not provide CUDA stub libraries, if we don't bring host libraries to the container. Need to check how to fix it without colliding if we bring host libraries tot he container.
- Bumped to CUDA 13.4 and ROCm 10.1, we can now bump to GCC 16.2.1
TPC / GPU Processing
- WIP: Use alignas() or find a better solution to fix alignment of monte carlo labels: https://its.cern.ch/jira/browse/O2-5314
- Need to check the problem with ONNX external memory allocator.
- Next high priority topic: Improvements for cluster sharing and cluster attachment at lower TPC pad rows. PR: https://github.com/AliceO2Group/AliceO2/pull/14542
- Check for unnecessary f64 instructions in GPU code.
- Added CPUDevice backend for debugging (Runs on CPU, but in the way we normally run on the device).
Other topics:
- GCC bump - Status?
EPN GPU Topics:
- Discuss with Ernst / Giulio how to do a build for SLC10 on EPNs with build mode.
- Need to test Maksim's SLC9 container with ROCm included on SLC10 node, Giada fixed CVMFS for me :).
-
10:20
→
10:25
TPC ML Clustering 5mSpeaker: Christian Sonnabend (CERN)
Neural network seeding vector (https://github.com/AliceO2Group/AliceO2/pull/15756)
- Found a bug in the implementation (missing setting during application, seeding vectors were never used) -> Fixed
- Results don't change much. 0-100% centrality is mostly unchanged, but one improvement can be seen in 0-5% centrality enforced simulation


Is it worth the complications? 🤷♂️
ROOT: TMVA SOFIE inference of .onnx models on the EPNs
- Discussion during ROOT Q3 meeting last week: ALICE has great interest to give SOFIE a try. Meeting with SOFIE devs (Lorenzo and Sanjban) was very productive
- One night of synergetic work between me and ChatGPT Codex: Huston, we have liftoff!
- SOFIE PR: https://github.com/ML4EP/SOFIE/pull/63
- O2 PR: https://github.com/AliceO2Group/AliceO2/pull/15913
- alidist PR: https://github.com/alisw/alidist/pull/6369

This brings...
- Support for cuda_ / hip_stream
- GPU code generation and on-the-fly compilation of .onnx models. Inflight-CCDB-loaded models are supported
- Support for FP16 and FP32
- Support for dense layers and ReLU operations
- Support for memory arena pointer (provided by reco framework as scratch allocation)
- Direct interface with SOFIE's generalised kernel generation code
Benchmarks (ONNX to be redone with new software. Old benchmark from over several months ago).
- SOFIE

- ONNX

10% improvement in TF reconstruction time on single GPU benchmark! -> probably 2-3x faster NN inference by SOFIE
AOB
- Brevitas presentation
- Slides attached, please have a look
- Presentation is next Monday at 14:30, 30 mins including questions
- New physics accurate event display rendering with blender and video support

- Testing kubernetes setup with Giada on test nodes
- O2 analysis tutorial
- Will now continue with Sandro's simulation project on the track pruning
-
10:25
→
10:30
GPU Parameter Optimizations 5mSpeaker: Gabriele Cimador (CERN, Università and INFN Torino)
-
10:30
→
10:35
Efficient Data Structures 5mSpeaker: Dr Oliver Gregor Rietmann (CERN)
Implement NGT SoA Code in O2 standalone benchmark
- Working on this fork of the AliceO2 repo, with a CI pipeline:
- Running on NGT hardware with 4 different GPUs (Nvidia and AMD)
- Extended CI-pipline to fail if GPU.out changes
- Implemented SoA in SectorTracker: GPUTPCBaseTrackParam, GPUTPCTrackParam, GPUTPCTracklet, GPUTPCTrack
- Make better use of SoA to improve performance
- ExtrapolationTracking
- I made TrackletSelector filter the Tracklets that cross sector boundaries, so that we extrapolate only those.
- Sorting extrapolation candidates by estimated length (row where it leaves the sector).
- Overall performance increased or decreased based on the GPU, but in any case it's just a small difference.
- Standart deviation of ExtrapolationTracking got much smaller.
- Tried both a custom bucket sort and
sortOnDevice.
- ExtrapolationTracking
metric old new newer active warps per thread 3 / 32 15 / 32 18 / 32 average SM occupancy 26% 10% 9.38% bytes utilized per loaded sector 4.4 / 32 12 / 32 11.6 / 32 register spill to lcoal memory (bytes) 1,402,723 19,308 18,003 The numbers are in us and computed using standalone benchmark. The number in brackets is the timing obtained from the profiler. These numbers were generated using
sortOnDevice.GPU new (us) baseline (us) sort overhead (us) Nvidia H100 28,402 (28,789) 35,827 (35,978) 2700 Nvidia L40S 19,183 (21,759) 13,641 (13,484) 4000 AMD MI300x 79,777 (71,977) 81,985 (73,602) 4400 AMD W7900 18,392 (24,250) 21,242 (20,555) ? - Next Steps
- Merge current dev branch into my branch and re-run benchmarks.
- Sort both
TrackletSelectorandExtrapolationTrackingand benchmark newest dev commit. - Then try again AoS vs SoA and deep-copy vs index sort
- Generate
GPUParameters.csvfile for the newest changes.
- Working on this fork of the AliceO2 repo, with a CI pipeline:
-
10:40
→
10:45
TPC Clusterization / OpenCL / Highly Ionizing Particles 5mSpeaker: Felix Weiglhofer (CERN)
-
10:45
→
10:50
ITS Tracking 5mSpeakers: Felix Schlepper (CERN, Heidelberg University (DE)), Gabriele Cimador (CERN, Università and INFN Torino), Matteo Concas (CERN)
Gabriele: Working on optimzing vertexer kernels
E.g.: fitPeaksKernels
Before: one thread per peak
- Each GPU thread handles one vertex candidate (peak)
- Collects the lines inside the peak's z-window that are time-compatible with the peak line
- Fits a seed vertex: the point closest to all those lines (least squares)
- Keeps only the lines passing within
pairCutof the seed - Refits the kept lines → final vertex; applies the cuts (min. contributors, distance to beam)
- Computes the vertex quality: mean residual components and mean squared distance of the lines to the vertex
- Problem: a busy peak (thousands of lines) is walked three times by a single thread, and the whole kernel waits for it
After: one warp per peak
- Each warp handles one peak
- The window's lines are split across the lanes: lane i takes lines i, i+64, i+128, …
- Each lane adds its lines into a partial sum of the fit equations (9 numbers)
- A warp reduction (register shuffles: no shared memory, no synchronisation) adds up the partial sums, so every lane ends up with the full total
- Every lane solves the same 3×3 system → the seed vertex is in every lane's registers
- Lines are reselected around the seed: again partial sums, reduce, solve → final vertex
- Residuals and distances to the vertex: partial sums per lane, reduce, divide by the number of lines
- Lane 0 writes the result
- The work per peak is split 64 ways, so busy peaks no longer hold up the whole kernel

Next:
- Optimise other kernels
- Tune
-
10:50
→
10:55
System Run Coordination Topics 5mSpeaker: Ernst Hellbar (CERN)
-
10:00
→
10:20