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:
        • RC did timeout tests. COSMICs tests were not helpful, since no backpressure at all in processing.
        • In tests in PHYSICS, saw that indeed there is a bug that calib data is not processed after the dataprocessing timeout (the feature not yet tested).
          • Fix for that is being worked on. Partial fix available, but seems still incomplete (Ernst and Giulio).
        • Need to test in production, and repeat the timeout tests in PHYSICS.
        • 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.
      • Dummy sink issue fixed.

       

      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.

       

      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.
      • Once we bump arrow (PR opened by Giulio), we can bump LLVM to 19.
      • Waiting for ROCm 6.4.1 with the synchronization fix, then we can test the 6.4 version.
      • GPU Deterministic mode for async fully fixed.
      • Tested new ROCm 6.4.1, which contains the fix/workaround for the serialization bug. Seems to work on both MI50/Mi100 without extra patches. Opened JIRA ticket for EPN to bump https://its.cern.ch/jira/browse/EPN-557
      • Memory error reported by Gabriele for certain parameters in parameter range scan unclear. Disappears with O2/dev, and disappears with ROCm 6.4. Bisection shows a cleanup PR getting rid of an unused variable. This is totally unrelated, and might hint to a compiler problem. Sent a reproducer to AMD.
      • Even with AMD’s fix for the ROCm runtime, seeing sporadic crashes on MI100. Very difficult to debug. Dumped one TF  that led to a crash online, but reprocessed it 500000 times locally without crash.
        • Not sure if we can get a good reproducer.
        • Fix by AMD is no real fix, but only a mitigation which makes it less likely. AMD is working on a better mitigation, but not clear to me if we will every get a full fix for deprecated MI100 GPUs.
        • Have a mitigation in O2, which might also make it less likely. But need to test online.
      • CUDA and ROCm can meanwhile use some std headers on GPU (like <array>, <type_traits>), for which we had custom reimplementations so far. Unfortunately, doesn’t work yet with OpenCL.
        • OpenCL in clang should support <type_traits>, but that fails compilation. Filed a bug report with clang.
        • Changed O2 such to use the custom GPU versions only for OpenCL, and the std headers otherwise. Can only clean up fully when OpenCL supports everything.
      • GPU check in gpu-systems.sh alidist package (bash script) and in O2 (CMake modules) do not match, so we can have that one finds a GPU feature and one doesn’t. Draft PR to improve this open here: https://github.com/alisw/alidist/pull/5896

       

      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.
      • 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:

       

    • 10:20 → 10:25
      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.
      • 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.
      • 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
      • new BEAMTYPE for oxygen period
        • https://its.cern.ch/jira/browse/O2-5797
        • beam types
          • p-O and O-O
          • Ne-Ne highly likely
        • RC asked for a synthetic OO dataset
          • Ruben uploaded respective GRPECS and CTPConfig (created by Roman) objects for unanchored MC to CCDB
        • O2 code to be checked for pp and PbPb specific variables

       

       

    • 10:25 → 10:30
      TPC ML Clustering 5m
      Speaker: Christian Sonnabend (CERN, Heidelberg University (DE))

      Framework

      • Investigated ONNX usage of memory allocation.
        • Onnx can make use of them in initialisation, but somehow not in layer-wise (temporary) allocations

       

       

      • MIGraphX failure on build
        • Requires: migraphx-devel hipblaslt-devel
      [533/1835] Scanning /scratch/csonnabe/MyO2/sw/SOURCES/ONNXRuntime/ee70b3049a/0/onnxruntime/core/providers/migraphx/migraphx_execution_provider_info.cc for CXX dependencies
