Alice Weekly Meeting: Software for Hardware Accelerators

Wednesday 1 Jul 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

• Ernst provided new reference data for 20, 30 and 50 kHz. Need to repeat 5090 GPU measurements for updated GPU requirement numbers.

 

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.
  
• Need to check ROCm 7.2 corrtecness.
• Need to understand deterministic mode issue on AMD Pro 9700 reported by Oliver - Status?

 

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.
• PR with changes for qTot-reading / dEdx code, to support >16 bit range for saturated signals: https://github.com/AliceO2Group/AliceO2/pull/15523, for now clamping it in TPC Track Interpolation to restrict it to 16 bit. Asked TPC to check. Will merge once it is green.

 

Other topics:

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

 

EPN GPU Topics:

 

10:20 → 10:25 TPC ML Clustering

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

Int8 ONNX evaluation

• Model internal computations are scaled to Int8 -> Faster on modern GPUs
• Caveats:
  • For float16: Input can be cast to float16 by rounding
  • Int8 has better coverage with scaling -> Input cannot be cast without knowing the scaling factors
  • Input to NN is therefore kept at float32 and model internally converts to int8
• Fusing of ONNX layers strongly improved performance of the graph eval

 

EPN000 (MI50 -> No speedup with int8)
(This shows how much the int8 internal conversion may slow down the graph evaluation compared to float16 computations)

GPU timeslice reconstruction excluding gaps

Default reco	2.06 ± 0.01
NN class. + default reg.	2.57 ± 0.02
NN full, float16, threshold 18%	2.93 ± 0.01
NN full, int8, threshold 18%	3.78 ± 0.01

 

RTX5080

ONNXRuntime error for INT8 execution

[E:onnxruntime:, sequential_executor.cc:572 ExecuteKernel]

Non-zero status code returned while running MatMulInteger node.

Name:'StaticINT8_Layer5__MatMulInteger' Status Message: CUBLAS failure 15: CUBLAS_STATUS_NOT_SUPPORTED ;

GPU=0 ; hostname=pdp-dev00.internal ; file=/home/csonnabe/MyO2/sw/BUILD/ce42118eb63c984cd8d7b1d3a7a79338b5e6a4dc/

ONNXRuntime/onnxruntime/core/providers/cuda/integer_gemm.cc ; line=52 ;

expr=cublasGemmEx( cublas, CUBLAS_OP_N, CUBLAS_OP_N, n, m, k, &alpha, ldb_aligned == ldb ? b : b_padded.get(),

CUDA_R_8I, ldb_aligned, lda_aligned == lda ? a : a_padded.get(), CUDA_R_8I, lda_aligned, &beta, c, CUDA_R_32I,

ldc, CUDA_R_32I, CUBLAS_GEMM_DFALT);

 

Tried the following

• Tried building with tensor-rt support, but Tensor-rt require gnu++17 -> c++ standards mismatch
• Fixed Tensor-RT support: https://github.com/alisw/alidist/pull/6266
• But finally worked only with CUDA EP

 

Benchmark result: 19.16 ± 0.14s -> Something is not right yet... but at least it runs. Baby steps.

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 (CERN)

 

NextGenTrigger Task 1.7

• Supervising a student to work an a CMS clustering algorithm. Onboarding complete.
• I have to give a presentation about my NGT cluster usage this Friday.

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
• Writing ACAT Proceedings
• Make better use of SoA to improve performance
• Next Steps:
  • ...

10:40 → 10:45 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)
• New issue: GPU server had issues with handling Vault secrets (will follow up with Giulio)

Highly Ionizing Particles

• vectorized CPU version of tail filter done
  • TODO: need to check results in deterministic mode
  • TODO: check performance impact of 128B cache line layout

10:45 → 10:50 ITS Tracking

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

10:50 → 10:55 System Run Coordination Topics

Speaker: Ernst Hellbar (CERN)