Alice Weekly Meeting: Software for Hardware Accelerators

Wednesday Sep 2, 2026, 10:00 AM → 11:30 AM Europe/Zurich

10:00 AM → 10:20 AM 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.
  
• Started work to add CUDA and ROCm recipies, and added ALIBUILD_O2_FORCE_GPU=build env variable option to use CUDA / ROCm from alidist.
  • Current problem is with MigraphX: With ROCm >= 7.14, AMD seems to no longer provide RPMs, but only pip modules, but they lack CMake config files.
  • For now, trying to get it to run with CUDA cuDNN, we can handle MigraphX later.

 

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.
• TPC Buffer Size exceeded issue in async reco fixed by Felix.

 

Other topics:

• Bumped ONNXRuintime to 1.29. ROCm execution provider no longer exists. Need to use MigraphX. For now, ONNXRuntime on AMD GPUs not working until this gets fixed.
• GCC bump - Status?

 

EPN GPU Topics:

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

10:20 AM → 10:25 AM TPC ML Clustering

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

Physics performance of ML cluster finder

D0 yields

K0S performance

Software on EPNs

New version of ONNXRuntime merged by Giulio (1.29.0)

-> ROCm execution provider removed, should now use MIGraphX execution provider: That requires MIGraphX with an fp4x2 header file, only available in ROCm 7.x

-> Need a node with a new ROCm for testing, in contact with Giada

10:25 AM → 10:30 AM GPU Parameter Optimizations

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

10:30 AM → 10:35 AM Efficient Data Structures

Speaker: Dr Oliver Gregor Rietmann (CERN)

 

NextGenTrigger Task 1.7

• Supervising a student to work an a CMS clustering algorithm
  • Integrated our idea in the CLUEstering code
    • The student gave a presetnation last Friday
  • Next steps:
    • See if she can continue working on the project. (She ist stall an ALICE USER until end of October.)

 

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
• CHEP2026 proceedings
• Make better use of SoA to improve performance
  • Sorted Tracklets by lastRow and (lastRow - firstRow) to counter warp divergence.
  • For TrackletSelector on AMD W7900:

  • Kernel	old (us)	new (us)
    GPUTPCTrackletSelector	34239	20250
    GPUTPCTrackletSelector_sort		10371

• Next Steps:
  • Fix worsening of the coalescing (due to TrackletRowHits).
  • Try different way of allocating tracklets.
  • Generate GPU parameters .csv file for SoA and to see if the optimal parameters are different

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

Speaker: Felix Weiglhofer (CERN)

OpenCL

No news.

GPU Servers

• CI Server:
  • GPU CI enabled for ROCm test PR

 

Highly Ionizing Particles

Started to implement changes from TPC discussion:

• Working on fixing the time position

 

Other 

Fixed uninitialized memory in cluster finder: https://github.com/AliceO2Group/AliceO2/pull/15707

• Weird issue where some buffers weren't initialized when first fragment was empty
• Hard to reproduce as sometimes memory is still zeroed by OS, ...
• ... likely that's why it was triggered by unrelated commit by Sandro

 

 Focusing on ML4Jets talk currently.

10:45 AM → 10:50 AM 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 = launch_bounds; shown here adhoc settings) 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 

 

 

Gabriele: presentation of GPU seeding vertexer at Combined WP12/WP13 meeting today at 15:30 (https://indico.cern.ch/event/1720056/)

 

10:50 AM → 10:55 AM System Run Coordination Topics

Speaker: Ernst Hellbar (CERN)