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

Wednesday 11 Sept 2024, 11:00 → 12:30 Europe/Zurich

11:00 → 11: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:

• Fix dropping lifetime::timeframe for good: Still pending: problem with CCDB objects getting lost by DPL leading to "Dropping lifetime::timeframe", saw at least one occation during SW validation.
  • Found yet another bogus message at end of run. Potential fix is in this PR https://github.com/AliceO2Group/AliceO2/pull/13481, but fails DPL tests since some service is not found. 
• 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:
      
• Stabilize calibration / fix EoS: New scheme: https://its.cern.ch/jira/browse/O2-4308: Status?.
  
• Fix problem with ccdb-populator: no idea yet - since Ole left, someone else will have to take care.
• Giulio fixed the second reproducer I obtained from the full topology at P2. Ernst deployed new SW test version and spurious cyclic channels. Waiting for RC to test expendable tasks with it?
  

 

Global calibration topics:

• TPC IDC and SAC workflow issues to be reevaluated with new O2 at restart of data taking. Cannot reproduce the problems any more.

 

Sync reconstruction

• Had a problem with COSMIC runs crashing due to bug in ITS code.
  • Waiting for RC to test COSMIC replay data set.
    
  • Again crashes in its-tracker, fixed by Ruben in O2/dev, deployed at P2.
    
• ITS reco had a bug with excessive logging in first SW version deployed at P2, was fixed by deploying nightly tag with fix.
  
• gpu-reconstruction crashing due to receiving bogus CCDB objects.
  
  • After investigating with Costin, Ruben, the problem was incorrect propagation of HTTP header Content-Size information, and then high CPU load due to extension of SHM PMR vector. Fixed.
    
• Some crashes of tpc-tracking yesterday in SYNTHETIC runs, to be investigated.
  
  • RC did a test run with CPU only to see if the problem is only on GPUs. Had some side effects on EPNs at STOP. To be checked. More news from EPN checks?
    
• STOP timeout discussion: Waiting for RC to test with extra env variable to find good timeout value.
  • Alice will add number of CTF orbit monitoring to GRID DAQ monitoring website.
• But in O2DPG/GenTopo set incorrect IS_SIMULATED_DATA flag for SYNTHETIC runs leading to EMCAL warnings, fixed.
  
• Reported a bug that ECS GUI sends the default string instead of an empty string if a GUI field is empty, fixed by ECS.
  

 

Async reconstruction

• Remaining oscilation problem: GPUs get sometimes stalled for a long time up to 2 minutes. Checking 2 things:
  • does the situation get better without GPU monitoring? --> Inconclusive
  • We can use increased GPU processes priority as a mitigation, but doesn't fully fix the issue.
• ḾI100 GPU stuck problem will only be addressed after AMD has fixed the operation with the latest official ROCm stack.
  
• 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.
• Problem reported by Max/Ruben that analysis fails, since sometimes at EOR we receive TFs with no data and bogus orbit. Started to implement a workaround to reject bogus TFs at DPL proxy: https://github.com/AliceO2Group/AliceO2/pull/13495. Will need a another PR on top by Giulio. Should also report this to EPN to fix the root cause in DD.
  

 

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
  
• Improve core dumps and time stamps: https://its.cern.ch/jira/browse/EPN-487
• Tentative time for ALMA9 deployment: december 2024.

 

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.

 

Full system test issues:

Topology generation:

• 2 Occucances where the git repository in the topology cache was corrupted. Not really clear how this can happen, also not reproducible. Was solved by wiping the cache. Will add a check to the topology scripts to check for a corrupt repository, and in that case delete it and check it out anew.
  

 

AliECS related topics:

• Extra env var field still not multi-line by default.

 

GPU ROCm / compiler topics:

• Compilation failure due to missing symbols when compiling with -O0. Similar problem found by Matteo, being debugged.. Sent a reproducer to AMD.
  
  • Still discussing with AMD, but basically the behavior with the missing symbols is not specified, and apparently clang (and thus ROCm) with the latest update deviates from what NVIDIA does. Will wait for AMD to reply again, otherwise we need to check how we can alter our code to not have missing unused symbols.
    
• Internal compiler error with LOG(...) macro: we have a workaround, AMD has a reproducer, waiting for a fix.
• New miscompilation for >ROCm 6.0
  
  • Waiting for AMD to fix the reproducer we provided.
