Alice Weekly Meeting: Software for Hardware Accelerators

Wednesday 2 Sept 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

 

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.
  
• As discussed with Giulio, want to attempt to move CUDA / ROCm from system installation to alidist.

 

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.
• Buffer overrun is real, not due to wrong estimation. Bug somewhere in TPC digit processing or clusterization. Still need to check.

 

Other topics:

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

 

EPN GPU Topics:

• Still need to do long-term stress test with ROCm 7.14.

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)

10:30 → 10:35 Efficient Data Structures

Speaker: Dr Oliver Gregor Rietmann (CERN)

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

Speaker: Felix Weiglhofer (CERN)

10:45 → 10:50 ITS Tracking

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

After some more tuning of the ITS GPU tracking code (note that this does not change the current output; we will see more impactful by retuning the parameters) this is the state:

First bin is always longer since first the buffer estimate sizes have to be established.

The other spikes is due to higher occupancy and conservative buffer provisioning.

Improved occupancy of the fitting steps by sorting candidates by the expected energy loss correction steps applied, less warp divergence.

Also sort the seeds in phi and length.

This is measured on RTX 5080.

I removed some unwanted FP64 instructions in the track model but some are intentional (e.g., to keep cov. matrix stable in update step). I changed it locally to a different formalism like Joseph form using only FP32 this improved this part by -30% (Ruben suggested square-root formalism, I found there exist also UD Bierman-Thornton factorisation). But this is just playing around and no serious test (maybe something for the future).

Additionally, I tried to optimised the MatLUT queries achieved -5% and reduced voxel map size by half 40kB -> 20Kb which fits better into cache (ITS uses only 1.8kB in the hot path).

There is some more optimisation potential we could achieve with maybe changing the data layout AoS -> SoA 

 

10:50 → 10:55 System Run Coordination Topics

Speaker: Ernst Hellbar (CERN)