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

Wednesday 13 Nov 2024, 10:00 → 11:30 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:

• Fix dropping lifetime::timeframe for good: No more instances spotted.
  
• 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: Reported 2 issues to GIulio, waiting for a fix.
  
• Fix problem with ccdb-populator: no idea yet - since Ole left, someone else will have to take care.

 

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

• Waiting for RC to test COSMIC replay data set.
• Waiting for RC to test STOP timeout impact.
• Again several software updates with minor fixes of detector code, currently one update pending with a fix for the event display.
• Had first fill with leveling at 50 kHz yesterday run 559933.
  
  • 255 MI50 nodes 100% busy, 64 MI100 nodes 80% busy. --> 17 MI50 equivalent nodes margin out of 340 in production (5%).
    
• Checked number of clusters in online processing (before CTF cluster removal compared to 2023 47 kHz Pb-Pb): Now 10% fewer clusters. Not clear why the processing requirements are still higher than with replay of last years data.
  

 

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.
  • Giulio will implement tf-status message (from https://github.com/AliceO2Group/AliceO2/pull/13495) sent by readout-proxy. Status?
    

 

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
• 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.
• Understand different time stamps: https://its.cern.ch/jira/browse/EPN-487

 

Full system test issues:

Topology generation:

• Deployed a workaround in O2DPG GenTopo, to delete repository and check out anew if this happens. To be seen if this fixes it fully.
  

 

AliECS related topics:

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

 

GPU ROCm / compiler topics:

• ROCm 6.2.2 available from AMD, old problems seem fixed, but we see 2 new types of crashes
  • GPU memory error : no update
  • Server crash: Serial console works, AMD got some debug info, work ongoing.
• New miscompilation for >ROCm 6.0
  
  • Waiting for AMD to fix the reproducer we provided (not yet fixed in 6.2.2, but we have a workaround).
    
• Try to find a better solution for the problem with __device__ inline functions leaking symbols in the host code.
  
• Created PR to bump GCC to 14.2 (just to test).
  • GCC14 PR now gree despite CUDA problem in FullCI, now have to wait for new CUDA version.

 

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 including average charge and occupancy map errors during seeding.
  
• Changing IFC inner radius filter to remove cluster instead of increasing its error. Requested by Ruben.
  

 

TPC processing performance regression:

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

 

General GPU Processing

 

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.

• Backpressure reporting when there is only 1 input channel: no progress

• 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
• 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.
• DPL sending SHM metrics for all devices, not only input proxy
• Some improvements to ease debugging with the GUI:
  • ROOT Messages in the output of a workflow should not generally be interpreted as error
    • done https://github.com/AliceO2Group/AliceO2/pull/13635
      • for unknown log level, check for typical ROOT messages (Info in , ...) at the beginning of the message string to assign the proper log level
      • assign color (YELLOW ) for ALARM messages
    • https://github.com/AliceO2Group/AliceO2/pull/13683
      • assignment of proper severity for ROOT logs in Utilities/EPNMonitoring/src/EPNstderrMonitor.cxx for EPN InfoLogger
      • assignment of proper severity for ROOT logs printed in terminal output (+GUI)
      • new variables DeviceInfo::logLevel set to --severity from command line arguments
        • now determines the minimum severity of printed logs for each device
        • independent of DeviceControl::logLevel which determines the filter severity in the GUI
• 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

 

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

10:25 → 10:30 TPC ML Clustering

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

Multi-class classification

• Adapted training data selection for multi-class classification
  • Try to predict class 0 (reject), class 1 (accept, don't split), class 2 (accept and split)
  • Try to split digit maxima which belong to two real tracks, but only predict class 1 if e.g. track + high incl. track (e.g. looper) create cluster 
• Both classification and regression data adjusted, but facing some limitations:
  • After downsampling from 5 x 50kHz, PbPb simulations I got ~3000 class 2 points
• Model seems to find more split clusters at smaller radii, but also sometimes splits clusters which shouldn't be split

 

 

Area overlap study

• Simple way of removing loopers: Make cut on the maximum area of a cluster with same MC label (if area < 200 -> accept)
• Now the trending is gone, interestingly

 

• Compare to GPU CF

 

Plans & To-Do's

Title and text; High priority, medium priority, low priority; short term, mid term, long term; other

• Neural networks
  • Cluster splitter network
    • (✓) N-class classifier network: Probably only going until split-level 2 or 3. Higher gets really sparse in training data
    • (✓) N-class regression network: Similar approach as the N-class classifier, but need to see how good performance is...
    • Pass momentum vector to downstream reconstruction
• GPU developments
  • Pull requests
    • O2
      • PR, ORT library integration: https://github.com/AliceO2Group/AliceO2/pull/13522 
      • PR, Full clusterization integration: https://github.com/AliceO2Group/AliceO2/pull/13610
    • alidist
      • PR, ORT GPU build: https://github.com/alisw/alidist/pull/5622
  • Issues & Feature requests
    • OnnxRuntime
      • Coming up...
• QA task & algorithmic developments
  • Include SC distortion simulation
  • Use black PbPb data to evaluate performance and for training
  • Redo 2D study with NN

10:30 → 10:35 ITS Tracking

Speaker: Matteo Concas (CERN)

ITS GPU tracking

• General priorities:
  • Focusing on porting all of what is possible on the device, extending the state of the art, and minimising computing on the host.
  • Optimizations via intelligent scheduling and multi-streaming can happen right after.
  • Kernel-level optimisations to be investigated.
• Tracklet finding:
  • Still WIP, it requires some more new developments to port the features we have for CPU (delta-rof, timeframe fractioning, per-vertex tracking...).
• TODO:
  
  • Reproducer for HIP bug on multi-threaded track fitting: no progress yet.
  • Move more of the track-finding tricky steps on GPU: no progress yet.
  • Fix possible execution issues and known discrepancies when using gpu-reco-workflow : no progress yet; will start after the tracklet finding is ported.

DCAFitterGPU

• Deterministic approach via using SMatrixGPU on the host, under particular configuration: no progress yet.  

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

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

Code duplicaiton

• Added new class TPCClusterDecompressionCore.inc
• Single entry point for decoding logic
  • decompressTrack() -> logic to decode attached cluster to a specific track
  • decompressTrackStore() -> logic to store clusters, either using vectors or plain arrays
  • decompressHits()
  • decompressHitsStore() -> same as before but for unattached
• The decoding classes include TPCClusterDecompressionCore.inc and uses the functions
• No significant difference for GPU decoding (before duplication and after)
• CPU version appears to be slower, need to check if inline declarations are correct
• Process slowed down by technical problem with no fast math option

10:45 → 10:55 Efficient Data Structures

Speaker: Dr Oliver Gregor Rietmann (CERN)

Efficient Data Structures

Context

• Create data structures for controling the data layout (AoS vs SoA)
• These data structures should hide the underlying data layout.
• We want to change the underlying data layout without affecting the code using it.

value, reference, const_reference

Consider the following structs to understand the problem.

struct point_soa { std::vector<int> x, y, x; };
point_soa p_soa = {{1, 2, 3}, {4, 5, 6}, {8, 9, 10}};

struct point { int x, y, z; };
point& p = { p_soa[1].x, p_soa[1].y, p_soa[1].z }; // error

struct point_reference { int& x, y, z; };
point_reference pr = { p_soa[1].x, p_soa[1].y, p_soa[1].z }; // works

pr = p_soa[1]; // we want something like this

 

We use the types "value", "reference", "const_reference" below, and a type "proxy_type" that converts between them.

 

#include <cstddef>
#include <iostream>
#include <vector>

template <typename T> using value = T;
template <typename T> using reference = T&;
template <typename T> using const_reference = const T&;

template <template <class> class F_in, template <template <class> class> class S>
struct proxy_type : S<F_in> {
    template<template <class> class F_out>
    operator S<F_out>() const {
        auto& [x, y, z] = *this;
        return {x, y, z};
    }
};

template <template <class> class F, template <template <class> class> class S>
struct wrapper {
    S<F> data;
    proxy_type<reference, S> operator[](std::size_t i) {
        auto& [x, y, z] = data;
        return {x[i], y[i], z[i]};
    }
    proxy_type<const_reference, S> operator[](std::size_t i) const {
        auto& [x, y, z] = data;
        return {x[i], y[i], z[i]};
    }
};

template<template <class> class F> struct S { F<int> x, y, z; };

int main() {
    wrapper<std::vector, S> my_array = {{{1, 2, 3}, {4, 5, 6}, {8, 9, 10}}};

    S<reference> r = my_array[1];
    r.x = 42; // my_array[1] changed

    S<value> v = my_array[1];
    v.x = 43; // v ist just a copy, my_array[1] remains unchanged

    S<const_reference> cr = my_array[1];
    cr.x = 43; // error: cannot assign to const-qualified type

    std::cout << cr.x << std::endl; // 42

    return 0;
}