Alice Weekly Meeting: Software for Hardware Accelerators

Wednesday 11 Nov 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.
  
• All gpu-ci related issues fixed, it is now compiling for the correct architectures and it is green.
• Added OpenCL build to gpu-system.sh in build mode.
• ROCm 10.1 seems as good as ROCm 10 on MI50.
• SLC10-FullCI added by Giulio/Vit: currently failing since gpu-system's build option does not provide CUDA stub libraries, if we don't bring host libraries to the container. Need to check how to fix it without colliding if we bring host libraries tot he container.
• Bumped to CUDA 13.4 and ROCm 10.1, we can now bump to GCC 16.2.1

 

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.
• Added CPUDevice backend for debugging (Runs on CPU, but in the way we normally run on the device).

 

Other topics:

• GCC bump - Status?

 

EPN GPU Topics:

• Discuss with Ernst / Giulio how to do a build for SLC10 on EPNs with build mode.
• Need to test Maksim's SLC9 container with ROCm included on SLC10 node, Giada fixed CVMFS for me :).

10:20 → 10:25 TPC ML Clustering

Speaker: Christian Sonnabend (CERN)

10:25 → 10:30 GPU Parameter Optimizations

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

On vacation, no news

Next step:

tuning without rtc on its kernels

10:30 → 10:35 Efficient Data Structures

Speaker: Dr Oliver Gregor Rietmann (CERN)

 

Next Generation Triggers

• Handed in an abstract to ACAT 2027 (Fourier Transform and CLUEstering).

O2 CI Pipeline

On every commit on the dev branch:

1. Build deterministic and check on a small dataset if we get the same tracks
2. Build non-deterministic and benchmark
3. Plot the benchmarks with the different commits on the x-axis

 

Possible implementation:

1. Do this still in the PR
2. Do this in the dev branch via a pipeline that dumps the benchmark results
3. Use a manually triggered pipeline that fetches the results from Step 2 and plots them.

 

O2 TPC Track Reconstruction

• 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

 

• Adressing feedback of CHEP 2026 review
• Merged newest dev branch into my branches
• Using newest version of dependencies from CVMFS
• Make better use of SoA to improve performance
  • TrackletConstructor
  • TrackletSelector 
  • ExtrapolationTracking
    • Select extrapolation candidates during TrackletSelector and process only those in ExtraploationTracking
    • Only lowering threads per block in GPUParameters.csv brings speedup
    • Storing left and right extrapolation candidates separately
    • Cleaned up the implementation: It seems that MI300 with wavefront size 64 profits the most.

 

GPU	ExtrapolationTracking (old / new)	TPC Sector Tracking (old / new)	Wall (old / new)
Nvidia H100	20,191 / 34,783 (x0.58)	229,166 / 243,667 (x0.94)	1,169,005 / 1,175,663 (x0.99)
Nvidia RTX6000	17,850 / 19,824 (x0.90)	142,867 / 145,001 (x0.99)	856,640 / 855,431 (x1.00)
Nvidia L40S	15,419 / 13,819 (x1.12)	302,630 / 291,954 (x1.04)	1,232,809 / 1,217,002 (x1.01)
AMD MI300	34,803 / 85,448 (x0.41)	312,679 / 360,710 (x0.87)	1,259,720 / 1,318,674 (0.96)
AMD W7900	13,241 / 20,572 (0.64)	321,676 / 331,073  (x0.97)	1,500,447 / 1,507,110 (x1.00)
CPU	197,498 / 205,857 (x0.95)	13,090,944 / 13,157,209 (x0.99)	50,646,580 / 50,301,550 (x1.01)

 

• Next Steps
  • Split ExtrapolationTracking into left and right, so tht we have to wait only for one Sector instead of two.
  • Generate GPUParameters.csv file for the newest changes.

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)

Gabriele: on vacation, will report next week

Felix: nothing

10:50 → 10:55 System Run Coordination Topics

Speaker: Ernst Hellbar (CERN)