Alice Weekly Meeting: Software for Hardware Accelerators

Wednesday 5 Aug 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.
  
• ROCm 7.14 solves the known miscompilations, the performance regression of 7.2, and perhaps the MI100 serialization issue. So a promissing candidate to update.
  • However, there is a problem that ROCm-provided LLVM and our LLVM for arrow collide with their symbols, leading to crashes and other random failures.
  • For now, mitigated by improving how we hide the LLVM symbols of the arrow gandiva library (which we want to do in any case).
  • Meanwhile AMD has acknowledged that this is a real issue on their side, and they are working to fix it by switching to statically linked LLVM.

 

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.

 

Other topics:

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

 

EPN GPU Topics:

• ROCm 7.2.4 did not look goot, but ROCm 7.14 looks promising. Should do some more long-term tests, then we can plan to update to ALMA 10 and the new ROCm.

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)

• Writing CHEP proceedings

10:30 → 10:35 Efficient Data Structures

Speaker: Dr Oliver Gregor Rietmann (CERN)

 

NextGenTrigger Task 1.7

• Supervising a student to work an a CMS clustering algorithm.
  • Integrating the standalone version in the actual library
  • Next step: Test and see if we get the right result

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
• 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 → 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)

Implemented seeding vertexer on for GPUs

Some tests with 50 Pb-Pb TFs

 

Computing results

• Implemented GPU seeding vertexer version
• Removed redundant h2d transfers
• Moved cluster sorting and preparation to GPUs (sorting into phi and z bins)

Run with CPU (20 threads) TF initialization + CPU (20 threads) vertexer + GPU tracking:

[1853565:its-gpu-tracker]: [09:28:42][INFO] === ITS CPU Seeding vertexer iteration 0 summary:
[1853565:its-gpu-tracker]: [09:28:43][INFO]  - Initialisation: completed in 44.02 ms
[1853565:its-gpu-tracker]: [09:28:43][INFO]  - Tracklet finding: found 10895769 | 7578426 tracklets in 187.70 ms
[1853565:its-gpu-tracker]: [09:28:43][INFO]  - Tracklet selection: selected 113082 tracklets in 101.88 ms
[1853565:its-gpu-tracker]: [09:28:43][INFO]  - Vertex finding: found 133 vertices (total 133) in 53.17 ms
[1853565:its-gpu-tracker]: [09:28:43][INFO]  + Vertex seeding total elapsed time: 386.77176 ms for 133 vertices found
[1853565:its-gpu-tracker]: [09:28:43][INFO]  + Beam position computed for the TF: -0.0005932135, 0.0019080117
[1853565:its-gpu-tracker]: [09:28:43][INFO] ==== ITS GPU Tracking iteration 0 summary ====
[1853565:its-gpu-tracker]: [09:28:43][INFO]  - TimeFrame initialisation completed in 66.79 ms
[1853565:its-gpu-tracker]: [09:28:43][INFO]  - Tracklet finding: 2214325 tracklets found in 271.70 ms
[1853565:its-gpu-tracker]: [09:28:43][INFO]  - Cell finding: 1728914 cells found in 25.46 ms
[1853565:its-gpu-tracker]: [09:28:43][INFO]  - Neighbours finding: 1042507 neighbours found in 41.93 ms
[1853565:its-gpu-tracker]: [09:28:43][INFO]  - Track finding: 76213 tracks found in 157.53 ms
[1853565:its-gpu-tracker]: [09:28:43][INFO] ==== ITS GPU Tracking iteration 1 summary ====
[1853565:its-gpu-tracker]: [09:28:43][INFO]  - TimeFrame initialisation completed in 0.90 ms
[1853565:its-gpu-tracker]: [09:28:44][INFO]  - Tracklet finding: 1985297 tracklets found in 261.85 ms
[1853565:its-gpu-tracker]: [09:28:44][INFO]  - Cell finding: 1626428 cells found in 23.94 ms
[1853565:its-gpu-tracker]: [09:28:44][INFO]  - Neighbours finding: 725931 neighbours found in 42.07 ms
[1853565:its-gpu-tracker]: [09:28:44][INFO]  - Track finding: 10329 tracks found in 30.78 ms
[1853565:its-gpu-tracker]: [09:28:44][INFO] ==== ITS GPU Tracking iteration 2 summary ====
[1853565:its-gpu-tracker]: [09:28:44][INFO]  - TimeFrame initialisation completed in 0.85 ms
[1853565:its-gpu-tracker]: [09:28:46][INFO]  - Tracklet finding: 6415815 tracklets found in 382.49 ms
[1853565:its-gpu-tracker]: [09:28:46][INFO]  - Cell finding: 16763194 cells found in 209.79 ms
[1853565:its-gpu-tracker]: [09:28:46][INFO]  - Neighbours finding: 13069894 neighbours found in 684.44 ms
[1853565:its-gpu-tracker]: [09:28:46][INFO]  - Track finding: 113246 tracks found in 1314.30 ms
[1853565:its-gpu-tracker]: [09:28:46][INFO]  + Tracking total elapse time: 3514.8193 ms for 199788 tracks found
[1853565:its-gpu-tracker]: [09:28:46][INFO] === TimeSlice 0 processing completed in: 3901.59 ms using 20/20 thread(s) ===
[1853565:its-gpu-tracker]: [09:28:46][INFO] ITSTracker pushed 199788 tracks in 192 rofs and 133 vertices
[1853565:its-gpu-tracker]: [09:28:46][INFO] GPU Reconstruction time for this TF 4.45 s (cpu), 5.31155 s (wall)

