.gitlab-ci.yml

@@ -16,7 +16,7 @@ variables:
   PROFILE_DEVICE: "a5000"
   RUN_THROUGHPUT_OPTIONS_CUDAPROF: "-n 500 -m 500 -r 1 -t 1"
   RUN_THROUGHPUT_OPTIONS_CUDA: "-n 500 -m 500 -r 1000 -t 16"
-  RUN_THROUGHPUT_OPTIONS_HIP: "-n 2800 --events-per-slice 2800 -m 2800 -t 10 -r 1000"
+  RUN_THROUGHPUT_OPTIONS_HIP: "-n 2800 --events-per-slice 2800 -m 2800 -t 10 -r 100"
   RUN_THROUGHPUT_OPTIONS_CPU: "-n 100 -m 100 -r 200"
 
   AVG_THROUGHPUT_DECREASE_THRESHOLD: "-2.5" # (%); fail throughput check if averaged throughput % change falls below -2.5%
@@ -152,4 +152,4 @@ test_pages:
     paths:
       - .cache/pip
 
-include: "scripts/ci/config/main.yaml"
\ No newline at end of file
+include: "scripts/ci/config/main.yaml"

CMakeLists.txt

@@ -10,7 +10,7 @@ if(STANDALONE AND NOT DEFINED CMAKE_TOOLCHAIN_FILE AND DEFINED ENV{CMAKE_TOOLCHA
   set(CMAKE_TOOLCHAIN_FILE $ENV{CMAKE_TOOLCHAIN_FILE})
 endif()
 
-project(Allen VERSION 2.0
+project(Allen VERSION 2.1
         LANGUAGES CXX)
 
 # ---------------
@@ -70,6 +70,16 @@ set_property(CACHE CUDA_ARCH PROPERTY STRINGS COMPATIBILITY MAX MIN 53 60 61 62
 set(HIP_ARCH gfx908 CACHE STRING "HIP target architecture")
 set_property(CACHE HIP_ARCH PROPERTY STRINGS gfx900 gfx902 gfx904 gfx906 gfx908 gfx909 gfx1010 gfx1011 gfx1012 gfx1030)
 
+# Forces a no separable compilation build
+# Note: For now, HIP *requires* a no separable compilation build
+option(SEPARABLE_COMPILATION "Sets the build to no separable compilation" ON)
+if(TARGET_DEVICE STREQUAL "HIP")
+  set(SEPARABLE_COMPILATION OFF)
+  message(STATUS "Separable compilation: OFF (required by HIP)")
+else()
+  message(STATUS "Separable compilation: " ${SEPARABLE_COMPILATION})
+endif()
+
 # --------------------
 # Unrestricted options
 # --------------------
@@ -135,7 +145,7 @@ if(ENABLE_CONTRACTS)
   add_compile_definitions(ENABLE_CONTRACTS)
 endif()
 
-# Chioce of profiling with callgrind
+# Choice of profiling with callgrind
 option(CALLGRIND_PROFILE "Set Callgrind macros around sequence" OFF)
 if(CALLGRIND_PROFILE)
   add_compile_definitions(CALLGRIND_PROFILE)
@@ -240,7 +250,7 @@ if(TARGET_DEVICE STREQUAL "CUDA")
       string(APPEND DEVICE_ARCHITECTURE " -gencode=arch=compute_${arch},code=sm_${arch}")
     endforeach()
     if(CUDA_ARCH STREQUAL "COMPATIBILITY")
-      string(APPEND DEVICE_ARCHITECTURE " -gencode=arch=compute_86,code=compute_86")
+      string(APPEND DEVICE_ARCHITECTURE " -gencode=arch=compute_86,code=sm_86")
     endif()
   endif()
   message(STATUS "Device architecture set to ${DEVICE_ARCHITECTURE}")
@@ -264,7 +274,7 @@ if(TARGET_DEVICE STREQUAL "CPU")
 
     add_library(${ARGV})
     add_library(Allen::${ARGV0} ALIAS ${ARGV0})
-    target_compile_definitions(${ARGV0} PRIVATE ${TARGET_DEFINITION})
+    target_compile_definitions(${ARGV0} PRIVATE ${TARGET_DEFINITION} ODIN_WITHOUT_GAUDI)
     install(TARGETS ${ARGV0}
       EXPORT Allen
       LIBRARY DESTINATION lib)
@@ -288,7 +298,7 @@ if(TARGET_DEVICE STREQUAL "CPU")
     endif()
 
     add_library(Allen::${ARGV0} ALIAS ${ARGV0})
-    target_compile_definitions(${ARGV0} PRIVATE ${TARGET_DEFINITION})
+    target_compile_definitions(${ARGV0} PRIVATE ${TARGET_DEFINITION} ODIN_WITHOUT_GAUDI)
     install(TARGETS ${ARGV0}
       EXPORT Allen
       DESTINATION ${destination})
@@ -313,7 +323,6 @@ if(TARGET_DEVICE STREQUAL "CPU")
    endif()
 
 elseif(TARGET_DEVICE STREQUAL "HIP")
-
   # Deal with build type (HIP)
   if(NOT CMAKE_BUILD_TYPE)
     set(HIP_RELEASE_OPTIONS "-O3 -g -DNDEBUG")
@@ -322,7 +331,7 @@ elseif(TARGET_DEVICE STREQUAL "HIP")
   elseif(${CMAKE_BUILD_TYPE} STREQUAL "Release")
     set(HIP_RELEASE_OPTIONS "-O3 -DNDEBUG")
   elseif(${CMAKE_BUILD_TYPE} STREQUAL "Debug")
-    set(HIP_RELEASE_OPTIONS "-g -DNDEBUG")
+    set(HIP_RELEASE_OPTIONS "-O0 -g -DALLEN_DEBUG")
   endif()
 
   if(GCC_TOOLCHAIN_ROOT)
@@ -362,7 +371,7 @@ elseif(TARGET_DEVICE STREQUAL "HIP")
 
     add_library(Allen::${ARGV0} ALIAS ${ARGV0})
     target_include_directories(${ARGV0} PRIVATE ${HIP_PATH}/include ${ROCM_PATH}/hsa/include)
-    target_compile_definitions(${ARGV0} PRIVATE ${TARGET_DEFINITION})
+    target_compile_definitions(${ARGV0} PRIVATE ${TARGET_DEFINITION} ODIN_WITHOUT_GAUDI)
 
     install(TARGETS ${ARGV0}
       EXPORT Allen
@@ -391,8 +400,11 @@ elseif(TARGET_DEVICE STREQUAL "CUDA")
     set(CUDA_VERBOSE_FLAGS "")
   endif()
 
-  string(REGEX REPLACE "toolchain/.*" "toolchain" CXX_COMPILER_DIR ${CMAKE_CXX_COMPILER})
-  string(APPEND CMAKE_CUDA_FLAGS " ${DEVICE_ARCHITECTURE} -ccbin=${CXX_COMPILER_DIR} ${CUDA_TOOLCHAIN_OPTIONS} -Xcudafe --display_error_number --use_fast_math --expt-relaxed-constexpr ${CUDA_VERBOSE_FLAGS}")
+  string(APPEND CMAKE_CUDA_FLAGS " ${DEVICE_ARCHITECTURE} ${CUDA_TOOLCHAIN_OPTIONS} -Xcudafe --display_error_number --use_fast_math --expt-relaxed-constexpr ${CUDA_VERBOSE_FLAGS}")
+  if (NOT CMAKE_CUDA_HOST_COMPILER)
+    string(REGEX REPLACE "toolchain/.*" "toolchain" CXX_COMPILER_DIR ${CMAKE_CXX_COMPILER})
+    string(APPEND CMAKE_CUDA_FLAGS " -ccbin=${CXX_COMPILER_DIR} ")
+  endif()
 
   message(STATUS "CMAKE CUDA FLAGS:" ${CMAKE_CUDA_FLAGS})
   message(STATUS "Detected CUDA include directory: " ${CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES})
@@ -407,8 +419,11 @@ elseif(TARGET_DEVICE STREQUAL "CUDA")
     endif()
 
     add_library(${ARGV})
+    if(SEPARABLE_COMPILATION)
+      set_property(TARGET ${ARGV0} PROPERTY CUDA_SEPARABLE_COMPILATION ON)
+    endif()
     target_include_directories(${ARGV0} PRIVATE ${CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES} ${PROJECT_BINARY_DIR}/code_generation)
-    target_compile_definitions(${ARGV0} PRIVATE ${TARGET_DEFINITION})
+    target_compile_definitions(${ARGV0} PRIVATE ${TARGET_DEFINITION} ODIN_WITHOUT_GAUDI)
     add_library(Allen::${ARGV0} ALIAS ${ARGV0})
     install(TARGETS ${ARGV0}
       EXPORT Allen
@@ -425,7 +440,9 @@ elseif(TARGET_DEVICE STREQUAL "CUDA")
     endif()
 
     add_library(${ARGV})
-
+    if(SEPARABLE_COMPILATION)
+      set_property(TARGET ${ARGV0} PROPERTY CUDA_SEPARABLE_COMPILATION ON)
+    endif()
     target_include_directories(${ARGV0} PRIVATE ${CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES})
     target_compile_definitions(${ARGV0} PRIVATE ${TARGET_DEFINITION} ODIN_WITHOUT_GAUDI)
     add_library(Allen::${ARGV0} ALIAS ${ARGV0})
@@ -436,6 +453,9 @@ elseif(TARGET_DEVICE STREQUAL "CUDA")
 
   function(allen_add_executable)
     add_executable(${ARGV})
+    if(SEPARABLE_COMPILATION)
+      set_property(TARGET ${ARGV0} PROPERTY CUDA_SEPARABLE_COMPILATION ON)
+    endif()
     target_include_directories(${ARGV0} PRIVATE ${CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES})
     target_compile_definitions(${ARGV0} PRIVATE ${TARGET_DEFINITION})
     install(TARGETS ${ARGV0} RUNTIME DESTINATION bin)
@@ -448,6 +468,50 @@ elseif(TARGET_DEVICE STREQUAL "CUDA")
   endif()
 endif()
 
+# Mechanism to mimic a UNITY build
+set(COUNTER 0)
+set(GROUPCOMP_DIR ${PROJECT_BINARY_DIR}/group_comp)
+file(MAKE_DIRECTORY ${GROUPCOMP_DIR})
+
+function(allen_add_device_library_unified)
+  set(destination lib)
+
+  foreach(arg IN LISTS ARGN)
+    if(${arg} MATCHES "\\.cu$")
+      list(APPEND LIBRARY_SOURCES ${arg})
+    else()
+      list(APPEND LIBRARY_SETTINGS ${arg})
+    endif()
+  endforeach()
+
+  set(UNIFIED_FILENAME "${GROUPCOMP_DIR}/group_${COUNTER}.cu")
+  add_custom_command(
+    OUTPUT "${UNIFIED_FILENAME}"
+    COMMAND ${CMAKE_COMMAND} -E remove "${UNIFIED_FILENAME}" && cat ${LIBRARY_SOURCES} >> "${UNIFIED_FILENAME}"
+    DEPENDS "${LIBRARY_SOURCES}")
+
+  if(TARGET_DEVICE STREQUAL "CPU")
+    add_library(${LIBRARY_SETTINGS} ${UNIFIED_FILENAME})
+    set_source_files_properties(${UNIFIED_FILENAME} PROPERTIES LANGUAGE CXX)
+  elseif(TARGET_DEVICE STREQUAL "HIP")
+    hip_add_library(${LIBRARY_SETTINGS} ${UNIFIED_FILENAME} HIPCC_OPTIONS ${HIPCC_OPTIONS})
+    target_include_directories(${ARGV0} PRIVATE ${HIP_PATH}/include ${ROCM_PATH}/hsa/include)
+  elseif(TARGET_DEVICE STREQUAL "CUDA")
+    add_library(${LIBRARY_SETTINGS} ${UNIFIED_FILENAME})
+    target_include_directories(${ARGV0} PRIVATE ${CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES})
+  endif()
+
+  target_sources(${ARGV0} PRIVATE ${UNIFIED_FILENAME} ${LIBRARY_SOURCES})
+  add_library(Allen::${ARGV0} ALIAS ${ARGV0})
+  target_compile_definitions(${ARGV0} PRIVATE ${TARGET_DEFINITION})
+
+  install(TARGETS ${ARGV0}
+    EXPORT Allen
+    LIBRARY DESTINATION ${destination})
+
+  MATH(EXPR COUNTER "${COUNTER}+1")
+endfunction()
+
 if(CMAKE_CXX_COMPILER_ID STREQUAL "Clang")
   string(APPEND CMAKE_CXX_FLAGS " -Wno-gnu-zero-variadic-macro-arguments")
 endif()
@@ -523,6 +587,10 @@ target_include_directories(AllenCommon
   $<INSTALL_INTERFACE:${CMAKE_INSTALL_INCLUDEDIR}>
   $<BUILD_INTERFACE:${PROJECT_SOURCE_DIR}/test/contracts/include>)
 
+if (STANDALONE)
+  target_compile_definitions(AllenCommon INTERFACE ALLEN_STANDALONE)
+endif()
+
 set(AllenHeaders)
 foreach(header
     Allen.h
@@ -540,6 +608,7 @@ foreach(header
     InputProvider.h
     OutputHandler.h
     SliceUtils.h
+    FileSystem.h
     TransposeTypes.h)
   list(APPEND AllenHeaders main/include/${header})
 endforeach()
@@ -568,10 +637,10 @@ endif()
 
 message(STATUS "CMAKE CXX FLAGS: " ${CMAKE_CXX_FLAGS})
 
+
 # Add host library
 allen_add_host_library(HostCommon STATIC
   main/src/BankTypes.cpp
-  main/src/CheckVP.cpp
   main/src/InputReader.cpp
   main/src/InputTools.cpp
   main/src/Logger.cpp
@@ -595,6 +664,8 @@ allen_add_host_library(AllenLib SHARED
   main/src/RegisterConsumers.cpp
   main/src/ZMQOutputSender.cpp)
 
+
+
 target_compile_definitions(AllenLib PUBLIC ${TARGET_DEFINITION})
 
 target_link_libraries(AllenLib PUBLIC
@@ -642,3 +713,22 @@ if (NOT STANDALONE)
   # Finalize configuration
   lhcb_finalize_configuration()
 endif()
+
+
+# Download parameter files if in STANDALONE
+if (STANDALONE)
+  find_package(Git REQUIRED)
+  file(MAKE_DIRECTORY "${CMAKE_CURRENT_SOURCE_DIR}/input/PARAM/")
+  SET(PARAMFILESROOT "${CMAKE_CURRENT_SOURCE_DIR}/input/PARAM/ParamFiles/")
+  add_custom_command(
+    OUTPUT "${PARAMFILESROOT}"
+    COMMENT "Checking out parameterization files from the LHCb stack"
+    COMMAND
+      ${CMAKE_COMMAND} -E env ${GIT_EXECUTABLE} clone https://gitlab.cern.ch/lhcb-datapkg/ParamFiles.git --no-checkout &&
+      ${CMAKE_COMMAND} -E env ${GIT_EXECUTABLE} --work-tree=ParamFiles --git-dir=ParamFiles/.git checkout origin/master -- data
+    WORKING_DIRECTORY "${CMAKE_CURRENT_SOURCE_DIR}/input/PARAM/")
+  add_custom_target(checkout_param_files DEPENDS "${PARAMFILESROOT}")
+  add_dependencies(AllenLib checkout_param_files)
+  add_compile_definitions(PARAMFILESROOTPATH=${PARAMFILESROOT})
+  message(STATUS "PARAMFILESROOT set to ${PARAMFILESROOT}")
+endif()

Dumpers/BinaryDumpers/CMakeLists.txt

@@ -11,6 +11,7 @@ gaudi_add_library(BinaryDumpers
                     src/lib/TestUTBoards.cpp
                     src/lib/Utils.cpp
                   LINK PUBLIC
+                    AllenCommon
                     Gaudi::GaudiKernel
                     LHCb::DAQEventLib
                     LHCb::MuonDAQLib)

Dumpers/BinaryDumpers/include/Dumpers/Utils.h

@@ -16,13 +16,7 @@
 #include <type_traits>
 #include <vector>
 
-template<class Tuple, class F>
-void for_each(Tuple&& tup, F&& f)
-{
-  std::apply(
-    [f = std::forward<F>(f)](auto&&... args) { (std::invoke(f, std::forward<decltype(args)>(args)), ...); },
-    std::forward<Tuple>(tup));
-}
+#include <Common.h>
 
 template<class T>
 void optional_resize(T&, size_t)

Dumpers/BinaryDumpers/options/allen.py

@@ -5,6 +5,7 @@
 import os
 from GaudiPython.Bindings import AppMgr, gbl
 from Configurables import LHCbApp, CondDB, ApplicationMgr
+from Configurables import Gaudi__RootCnvSvc as RootCnvSvc
 from Allen.config import setup_allen_non_event_data_service
 from threading import Thread
 from time import sleep
@@ -25,7 +26,7 @@ interpreter.Declare("#include <Allen/Provider.h>")
 interpreter.Declare("#include <Dumpers/PyAllenHelper.h>")
 
 sequence_default = os.path.join(os.environ['ALLEN_INSTALL_DIR'], 'constants',
-                                'hlt1_pp_default')
+                                'hlt1_pp_default.json')
 
 
 def cast_service(return_type, svc):
@@ -42,10 +43,6 @@ parser.add_argument(
     "-g",
     dest="det_folder",
     default=os.path.join(allen_dir, "input", "detector_configuration", "down"))
-parser.add_argument(
-    "--params",
-    dest="param_folder",
-    default=os.path.join(allen_dir, "input", "parameters"))
 parser.add_argument("-n", dest="n_events", default=0)
 parser.add_argument("-t", dest="threads", default=1)
 parser.add_argument("-r", dest="repetitions", default=1)
@@ -59,8 +56,13 @@ parser.add_argument(
     dest="mdf",
     default=os.path.join(
         allen_dir, "input", "minbias", "mdf",
-        "MiniBrunel_2018_MinBias_FTv4_DIGI_retinacluster.mdf"))
+        "MiniBrunel_2018_MinBias_FTv4_DIGI_retinacluster_v1.mdf"))
 parser.add_argument("--mep", dest="mep", default=None)
+parser.add_argument(
+    "--mep-mask-source-id-top-5",
+    action="store_true",
+    dest="mask_top5",
+    default=False)
 parser.add_argument(
     "--reuse-meps",
     action="store_true",
@@ -83,6 +85,8 @@ parser.add_argument(
     help="Add profiler start and stop calls",
 )
 parser.add_argument("--output-file", dest="output_file", default=None)
+parser.add_argument(
+    "--output-batch-size", dest="output_batch_size", default=10)
 parser.add_argument(
     "--monitoring-save-period", dest="mon_save_period", default=0)
 parser.add_argument(
@@ -110,6 +114,8 @@ parser.add_argument(
     default=300,
     help="How long to run when reusing MEPs",
 )
+parser.add_argument(
+    "--tags", dest="tags", default="dddb-20171122,sim-20180530-vc-md100")
 
 args = parser.parse_args()
 
@@ -120,14 +126,14 @@ runtime_lib = None
 if args.profile == "CUDA":
     runtime_lib = ctypes.CDLL("libcudart.so")
 
+dddb_tag, conddb_tag = args.tags.split(',')
+
 app = LHCbApp(
     DataType="Upgrade",
     EvtMax=1000,
     Simulation=True,
-    DDDBtag="dddb-20171122",
-    CondDBtag="sim-20180530-vc-md100")
-# DDDBtag="dddb-20210617",  # tags for FEST sample from 10/2021
-# CondDBtag="sim-20210617-vc-md100")
+    DDDBtag=dddb_tag,  # tags for FEST sample from 10/2021
+    CondDBtag=conddb_tag)
 
 # Upgrade DBs
 CondDB().Upgrade = True
@@ -138,24 +144,29 @@ setup_allen_non_event_data_service()
 ApplicationMgr().EvtSel = "NONE"
 ApplicationMgr().ExtSvc += ["ToolSvc", "AuditorSvc", "ZeroMQSvc"]
 
+# RootCnvSvc becauce it sets up a bunch of ROOT IO stuff
+rootSvc = RootCnvSvc("RootCnvSvc", EnableIncident=1)
+ApplicationMgr().ExtSvc += [rootSvc]
+
 if args.mep is not None:
     ApplicationMgr().ExtSvc += ["AllenConfiguration", "MEPProvider"]
     from Configurables import MEPProvider, AllenConfiguration
 
     allen_conf = AllenConfiguration("AllenConfiguration")
     allen_conf.JSON = args.sequence
-    allen_conf.OutputLevel = 2
+    allen_conf.OutputLevel = 3
 
     mep_provider = MEPProvider()
     mep_provider.NSlices = args.slices
     mep_provider.EventsPerSlice = args.events_per_slice
-    mep_provider.OutputLevel = 2
+    mep_provider.OutputLevel = 3
     # Number of MEP buffers and number of transpose/offset threads
     mep_provider.BufferConfig = (10, 8)
     mep_provider.TransposeMEPs = False
     mep_provider.SplitByRun = False
     mep_provider.Source = "Files"
-    mep_dir = args.mep
+    mep_provider.MaskSourceIDTop5 = args.mask_top5
+    mep_dir = os.path.expandvars(args.mep)
     if os.path.isdir(mep_dir):
         mep_provider.Connections = sorted([
             os.path.join(mep_dir, mep_file) for mep_file in os.listdir(mep_dir)
@@ -181,13 +192,15 @@ updater = cast_service(gbl.Allen.NonEventData.IUpdater, svc)
 
 # options map
 options = gbl.std.map("std::string", "std::string")()
-for flag, value in [("g", args.det_folder), ("params", args.param_folder),
+for flag, value in [("g", args.det_folder),
+                    ("params", os.getenv("PARAMFILESROOT")),
                     ("n", args.n_events), ("t", args.threads),
                     ("r", args.repetitions), ("output-file", args.output_file),
+                    ("output-batch-size", args.output_batch_size),
                     ("m", args.reserve), ("v", args.verbosity),
                     ("p", args.print_memory),
                     ("sequence", os.path.expandvars(args.sequence)),
-                    ("s", args.slices), ("mdf", args.mdf),
+                    ("s", args.slices), ("mdf", os.path.expandvars(args.mdf)),
                     ("cpu-offload", args.cpu_offload),
                     ("disable-run-changes", int(not args.enable_run_changes)),
                     ("monitoring-save-period", args.mon_save_period),

Dumpers/BinaryDumpers/src/DumpFTGeometry.cpp

@@ -8,6 +8,7 @@
 #include <vector>
 
 #include "DumpFTGeometry.h"
+#include "FTDAQ/FTReadoutMap.h"
 #include "Detector/FT/FTChannelID.h"
 
 namespace {
@@ -109,55 +110,51 @@ DumpUtils::Dumps DumpFTGeometry::dumpGeometry() const
   }
 #endif
   // Raw bank layout (from FTReadoutTool)
-  vector<uint32_t> bank_first_channel;
   string conditionLocation = "/dd/Conditions/ReadoutConf/FT/ReadoutMap";
   Condition* rInfo = getDet<Condition>(conditionLocation);
-
-  std::vector<int> stations_ = rInfo->param<std::vector<int>>("FTBankStation");
-  std::vector<int> layers = rInfo->param<std::vector<int>>("FTBankLayer");
-  std::vector<int> quarters = rInfo->param<std::vector<int>>("FTBankQuarter");
-  std::vector<int> firstModules = rInfo->param<std::vector<int>>("FTBankFirstModule");
-  std::vector<int> firstMats = rInfo->param<std::vector<int>>("FTBankFirstMat");
-
-  // Construct the first channel attribute
-  uint32_t number_of_tell40s = stations_.size();
-  bank_first_channel.reserve(number_of_tell40s);
-  for (unsigned int i = 0; i < number_of_tell40s; i++) {
-    bank_first_channel.push_back(
-      static_cast<uint32_t>(LHCb::Detector::FTChannelID {LHCb::Detector::FTChannelID::StationID(stations_[i]),
-                                                         LHCb::Detector::FTChannelID::LayerID(layers[i]),
-                                                         LHCb::Detector::FTChannelID::QuarterID(quarters[i]),
-                                                         LHCb::Detector::FTChannelID::ModuleID(firstModules[i]),
-                                                         LHCb::Detector::FTChannelID::MatID(firstMats[i]),
-                                                         0u,
-                                                         0u}));
-  }
+  auto readoutMap = FTReadoutMap {this, *rInfo};
   DumpUtils::Writer output {};
-  output.write(
-    number_of_stations,
-    number_of_layers_per_station,
-    number_of_layers,
-    number_of_quarters_per_layer,
-    number_of_quarters,
-    number_of_modules,
-    number_of_mats_per_module,
-    number_of_mats,
-    number_of_tell40s,
-    bank_first_channel,
-    max_uniqueMat,
-    mirrorPointX,
-    mirrorPointY,
-    mirrorPointZ,
-    ddxX,
-    ddxY,
-    ddxZ,
-    uBegin,
-    halfChannelPitch,
-    dieGap,
-    sipmPitch,
-    dxdy,
-    dzdy,
-    globaldy);
 
+  auto comp = readoutMap.compatibleVersions();
+  if (comp.count(4)) {
+    auto number_of_tell40s = readoutMap.nBanks();
+    // Decoding v6
+    vector<uint32_t> bank_first_channel;
+    bank_first_channel.reserve(number_of_tell40s);
+    for (unsigned int i = 0; i < number_of_tell40s; i++) {
+      bank_first_channel.push_back(readoutMap.channelIDShift(i));
+    }
+    output.write(
+      number_of_stations,
+      number_of_layers_per_station,
+      number_of_layers,
+      number_of_quarters_per_layer,
+      number_of_quarters,
+      number_of_modules,
+      number_of_mats_per_module,
+      number_of_mats,
+      number_of_tell40s,
+      bank_first_channel,
+      max_uniqueMat,
+      mirrorPointX,
+      mirrorPointY,
+      mirrorPointZ,
+      ddxX,
+      ddxY,
+      ddxZ,
+      uBegin,
+      halfChannelPitch,
+      dieGap,
+      sipmPitch,
+      dxdy,
+      dzdy,
+      globaldy);
+  }
+  else {
+    std::stringstream s;
+    Gaudi::Utils::toStream(comp, s);
+    throw GaudiException {"Unsupported conditions compatible with " + s.str(), __FILE__, StatusCode::FAILURE};
+  }
+  return {{tuple {output.buffer(), "scifi_geometry", Allen::NonEventData::SciFiGeometry::id}}};
   return {{tuple {output.buffer(), "scifi_geometry", Allen::NonEventData::SciFiGeometry::id}}};
 }

Dumpers/BinaryDumpers/src/TransposeRawBanks.cpp

@@ -6,14 +6,20 @@
 #include <string>
 #include <unordered_map>
 #include <vector>
-#include <GaudiAlg/MergingTransformer.h>
+
 #include <AIDA/IHistogram1D.h>
+
+#include <GaudiAlg/MergingTransformer.h>
+#include <GaudiAlg/GaudiHistoAlg.h>
+#include <GaudiKernel/GaudiException.h>
+
 #include <Event/ODIN.h>
 #include <Event/RawBank.h>
 #include <Event/RawEvent.h>
-#include <GaudiAlg/GaudiHistoAlg.h>
+
 #include <Dumpers/Utils.h>
-#include <GaudiKernel/GaudiException.h>
+
+#include <BankTypes.h>
 
 template<typename T>
 using VOC = Gaudi::Functional::vector_of_const_<T>;
@@ -41,17 +47,16 @@ using VOC = Gaudi::Functional::vector_of_const_<T>;
  *  @author Roel Aaij
  *  @date   2018-08-27
  */
-class TransposeRawBanks
-  : public Gaudi::Functional::MergingTransformer<
-      std::array<std::tuple<std::vector<char>, int>, LHCb::RawBank::types().size()>(VOC<LHCb::RawEvent*> const&),
-      Gaudi::Functional::Traits::BaseClass_t<GaudiHistoAlg>> {
+class TransposeRawBanks : public Gaudi::Functional::MergingTransformer<
+                            std::array<TransposedBanks, LHCb::RawBank::types().size()>(VOC<LHCb::RawEvent*> const&),
+                            Gaudi::Functional::Traits::BaseClass_t<GaudiHistoAlg>> {
 public:
   /// Standard constructor
   TransposeRawBanks(const std::string& name, ISvcLocator* pSvcLocator);
 
   StatusCode initialize() override;
 
-  std::array<std::tuple<std::vector<char>, int>, LHCb::RawBank::types().size()> operator()(
+  std::array<TransposedBanks, LHCb::RawBank::types().size()> operator()(
     VOC<LHCb::RawEvent*> const& rawEvents) const override;
 
 private:
@@ -86,11 +91,11 @@ StatusCode TransposeRawBanks::initialize()
   return StatusCode::SUCCESS;
 }
 
-std::array<std::tuple<std::vector<char>, int>, LHCb::RawBank::types().size()> TransposeRawBanks::operator()(
+std::array<TransposedBanks, LHCb::RawBank::types().size()> TransposeRawBanks::operator()(
   VOC<LHCb::RawEvent*> const& rawEvents) const
 {
 
-  std::array<std::tuple<std::vector<char>, int>, LHCb::RawBank::types().size()> output;
+  std::array<TransposedBanks, LHCb::RawBank::types().size()> output;
   std::array<LHCb::RawBank::View, LHCb::RawBank::types().size()> rawBanks;
 
   for (auto const* rawEvent : rawEvents) {
@@ -109,12 +114,16 @@ std::array<std::tuple<std::vector<char>, int>, LHCb::RawBank::types().size()> Tr
     auto const& banks = rawBanks[bt];
     if (banks.empty()) continue;
 
-    const uint32_t number_of_rawbanks = banks.size();
+    const uint32_t nBanks = banks.size();
     uint32_t offset = 0;
 
     std::vector<uint32_t> bankOffsets;
-    std::vector<uint32_t> bankData;
     bankOffsets.push_back(0);
+    std::vector<uint32_t> bankData;
+    std::vector<uint16_t> bankSizes;
+    bankSizes.reserve(nBanks);
+    std::vector<uint8_t> bankTypes;
+    bankTypes.reserve(nBanks);
 
     for (auto& bank : banks) {
       const uint32_t sourceID = static_cast<uint32_t>(bank->sourceID());
@@ -141,12 +150,15 @@ std::array<std::tuple<std::vector<char>, int>, LHCb::RawBank::types().size()> Tr
         offset++;
       }
       bankOffsets.push_back(offset * sizeof(uint32_t));
+      bankSizes.push_back(bank->size());
+      bankTypes.push_back(static_cast<uint8_t>(bank->type()));
     }
 
     // Dumping number_of_rawbanks + 1 offsets!
     DumpUtils::Writer bank_buffer;
-    bank_buffer.write(number_of_rawbanks, bankOffsets, bankData);
-    output[bt] = std::tuple {bank_buffer.buffer(), banks[0]->version()};
+    bank_buffer.write(nBanks, bankOffsets, bankData);
+    output[bt] =
+      TransposedBanks {bank_buffer.buffer(), std::move(bankSizes), std::move(bankTypes), banks[0]->version()};
   }
   return output;
 }

Rec/Allen/Allen/AlgorithmConversionTools.h

@@ -10,6 +10,7 @@
 \*****************************************************************************/
 #pragma once
 
+#include <algorithm>
 #include <cstdio>
 #include <BackendCommon.h>
 #include <Argument.cuh>
@@ -20,6 +21,9 @@
 #include <GaudiKernel/ParsersFactory.h>
 #include <GaudiKernel/StdArrayAsProperty.h>
 
+// Needed for the Gaudi wrappers when doing a device build in the stack
+using std::min;
+
 namespace Allen {
   // Shortcut for type used in input / outputs of Allen - Gaudi wrappers
   template<typename T>

Rec/Allen/python/Allen/qmtest/exclusions.py

@@ -16,6 +16,7 @@ remove_throughput = LineSkipper(regexps=[
     # Throughput messages
     r"\s*(\d+\.\d+)\s+events/s",
     r"Ran test for (\d+\.\d+)\s+seconds",
+    r"Providing banks for",
     r"Providing events in.*",
     r"Opened\s.*",
     r"Cannot read more data.*"

Rec/Allen/python/Allen/qmtest/validators.py

+###############################################################################
+# (c) Copyright 2022 CERN for the benefit of the LHCb Collaboration           #
+#                                                                             #
+# This software is distributed under the terms of the GNU General Public      #
+# Licence version 3 (GPL Version 3), copied verbatim in the file "COPYING".   #
+#                                                                             #
+# In applying this licence, CERN does not waive the privileges and immunities #
+# granted to it by virtue of its status as an Intergovernmental Organization  #
+# or submit itself to any jurisdiction.                                       #
+###############################################################################
+def checkSizes(reference_file, causes, result):
+    import os
+
+    ref_sizes = {}
+    with open(os.path.expandvars(reference_file)) as ref:
+        for line in ref:
+            f, s = line.strip().split()
+            ref_sizes[f] = int(s)
+
+    missing = []
+    bad_size = []
+    for f, s in ref_sizes.items():
+        if not os.path.exists(f):
+            missing.append(f)
+        else:
+            size = os.stat(f).st_size
+            if size != s:
+                bad_size.append((f, size, s))
+
+    result_str = ''
+    if not missing and not bad_size:
+        result_str += 'Files and sizes match\n'
+    if missing:
+        causes.append("missing files")
+        result_str += 'Missing:\n\t' + '\n\t'.join(missing) + '\n'
+    if bad_size:
+        pat = "%s is %d bytes instead of %d"
+        causes.append("file size mismatch(es)")
+        result_str += (
+            'Wrong size:\n\t' + '\n\t'.join([pat % e
+                                             for e in bad_size]) + '\n')
+
+    result["validator.output"] = result.Quote(result_str)

Rec/Allen/src/ProvideRuntimeOptions.cpp

@@ -21,8 +21,8 @@
 #include <Constants.cuh>
 #include <Logger.h>
 
-class ProvideRuntimeOptions final : public Gaudi::Functional::Transformer<RuntimeOptions(
-                                      std::array<std::tuple<std::vector<char>, int>, LHCb::RawBank::LastType> const&)> {
+class ProvideRuntimeOptions final
+  : public Gaudi::Functional::Transformer<RuntimeOptions(std::array<TransposedBanks, LHCb::RawBank::LastType> const&)> {
 
 public:
   /// Standard constructor
@@ -31,8 +31,7 @@ public:
   StatusCode initialize() override;
 
   /// Algorithm execution
-  RuntimeOptions operator()(
-    std::array<std::tuple<std::vector<char>, int>, LHCb::RawBank::LastType> const& allen_banks) const override;
+  RuntimeOptions operator()(std::array<TransposedBanks, LHCb::RawBank::LastType> const& allen_banks) const override;
 
 private:
   Gaudi::Property<std::string> m_monitorFile {this, "MonitorFile", "allen_monitor.root"};
@@ -57,7 +56,7 @@ StatusCode ProvideRuntimeOptions::initialize()
 }
 
 RuntimeOptions ProvideRuntimeOptions::operator()(
-  std::array<std::tuple<std::vector<char>, int>, LHCb::RawBank::LastType> const& allen_banks) const
+  std::array<TransposedBanks, LHCb::RawBank::LastType> const& allen_banks) const
 {
   const unsigned number_of_repetitions = 1;
   const bool do_check = false;

Rec/Allen/src/RunAllen.cpp

@@ -62,15 +62,15 @@ StatusCode RunAllen::initialize()
   }
 
   // get constants
-  const std::string params_path = resolveEnvVars(m_paramDir);
+  const std::string params_path = resolveEnvVars(m_paramDir) + "/data";
 
   std::vector<float> muon_field_of_interest_params;
-  read_muon_field_of_interest(muon_field_of_interest_params, params_path + "/field_of_interest_params.bin");
+  read_muon_field_of_interest(muon_field_of_interest_params, params_path + "/allen_muon_field_of_interest_params.bin");
 
-  m_constants.reserve_and_initialize(muon_field_of_interest_params, params_path + "/params_kalman_FT6x2/");
+  m_constants.reserve_and_initialize(muon_field_of_interest_params, params_path);
 
   std::unique_ptr<CatboostModelReader> muon_catboost_model_reader =
-    std::make_unique<CatboostModelReader>(params_path + "/muon_catboost_model.json");
+    std::make_unique<CatboostModelReader>(params_path + "/allen_muon_catboost_model.json");
   m_constants.initialize_muon_catboost_model_constants(
     muon_catboost_model_reader->n_trees(),
     muon_catboost_model_reader->tree_depths(),
@@ -81,7 +81,7 @@ StatusCode RunAllen::initialize()
     muon_catboost_model_reader->split_feature());
 
   std::unique_ptr<TwoTrackMVAModelReader> two_track_mva_model_reader =
-    std::make_unique<TwoTrackMVAModelReader>(params_path + "/two_track_mva_model.json");
+    std::make_unique<TwoTrackMVAModelReader>(params_path + "/allen_two_track_mva_model.json");
   m_constants.initialize_two_track_mva_model_constants(
     two_track_mva_model_reader->weights(),
     two_track_mva_model_reader->biases(),
@@ -181,7 +181,7 @@ StatusCode RunAllen::finalize()
 /** Calls Allen for one event
  */
 std::tuple<bool, HostBuffers> RunAllen::operator()(
-  const std::array<std::tuple<std::vector<char>, int>, LHCb::RawBank::types().size()>& allen_banks,
+  const std::array<TransposedBanks, LHCb::RawBank::types().size()>& allen_banks,
   const LHCb::ODIN&) const
 {
   // Initialize input provider

Rec/Allen/src/RunAllen.h

@@ -31,10 +31,9 @@
 // STL includes
 #include <deque>
 
-class RunAllen final
-  : public Gaudi::Functional::MultiTransformerFilter<std::tuple<HostBuffers>(
-      const std::array<std::tuple<std::vector<char>, int>, LHCb::RawBank::types().size()>& allen_banks,
-      const LHCb::ODIN& odin)> {
+class RunAllen final : public Gaudi::Functional::MultiTransformerFilter<std::tuple<HostBuffers>(
+                         const std::array<TransposedBanks, LHCb::RawBank::types().size()>& allen_banks,
+                         const LHCb::ODIN& odin)> {
 public:
   /// Standard constructor
   RunAllen(const std::string& name, ISvcLocator* pSvcLocator);
@@ -47,7 +46,7 @@ public:
 
   /// Algorithm execution
   std::tuple<bool, HostBuffers> operator()(
-    const std::array<std::tuple<std::vector<char>, int>, LHCb::RawBank::types().size()>& allen_banks,
+    const std::array<TransposedBanks, LHCb::RawBank::types().size()>& allen_banks,
     const LHCb::ODIN& odin) const override;
 
 private:
@@ -63,11 +62,9 @@ private:
   std::unique_ptr<ROOTService> m_root_service;
 
   Gaudi::Property<std::string> m_sequence {this, "Sequence", "hlt1_pp_default"};
-
   Gaudi::Property<std::string> m_updaterName {this, "UpdaterName", "AllenUpdater"};
-
+  Gaudi::Property<std::string> m_paramDir {this, "ParamDir", "${PARAMFILESROOT}"};
   Gaudi::Property<std::string> m_json {this, "JSON", "${ALLEN_INSTALL_DIR}/constants/hlt1_pp_default.json"};
-  Gaudi::Property<std::string> m_paramDir {this, "ParamDir", "${ALLEN_PROJECT_ROOT}/input/parameters"};
   Gaudi::Property<std::string> m_monitorFile {this, "MonitorFile", "allen_monitor.root"};
 
   // If set to false, events are only filtered by the GEC

Rec/Allen/tests/qmtest/mdf_input.qmt

@@ -16,7 +16,7 @@
   <argument name="args"><set>
     <text>${BINARYDUMPERSROOT}/options/allen.py</text>
     <text>--monitoring-filename</text><text>mdf_input_hists.root</text>
-    <text>--mdf</text><text>${ALLEN_PROJECT_ROOT}/input/minbias/mdf/MiniBrunel_2018_MinBias_FTv4_DIGI_retinacluster.mdf</text>
+    <text>--mdf</text><text>${ALLEN_PROJECT_ROOT}/input/minbias/mdf/MiniBrunel_2018_MinBias_FTv4_DIGI_retinacluster_v1.mdf</text>
   </set></argument>
   <argument name="reference"><text>../refs/allen_10ev.ref</text></argument>
   <argument name="timeout"><integer>600</integer></argument>

Rec/Allen/tests/refs/allen_10ev.ref

@@ -6,6 +6,7 @@ DetectorPersist...   INFO Added successfully Conversion service:XmlCnvSvc
 DetectorDataSvc   SUCCESS Detector description database: git:/lhcb.xml
 EventClockSvc.F...   INFO Event times generated from 0 with steps of 0
 DeFTDetector         INFO Current FT geometry version =   63
+DumpFTGeometry       INFO Conditions DB is compatible with FT bank version 4, 5, 6.
 DumpMuonGeometry     INFO Registering /dd/Conditions/ReadoutConf/Muon/Cabling/M2/Cabling
 DumpMuonGeometry     INFO Registering /dd/Conditions/ReadoutConf/Muon/Cabling/M3/Cabling
 DumpMuonGeometry     INFO Registering /dd/Conditions/ReadoutConf/Muon/Cabling/M4/Cabling

ReleaseNotes/v2r1.md

+2022-04-14 Allen v2r1
+===
+
+This version uses
+Rec [v34r2](../../../../Rec/-/tags/v34r2),
+Lbcom [v33r8](../../../../Lbcom/-/tags/v33r8),
+LHCb [v53r8](../../../../LHCb/-/tags/v53r8),
+Gaudi [v36r5](../../../../Gaudi/-/tags/v36r5),
+Detector [v1r1](../../../../Detector/-/tags/v1r1) and
+LCG [101](http://lcginfo.cern.ch/release/101/) with ROOT 6.24.06.
+
+This version is released on `master` branch.
+Built relative to Allen [v2r0](/../../tags/v2r0), with the following changes:
+
+### New features ~"new feature"
+
+- Beam Gas Line for VELO Alignment, !819 (@hyeung)
+
+
+### Enhancements ~enhancement
+
+- ~Tracking | Follow streamline of PR hits, !818 (@gunther)
+- Create and use custom is_trivially_copyable, !832 (@dcampora)
+
+
+### Code cleanups and changes to tests ~modernisation ~cleanup ~testing
+
+- ~Build | Use CMAKE_TOOLCHAIN_FILE, various cmake changes, !797 (@dcampora) [LBCORE-2051]
+- Use rocm-smi to print state of HIP device usage in CI, !833 (@roneil) [#314]
+- CI: whitelisted mi100-full: [run_throughput, RelWithDebInfo, hlt1_pp_veloSP, upgrade_mc_minbias_scifi_v5_000], !830 (@roneil)
+- Prevent full run jobs from starting when they shouldn't, !829 (@roneil)
+- Update build.rst, !828 (@dcampora)
+- Removed duplication of channelIDs between LHCb and Detector. Kept Detector ones, !825 (@sponce)
+- Explicitly enforce all Allen datatypes be trivially copyable, !821 (@dcampora)
+- Remove PatPV, !820 (@dcampora)

backend/include/AllenTypeTraits.h

@@ -97,23 +97,8 @@ namespace details {
 template<typename Tuple>
 using reverse_tuple_t = typename details::ReverseTuple<Tuple, std::make_index_sequence<std::tuple_size_v<Tuple>>>::type;
 
-namespace details {
-  template<typename...>
-  struct ConcatTuple;
-
-  template<typename... First, typename... Second>
-  struct ConcatTuple<std::tuple<First...>, std::tuple<Second...>> {
-    using type = std::tuple<First..., Second...>;
-  };
-
-  template<typename T1, typename T2, typename... Ts>
-  struct ConcatTuple<T1, T2, Ts...> {
-    using type = typename ConcatTuple<typename ConcatTuple<T1, T2>::type, Ts...>::type;
-  };
-} // namespace details
-
-template<typename... Tuples>
-using cat_tuples_t = typename details::ConcatTuple<Tuples...>::type;
+template<typename... input_t>
+using tuple_cat_t = decltype(std::tuple_cat(std::declval<input_t>()...));
 
 namespace details {
   template<typename T>
@@ -126,12 +111,12 @@ namespace details {
 
   template<typename... InTuple, typename... Ts>
   struct FlattenTuple<std::tuple<std::tuple<InTuple...>, Ts...>> {
-    using type = cat_tuples_t<std::tuple<InTuple...>, typename FlattenTuple<std::tuple<Ts...>>::type>;
+    using type = tuple_cat_t<std::tuple<InTuple...>, typename FlattenTuple<std::tuple<Ts...>>::type>;
   };
 
   template<typename T, typename... Ts>
   struct FlattenTuple<std::tuple<T, Ts...>> {
-    using type = cat_tuples_t<std::tuple<T>, typename FlattenTuple<std::tuple<Ts...>>::type>;
+    using type = tuple_cat_t<std::tuple<T>, typename FlattenTuple<std::tuple<Ts...>>::type>;
   };
 } // namespace details
 
@@ -231,6 +216,13 @@ namespace Allen {
           is_template_base_of_v<Base1, typename T::dev_particle_container_t>> {
   };
 
+  template<typename T, typename = void>
+  struct has_enable_monitoring : std::false_type {
+  };
+  template<typename T>
+  struct has_enable_monitoring<T, std::void_t<typename T::enable_monitoring_t>> : std::true_type {
+  };
+
   template<typename T>
   void initialize_algorithm(T& alg)
   {
@@ -245,4 +237,7 @@ namespace Allen {
 
   template<typename T>
   inline constexpr bool is_trivially_copyable_v = is_trivially_copyable<T>::value;
+
+  template<typename T>
+  using func_t = T (*)();
 } // namespace Allen

backend/include/BackendCommon.h

@@ -18,6 +18,7 @@
 #include <string>
 #include <cassert>
 #include <array>
+#include <gsl/gsl>
 #include "AllenTypeTraits.h"
 #include "BackendCommonInterface.h"
 
@@ -96,9 +97,10 @@ namespace Allen {
 #include <cmath>
 #endif
 
-// Replacement for gsl::span in device code
+// // Replacement for gsl::span in device code
 namespace Allen {
   namespace device {
+#if defined(TARGET_DEVICE_HIP)
     template<class T>
     struct span {
       T* __ptr = nullptr;
@@ -123,6 +125,9 @@ namespace Allen {
       constexpr __device__ __host__ T& operator[](int i) { return __ptr[i]; }
       constexpr __device__ __host__ const T& operator[](int i) const { return __ptr[i]; }
     };
+#else
+    using gsl::span;
+#endif
   } // namespace device
 } // namespace Allen
 

backend/include/CPUBackend.h

@@ -40,6 +40,8 @@ using std::signbit;
 #define fminf_impl fminf
 constexpr int warp_size = 1;
 
+unsigned inline __ballot_sync(unsigned mask, int predicate) { return predicate & mask; }
+
 // Support for dynamic shared memory buffers
 #define DYNAMIC_SHARED_MEMORY_BUFFER(_type, _instance, _config)                                                  \
   auto _dynamic_shared_memory_buffer = std::vector<_type>(_config.dynamic_shared_memory_size() / sizeof(_type)); \

backend/include/HIPBackend.h

@@ -51,6 +51,7 @@
 constexpr int warp_size = 64;
 #define __shfl_down_sync(_, val, offset) __shfl_down(val, offset)
 #define __shfl_sync(_, val, offset) __shfl(val, offset)
+#define __ballot_sync(_, predicate) __ballot(predicate)
 
 // Support for dynamic shared memory buffers
 #define DYNAMIC_SHARED_MEMORY_BUFFER(_type, _instance, _config) HIP_DYNAMIC_SHARED(_type, _instance)