• Fixed missing omp.h header in O2CodeChecker.
• Problems of Matteo in DCAFitter come from bug / unexpected or undocumented behavior of CUDA compiler putting bogus host symbols in the object file that kill the application.
  • Filed a bug report to CUDA, but NVIDIA explains that they cannot fix it due to how their compiler works, and since they say it is anyway illegal to have the same device and host functions in different libraries, they closed the bug report as "won't fix" : https://developer.nvidia.com/bugs/4853553
    
    • (Note, the problem comes from an inline __device__ function in the CUDA code that has the same name as an inline function in a different shared object in the host code, which collide at linking. IMHO such a condition is very difficult to prevent a priori.)
      
  • For now, found a workaround, by adding additional inline keywords to GPUd() function definitions during host compile pass and disabling the CUDA warnings this causes, and for templated classes adding some tricky #ifdef for all explicit class instantiations.
    
    • Don't really like this, since I don't want to disable warnings, and the #ifdef fix needs to be repeated for all future class instatiations we'll do, which can cause undefined behavior again that is not easy to understand.
      
  • Attempting a different idea to use a linker version-script to hide the symbols. But so far seems not to work, since we cannot hide everything in o2::..., and it is not easily possible to distinguish what to hide...
    
  • Other possibility would be namespaces, but namespacing all GPU code would be a mess, since we'd need different namespaces for host and GPU compilation, i.e. use a macro definition everywhere we have code used on GPU, which include propgagator, matlut, mathutils, etc...
    
  • Still trying to think of a better solution.
    
• Created PR to bump GCC to 14.2 (just to test). Problem with JSON-C already fixed by Sergio, Giulio, David. Investigated another problem in binutils (Giulio, David), will probably just downgrade to a working version. Now failing due to compile errors in AlfFred, filed a bug report https://its.cern.ch/jira/browse/ALIDCS-1496. Note that it will eventually fail until we have new CUDA version anyway.
  

 

TPC GPU Processing

• GPU reconstruction crashing in the 50khz data REPLAY, SYNTHETIC seems ok, under investigation.
  
• Bug in TPC QC with MC embedding, TPC QC does not respect sourceID of MC labels, so confuses tracks of signal and of background events.
  • All the TPC MC/QC problems fixed. 
  • Tests with TPC MC processing revealed problem with incorrectly aligned labels, leading to undefined behavior.
    • Simple proper fix using alignas() statement works, but makes old MC label data unreadable.
      
    • Tried several approaches to fix this in transparent way, but all failed. Documented in this JIRA: https://its.cern.ch/jira/browse/O2-5314
      
    • Only solution that still comes to my mind is to change the anyway custom ROOT streamer for the MC labels to use the FlatHeader version information (which is anyway available) and do a copy of the data inserting the necessary padding.
• New problem with bogus values in TPC fast transformation map still pending.
  
  • Fixed by adding a check in TPCFastTransform.h every time the map is querries, while we wait for a proper fix from TPC.
    
• Waiting for TPC to check PR which uses full cluster including average charge and occupancy map errors during seeding.
• Want to switch to int8 ... uint64 types instead of char, short, ...
  

 

TPC processing performance regression:

• Final solution: merging transformation maps on the fly into a single flat object: Still WIP

 

General GPU Processing

 

 

11:20 → 11:25 Following up JIRA tickets

Speaker: Ernst Hellbar (CERN)

11:25 → 11:30 TPC ML Clustering

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

Fixed the issues from last week and added momentum vector selection for training data selection (p_x / p_y < tan(45°) (= 1), maybe this can be adjusted). Need to bug-fix some issues for submission / training of NN's but should be doable within 1-2 days and then submit NN's -> Actually just fixed it before the meeting. Submission works now

11:30 → 11:35 ITS Tracking

Speaker: Matteo Concas (CERN)

• ITS tracking: back to port the usage of external allocator on in the vertexer
  • The "flexibility" in the memory management (not a lot of gpu mem allocations, but just computed pointers) now drives a simplification process of the vertexing routines.
• GPU DCA fitter: now the test seems able to run at least one of the many tested cases (different configurations -> different functions DCA, trackparametrization)
  • Now investigating the second configuration that uses chi2 weighted re-calculation.
  • No weird behaviours spotted wrt fitter's "statefulness" (rerun the same test back to back succeeds).

11:35 → 11:45 TPC Track Model Decoding on GPU

Speaker: Gabriele Cimador (Universita e INFN Trieste (IT))

11:45 → 11:55 Efficient Data Structures

Speaker: Dr Oliver Gregor Rietmann (CERN)