FAILED: CMakeFiles/onnxruntime_providers_migraphx.dir/scratch/csonnabe/MyO2/sw/SOURCES/ONNXRuntime/ee70b3049a/0/onnxruntime/core/providers/migraphx/migraphx_execution_provider_info.cc.o.ddi 
/scratch/csonnabe/MyO2/sw/slc9_x86-64/GCC-Toolchain/v14.2.0-alice2-1/bin/c++ -DCPUINFO_SUPPORTED_PLATFORM=1 -DEIGEN_MPL2_ONLY -DEIGEN_USE_THREADS -DENABLE_CPU_FP16_TRAINING_OPS -DHAS_STRING_VIEW=1 -DMIOPEN_VERSION=30300 -DONLY_C_LOCALE=0 -DONNXIFI_BUILD_LIBRARY=1 -DONNX_ML=1 -DONNX_NAMESPACE=onnx -DORT_ENABLE_STREAM -DPLATFORM_POSIX -DROCM_VERSION=60302 -DUSE_MIGRAPHX=1 -DUSE_PROF_API=1 -DUSE_ROCM=1 -D_GNU_SOURCE -D__HIP_PLATFORM_AMD__=1 -D__ONNX_DISABLE_STATIC_REGISTRATION -Donnxruntime_providers_migraphx_EXPORTS -I/scratch/csonnabe/MyO2/sw/SOURCES/ONNXRuntime/ee70b3049a/0/include/onnxruntime -I/scratch/csonnabe/MyO2/sw/SOURCES/ONNXRuntime/ee70b3049a/0/include/onnxruntime/core/session -I/scratch/csonnabe/MyO2/sw/slc9_x86-64/pytorch_cpuinfo/alice1-3/include -I/scratch/csonnabe/MyO2/sw/BUILD/ff69defdb056aee197ee5b00d887c93e809528b3/ONNXRuntime -I/scratch/csonnabe/MyO2/sw/SOURCES/ONNXRuntime/ee70b3049a/0/onnxruntime -I/scratch/csonnabe/MyO2/sw/slc9_x86-64/safe_int/v3.0.28a-1/include -I/scratch/csonnabe/MyO2/sw/slc9_x86-64/ms_gsl/4.0.0-19/include -I/scratch/csonnabe/MyO2/sw/slc9_x86-64/date/v3.0.3-2/include -I/scratch/csonnabe/MyO2/sw/BUILD/ff69defdb056aee197ee5b00d887c93e809528b3/ONNXRuntime/amdgpu/onnxruntime -isystem /scratch/csonnabe/MyO2/sw/slc9_x86-64/abseil/20240722.0-2/include -isystem /scratch/csonnabe/MyO2/sw/slc9_x86-64/onnx/v1.17.0-alice2-6/include -isystem /scratch/csonnabe/MyO2/sw/slc9_x86-64/protobuf/v29.3-2/include -isystem /scratch/csonnabe/MyO2/sw/slc9_x86-64/flatbuffers/v24.3.25-10/include -isystem /scratch/csonnabe/MyO2/sw/slc9_x86-64/boost/v1.83.0-alice2-38/include -isystem /opt/rocm/include -isystem /opt/rocm-6.3.2/include -fPIC -O2 -std=c++20 -Wno-unknown-warning -Wno-unknown-warning-option -Wno-pass-failed -Wno-error=unused-but-set-variable -Wno-pass-failed=transform-warning -Wno-error=deprecated -Wno-error=maybe-uninitialized -Wno-error=deprecated-enum-enum-conversion -Wno-error -Wno-error=missing-requires -w -ffunction-sections -fdata-sections -DCPUINFO_SUPPORTED -Wno-unused-parameter -O3 -DNDEBUG -fPIC -Wno-deprecated-declarations -Wall -Wextra -Wno-deprecated-copy -Wno-nonnull-compare -w -Wno-error=sign-compare -Werror -E -x c++ /scratch/csonnabe/MyO2/sw/SOURCES/ONNXRuntime/ee70b3049a/0/onnxruntime/core/providers/migraphx/migraphx_execution_provider_info.cc -MT CMakeFiles/onnxruntime_providers_migraphx.dir/scratch/csonnabe/MyO2/sw/SOURCES/ONNXRuntime/ee70b3049a/0/onnxruntime/core/providers/migraphx/migraphx_execution_provider_info.cc.o.ddi -MD -MF CMakeFiles/onnxruntime_providers_migraphx.dir/scratch/csonnabe/MyO2/sw/SOURCES/ONNXRuntime/ee70b3049a/0/onnxruntime/core/providers/migraphx/migraphx_execution_provider_info.cc.o.ddi.d -fmodules-ts -fdeps-file=CMakeFiles/onnxruntime_providers_migraphx.dir/scratch/csonnabe/MyO2/sw/SOURCES/ONNXRuntime/ee70b3049a/0/onnxruntime/core/providers/migraphx/migraphx_execution_provider_info.cc.o.ddi -fdeps-target=CMakeFiles/onnxruntime_providers_migraphx.dir/scratch/csonnabe/MyO2/sw/SOURCES/ONNXRuntime/ee70b3049a/0/onnxruntime/core/providers/migraphx/migraphx_execution_provider_info.cc.o -fdeps-format=p1689r5 -o CMakeFiles/onnxruntime_providers_migraphx.dir/scratch/csonnabe/MyO2/sw/SOURCES/ONNXRuntime/ee70b3049a/0/onnxruntime/core/providers/migraphx/migraphx_execution_provider_info.cc.o.ddi.i
In file included from /scratch/csonnabe/MyO2/sw/SOURCES/ONNXRuntime/ee70b3049a/0/onnxruntime/core/providers/migraphx/migraphx_inc.h:8,
                 from /scratch/csonnabe/MyO2/sw/SOURCES/ONNXRuntime/ee70b3049a/0/onnxruntime/core/providers/migraphx/migraphx_execution_provider_info.cc:9:
