Alice Weekly Meeting: Software for Hardware Accelerators

Wednesday 24 Jun 2026, 10:00 → 11:30 Europe/Zurich

10:00 → 10:20 Discussion

Speaker: David Rohr (CERN)

Color code: (critical, news from this week: blue, news from last week: purple, no news: black)

Sync reconstruction

• Ernst provided new reference data for 20, 30 and 50 kHz. Need to repeat 5090 GPU measurements for updated GPU requirement numbers.

 

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.
  
• Need to check ROCm 7.2 corrtecness.
• Need to understand deterministic mode issue on AMD Pro 9700 reported by Oliver - Status?

 

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.
• PR with changes for qTot-reading / dEdx code, to support >16 bit range for saturated signals: https://github.com/AliceO2Group/AliceO2/pull/15523 (to be verified by Ruben and Jens)

 

Other topics:

• Need to bump ONNXRuntime to 1.24, Giulio is checking, needed for ROCm 7.2 - Status?

 

EPN GPU Topics:

• Test of MI50-EPN equipped with 1 RTX Pro 6000 quite successful.
  • GPU temperature saturates at 78°C under full load.
  • At high temperature, seeing up to 2% degradation in GPU clocks and GPU performance.
  • 1 RTX Pro 6000 can run at full speed in FST shared by both CPU sockets / NUMA domains.
  • Right now, EPN with 1 RTX Pro 6000 Max-Q (300W TDP= reaches 5/6th of the performance of 1 8-MI50-EPN.
    • (From CPU side, would need ~280 64-core EPNs, but then GPU performance should be slightly above 8*MI50).
    • Perhaps 1 RTX Pro Server GPU (600W) might already be enough, 2 * RTX Pro 5000 might be another option. Almost certainly, next gaming GPU generation will be fast enough.

 

10:20 → 10:25 TPC ML Clustering

Speaker: Christian Sonnabend (CERN, Heidelberg University (DE))

10:25 → 10:30 GPU Parameter Optimizations

Speaker: Gabriele Cimador (CERN, Università and INFN Torino)

TPC Processing estimates for EPN farm Upgrade

• Received TPC GPU dumps from Ernst
• Will benchmark usual GPU set on three different datasets
  • First will check improvement of opt parameters over default
  • If gain is constant among datasets, use the opt params
  • Otherwise, retune based on that IR

Opt Params PR

• @David have you tried the retuned params for Blackwells on 5090?

10:30 → 10:35 Efficient Data Structures

Speaker: Dr Oliver Gregor Rietmann (CERN)

 

NextGenTrigger Task 1.7

• Supervising a student to work an a CMS clustering algorithm. Currently onboarding her.

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
• Writing ACAT Proceedings
• Next Steps:
  • Make better use of SoA to improve performance

10:40 → 10:45 TPC Clusterization / OpenCL / Highly Ionizing Particles

Speaker: Felix Weiglhofer (CERN)

OpenCL

No news.

GPU Servers

• CI Server: Working on alidist recipe for GPU CI. 
  • Taking the Full CI as template and adding recipe for GPU specific tests

Highly Ionizing Particles

• Working on (vectorized) CPU version of tail filter
• Found an easy way to increase performance (on GPU) by adjusting to memory layout to 128B cache lines:

  • On RTX 5080: 221 GB/s to 379 GB/s.

     

  • On Radeon VII: 58 GB/s to 81 GB/s.

10:45 → 10:50 ITS Tracking

Speakers: Felix Schlepper (CERN, Heidelberg University (DE)), Gabriele Cimador (CERN, Università and INFN Torino), Matteo Concas (CERN)

GPU Seeding Vertexer

• Implemented new version which reuses trackleting and celling from tracker
• Remove semi-duplicated code of tracklet finding and tracklet selection from ITS code
• Tests from 4 simulated TFs:

 

 

TF / Vertex seeding total elapsed time [ms]	O2/dev - Vertex seeding total elapsed time [s]	new Vertex seeding total elapsed time [s]
0	481.84 ms for 112 vertices found	477.92 ms for 118 vertices found
1	703.49 ms for 123 vertices found	594.70 ms for 131 vertices found
2	537.67 ms for 132 vertices found	513.21 ms for 140 vertices found
3	481.95 ms for 122 vertices found	458.45 ms for 125 vertices found

 

Cons:

O2/dev  fake rate 0/489 = 0.000

new fake rate 6/514 = 0.012

Next steps:

• Validate with more TFs
• Validate with all reconstruction step
• Port to GPU

10:50 → 10:55 System Run Coordination Topics

Speaker: Ernst Hellbar (CERN)

I produced the reference data for the GPU benchmarks with 20 kHz, 30 kHz and 50 kHz PbPb data. I did the following to obtain the data:

1) reconstruct 100 rawtf files each for 3 runs with an average IR close to the target rates

2) calculate the average IR of the reconstructed TFs of the rawtf files (might be more than 100 if >1 TF per file) using CTP digits (pile-up corrected) -> IR_100tf
- for the 50 kHz sample and 20 kHz sample, the IRs per TF were distributed quite closely around the average since for these rawtfs we were still levelled
- for the 30 kHz sample, I have a more or less uniform sample between 35 and 25 kHz

3) for all reconstructed TFs, calculate the average number of TPC clusters from the gpu-reco logs -> NCl_100tf

4) get the target number of TPC clusters for the target rates by scaling the average NCl obtained in 3) to the target rate using the average rate obtained in 2)

5) select the 10 TFs for which the reconstructed number of TPC clusters is closest to the target number calculated in 4)

6) then I ran over all 100 rawtf files again, but only selecting and reconstructing the TFs obtained by 5) (was too lazy to handpick the rawtf files which becomes non-trivial if there are more than 1 TF per rawtf file), I cross-checked that the number of reconstructed TPC clusters (NCl_10tf_selected and the std. dev.) matches the target average I calculated in 4) and I dumped the GPU data



Here’s the final results for the different IRs:

20 kHz, LHC26ak, 572397:
IR_100tf: 22032 Hz, scaled IR: 20000 Hz, NCl_100tf: 42755606, NCl_10tf_selected: 38604835, std. dev.: 693143
gpu event data: /scratch/ehellbar/GPUbenchmark/20kHz/events/dump

30 kHz, LHC26ar, 560089:
IR_100tf: 31149 Hz, scaled IR: 30000 Hz, NCl_100tf: 62084016, NCl_10tf_selected: 59228746, std. dev.: 870193
gpu event data: /scratch/ehellbar/GPUbenchmark/30kHz/events/dump

50 kHz, LHC24ar, 560161:
IR_100tf: 48175 Hz, scaled IR: 50000 Hz, NCl_100tf: 89397531, NCl_10tf_selected: 93076232, std. dev.: 1114990
gpu event data: /scratch/ehellbar/GPUbenchmark/50kHz/events/dump