Alice Weekly Meeting: Software for Hardware Accelerators

Wednesday 1 Apr 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.
• Test with GPU GRID jobs at NERSC pending.
• 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.
  
• Need to check ROCm 7.2 corrtecness.
• Need to understand and fix crash on RTX Pro 6000.
• Understand deterministic mode issue on NVIDIA Blackwell.

 

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
• Waiting for TPC to fix bogus TPC transformations for good, then we can revert the workaround.
• Final solution: merging transformation maps on the fly into a single flat object:
  • Maps now yielding correct results, but 1.5x performance regression running on GPUs.
  • 50% of performance regression from use of std::array in GPU code (not clear why), recovered by using references.
  • Beyond, very difficult to understand. Slowdown seems concentrated in inverse transformation. But not a single code path is slow, but the combination of multiple make the compiler optimize poorly.
  • Not sure if it is meaningful / possible, to recover the full performance loss in current implementation, since it will anyway not run on GPUs. But right now we still cannot merge the PR since it will break online operations.
  • Perhaps, at some point should try to test the new POD version directly:
    • Rebase Ruben's PR once more.
    • Merge maps and convert inside gpu-reco-workflow, for simple testing.
    • And compare performance of the whole thing. If that is faster, or equally fast, we can merge it.
  • However, current version is slower even when I reduce it to a single map ignoring the reference map (meaning slower than old version with reference + differential map).
• 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
• TPC decoding can now handle data that was encoded with wrong b field, memory scaling factors parameters are configurable as --configKeyValue.
• 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?
• Status of bumping CMake and boost (https://github.com/alisw/alidist/pull/6135):
  • Merged
• Test at NERSC still ongoing, all jobs so far failed for non-gpu related reasons.
  • Irakli needs to investigate current issue, but was in vacation last week. Pinged him already this week, but no reply yet.

 

EPN GPU Topics:

 

10:20 → 10:25 TPC ML Clustering

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

Cluster error parameterization

• Tried the following metrics:
  1. ((Correctly attached non-fake) - (fake attached)) * avg_eff_NN/avg_eff_ref *((avg_fake_ref / avg_fake_NN * avg_clone_ref / avg_clone_NN) if (avg_eff_NN > avg_eff_ref) else 1)
  2. avg_eff_NN/avg_eff_ref *((avg_fake_ref / avg_fake_NN * avg_clone_ref / avg_clone_NN) if (avg_eff_NN > avg_eff_ref) else 1)
• Side note: Calculation of averages: sum_{i in bins}(mean_i / std_i^2) / sum_{i in bins}(1 / std_i^2)
• Optimise all parameters: 
  [

  "nnScaleClusterErrorY",

  "nnScaleClusterErrorZ",

  "scaleChiY1",

  "scaleChiY2",

  "scaleChiY3",

  "scaleChiZ1",

  "scaleChiZ2",

  "scaleChiZ3"

  ]
• Optimize 2 cases:
  1. Regular cluster errors
  2. NN cluster errors

 

Metric 1 leads to strong reduction in fake-rate and mild reduction in clone rate ~1% efficiency loss

For metric 2, none of the sampled points over 6 iterations surpassed the default settings for both NN and default cluster errors with scaling parameters.

Noticed a mistake in the execution of the task for the default case, need to rerun.

10:25 → 10:30 GPU Parameter Optimizations

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

GPU Parameters news

Tuned MI210 without RTC for simulated PbPb

3.25 per event --> 2.50 per event (-23,17%)

10:30 → 10:35 Efficient Data Structures

Speaker: Dr Oliver Gregor Rietmann (CERN)

 

NextGenTrigger Task 1.7

• CHEP talk together with Jolly got accepted.
• Co-Supervision of a summer student. Topic: Imrpove clustering algorithm CLUE.

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
• Benchmarking:
  • Two independent measurements for each of the 4 NGT GPUs (standlaone + profiler)
  • Integrated in CI-pipeline with a comparison to the unaltered code
  • Results are now much more consistent, at least with the profiler
• Next Steps:
  • Build and run O2 with C++26 compiler to integrate reflections (needed for CHEP benchmark) 
  • Write minimal reproducer for the (weird) behavior I observed with AMD W7900 and custom .par file.
  • Make better use of SoA to improve performance
  • Try David's suggestion

Comparison Standalone / rocprofv2

Nvidia H100

name	time (nsight compute)	time (standalone)	standard deviation
GPUTPCTrackletConstructor	205950	206419	606.93
GPUTPCTrackletSelector	13604	13800	45.46

 

AMD MI300

name	time (rocprofv2)	time (standalone)	standard deviation
GPUTPCTrackletConstructor	374527	381736	4231
GPUTPCTrackletSelector	27732	34623	1032

10:35 → 10:40 Following up GPU to-dos

Speaker: Dr Vikas Singhal (Department of Atomic Energy (IN))

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:

Tried different approach

• Amanatides & Woo (1987), A Fast Voxel Traversal Algorithm for Ray Tracing
• Instead of sampling the lines at fixed t steps, calculate at each step t_exit s.t. the line exits the voxel and enters the next. Use t_exit - t_enter as a weight
• First implementation seems slower than sampling
  [2125556:its-tracker]: [17:32:25][INFO] Voxel filling with sampling took 140 ms, 0.7291667 ms per ROF
  [2125556:its-tracker]: [17:32:25][INFO] Voxel filling with traversal took 229 ms, 1.1927084 ms per ROF
• Need to investigate why, might be:
  • More branchy code
  • Less cache friendly
  • In RT work per voxel is higher, here just bin update
• Will investigate further on this slowdown

 

Migration from trackleting to celling vertexer

• So far used tracklets from first 3 ITS layers given by old vertexer
• Trying to migrate to cell seeds used in tracking:
  • Made from two tracklets that share one cluster on the middle layer
  • Kalman filtering applied
  • A full TrackParametrizationWithError object
• Trying to understand this weird behavior: video

10:50 → 10:55 System Run Coordination Topics

Speaker: Ernst Hellbar (CERN)