/opt/rocm/include/migraphx/migraphx.hpp:1234:5: error: module control-line cannot be in included file
 1234 |     module get_main_module()
      |     ^~~~~~
/opt/rocm/include/migraphx/migraphx.hpp:1248:5: error: module control-line cannot be in included file
 1248 |     module create_module(const std::string& name)
      |     ^~~~~~
cc1plus: note: unrecognized command-line option '-Wno-pass-failed=transform-warning' may have been intended to silence earlier diagnostics
cc1plus: note: unrecognized command-line option '-Wno-pass-failed' may have been intended to silence earlier diagnostics
cc1plus: note: unrecognized command-line option '-Wno-unknown-warning-option' may have been intended to silence earlier diagnostics
cc1plus: note: unrecognized command-line option '-Wno-unknown-warning' may have been intended to silence earlier diagnostics

       

      Physics

      • Enlarging the data volume and rerunning QA: 5x50Ev 0-5% centrality
      • Trying to combine digit files -> Would make QA a lot easier (wrote a task, but its quite memory heavy)
      • Implementing full PID calibration chain -> As per Silvias request for the PhD thesis. Will do it on raw data of e.g. 24as

      Other

      • Potential bug in ppl-workflow.sh?
        export GPUTYPE=CPU
export RAWINPUTDIR="/scratch/csonnabe/PhD/jobs/simulation/output/test_sim_2"
export WORKFLOW_DETECTORS=ITS,TPC,TRD,TOF,FT0,CTP
export TRD_SOURCES=ITS-TPC
export TOF_SOURCES=ITS-TPC,ITS-TPC-TRD
export HMP_SOURCES=
export TRACK_SOURCES=ITS-TPC,ITS-TPC-TRD,ITS-TPC-TOF,ITS-TPC-TRD-TOF,ITS,TPC,TOF,FT0,TRD,CTP
export VERTEXING_SOURCES=ITS-TPC,ITS-TPC-TRD,ITS-TPC-TOF,ITS-TPC-TRD-TOF,ITS,TPC,TOF,FT0,TRD,CTP
export VERTEX_TRACK_MATCHING_SOURCES=ITS-TPC,ITS-TPC-TRD,ITS-TPC-TOF,ITS-TPC-TRD-TOF,ITS,TPC,TOF,FT0,TRD,CTP
export SVERTEXING_SOURCES=ITS-TPC,ITS-TPC-TRD,ITS-TPC-TOF,ITS-TPC-TRD-TOF,ITS,TPC,TOF,FT0,TRD,CTP
export AOD_SOURCES=ITS-TPC,ITS-TPC-TRD,ITS-TPC-TOF,ITS-TPC-TRD-TOF,ITS,TPC,TOF,FT0,TRD,CTP
export WORKFLOW_PARAMETERS="AOD,CALIB,CTF,EVENT_DISPLAY"
export WORKFLOW_EXTRA_PROCESSING_STEPS="MATCH_SECVTX,ZDC_RECO"
export WORKFLOWMODE="print"

