Alice Weekly Meeting: Software for Hardware Accelerators

Wednesday Aug 12, 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.
  
• 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.
• Problem running out of buffer size during clusterization in raw reco of MC, 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 AM → 10:25 AM TPC ML Clustering

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

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

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

Finished CHEP 2026 proceedings

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. We are still fixing some bugs.
  • Next step: Fix bugs and improve performance.

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
  • Changed a few things about GPUTPCBaseTrackParam, but performance got only worse...
• Next Steps:
  • 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: alidist recipe for a simple GPU smoke test done (https://github.com/alisw/alidist/pull/6264)
• Fixed: Errors between Nomad, Consul and Vault due to missing DNS configuration for Consul
• Node appears healthy now, but need Nomad token to submit test jobs to fully verify

Highly Ionizing Particles

Started to implement changes from TPC discussion:

• Working on fixing the time position

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)

Current pattern of ITS kernels:

1. Execute the kernel (e.g. trackleting), count how many objects are produced
2. Allocate buffers
3. Execute the kernel again, store the object in the buffer

I.e. double computation.

Currently experimenting new ways to do it one-shot.

1. Build some density estimates of the output of each kernel based on the input (e.g. clusters -> tracklets)
2. Allocate buffers using the estimates
3. Execute the kernel, storing the objects using atomic counters
4. If at the end counter > capacity, reallocate buffers and rerun the kernel
5. Update buffer estimation based on last processed TF

 

This way, kernel is executed twice only if the buffer estimation was wrong. First experiment with tracklet kernel only:

Next steps:

• Try with other kernels
• Received test simulation from Maximiliano -> test vertexer quality

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

Speaker: Ernst Hellbar