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

Europe/Zurich
Zoom Meeting ID
61230224927
Host
David Rohr
Useful links
Join via phone
Zoom URL
    • 10:00 AM 10:20 AM
      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: Reported 2 issues to GIulio, waiting for a fix. Status?
      • 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.
      • Smooth data taking till end of run. Minor errors were still fixed, but fixes not deployed since no time and issues not critical, e.g. MFT corrupt data leading to crashes at start of run, and TPC calibration causing trouble in CCDB when older objects arrive later.

       

      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.
      • Changed MI100 vobox to run default container, and disabled matching for 96 core GPU jobs, so should not run normal 8 core async jobs CPU-only.

       

      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:

       

      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
        • No update
      • 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.
      • Merged GPUDataTypes and GPUDataTypeHeaders libraries, since it was causing trouple for Jens with ROOT DebugStreamers, since I had some hack in there with respect to ROOT dictionaries for the TRD GPU track model.
      • A lot of progress for C++ for OpenCL / OpenCL 3.0 backend, using POCL (Portable OpenCL C++ runtime) and Clang 19, we can run our SPIR-V IL code on CPU. Clusterization works, but then TPC tracking crashes (under investigation).
      • Plan is to fix this, and then remove all the obsolete support for ROOT 5 / AliRoot build / OpenCL 1.2 / C++ < C++17. Run 2 data will still be supported by using the existing raw dump in AliRoot and loading it in the standalone benchmark.

       

      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.
      • Next PR will allow to remove multiple clusters in the same pad row and keep only the best one.
      • Added feature to cut TPC clusters above a certain time bin in TF to cut away bogus data due to altro sync signal.
      • Fixed a regression that lead to RTC compilation failure due to using system headers.
      • Fixed a bug in TPC GPU track model decoding when TF was truncated during sync processing due to buffer overflows / TPC trips.

       

      TPC processing performance regression:

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

       

      General GPU Processing

       

    • 10:20 AM 10:25 AM
      Following up JIRA tickets 5m
      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 https://its.cern.ch/jira/browse/O2-4237
        • using testworkflows
          • simple-source
          • simplified dummy-workflow
            • one input, one output channel
            • taking simple-source output as input
          • putting one processor to sleep for 1s to create backpressure
            • backpressure correctly reported by processers in front
      • 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
        • moved checking for readout-proxy name from when sending the metric to when registering the metric 
          • sending the metric only when registered
          • tested locally with GUI
        • calibrations on aggregator nodes have different proxy names
        • standalone calibration sometimes have different input proxy names
        • tested locally with raw-proxy 
        • merged
      • 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 AM 10:30 AM
      TPC ML Clustering 5m
      Speaker: Christian Sonnabend (CERN, Heidelberg University (DE))

      Centre of gravity checks

      Previously:

       

      Problem

      • Pad size varies over rows, so plotting in units of pad makes no sense when compressing over all rows -> switch to cm as unit of measurement
      • Plot pad and time resolution vs row to see how much it is impacted

       

      GPU CF vs. NN:

      To be taken with a grain of salt: GPU CF plot has 15.6 mio. entries, NN plot has 11.6 mio. -> Need to check where I loose them in the assignment

      dE/dx from reconstruciton (with 2 class classification network and regression networks)

      • GPU CF:

      • NN:

       

      Other topics

      • PR on alidist modified to use [[ ... ]] syntax and works now even if variables are not defined
      • ort_interface.cxx / .h are now merged. Might need to some work for the GPU part

       

    • 10:30 AM 10:35 AM
      ITS Tracking 5m
      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:
        • WIP: Draft PR. The largest kernel has been ported, now testing it. ToDo: porting of duplicate removal kernel.
      • 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 AM 10:45 AM
      TPC Track Model Decoding on GPU 10m
      Speaker: Gabriele Cimador (Universita e INFN Trieste (IT))

      TPC GPU Decoding

      Filter TPC clusters at decoding time (timebin cutoff) [work in progress]

      • Makes the final #clusters unknown
      • Requires small redesign of buffer allocations in the workflow
      • If parameter for cutoff is not set, normal workflow
      • If it is set, allocate tmp buffer and tmp ClusterNativeAccess
        1. Decode all clusters in tmp buffer
        2. Scan all clusters, count how many remains after filtering (done in GPU kernel)
        3. Allocate final buffer with correct size
        4. Move the filtered clusters to the final buffer (done in GPU kernel)

      New projects

      Familiar with simulation script of FST, will start to simulate datasets with different beams and interaction rates

    • 10:45 AM 10:55 AM
      Efficient Data Structures 10m
      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.

      Using std::span

      We can now create and manage instances of S with F == std::span.

      struct Point2D { double x, y; };

      template <template <class> class F>
      struct S {
      F<int> x;
      F<int> y;
      F<Point2D> point;
      F<std::string> identifier;

      int abs2() const { return x * x + y * y; }
      int& getX() { return x; }
      const int& getX() const { return x; }
      void setX(auto x_new) { x = x_new; }
      };

      Using Memory of an Existing Buffer

      Problem: Given an existing buffer, we want to allocate our data structures within that buffer.

      Solution: We use std::pmr::memory_resource that takes a pointer to the buffer and its size in bytes.

      struct BufferResource : public std::pmr::memory_resource {
      char* buffer_begin;
      char* buffer_end;
      char* current;

      BufferResource(char* buffer, std::size_t bytes)
          : buffer_begin(buffer), buffer_end(buffer + bytes), current(buffer) {}

      protected:

      void* do_allocate(std::size_t bytes, std::size_t alignment) override {
          std::size_t space = buffer_end - current;
      void* aligned = std::align(
      alignment,
      bytes,
      reinterpret_cast<void*&>(current),
      space
      );

      if (!aligned) throw std::bad_alloc();
          current += bytes;
      return aligned;
      }

      void do_deallocate(void*, std::size_t, std::size_t) override {}

      bool do_is_equal(const std::pmr::memory_resource& other) const noexcept override {
      return this == &other;
      }
      };