Alice Weekly Meeting: Software for Hardware Accelerators / PDP-SRC

Wednesday 18 Jun 2025, 10:00 → 11:00 Europe/Zurich

10:00 → 10:20 Discussion

Speakers: David Rohr (CERN), Giulio Eulisse (CERN)

Color code: (critical, news during the meeting: green, news from this week: blue, news from last week: purple, no news: black)

High priority Framework issues:

• Stabilize calibration / fix EoS: New scheme: https://its.cern.ch/jira/browse/O2-4308:
  • EoS bug with calibration data fixed.
  • RC is doing timout tests in PHYSICS, should have the result in the next days. Then have to decide on timeouts. 
  • PR with InfoLogger improvements still WIP - Status?
  • Processes crashing at shutdown if STOP timeout was short, not clear if related, but should be checked.
• Fix problem with ccdb-populator and cpv calib: no idea yet - since Ole left, someone else will have to take care.
• TF-status message (from https://github.com/AliceO2Group/AliceO2/pull/13495) sent by readout-proxy.

 

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.

 

AliECS related topics:

• Extra env var field still not multi-line by default., created https://its.cern.ch/jira/browse/OGUI-1624 to follow this up separately from other tickets.

 

Other Topics

• Usage of TPC GPU Standalone tracking in HS23 benchmark suite.

GPU ROCm / compiler topics:

• List of important issues with AMD:
  • Issues that disappeared but not yet understood: random server reboot with alma 9.4, miscompilation with ROCm 6.2, GPU getting stuck when DMA engine turned off, MI100 stalling with ROCm 5.5.
• Problem with building ONNXRuntime with MigraphX support, to be checked.
• 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.
  
• Arrow bumped to 20, should be able to bump LLVM to 20.
• ROCm 6.4.1 status:
  • Installed on EPNs in parallel to 6.3.2.
  • New build container with 6.4.1. available.
  • Builds in the container were failing due to incorrect rpm dependencies of new ROCm, missing libdrm-devel. Container fixed by installing this manually and sent a bug report to AMD. This problem slipped our tests, since I only updated ROCm to save time for the tests, so the package was already there.
  • Unfortunately, new regression with ROCm 6.4.1. Random data-driven crashes seen at P2. Sent a reproducer to AMD. Issue happens only with 6.4.1 and -ffast-match flag (without the flag, with 6.3.2, on NVIDIA and on CPU we cannot reproduce it). Most likely compiler bug again. For now, reverted to 6.3.2.
• New mitigation for MI100 serialization problem seems to work well.
• New gpu-system package:
  • Package is ready, but needs a new version of aliBuild, with some features still to be merged.
• Bumped CUDA to 12.9 in build container, and added CUDA ORT ONNX requirements for TensorRT build.
• With new clang in new ROCm, receiving warnings about NaN being disabled with --fast-math. Coming e.g. from libfmt. Suppressed this warning.

 

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.
  
• Waiting for TPC to check PR which uses full cluster errors including average charge and occupancy map errors during seeding.
  
• Final solution: merging transformation maps on the fly into a single flat object: Still WIP
• Pending OpenCL2 issues:
  • printf not working due to confirmed bug in clang, fix is being prepared. Prevents further debugging for now.
    
  • Crash in merger, which can be worked around by disabling clang SPIRV optimization. Probably bug in clang, but need to fix printf first to debug.
    
  • Also with optimization disabled, crashing later in TPC merging, need printf to debug.
    
  • Felix debugged the OpenCL clusterization problem to be due to off-by-one offset in NoiseSuppression. Need to check how that can happen only in OpenCL.
• Next high priority topic: Improvements for cluster sharing and cluster attachment at lower TPC pad rows.
  
• Adding a GPU Standalone Benchmark build to the FullCI, to avoid breaking it in the future: Done.
• Need to check the problem with ONNX external memory allocator, and the build failure with MigraphX.

Other Topics

 

EPN major topics:

• Fast movement of nodes between async / online without EPN expert intervention.
  • 2 goals I would like to set for the final solution:
    • It should not be needed to stop the SLURM schedulers when moving nodes, there should be no limitation for ongoing runs at P2 and ongoing async jobs.
    • We must not lose which nodes are marked as bad while moving.
• Interface to change SHM memory sizes when no run is ongoing. Otherwise we cannot tune the workflow for both Pb-Pb and pp: https://alice.its.cern.ch/jira/browse/EPN-250
  • Lubos to provide interface to querry current EPN SHM settings - ETA July 2023, Status?
• Improve DataDistribution file replay performance, currently cannot do faster than 0.8 Hz, cannot test MI100 EPN in Pb-Pb at nominal rate, and cannot test pp workflow for 100 EPNs in FST since DD injects TFs too slowly. https://alice.its.cern.ch/jira/browse/EPN-244 NO ETA.
• DataDistribution distributes data round-robin in absense of backpressure, but it would be better to do it based on buffer utilization, and give more data to MI100 nodes. Now, we are driving the MI50 nodes at 100% capacity with backpressure, and then only backpressured TFs go on MI100 nodes. This increases the memory pressure on the MI50 nodes, which is anyway a critical point. https://alice.its.cern.ch/jira/browse/EPN-397
• TfBuilders should stop in ERROR when they lose connection.
  
• Allow epn user and grid user to set nice level of processes: https://its.cern.ch/jira/browse/EPN-349
• Slurm bump

 

Other EPN topics:

• Check NUMA balancing after SHM allocation, sometimes nodes are unbalanced and slow: https://alice.its.cern.ch/jira/browse/EPN-245
• Fix problem with SetProperties string > 1024/1536 bytes: https://alice.its.cern.ch/jira/browse/EPN-134 and https://github.com/FairRootGroup/DDS/issues/440
• After software installation, check whether it succeeded on all online nodes (https://alice.its.cern.ch/jira/browse/EPN-155) and consolidate software deployment scripts in general.
• Improve InfoLogger messages when environment creation fails due to too few EPNs / calib nodes available, ideally report a proper error directly in the ECS GUI: https://alice.its.cern.ch/jira/browse/EPN-65
• Create user for epn2eos experts for debugging: https://alice.its.cern.ch/jira/browse/EPN-383
• EPNs sometimes get in a bad state, with CPU stuck, probably due to AMD driver. To be investigated and reported to AMD.
• Understand different time stamps: https://its.cern.ch/jira/browse/EPN-487

 

10:20 → 10:25 Following up JIRA tickets

Speaker: Ernst Hellbar (CERN)

Low-priority framework issues https://its.cern.ch/jira/browse/O2-5226

• Grafana metrics: Might want to introduce additional rate metrics that subtract the header overhead to have the pure payload: low priority.

• Merged workflow fails if outputs defined after being used as input
  • needs to be implemented by Giulio
• Cannot override options for individual processors in a workflow
  • requires development by Giulio first 
• Problem with 2 devices of the same name
  • https://github.com/AliceO2Group/AliceO2/pull/13963
• Usage of valgrind in external terminal: The testcase is currently causing a segfault, which is an unrelated problem and must be fixed first. Reproduced and investigated by Giulio.
• Run getting stuck when too many TFs are in flight.
• Do not use string comparisons to derrive processor type, since DeviceSpec.name is user-defined.
• Support in DPL GUI to send individual START and STOP commands.
• Add additional check on DPL level, to make sure firstOrbit received from all detectors is identical, when creating the TimeFrame first orbit.
• Implement a proper solution to detect wheter a device is firstInChain 
• Deploy topology with DPL driver
• Start / Stop / Start: 2 problems on O2 side left:
  • • All processes are crashing randomly (usually ~2 out of >10k) when restarting. Stack trace hints to FMQ. https://its.cern.ch/jira/browse/O2-4639
      
    • TPC ITS matching QC crashing accessing CCDB objects. Not clear if same problem as above, or a problem in the task itself

 

PDP-SRC issues

• Check if we can remove dependencies on /home/epn/odc/files in DPL workflows to remove the dependency on the NFS
  • reading / writing already disabled
  • remaining checks for file existence?
  • check after Pb-Pb by removing files and find remaining dependencies
• logWatcher.sh and logFetcher scripts modified by EPN to remove dependencies on epnlog user
  • node access privileges fully determined by e-groups
  • new log_access role to allow access in logWatcher mode to retrieve log files, e.g. for on-call shifters
  • to be validated on STG
  • waiting for EPN for further feedback and modifications of the test setup
• new BEAMTYPE for oxygen period
  • https://its.cern.ch/jira/browse/O2-5797
  • RC asked for a synthetic OO dataset
    • tested and working fine
  • RC request for a pO dataset
  • O2 code to be checked for pp and PbPb specific variables

10:25 → 10:30 TPC ML Clustering

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

10:30 → 10:35 ITS Tracking

Speaker: Matteo Concas (CERN)

10:35 → 10:40 TPC Track Model Decoding on GPU

Speaker: Gabriele Cimador (Universita e INFN Torino (TO))

10:40 → 10:45 Efficient Data Structures

Speaker: Dr Oliver Gregor Rietmann (CERN)

10:45 → 10:50 Following up GPU to-dos

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

10:50 → 10:55 TPC Clusterization / OpenCL / Highly Ionizing Particles

Speaker: Felix Weiglhofer (Goethe University Frankfurt (DE))