Login

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

→ Europe/Zurich
ALICE GPU Meeting
Zoom Meeting ID
61230224927
Host
David Rohr
Useful links
Join via phone
Zoom URL
    • 10:00 → 10:20
      Discussion 20m
      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:

      • 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:
        • Test last week was wrong due to misunderstanding between GIulio and me about which software to use :(.
        • New version seems to work well, except one feature we could not test yet: to keep running the calibration after the data processing timeout.
        • Some improvements for infoloogger messages to be done, Giulio will do in follow-up PRs.
      • Fix problem with ccdb-populator: 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. Status?

       

      Sync reconstruction

      • Waiting for RC to test COSMIC replay data set.
      • Waiting for RC to test STOP timeout impact.
      • Problem that EPN2EOS is not working, could be due to CTF files now written with wrong permission?
      • New SW version deployed at P2, running on Alma9, and has the new CRITICAL FairLogger severity.

       

      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.
        • Ready now to test on MI100 with new setup if problem persists, asked Catalin to start a test run.
      • 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 seperately from other tickets.
        • FLP working on this, but not yet deployed.

       

      GPU ROCm / compiler topics:

      • List of important issues with AMD:
        • Random server reboots on MI100: Tried several workarounds, but no solution found so far. Giada spotted some weird FairMQ problems in the large scale test, which could probably be due to some memory corruption happening.
          • Disappeared with Alma 9.5, AMD might still check with our Alma 9.4 / ROCm 6.2 server to understand root cause.
        • Random crashes on MI100 due to memory error, can be worked around by serializing all kernel and DMA transfers, which has 20% performance degradation.
          • Still there with our latest setup. Only persisting problem. Deployed my serialization workaround automatically for online.
          • AMD proposed a different workaround to disable the DMA engine. It works for working around this bug, but has another problem that GPUs get stuck. Sent a reproducer to AMD, although not important for us. We want a proper fix anyway.
          • Debugged in detail what is going on: multiple commands enqueued in a command queue get executed in parallel or in the wrong order, so that data is read before it is written. Clearly synchronization problem on AMD side.
        • Miscompilation leading to crashes, worked around by changing our code, but compiler bug still there.
          • Not appearing any more with ROCm 6.3.2, not clear if fixed, AMD might check the reproducer with the old ROCm.
        • Provide an RPM ROCm version with all fixes, so that we don't need to compile clang manually with custom patches.
          • No compiler patch or HIP patch necessary any more for ROCm 6.3.2, big step ahead. Still need new version which proper fix for the memory error.
        • Proper way to enable amdgpu-function-calls instead of hacking AMD scripts and binaries.
          • Can now do everything with CMake HIP language and Clang. Got rid of the hipcc binary, and thus of the hacks to enable function calls.
        • hipHostRegister has become very slow when more than 1 GPU visible (via ROCR_VISIBLE_DEVICES).
          • Fixed with ROCm 6.3.2
        • Discussed with AMD the meaning of a couple of workarounds we are applying. Once we are stable, should remove some and see if it is still needed.
      • Damon from AMD is back working for us, and will look at the GPU memory error next.
      • SLC9 GPU container in Jenkins working and RPM publishing fixed. Set up as secondary container for online builds (since we still need the old container for the infra node) and switched to a unified SLC9 container for offline builds (can run on the GRID and on the EPNs with GPU).
      • Updated vobox to use el9 container, reported a bug to Max about a missing library that was fixed. MI50s now working for offline in the new setup, MI100 to be tested.
      • Try to find a better solution for the problem with __device__ inline functions leaking symbols in the host code.
      • Added some functionality to CAMath class that was needed for BFloat16 for NN clusterizer.
      • Can now bump to GCC 14, PR reopened. Suppressed one set of bogus compiler warnings due to GCC regressions. Now some other GC14 warnings in simulation code, pinged Sandro.

       

      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.
        • GPU MemClean not working in TPC clusterization, need to debug.
        • 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.
      • Solved memset issue with OpenCL, but Clusterizer still gives slightly different clusters running on OpenCL.
        • Felix reported the problem is 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.

       

      Other Topics

      • Hiring new fellow: open until 17.3., 7 applications so far, 4 of them declined by HR for formal reasons.

       

      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
      • Online and Async updated to Alma9.5, Infrastructure nodes to follow.

       

      Other EPN topics:

       

    • 10:20 → 10:25
      Following up JIRA tickets 5m
      Speaker: Ernst Hellbar (CERN)
    • 10:25 → 10:30
      TPC ML Clustering 5m
      Speaker: Christian Sonnabend (CERN, Heidelberg University (DE))
    • 10:30 → 10:35
      ITS Tracking 5m
      Speaker: Matteo Concas (CERN)
    • 10:35 → 10:45
      TPC Track Model Decoding on GPU 10m
      Speaker: Gabriele Cimador (Universita e INFN Torino (TO))
    • 10:45 → 10:55
      Efficient Data Structures 10m
      Speaker: Dr Oliver Gregor Rietmann (CERN)
Powered by Indico v3.3.6-pre