Run with GPU TF initialization + GPU vertexer + GPU tracking:

[1851852:its-gpu-tracker]: [09:27:20][INFO] ==== ITS GPU Seeding-vertex pass ====
[1851852:its-gpu-tracker]: [09:27:20][INFO]  - TimeFrame initialisation completed in: 70.9118 ms
[1851852:its-gpu-tracker]: [09:27:20][INFO]  - Tracklet finding completed in: 183.922 ms
[1851852:its-gpu-tracker]: [09:27:20][INFO]  - Cell finding completed in: 31.0391 ms
[1851852:its-gpu-tracker]: [09:27:20][INFO]  - Cell linearisation completed in: 69.0919 ms
[1851852:its-gpu-tracker]: [09:27:20][INFO]  - Seeding vertices completed in: 0.348759 ms
[1851852:its-gpu-tracker]: [09:27:20][INFO]  + Vertex seeding total elapsed time: 355.43658 ms for 141 vertices found
[1851852:its-gpu-tracker]: [09:27:20][INFO]  + Beam position computed for the TF: -0.0005310316, 0.0012643747
[1851852:its-gpu-tracker]: [09:27:20][INFO] ==== ITS GPU Tracking iteration 0 summary ====
[1851852:its-gpu-tracker]: [09:27:20][INFO]  - TimeFrame initialisation completed in 1.11 ms
[1851852:its-gpu-tracker]: [09:27:21][INFO]  - Tracklet finding: 2269988 tracklets found in 270.52 ms
[1851852:its-gpu-tracker]: [09:27:21][INFO]  - Cell finding: 1770213 cells found in 23.38 ms
[1851852:its-gpu-tracker]: [09:27:21][INFO]  - Neighbours finding: 1060005 neighbours found in 42.70 ms
[1851852:its-gpu-tracker]: [09:27:21][INFO]  - Track finding: 76560 tracks found in 145.15 ms
[1851852:its-gpu-tracker]: [09:27:21][INFO] ==== ITS GPU Tracking iteration 1 summary ====
[1851852:its-gpu-tracker]: [09:27:21][INFO]  - TimeFrame initialisation completed in 0.01 ms
[1851852:its-gpu-tracker]: [09:27:21][INFO]  - Tracklet finding: 2039025 tracklets found in 263.11 ms
[1851852:its-gpu-tracker]: [09:27:21][INFO]  - Cell finding: 1689835 cells found in 24.11 ms
[1851852:its-gpu-tracker]: [09:27:21][INFO]  - Neighbours finding: 750328 neighbours found in 41.79 ms
[1851852:its-gpu-tracker]: [09:27:21][INFO]  - Track finding: 10395 tracks found in 34.41 ms
[1851852:its-gpu-tracker]: [09:27:21][INFO] ==== ITS GPU Tracking iteration 2 summary ====
[1851852:its-gpu-tracker]: [09:27:21][INFO]  - TimeFrame initialisation completed in 0.01 ms
[1851852:its-gpu-tracker]: [09:27:24][INFO]  - Tracklet finding: 6519801 tracklets found in 383.47 ms
[1851852:its-gpu-tracker]: [09:27:24][INFO]  - Cell finding: 17134595 cells found in 216.87 ms
[1851852:its-gpu-tracker]: [09:27:24][INFO]  - Neighbours finding: 13463180 neighbours found in 708.91 ms
[1851852:its-gpu-tracker]: [09:27:24][INFO]  - Track finding: 113221 tracks found in 1360.07 ms
[1851852:its-gpu-tracker]: [09:27:24][INFO]  + Tracking total elapse time: 3515.6353 ms for 200176 tracks found
[1851852:its-gpu-tracker]: [09:27:24][INFO] === TimeSlice 0 processing completed in: 3871.07 ms using 20/20 thread(s) ===
[1851852:its-gpu-tracker]: [09:27:24][INFO] ITSTracker pushed 200176 tracks in 192 rofs and 141 vertices
[1851852:its-gpu-tracker]: [09:27:24][INFO] GPU Reconstruction time for this TF 1.98 s (cpu), 5.29232 s (wall)

Next Steps:

• Wait for Max simulations to assess vertexing quality

 

 

10:50 → 10:55 System Run Coordination Topics

Speaker: Ernst Hellbar