Alice Weekly Meeting: Software for Hardware Accelerators

Wednesday 9 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.
  
• First version of ALIBUILD_O2_FORCE_GPU=build env variable option to use CUDA / ROCm from alidist merged.
  • Want to add OpenCL as well.
  • At the moment,  ALIBUILD_O2_FORCE_GPU_CUDA and ALIBUILD_O2_FORCE_GPU_HIP are ignored in the build option, required new alibuild features, which are available now. Will add that with the next PR.
  • Caused OOM issues with ONNXRuntime CUDA build. Giulio reduced the number of threads for CUDA, if I understand correctly. Could we discuss what exactly was done?
  • MigraphX: AMD does not provide RPMs from the normal repository. After checking with AMD, they provide tarballs and RPMs online, but the URL is not documented. AMD will fix the documentation, and I have added MigraphX for the new ROCm. Seems to work with ONNXRuntime. @Christian: could you confirm?
  • 2 bugs in ONNXRuntime found: CUDA architecture 75-vitual leads to a CMake error, since their build script is buggy. Mixing architectures 75 and >80 in the same build leads to a linker issue most like due to bad usage of architecture-dependent #ifdefs.
    • For now, implemented workarounds in our ONNXRuntime recipe. File 2 bug reports to ONNXRuntime, they are working on it.
  • ONNXRuntime with CUDA and cuDNN working, cuDNN is installed together with CUDA in the "build" mode. For now, not installing TensorRT, since I am not sure about the license. A free NVIDIA developer account is needed to get the URL for the download of the tarball (even though their RPM repository contains the RPMs freely...). @Christian: do we need TensorRT? Should I investigate?
  • New version causes trouble building on EPNs due to the version string with environment modules. (Actually, the version string did not change, but now it contains the architecturs also on EPNs, which it did not do before. With the new aliDist, we can rework this fully, and avoid extracting the needed information from the version string, and with that we can simplify it. Will work on that today.
  • One missing thing: RTC does not work if the installation PATH is relocated for now. Will also fix that.
• With the new "build" mode, the GPU CI is working on SLC10 with CUDA 13.3 and ROCm 10.0.

 

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.
• GCC bump - Status?

 

EPN GPU Topics:

• Delayed long-term test of ROCm 7.14: first want to finish the new "build" option, and then we can test ROCm 10.0 directly.
• Giada provided test node with only kernel module, but no ROCm installation, for ROCm 10.0 test with "build" option.
• For the future, I suggest we do not use the system ROCm on EPNs at all, but:
  • We use the "build" option for all online builds.
  • We use the "build" option for async builds based on new O2 / alidist and using SLC10.
  • For old builds, we can add the full /opt/rocm folder to the container, and use that together with the new kernel module. I verified that this works,
• Need to discuss with DPG and the GRID group how to make this change as transparent as possible.

10:20 → 10:25 TPC ML Clustering

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

PRs

• alidist
  • Problem found in build of O2 on epn000 due to usage of @ characters in gpu-systems package, PR open, currently under discussion
• O2Physics
  • Problem found in memory allocation of ORT model output (simple heap pointer instead of std::vector). Fixed by Claude, PR open: https://github.com/AliceO2Group/O2Physics/pull/17836 (iterations needed)
• O2
  • Momentum vector usage in seeding: https://github.com/AliceO2Group/AliceO2/pull/15756 -> See next results

 

Momentum vector in seeding

Managed to get the NN momentum vector into the seeding algorithm, benchmarks on MC.

 

• 0-5% centrality enforced, 50 Events, 38 kHz PbPb, LHC24ar anchored
  • Primaries, absolute

• • Secondaries, absolute
    
  • All tracks, ratio

• Similar plots for 0-100% but no significant improvement of regression observed

 

MIGraphX execution successful with ORT 1.29 on ROCm 6.4 (epn000), but very very slow as model is newly compiled by ORT at every inference call. Need to see how to fix it.

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

 

NextGenTrigger Task 1.7

• Nothing

 

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.
  • Tried grouping Tracklets into non-overlaping blocks and run each Selector thread over all rows.

• 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 → 10:45 TPC Clusterization / OpenCL / Highly Ionizing Particles

Speaker: Felix Weiglhofer (CERN)

No news... Finishing ML4Jets talk.

10:45 → 10:50 ITS Tracking

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

Gabriele: on vacation, seeding vertexer PR at https://github.com/AliceO2Group/AliceO2/pull/15733

10:50 → 10:55 System Run Coordination Topics

Speaker: Ernst Hellbar (CERN)