$O2_ROOT/prodtests/full-system-test/dpl-workflow.sh
      • Duplicate run-number option in aod-producer and raw file reader
    • 10:30 → 10:35
      ITS Tracking 5m
      Speaker: Matteo Concas (CERN)
    • 10:35 → 10:40
      TPC Track Model Decoding on GPU 5m
      Speaker: Gabriele Cimador (Universita e INFN Torino (TO))

      News from GPU parameters tuning

      External framework integration

      • Tested Optuna optimization framework via O2tuner
      • Tree-structured Parzen Estimator (TPE) algorithm
        • On each trial, for each parameter, TPE fits one Gaussian Mixture Model (GMM) l(x) to the set of parameter values associated with the best objective values
        • And another GMM g(x) to the remaining parameter values
        • It chooses the parameter value x that maximizes the ratio l(x)/g(x)

      Results on FollowLoopers + CompressionStep0 (4 dimensions)

      Dataset Default step time LHS algorithm solution time TPE solution time
      pp 2MHz 1097.26 ms ± 6.82 ms 971.22 ms ± 5.55 ms (-11.49%) 936.41 ms ± 5.93 ms (-14.66%)
      PbPb 5kHz 224.94 ms ± 4.34 ms 195.63 ms ± 1.36 ms (-13.03%) 186.76 ms ± 3.08 ms (-16.97%)
      PbPb 50kHz 1895.49 ms ± 25.84 ms 1736.86 ms ± 29.63 ms (-8.37%) 1881.87 ms ± 11.22 ms (-0.72%)

       

      Results on SectorTracker step (16 dimensions)

      Dataset Default step time LHS algorithm solution time TPE solution time
      pp 2MHz 2333.71 ms ± 20.06 ms 2206.02 ms ± 2.23 ms (-5.47%) 2094.66 ms ± 3.23 ms (-10.24%)
      PbPb 5kHz 446.31 ms ± 1.87 ms 435.43 ms ± 1.66 ms (-2.44%) 435.20 ms ± 2.01 ms (-2.49%)
      PbPb 50kHz 4519.54 ms ± 39.57 ms 4230.26 ms ± 5.40 ms (-6.40%) 4219.58 ms ± 4.71 ms (-6.64%)

      Other parameters tuning (aside from grid and block size)

      • Tuned SectorTracker step including 5 additional parameters:
        • PAR_NEIGHBOURS_FINDER_MAX_NNEIGHUP
        • PAR_NEIGHBOURS_FINDER_UNROLL_GLOBAL
        • PAR_NEIGHBOURS_FINDER_UNROLL_SHARED
        • PAR_TRACKLET_SELECTOR_HITS_REG_SIZE
        • PAR_SORT_STARTHITS
      • Total dimension: 21 parameters
      • Utilized o2tuner since it can handle all types of parameters
      • Found solution with step time: 4170.62 ms ± 3.07 ms
      • Step is 7.72% faster (previous solutions for this step: 6.40%, 6.64%)

      Robustness of solutions

      • Simulated 2 more PbPb 50kHz datasets
      • Tested best parameters for SectorTracker step on these new datasets
      • So same parameters, different datasets with same "nature"

      Dataset Default step time Best parameters time
      PbPb_50kHz 4526.31 ms ± 33.60 ms 4163.49 ms ± 3.02 ms (-8.02%)
      PbPb_50KHz_1 4765.86 ms ± 32.32 ms 4375.31 ms ± 2.76 ms (-8.19%)
      PbPb_50KHz_2 4784.34 ms ± 23.60 ms 4404.51 ms ± 3.61 ms (-7.94%)

      Overall improvement

      • Created dumps with best parameters for every step tested so far, per dataset
      • Measured sync time of standalone benchmark
      • Caveat: sync time fluctuates a bit, so difficult to have a precise measurement (single steps much more stable)
      Dataset Default sync time Optimized sync time
      pp 2MHz 10216.74 ms ± 71.58 ms 9874.28 ms ± 8.16 m (-3.35%)
      PbPb 5kHz 2540.28 ms ± 3.45 ms 2440.55 ms ± 2.08 ms (-3.93%)
      PbPb 50kHz 19679.88 ms ± 12.88 ms 18586.67 ms ± 30.42 ms (-5.55%)

      To-Dos

      • Investigate more on robustness
        • Measure other steps
        • Measure overall improvement on sync time
      • Tune rest of the steps
      • Test for other IRs
      • Try other algorithms
      • ...?
    • 10:40 → 10:45
      Efficient Data Structures 5m
      Speaker: Dr Oliver Gregor Rietmann (CERN)

      SoA Benchmarks

      Benchmarks comparing run times with hardcoded SoA layouts revealed the following issue:

      #include <vector>

      struct SoA_Points { std::vector<float> x, y, z; };

      struct S_ref { float &x, &y, &z; };

      struct wrapper {
          SoA_Points soa_points;
          S_ref operator[]( int i ) {
              return { soa_points.x[i], soa_points.y[i], soa_points.z[i] };
          }
      };

      int main() {
          int N = 100;
          wrapper w {{ std::vector<float>(N), std::vector<float>(N), std::vector<float>(N) }};
          w[42].x = 3;  // Accesses all data members of soa_points, although only x is needed
          return 0;
      }

      I am trying to find a solution for this problem.

      NGT Kubernetes Cluster

      • We can now develop O2 code on the cluster (I am already developing like this).
      • I am ready to share the corresponding yaml files and present the setup.
      • Setup can be used with any IDE that allows remote development via ssh.
      • Still waiting for non-Nvidia GPUs, but they are coming.
    • 10:45 → 10:50
      Following up GPU to-dos 5m
      Speaker: Dr Vikas Singhal (Department of Atomic Energy (IN))

       

       

    • 10:50 → 10:55
      TPC Clusterization / OpenCL / Highly Ionizing Particles 5m
      Speaker: Felix Weiglhofer (Goethe University Frankfurt (DE))

      1. Orchestrate the access of ALICE to the CERN NGT benchmark/development hardware, and alternatively to other CERN resources, to benchmark the ALICE GPU code on modern / new GPUs.
      2. Establish a software infrastructure / scripts, to run the ALICE GPU processing benchmarks on CPUs and GPUs over different data set, and provide performance results in a way that enables further analysis.
        - In principle the standalone benchmark should be enough, but pp and Pb-Pb data sets of different interaction rates should be tested.
      3. Extend the Alice O2 CI infrastructure to run actual code on the GPU, check for correctness and performance, for pull requests.
        - Should be a fast test for the CI. The standalone benchmark should be enough. Ideally we can perhaps use the deterministic mode, to validate results between CPU and GPU.
      4. Establish automated continuous benchmarking, to monitor the evolution of the GPU processing performance per O2 version.
        - We already have/had such  benchmarking on the alibi server with some old GPUs. But this should be extended, we need to use more modern GPUs, and we need monitoring and alerting. 

Powered by Indico v3.3.7-pre