diff --git a/CMakeLists.txt b/CMakeLists.txt
index cdfd53561754fa50b79f18892d6c4554d572dfd0..4a8a6a2dda9e37f2c6e35edb82beb4119aa57e3c 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -29,16 +29,64 @@ if(TARGET_DEVICE STREQUAL "CUDA" AND TENSOR)
   add_compile_definitions(TENSOR_CORES_ENABLED)
 endif()
 
-# Workaround 1/2 for shared cuda runtime
-set(CMAKE_CUDA_FLAGS "" CACHE STRING "")
-if(CMAKE_CUDA_FLAGS)
-	list(REMOVE_ITEM CMAKE_CUDA_FLAGS "-cudart static")
-endif()
-string(APPEND CMAKE_CUDA_FLAGS "-cudart shared")
+list(APPEND CMAKE_PREFIX_PATH ${CMAKE_INSTALL_PREFIX}) # for find_package
+list(APPEND CMAKE_MODULE_PATH ${CMAKE_SOURCE_DIR}/cmake) # for find_package
+
+option(BUILD_TESTS "Build test programs" OFF)
+
+# Setup the project to build as a Gaudi project or else standalone
+macro(allen_project)
+  # Gaudi build
+  if (DEFINED ENV{LHCBRELEASES})
+	find_package(GaudiProject REQUIRED)
+    # Fake extra package to get CUDA library dir included in the
+    # environment
+    find_package(CUDAPaths REQUIRED)
+    # Declare project name and version
+    gaudi_project(Allen v0r7
+      USE Rec HEAD Online HEAD)
+
+    # Use the lowest-common-denominator cpu arch
+    set(CPU_ARCH core2)
+
+    # Always enable tests for gaudi builds
+    set(BUILD_TESTS ON)
+
+    # Find the CUDA compiler if it's not set
+    # FIXME: this assumes a standard cuda install: re-evaluate if
+    # LCG_XXcudaYY is used
+    if (TARGET_DEVICE STREQUAL "CUDA")
+      if (NOT CMAKE_CUDA_COMPILER)
+        find_program(CMAKE_CUDA_COMPILER nvcc
+                     HINTS /usr/local/cuda/bin)
+        if (CMAKE_CUDA_COMPILER)
+          message(STATUS "Found cuda compiler ${CMAKE_CUDA_COMPILER}")
+        endif()
+      else()
+        message(STATUS "Using cuda compiler ${CMAKE_CUDA_COMPILER}")
+      endif()
+      # Make sure the lcg compiler wrapper scripts are used to call
+      # the host compiler
+      set(CMAKE_CUDA_HOST_COMPILER ${CMAKE_CXX_COMPILER})
+
+      # Build for maximum compatibility as the build host is different
+      # from the the machine running it and may not even have a GPU
+      foreach(arch 30 52 61 70 75)
+        set(OVERRIDE_ARCH_FLAG "${OVERRIDE_ARCH_FLAG} -gencode arch=compute_${arch},code=sm_${arch}")
+      endforeach()
+    endif()
+
+  else()
+    project(Allen C CXX)
+    if (BUILD_TESTS)
+      enable_testing()
+    endif()
+  endif()
+endmacro()
 
 # Setup project for the configured target
 if(TARGET_DEVICE STREQUAL "CPU")
-  project(Allen C CXX)
+  allen_project()
   add_compile_definitions(CPU)
 
   # This seems to be needed across compilers
@@ -46,10 +94,12 @@ if(TARGET_DEVICE STREQUAL "CPU")
 
   function(allen_add_library)
     add_library(${ARGV})
+    install(TARGETS ${ARGV0} DESTINATION lib OPTIONAL)
   endfunction()
 
   function(allen_add_executable)
     add_executable(${ARGV})
+    install(TARGETS ${ARGV0} RUNTIME DESTINATION bin OPTIONAL)
   endfunction()
 elseif(TARGET_DEVICE STREQUAL "HIP")
   add_compile_definitions(HIP)
@@ -71,7 +121,7 @@ elseif(TARGET_DEVICE STREQUAL "HIP")
   include_directories(${HIP_PATH}/include)
   include_directories(${HIP_PATH}/../hsa/include)
 
-  project(Allen C CXX)
+  allen_project()
 
   find_package(HIP QUIET REQUIRED)
   if(HIP_FOUND)
@@ -84,30 +134,45 @@ elseif(TARGET_DEVICE STREQUAL "HIP")
 
   set(CMAKE_CXX_FLAGS_ADDITIONAL "-Wno-unused-result")
 else()
-  project(Allen C CXX CUDA)
+
+  # Workaround 1/2 for shared cuda runtime
+  set(CMAKE_CUDA_FLAGS "" CACHE STRING "")
+  if(CMAKE_CUDA_FLAGS)
+	list(REMOVE_ITEM CMAKE_CUDA_FLAGS "-cudart static")
+  endif()
+  string(APPEND CMAKE_CUDA_FLAGS "-cudart shared")
+
+  allen_project()
+  enable_language(CUDA)
+
+  # Workaround 2/2 for shared cuda runtime
+  if(CMAKE_CUDA_HOST_IMPLICIT_LINK_LIBRARIES)
+	list(REMOVE_ITEM CMAKE_CUDA_HOST_IMPLICIT_LINK_LIBRARIES "cudart_static")
+  endif()
+  if(CMAKE_CUDA_IMPLICIT_LINK_LIBRARIES)
+	list(REMOVE_ITEM CMAKE_CUDA_IMPLICIT_LINK_LIBRARIES "cudart_static" )
+  endif()
+
+  # Determine whether to use Tensor Cores
+  # Note: Tensor cores are only supported in CUDA
+  option(TENSOR "Produce CUDA code with Tensor Cores enabled if available" OFF)
+  if(TARGET_DEVICE STREQUAL "CUDA" AND TENSOR)
+	add_compile_definitions(TENSOR_CORES_ENABLED)
+  endif()
 
   function(allen_add_library)
     add_library(${ARGV})
     set_property(TARGET ${ARGV0} PROPERTY
       CUDA_SEPARABLE_COMPILATION ON)
+    install(TARGETS ${ARGV0} DESTINATION lib OPTIONAL)
   endfunction()
 
   function(allen_add_executable)
     add_executable(${ARGV})
+    install(TARGETS ${ARGV0} RUNTIME DESTINATION bin OPTIONAL)
   endfunction()
 endif()
 
-# Workaround 2/2 for shared cuda runtime
-if(CMAKE_CUDA_HOST_IMPLICIT_LINK_LIBRARIES)
-  list(REMOVE_ITEM CMAKE_CUDA_HOST_IMPLICIT_LINK_LIBRARIES "cudart_static")
-endif()
-if(CMAKE_CUDA_IMPLICIT_LINK_LIBRARIES)
-  list(REMOVE_ITEM CMAKE_CUDA_IMPLICIT_LINK_LIBRARIES "cudart_static" )
-endif()
-
-list(APPEND CMAKE_PREFIX_PATH ${CMAKE_INSTALL_PREFIX}) # for find_package
-list(APPEND CMAKE_MODULE_PATH ${CMAKE_SOURCE_DIR}/cmake) # for find_package
-
 # Deal with build type
 if(NOT CMAKE_BUILD_TYPE)
   set(CMAKE_BUILD_TYPE RelWithDebInfo)
@@ -170,10 +235,13 @@ if(TARGET_DEVICE STREQUAL "CUDA")
     set(CUDA_ARCH "MAX" CACHE STRING "Cuda architecture")
   endif()
 
-  if (CUDA_ARCH STREQUAL "MIN" OR CUDA_ARCH STREQUAL "MAX" OR CUDA_ARCH STREQUAL "COMP")
-    set(OUTPUTFILE ${PROJECT_BINARY_DIR}/cuda_arch) # No suffix required
+  if(DEFINED OVERRIDE_ARCH_FLAG)
+    message(STATUS "Cuda architecture set to ${OVERRIDE_ARCH_FLAG}")
+    set(ARCH_FLAG ${OVERRIDE_ARCH_FLAG})
+  elseif (CUDA_ARCH STREQUAL "MIN" OR CUDA_ARCH STREQUAL "MAX" OR CUDA_ARCH STREQUAL "COMP")
+    set(OUTPUTFILE ${CMAKE_BINARY_DIR}/cuda_arch) # No suffix required
     set(CUDAFILE ${CMAKE_CURRENT_SOURCE_DIR}/cmake/utils/cuda_arch.cu)
-    execute_process(COMMAND ${CMAKE_CUDA_COMPILER} -lcuda ${CUDAFILE} -o ${OUTPUTFILE})
+    execute_process(COMMAND ${CMAKE_CUDA_COMPILER} -ccbin=${CMAKE_CXX_COMPILER} -std=c++14 -lcuda ${CUDAFILE} -o ${OUTPUTFILE})
     if(CUDA_ARCH STREQUAL "MAX")
       set(CHECK_ARGS "-l")
     endif()
@@ -197,9 +265,6 @@ if(TARGET_DEVICE STREQUAL "CUDA")
     message(STATUS "Cuda architecture manually set to ${CUDA_ARCH}")
     set(ARCH_FLAG "-arch=${CUDA_ARCH}")
   endif()
-  if(DEFINED OVERRIDE_ARCH_FLAG)
-    set(ARCH_FLAG ${OVERRIDE_ARCH_FLAG})
-  endif()
 
   # Cuda: Deal with build type
   if(${CMAKE_BUILD_TYPE} STREQUAL RelWithDebInfo)
@@ -219,7 +284,6 @@ if(TARGET_DEVICE STREQUAL "CUDA")
     set(CUDA_VERBOSE_FLAGS "")
   endif()
 
-  set(CMAKE_CUDA_STANDARD 14)
   set(CMAKE_CUDA_FLAGS " ${ARCH_FLAG} --ftemplate-depth=300 --use_fast_math --expt-relaxed-constexpr ${CUDA_VERBOSE_FLAGS}")
 elseif(TARGET_DEVICE STREQUAL "HIP")
   # HCC Options (HIPCC compiler)
@@ -250,6 +314,7 @@ elseif(TARGET_DEVICE STREQUAL "HIP")
     hip_add_library(${ARGV} HIPCC_OPTIONS ${HIPCC_OPTIONS})
     set_property(TARGET ${ARGV0} PROPERTY
       HIP_SEPARABLE_COMPILATION ON)
+    install(TARGETS ${ARGV0} DESTINATION lib OPTIONAL)
   endfunction()
 
   function(allen_add_executable)
@@ -257,14 +322,10 @@ elseif(TARGET_DEVICE STREQUAL "HIP")
     set_target_properties(${ARGV0}
          PROPERTIES HIP_SEPARABLE_COMPILATION ON
                   HIP_RESOLVE_DEVICE_SYMBOLS ON)
+    install(TARGETS ${ARGV0} RUNTIME DESTINATION bin OPTIONAL)
   endfunction()
 endif()
 
-option(BUILD_TESTS "Build test programs" OFF)
-if (BUILD_TESTS)
-  enable_testing()
-endif()
-
 find_package(ZLIB REQUIRED)
 
 option(USE_LZMA "build with lzma" OFF)
@@ -284,8 +345,6 @@ endif(USE_LZ4)
 
 option(USE_KALMAN_SINGLE_PRECISION "Use single precision Kalman filter" ON)
 
-# find_package(CUDA REQUIRED)
-#set(CUDA_HOST_COMPILER "g++")
 set(CUDA_SEPARABLE_COMPILATION ON)
 option(CUDA_PROPAGATE_HOST_FLAGS "Propagate CUDA host flags" OFF)
 
@@ -297,7 +356,7 @@ if ((EXISTS $ENV{ROOTSYS}) AND (USE_ROOT))
      set(ALLEN_ROOT_CMAKE $ENV{ROOTSYS}/etc)
    endif()
    find_package(ROOT QUIET HINTS ${ALLEN_ROOT_CMAKE} NO_DEFAULT_PATH COMPONENTS Core Hist Tree)
-   if ( ROOT_FOUND )
+   if (ROOT_FOUND)
       message(STATUS "Compiling with ROOT: " ${ROOT_INCLUDE_DIRS})
 
       # If ROOT is built with C++17 support, everything that includes ROOT
@@ -436,16 +495,13 @@ endif()
 
 target_link_libraries(AllenLib PUBLIC
   Stream
+  Common
   TrackChecking
   PVChecking
   CheckClustering
   SelChecking
-  x86MuonDecoding
-  Common
-  CudaCommon
   NonEventData
-  AllenZMQ
-  x86Clustering)
+  AllenZMQ)
 
 if (ROOT_FOUND)
   target_compile_definitions(AllenLib PUBLIC ${ALLEN_ROOT_DEFINITIONS})
@@ -456,8 +512,6 @@ endif()
 allen_add_executable(Allen ${main_cpp_path})
 target_link_libraries(Allen PRIVATE AllenLib)
 
-if(TARGET_DEVICE STREQUAL "CUDA")
-  if(BUILD_TESTS)
-    add_subdirectory(tests)
-  endif(BUILD_TESTS)
+if(TARGET_DEVICE STREQUAL "CUDA" AND BUILD_TESTS)
+  add_subdirectory(tests)
 endif()
diff --git a/Online/AllenIntegration/CMakeLists.txt b/Online/AllenIntegration/CMakeLists.txt
new file mode 100644
index 0000000000000000000000000000000000000000..0156a118c926cc7121eb9b89a5430085179f575d
--- /dev/null
+++ b/Online/AllenIntegration/CMakeLists.txt
@@ -0,0 +1,20 @@
+################################################################################
+# Package: AllenIntegration
+################################################################################
+gaudi_subdir(AllenIntegration v1r0)
+
+gaudi_depends_on_subdirs(Online/OnlineKernel)
+
+find_package(Boost REQUIRED)
+find_package(Rangev3 REQUIRED)
+include_directories(SYSTEM ${Boost_INCLUDE_DIRS} ${Rangev3_INCLUDE_DIRS})
+
+include_directories(${CMAKE_SOURCE_DIR}/main/include)
+
+# gaudi_install_headers()
+
+gaudi_add_library(AllenIntegration
+                  src/*.cpp
+                  PUBLIC_HEADERS Integration
+                  INCLUDE_DIRS Online/OnlineKernel
+                  LINK_LIBRARIES OnlineKernel)
diff --git a/Online/AllenIntegration/Integration/placeholder.h b/Online/AllenIntegration/Integration/placeholder.h
new file mode 100644
index 0000000000000000000000000000000000000000..40f403a62ca7e1ca20a2536253ad3da5452e9e05
--- /dev/null
+++ b/Online/AllenIntegration/Integration/placeholder.h
@@ -0,0 +1,3 @@
+#pragma once
+
+void test_integration();
diff --git a/Online/AllenIntegration/src/placeholder.cpp b/Online/AllenIntegration/src/placeholder.cpp
new file mode 100644
index 0000000000000000000000000000000000000000..100455e4c2578775f26e1ff6b896baec7b3f7ae8
--- /dev/null
+++ b/Online/AllenIntegration/src/placeholder.cpp
@@ -0,0 +1,5 @@
+#include <Integration/placeholder.h>
+
+void test_integration() {
+
+}
diff --git a/bindings/.gaudi_project_ignore b/bindings/.gaudi_project_ignore
new file mode 100644
index 0000000000000000000000000000000000000000..e69de29bb2d1d6434b8b29ae775ad8c2e48c5391
diff --git a/checker/.gaudi_project_ignore b/checker/.gaudi_project_ignore
new file mode 100644
index 0000000000000000000000000000000000000000..e69de29bb2d1d6434b8b29ae775ad8c2e48c5391
diff --git a/checker/pv/CMakeLists.txt b/checker/pv/CMakeLists.txt
index aa408aa8f2c025fe99060bd0957feaf321c1ff17..8a97ae31b6aa2dfda7fcbf029fd29b43ceac08c0 100644
--- a/checker/pv/CMakeLists.txt
+++ b/checker/pv/CMakeLists.txt
@@ -17,7 +17,7 @@ add_library(PVChecking STATIC
   ${pv_checker_sources}
 )
 
-if(ROOT_FOUND)
+if(USE_ROOT AND ROOT_FOUND)
   target_compile_definitions(PVChecking PRIVATE WITH_ROOT)
   target_include_directories(PVChecking BEFORE PRIVATE
     ${ROOT_INCLUDE_DIRS}
diff --git a/checker/tracking/CMakeLists.txt b/checker/tracking/CMakeLists.txt
index 7ff1dcecfc3d0fa08e13846443e0482cc73dfbf7..c20bb54ea33503d325fa2adf0f2e992947b034d1 100644
--- a/checker/tracking/CMakeLists.txt
+++ b/checker/tracking/CMakeLists.txt
@@ -12,7 +12,7 @@ include_directories(${CMAKE_SOURCE_DIR}/cuda/SciFi/common/include)
 include_directories(${CMAKE_SOURCE_DIR}/cuda/kalman/ParKalman/include)
 include_directories(${CMAKE_SOURCE_DIR}/cuda/PV/common/include)
 include_directories(${CMAKE_SOURCE_DIR}/cuda/PV/beamlinePV/include)
-include_directories(${ROOT_INCLUDE_DIRS})
+include_directories(SYSTEM ${ROOT_INCLUDE_DIRS})
 
 file(GLOB tracking_checker_sources "src/*cpp")
 
diff --git a/configuration/.gaudi_project_ignore b/configuration/.gaudi_project_ignore
new file mode 100644
index 0000000000000000000000000000000000000000..e69de29bb2d1d6434b8b29ae775ad8c2e48c5391
diff --git a/cuda/.gaudi_project_ignore b/cuda/.gaudi_project_ignore
new file mode 100644
index 0000000000000000000000000000000000000000..e69de29bb2d1d6434b8b29ae775ad8c2e48c5391
diff --git a/cuda/PV/beamlinePV/CMakeLists.txt b/cuda/PV/beamlinePV/CMakeLists.txt
index 0169420c3a6df4b3d3c7066faebca637c333f2bc..1ca7a51f486980547efda51ef0a4ebfbe003d8f3 100644
--- a/cuda/PV/beamlinePV/CMakeLists.txt
+++ b/cuda/PV/beamlinePV/CMakeLists.txt
@@ -20,10 +20,10 @@ endif()
 allen_add_library(PV_beamline STATIC
   ${blPV})
 
-if(ROOT_FOUND)
+if(USE_ROOT AND ROOT_FOUND)
   add_library(PV_monitoring STATIC src/pv_beamline_monitoring.cpp)
   target_compile_definitions(PV_monitoring PRIVATE ${ALLEN_ROOT_DEFINITIONS})
-  target_include_directories(PV_monitoring BEFORE PRIVATE
+  target_include_directories(PV_monitoring SYSTEM BEFORE PRIVATE
     ${ROOT_INCLUDE_DIRS})
   target_link_libraries(PV_beamline PRIVATE PV_monitoring)
 endif()
diff --git a/cuda/PV/patPV/CMakeLists.txt b/cuda/PV/patPV/CMakeLists.txt
index 5419cc5e010bd251bf060217857f2032fde765f7..57d9c70411079bab8290101473bc221c81914747 100644
--- a/cuda/PV/patPV/CMakeLists.txt
+++ b/cuda/PV/patPV/CMakeLists.txt
@@ -18,6 +18,6 @@ if(TARGET_DEVICE STREQUAL "CPU")
   set_source_files_properties(${pat_pv} PROPERTIES LANGUAGE CXX)
 endif()
 
-allen_add_library(PatPV STATIC
+allen_add_library(AllenPatPV STATIC
   ${pat_pv}
 )
diff --git a/cuda/UT/CMakeLists.txt b/cuda/UT/CMakeLists.txt
index 3ab7cfd26cc9f9fa9e20736bcd48afabddd3d849..efa8770c3a18ecbc747c6eb1d54e7cc70b55c514 100644
--- a/cuda/UT/CMakeLists.txt
+++ b/cuda/UT/CMakeLists.txt
@@ -23,9 +23,14 @@ if(TARGET_DEVICE STREQUAL "CPU")
   endforeach(source_file)
 endif()
 
-allen_add_library(UT STATIC
+allen_add_library(UTCommon STATIC
   ${UT_common}
+)
+
+allen_add_library(UT STATIC
   ${UT_decoding}
   ${UT_consolidate}
   ${CompassUT_tracking}
 )
+
+target_link_libraries(UT UTCommon)
diff --git a/cuda/muon/CMakeLists.txt b/cuda/muon/CMakeLists.txt
index 204547b97fdbb349518fff86c6c6b2980d17bd1a..956bcbbd9c51a931c04a5932d9ad968ddb3566c9 100644
--- a/cuda/muon/CMakeLists.txt
+++ b/cuda/muon/CMakeLists.txt
@@ -36,3 +36,5 @@ allen_add_library(Muon STATIC
   ${muon_decoding}
   ${muon_decoding_steps}
 )
+
+target_link_libraries(Muon Utils)
diff --git a/cuda/muon/common/include/MuonBase.cuh b/cuda/muon/common/include/MuonBase.cuh
index 0c5e73384c35fbfced1280c117c6da96e64c0e74..ed5a3c56add2322eb2a458e55f30b9783a985938 100644
--- a/cuda/muon/common/include/MuonBase.cuh
+++ b/cuda/muon/common/include/MuonBase.cuh
@@ -50,5 +50,5 @@ namespace Muon {
     static constexpr int RIGHT = 1;
     static constexpr int LEFT = -1;
 
-  }; // namespace MuonBase
-};   // namespace Muon
+  } // namespace MuonBase
+}   // namespace Muon
diff --git a/cuda/utils/CMakeLists.txt b/cuda/utils/CMakeLists.txt
index c8963eca5642d56306febaafcff113d4c9f34af8..f8aebf2427d9795df05ee6c17f43f8c95f03cd8b 100644
--- a/cuda/utils/CMakeLists.txt
+++ b/cuda/utils/CMakeLists.txt
@@ -29,8 +29,3 @@ allen_add_library(Utils STATIC
   ${prefix_sum}
   ${binary_search}
 )
-
-target_link_libraries(Utils PRIVATE
-  UT
-  SciFi
-)
diff --git a/cuda/velo/CMakeLists.txt b/cuda/velo/CMakeLists.txt
index 15ef5dddcb3246adb6651c1bf864f03192fa30c9..53139dc26cda21ea24021a2bbdecf27382e0db4d 100644
--- a/cuda/velo/CMakeLists.txt
+++ b/cuda/velo/CMakeLists.txt
@@ -31,8 +31,11 @@ if(TARGET_DEVICE STREQUAL "CPU")
   endforeach(source_file)
 endif()
 
-allen_add_library(Velo STATIC
+allen_add_library(VeloCommon STATIC
   ${velo_common}
+)
+
+allen_add_library(Velo STATIC
   ${velo_clustering_sources}
   ${velo_prefix_sum}
   ${velo_phi_and_sort}
@@ -40,3 +43,5 @@ allen_add_library(Velo STATIC
   ${velo_simplified_kalman_filter}
   ${velo_consolidate_tracks}
 )
+
+target_link_libraries(Velo VeloCommon)
diff --git a/integration/.gaudi_project_ignore b/integration/.gaudi_project_ignore
new file mode 100644
index 0000000000000000000000000000000000000000..e69de29bb2d1d6434b8b29ae775ad8c2e48c5391
diff --git a/integration/non_event_data/CMakeLists.txt b/integration/non_event_data/CMakeLists.txt
index bc2ce24b4ae253a636eb2132acc7304dca07b7a3..1f7bc109133e695bf4fa0651cd2b0cdc20a548f1 100644
--- a/integration/non_event_data/CMakeLists.txt
+++ b/integration/non_event_data/CMakeLists.txt
@@ -1,8 +1,12 @@
 file(GLOB SOURCES "src/*.cpp")
 
-add_library(NonEventData ${SOURCES})
+allen_add_library(NonEventData ${SOURCES})
 
 target_link_libraries(NonEventData
+  x86Clustering
+  VeloCommon
+  UTCommon
+  CudaCommon
   Common)
 
 target_include_directories(NonEventData PUBLIC ${CMAKE_CURRENT_SOURCE_DIR}/include)
diff --git a/main/.gaudi_project_ignore b/main/.gaudi_project_ignore
new file mode 100644
index 0000000000000000000000000000000000000000..e69de29bb2d1d6434b8b29ae775ad8c2e48c5391
diff --git a/main/include/BinaryProvider.h b/main/include/BinaryProvider.h
index 173c6ee3cbc11b51e743f5520e2da350addc3f2f..38ede5d3b57cca565fd73da70c3d93c2d820dd27 100644
--- a/main/include/BinaryProvider.h
+++ b/main/include/BinaryProvider.h
@@ -162,7 +162,7 @@ public:
    *
    * @return     event IDs in slice
    */
-  std::vector<std::tuple<unsigned int, unsigned long>> const& event_ids(size_t slice_index) const
+  std::vector<std::tuple<unsigned int, unsigned long>> const& event_ids(size_t slice_index) const override
   {
     return m_event_ids[slice_index];
   }
@@ -232,7 +232,7 @@ public:
    *
    * @return     Banks and their offsets
    */
-  BanksAndOffsets banks(BankTypes bank_type, size_t slice_index) const
+  BanksAndOffsets banks(BankTypes bank_type, size_t slice_index) const override
   {
     auto ib = to_integral<BankTypes>(bank_type);
     auto const& [banks, offsets, offsets_size] = m_slices[ib][slice_index];
@@ -251,9 +251,8 @@ private:
    */
   void prefetch() {
 
-    bool eof = false, prefetch_done = false;
-    size_t bytes_read = 0, n_reps = 0;
-
+    bool prefetch_done = false;
+    size_t n_reps = 0;
     size_t eps = this->events_per_slice();
 
     size_t n_files = std::get<1>(m_files.front()).size();
@@ -281,7 +280,6 @@ private:
         *it = false;
       }
       size_t slice_index = distance(m_slice_free.begin(), it);
-      auto& slice = m_slices[slice_index];
       this->debug_output("Got slice index " + std::to_string(slice_index), 1);
 
       // "Reset" the slice
diff --git a/main/include/Common.h b/main/include/Common.h
index eb6d4907c74c0a126d9f33e2b5a3cfce5b437931..7aba99cf2e042bce1f22b23eb4d6de840f660b30 100644
--- a/main/include/Common.h
+++ b/main/include/Common.h
@@ -16,7 +16,7 @@ struct StrException : public std::exception {
   std::string s;
   StrException(std::string ss) : s(ss) {}
   ~StrException() throw() {} // Updated
-  const char* what() const throw() { return s.c_str(); }
+  const char* what() const throw() override { return s.c_str(); }
 };
 
 using EventID = std::tuple<unsigned int, unsigned long>;
diff --git a/main/include/InputProvider.h b/main/include/InputProvider.h
index 709313817debf047787258e57af8d4d23b1919eb..64f7166544a999e8adc9453fc6e00ccaf469a1cf 100644
--- a/main/include/InputProvider.h
+++ b/main/include/InputProvider.h
@@ -44,6 +44,10 @@ struct IInputProvider {
    * @return     spans spanning bank and offset memory
    */
   virtual BanksAndOffsets banks(BankTypes bank_type, size_t slice_index) const = 0;
+
+  /// Descturctor
+  virtual ~IInputProvider() {};
+
 };
 
 // InputProvider
@@ -57,6 +61,7 @@ public:
     m_nslices {n_slices}, m_events_per_slice {events_per_slice}, m_nevents {n_events}, m_types {banks_set<Banks...>()}
   {}
 
+  /// Descturctor
   virtual ~InputProvider() {};
 
   /**
diff --git a/main/include/InputReader.h b/main/include/InputReader.h
index 6a08410351da0079cc612b30288c82bd37a4b90b..99920964c2be82ec61de1e8f959c492c8a3759ac 100644
--- a/main/include/InputReader.h
+++ b/main/include/InputReader.h
@@ -34,6 +34,8 @@ using FolderMap = std::map<BankTypes, std::string>;
 struct EventReader : public Reader {
   EventReader(FolderMap folders) : Reader(begin(folders)->second), m_folders {std::move(folders)} {}
 
+  virtual ~EventReader() = default;
+
   gsl::span<char> events(BankTypes type)
   {
     auto it = m_events.find(type);
diff --git a/main/include/Logger.h b/main/include/Logger.h
index 0a9f84b01e52f74f599c228dd00a203cfca1879a..1f0d1c415ab297d2d3cdd27696fd90a4ee340773 100644
--- a/main/include/Logger.h
+++ b/main/include/Logger.h
@@ -38,7 +38,7 @@ public:
 
   ~MessageLogger() { _file_std_io->rdbuf(_old); }
 
-  int overflow(int c)
+  int overflow(int c) override
   {
     if (c == '\n') {
       std::cout << _buf << std::endl;
@@ -61,7 +61,7 @@ public:
 
   ~VoidLogger() { _void_stream->rdbuf(_old); }
 
-  int overflow(int c)
+  int overflow(int c) override
   {
     // Just don't do anything
     return c;
diff --git a/main/include/MDFProvider.h b/main/include/MDFProvider.h
index d3c30efb70e2ce03128b2105bd08e56940cc0f65..8e2669f6afe9c4af79b959b3547124f8c2ad75a4 100644
--- a/main/include/MDFProvider.h
+++ b/main/include/MDFProvider.h
@@ -22,13 +22,13 @@
 #include <read_mdf.hpp>
 #include <raw_bank.hpp>
 
+#include "Transpose.h"
+
 #ifndef NO_CUDA
 #include <CudaCommon.h>
 #endif
 
 namespace {
-  constexpr auto header_size = sizeof(LHCb::MDFHeader);
-
   using namespace Allen::Units;
 } // namespace
 
@@ -67,283 +67,6 @@ struct MDFProviderConfig {
   size_t n_loops = 0;
 };
 
-//
-/**
- * @brief      read events from input file into prefetch buffer
- *
- * @details    NOTE: It is assumed that the header has already been
- *             read, calling read_events will read the subsequent
- *             banks and then header of the next event.
- *
- * @param      input stream
- * @param      prefetch buffer to read into
- * @param      storage for the MDF header
- * @param      buffer for temporary storage of the compressed banks
- * @param      number of events to read
- * @param      check the MDF checksum if it is available
- *
- * @return     (eof, error, full, n_bytes)
- */
-std::tuple<bool, bool, bool, size_t> read_events(int input, ReadBuffer& read_buffer,
-                                                 LHCb::MDFHeader& header,
-                                                 std::vector<char> compress_buffer,
-                                                 size_t n_events, bool check_checksum)
-{
-  auto& [n_filled, event_offsets, buffer] = read_buffer;
-
-  // Keep track of where to write and the end of the prefetch buffer
-  auto* write = &buffer[0 + event_offsets[n_filled]];
-  auto* buffer_start = &buffer[0];
-  auto const* buffer_end = buffer.data() + buffer.size();
-  size_t n_bytes = 0;
-  bool eof = false, error = false, full = false;
-  gsl::span<char> bank_span;
-
-  // Loop until the requested number of events is prefetched, the
-  // maximum number of events per prefetch buffer is hit, an error
-  // occurs or eof is reached
-  while (!eof && !error && n_filled < event_offsets.size() - 1 && n_filled < n_events) {
-    // It is
-
-    // Read the banks
-    gsl::span<char> buffer_span{buffer_start + event_offsets[n_filled], buffer.size() - event_offsets[n_filled]};
-    std::tie(eof, error, bank_span) = MDF::read_banks(input, header, std::move(buffer_span),
-                                                      compress_buffer, check_checksum);
-    // Fill the start offset of the next event
-    event_offsets[++n_filled] = bank_span.end() - buffer_start;
-    n_bytes += bank_span.size();
-
-    // read the next header
-    ssize_t n_bytes = ::read(input, reinterpret_cast<char*>(&header), header_size);
-    if (n_bytes != 0) {
-      // Check if there is enough space to read this event
-      int compress = header.compression() & 0xF;
-      int expand = (header.compression() >> 4) + 1;
-      int event_size = (header.recordSize() + header_size +
-                        2 * (sizeof(LHCb::RawBank) + sizeof(int)) +
-                        (compress ? expand * (header.recordSize() - header_size) : 0));
-      if (event_offsets[n_filled] + event_size > buffer.size()) {
-        full = true;
-        break;
-      }
-    } else if (n_bytes == 0) {
-      info_cout << "Cannot read more data (Header). End-of-File reached.\n";
-      eof = true;
-    } else {
-      error_cout << "Failed to read header " << strerror(errno) << "\n";
-      error = true;
-    }
-  }
-  return {eof, error, full, n_bytes};
-}
-
-/**
- * @brief      Fill the array the contains the number of banks per type
- *
- * @details    detailed description
- *
- * @param      prefetched buffer of events (a single event is needed)
- *
- * @return     (success, number of banks per bank type; 0 if the bank is not needed)
- */
-std::tuple<bool, std::array<unsigned int, NBankTypes>> fill_counts(ReadBuffer const& read_buffer)
-{
-  auto& [n_filled, event_offsets, buffer] = read_buffer;
-
-  std::array<unsigned int, NBankTypes> count{0};
-
-  // Care only about the first event
-  size_t i_event = 0;
-
-  // Offsets are to the start of the event, which includes the header
-  auto const* bank = buffer.data() + event_offsets[i_event];
-  auto const* bank_end = buffer.data() + event_offsets[i_event + 1];
-
-  // Loop over all the bank data
-  while (bank < bank_end) {
-    const auto* b = reinterpret_cast<const LHCb::RawBank*>(bank);
-
-    if (b->magic() != LHCb::RawBank::MagicPattern) {
-      error_cout << "Magic pattern failed: " << std::hex << b->magic() << std::dec << "\n";
-      return {false, count};
-    }
-
-    // Check if Allen processes this bank type, count bank types that
-    // are wanted
-    auto bank_type_it = Allen::bank_types.find(b->type());
-    if (bank_type_it != Allen::bank_types.end()) {
-      auto bank_type_index = to_integral(bank_type_it->second);
-      ++count[bank_type_index];
-    }
-
-    // Increment overall bank pointer
-    bank += b->totalSize();
-  }
-
-  return {true, count};
-}
-
-/**
- * @brief      Transpose events to Allen layout
- *
- * @param      ReadBuffer containing events to be transposed
- * @param      slices to fill with transposed banks, slices are addressed by bank type
- * @param      index of bank slices
- * @param      event ids of banks in this slice
- * @param      number of banks per event
- * @param      number of events to transpose
- *
- * @return     (success, slice full for one of the bank types, number of events transposed)
- */
-template <BankTypes... Banks>
-std::tuple<bool, bool, size_t> transpose_events(const ReadBuffer& read_buffer,
-                                                Slices& slices, int const slice_index,
-                                                EventIDs& event_ids,
-                                                std::vector<int> const& bank_ids,
-                                                std::array<unsigned int, NBankTypes> const& banks_count,
-                                                size_t n_events) {
-  auto const& [n_filled, event_offsets, buffer] = read_buffer;
-
-  unsigned int* banks_offsets = nullptr;
-  // Number of offsets
-  size_t* n_banks_offsets = nullptr;
-
-  // Where to write the transposed bank data
-  uint32_t* banks_write = nullptr;
-
-  // Where should offsets to individual banks be written
-  uint32_t* banks_offsets_write = nullptr;
-
-  unsigned int bank_offset = 0;
-  unsigned int bank_counter = 1;
-  unsigned int bank_type_index = 0;
-
-  bool full = false;
-
-  // "Reset" the slice
-  for (auto bank_type : {Banks...}) {
-    auto ib = to_integral<BankTypes>(bank_type);
-    std::get<1>(slices[ib][slice_index])[0] = 0;
-    std::get<2>(slices[ib][slice_index]) = 1;
-  }
-  event_ids.clear();
-
-  // L0Calo doesn't exist in the upgrade
-  LHCb::RawBank::BankType prev_type = LHCb::RawBank::L0Calo;
-
-  // Loop over events in the prefetch buffer
-  size_t i_event = 0;
-  for (; i_event < n_filled && i_event < n_events; ++i_event) {
-    // Offsets are to the start of the event, which includes the header
-    auto const* bank = buffer.data() + event_offsets[i_event];
-    auto const* bank_end = buffer.data() + event_offsets[i_event + 1];
-
-    // Loop over all bank data of this event
-    while (bank < bank_end) {
-      const auto* b = reinterpret_cast<const LHCb::RawBank*>(bank);
-
-      if (b->magic() != LHCb::RawBank::MagicPattern) {
-        error_cout << "Magic pattern failed: " << std::hex << b->magic() << std::dec << "\n";
-        return {false, false, i_event};
-        // Decode the odin bank
-      }
-
-      // Check what to do with this bank
-      auto bt = b->type();
-      if (bt == LHCb::RawBank::ODIN) {
-        // decode ODIN bank to obtain run and event numbers
-        auto odin = MDF::decode_odin(b);
-        event_ids.emplace_back(odin.run_number, odin.event_number);
-        bank += b->totalSize();
-        continue;
-      } else if (bt >= LHCb::RawBank::LastType || bank_ids[bt] == -1) {
-        // This bank is not required: skip it
-        bank += b->totalSize();
-        continue;
-      } else if (bt != prev_type) {
-        // Switch to new type of banks
-        bank_type_index = bank_ids[b->type()];
-        auto& slice = slices[bank_type_index][slice_index];
-        prev_type = bt;
-
-        bank_counter = 1;
-        banks_offsets = std::get<1>(slice).data();
-        n_banks_offsets = &std::get<2>(slice);
-
-        // Calculate the size taken by storing the number of banks
-        // and offsets to all banks within the event
-        auto preamble_words = 2 + banks_count[bank_type_index];
-
-        // Initialize offset to start of this set of banks from the
-        // previous one and increment with the preamble size
-        banks_offsets[*n_banks_offsets] = (banks_offsets[*n_banks_offsets - 1]
-                                           + preamble_words * sizeof(uint32_t));
-
-        // Three things to write for a new set of banks:
-        // - number of banks/offsets
-        // - offsets to individual banks
-        // - bank data
-
-        // Initialize point to write from offset of previous set
-        banks_write = reinterpret_cast<uint32_t*>(std::get<0>(slice).data() + banks_offsets[*n_banks_offsets - 1]);
-
-        // New offset to increment
-        ++(*n_banks_offsets);
-
-        // Write the number of banks
-        banks_write[0] = banks_count[bank_type_index];
-
-        // All bank offsets are uit32_t so cast to that type
-        banks_offsets_write = banks_write + 1;
-        banks_offsets_write[0] = 0;
-
-        // Offset in number of uint32_t
-        bank_offset = 0;
-
-        // Start writing bank data after the preamble
-        banks_write += preamble_words;
-      } else {
-        ++bank_counter;
-      }
-
-      // Write sourceID
-      banks_write[bank_offset] = b->sourceID();
-
-      // Write bank data
-      ::memcpy(banks_write + bank_offset + 1, b->data(), b->size());
-
-      auto n_word = b->size() / sizeof(uint32_t);
-      bank_offset += 1 + n_word;
-
-      // Write next offset in bytes
-      banks_offsets_write[bank_counter] = bank_offset * sizeof(uint32_t);
-
-      // Update "event" offset (in bytes)
-      banks_offsets[*n_banks_offsets - 1] += sizeof(uint32_t) * (1 + n_word);
-
-      // Increment overall bank pointer
-      bank += b->totalSize();
-    }
-
-    // Check if any of the per-bank-type slices potentially has too
-    // little space to fit the next event
-    for (auto bank_type : {Banks...}) {
-      auto ib = to_integral<BankTypes>(bank_type);
-      const auto& [slice, slice_offsets, offsets_size] = slices[ib][slice_index];
-      // Use the event size of the next event here instead of the
-      // per bank size because that's not yet known for the next
-      // event
-      auto const event_size = event_offsets[i_event + 1] - event_offsets[i_event];
-      if ((slice_offsets[offsets_size - 1] + event_size) > slice.size()) {
-        full = true;
-        goto transpose_end;
-      }
-    }
-  }
- transpose_end:
-  return {true, full, i_event};
-}
-
 /**
  * @brief      Provide transposed events from MDF files
  *
@@ -372,9 +95,10 @@ public:
   MDFProvider(size_t n_slices, size_t events_per_slice, std::optional<size_t> n_events,
               std::vector<std::string> connections, MDFProviderConfig config = MDFProviderConfig{}) :
     InputProvider<MDFProvider<Banks...>>{n_slices, events_per_slice, n_events},
-    m_event_ids{n_slices}, m_banks_count{0},
     m_buffer_writable(config.n_buffers, true),
     m_slice_free(n_slices, true),
+    m_banks_count{0},
+    m_event_ids{n_slices},
     m_connections {std::move(connections)},
     m_config{config}
   {
@@ -403,8 +127,6 @@ public:
 
       // Fudge with extra 20% memory
       size_t n_bytes = std::lround(it->second * events_per_slice * bank_size_fudge_factor * kB);
-      m_banks_data[ib].reserve(n_bytes / sizeof(uint32_t));
-      m_banks_offsets[ib].reserve(events_per_slice);
       auto& slices = m_slices[ib];
       slices.reserve(n_slices);
       for (size_t i = 0; i < n_slices; ++i) {
@@ -459,11 +181,13 @@ public:
       // Wait for first read buffer to be full
       m_prefetch_cond.wait(lock, [this] { return !m_prefetched.empty() || m_read_error; });
       if (!m_read_error) {
-        size_t i_read = m_prefetched.front();
-
         // Count number of banks per flavour
         bool count_success = false;
-        std::tie(count_success, m_banks_count) = fill_counts(m_buffers[i_read]);
+
+        // Offsets are to the start of the event, which includes the header
+        auto i_read = m_prefetched.front();
+        auto& [n_filled, event_offsets, buffer] = m_buffers[i_read];
+        std::tie(count_success, m_banks_count) = fill_counts({buffer.data(), event_offsets[1]});
         if (!count_success) {
           error_cout << "Failed to determine bank counts\n";
           m_read_error = true;
@@ -530,7 +254,7 @@ public:
    *
    * @return     EventIDs of events in given slice
    */
-  std::vector<std::tuple<unsigned int, unsigned long>> const& event_ids(size_t slice_index) const
+  std::vector<std::tuple<unsigned int, unsigned long>> const& event_ids(size_t slice_index) const override
   {
     return m_event_ids[slice_index];
   }
@@ -543,7 +267,7 @@ public:
    *
    * @return     Banks and their offsets
    */
-  BanksAndOffsets banks(BankTypes bank_type, size_t slice_index) const
+  BanksAndOffsets banks(BankTypes bank_type, size_t slice_index) const override
   {
     auto ib = to_integral<BankTypes>(bank_type);
     auto const& [banks, offsets, offsets_size] = m_slices[ib][slice_index];
@@ -682,12 +406,17 @@ private:
         }
       }
 
+      // Reset the slice
+      auto& event_ids = m_event_ids[*slice_index];
+      reset_slice<Banks...>(m_slices, *slice_index, event_ids);
+
       // Transpose the events in the read buffer into the slice
       std::tie(good, transpose_full, n_transposed) = transpose_events<Banks...>(m_buffers[i_read],
-                                                                                m_slices, *slice_index,
-                                                                                m_event_ids[*slice_index],
+                                                                                m_slices,
+                                                                                *slice_index,
                                                                                 m_bank_ids,
                                                                                 m_banks_count,
+                                                                                event_ids,
                                                                                 this->events_per_slice());
       this->debug_output("Transposed " + std::to_string(*slice_index) + " " + std::to_string(good)
                          + " " + std::to_string(transpose_full) + " " + std::to_string(n_transposed),
@@ -912,12 +641,6 @@ private:
   // Allen IDs of LHCb raw banks
   std::vector<int> m_bank_ids;
 
-  // Offsets to all the banks in the single event that was read from file
-  mutable std::array<std::vector<uint32_t>, NBankTypes> m_banks_offsets;
-
-  // Raw bank data by subdetector as read from file
-  mutable std::array<std::vector<uint32_t>, NBankTypes> m_banks_data;
-
   // Memory slices, N for each raw bank type
   Slices m_slices;
 
diff --git a/main/include/MDFReader.h b/main/include/MDFReader.h
deleted file mode 100644
index 6f54e2e56e5183186e029700465af5d8a2486c43..0000000000000000000000000000000000000000
--- a/main/include/MDFReader.h
+++ /dev/null
@@ -1,20 +0,0 @@
-#ifndef MDFREADER_H
-#define MDFREADER_H 1
-
-#include <map>
-#include <string>
-
-#include "InputReader.h"
-
-struct MDFReader : public EventReader {
-
-  MDFReader(FolderMap folders) : EventReader(std::move(folders)) {}
-
-  /**
-   * @brief Reads files from the specified folder, starting from an event offset.
-   */
-  std::vector<std::tuple<unsigned int, unsigned long>> read_events(
-    uint number_of_events_requested = 0,
-    uint start_event_offset = 0) override;
-};
-#endif
diff --git a/main/include/Transpose.h b/main/include/Transpose.h
new file mode 100644
index 0000000000000000000000000000000000000000..d5bc3b7c63239aad3b5160abfd677f536fa455a2
--- /dev/null
+++ b/main/include/Transpose.h
@@ -0,0 +1,358 @@
+#pragma once
+
+#include <thread>
+#include <vector>
+#include <array>
+#include <deque>
+#include <mutex>
+#include <atomic>
+#include <chrono>
+#include <algorithm>
+#include <numeric>
+#include <condition_variable>
+
+#include <unistd.h>
+#include <sys/types.h>
+#include <sys/stat.h>
+#include <fcntl.h>
+
+#include <Logger.h>
+#include <mdf_header.hpp>
+#include <read_mdf.hpp>
+#include <raw_bank.hpp>
+
+#ifndef NO_CUDA
+#include <CudaCommon.h>
+#endif
+
+namespace {
+  constexpr auto header_size = sizeof(LHCb::MDFHeader);
+
+  using namespace Allen::Units;
+} // namespace
+
+// Read buffer containing the number of events, offsets to the start
+// of the event and the event data
+using ReadBuffer = std::tuple<size_t, std::vector<unsigned int>, std::vector<char>>;
+using ReadBuffers = std::vector<ReadBuffer>;
+
+// A slice contains transposed bank data, offsets to the start of each
+// set of banks and the number of sets of banks
+using Slice = std::tuple<gsl::span<char>, gsl::span<unsigned int>, size_t>;
+using BankSlices = std::vector<Slice>;
+using Slices = std::array<BankSlices, NBankTypes>;
+
+//
+/**
+ * @brief      read events from input file into prefetch buffer
+ *
+ * @details    NOTE: It is assumed that the header has already been
+ *             read, calling read_events will read the subsequent
+ *             banks and then header of the next event.
+ *
+ * @param      input stream
+ * @param      prefetch buffer to read into
+ * @param      storage for the MDF header
+ * @param      buffer for temporary storage of the compressed banks
+ * @param      number of events to read
+ * @param      check the MDF checksum if it is available
+ *
+ * @return     (eof, error, full, n_bytes)
+ */
+std::tuple<bool, bool, bool, size_t> read_events(
+  int input,
+  ReadBuffer& read_buffer,
+  LHCb::MDFHeader& header,
+  std::vector<char> compress_buffer,
+  size_t n_events,
+  bool check_checksum)
+{
+  auto& [n_filled, event_offsets, buffer] = read_buffer;
+
+  // Keep track of where to write and the end of the prefetch buffer
+  auto* buffer_start = &buffer[0];
+  size_t n_bytes = 0;
+  bool eof = false, error = false, full = false;
+  gsl::span<char> bank_span;
+
+  // Loop until the requested number of events is prefetched, the
+  // maximum number of events per prefetch buffer is hit, an error
+  // occurs or eof is reached
+  while (!eof && !error && n_filled < event_offsets.size() - 1 && n_filled < n_events) {
+    // It is
+
+    // Read the banks
+    gsl::span<char> buffer_span {buffer_start + event_offsets[n_filled], buffer.size() - event_offsets[n_filled]};
+    std::tie(eof, error, bank_span) =
+      MDF::read_banks(input, header, std::move(buffer_span), compress_buffer, check_checksum);
+    // Fill the start offset of the next event
+    event_offsets[++n_filled] = bank_span.end() - buffer_start;
+    n_bytes += bank_span.size();
+
+    // read the next header
+    ssize_t n_bytes = ::read(input, &header, header_size);
+    if (n_bytes != 0) {
+      // Check if there is enough space to read this event
+      int compress = header.compression() & 0xF;
+      int expand = (header.compression() >> 4) + 1;
+      int event_size =
+        (header.recordSize() + header_size + 2 * (sizeof(LHCb::RawBank) + sizeof(int)) +
+         (compress ? expand * (header.recordSize() - header_size) : 0));
+      if (event_offsets[n_filled] + event_size > buffer.size()) {
+        full = true;
+        break;
+      }
+    }
+    else if (n_bytes == 0) {
+      info_cout << "Cannot read more data (Header). End-of-File reached.\n";
+      eof = true;
+    }
+    else {
+      error_cout << "Failed to read header " << strerror(errno) << "\n";
+      error = true;
+    }
+  }
+  return {eof, error, full, n_bytes};
+}
+
+/**
+ * @brief      Fill the array the contains the number of banks per type
+ *
+ * @details    detailed description
+ *
+ * @param      prefetched buffer of events (a single event is needed)
+ *
+ * @return     (success, number of banks per bank type; 0 if the bank is not needed)
+ */
+std::tuple<bool, std::array<unsigned int, NBankTypes>> fill_counts(gsl::span<char const> bank_data)
+{
+
+  std::array<unsigned int, NBankTypes> count {0};
+
+  auto const* bank = bank_data.data();
+
+  // Loop over all the bank data
+  while (bank < bank_data.end()) {
+    const auto* b = reinterpret_cast<const LHCb::RawBank*>(bank);
+
+    if (b->magic() != LHCb::RawBank::MagicPattern) {
+      error_cout << "Magic pattern failed: " << std::hex << b->magic() << std::dec << "\n";
+      return {false, count};
+    }
+
+    // Check if Allen processes this bank type, count bank types that
+    // are wanted
+    auto bank_type_it = Allen::bank_types.find(b->type());
+    if (bank_type_it != Allen::bank_types.end()) {
+      auto bank_type_index = to_integral(bank_type_it->second);
+      ++count[bank_type_index];
+    }
+
+    // Increment overall bank pointer
+    bank += b->totalSize();
+  }
+
+  return {true, count};
+}
+
+/**
+ * @brief      Transpose events to Allen layout
+ *
+ * @param      slices to fill with transposed banks, slices are addressed by bank type
+ * @param      index of bank slices
+ * @param      number of banks per event
+ * @param      event ids of banks in this slice
+ * @param      start of bank data for this event
+ *
+ * @return     tuple of: (success, slice is full)
+ */
+template<BankTypes... Banks>
+std::tuple<bool, bool> transpose_event(
+  Slices& slices,
+  int const slice_index,
+  std::vector<int> const& bank_ids,
+  std::array<unsigned int, NBankTypes> const& banks_count,
+  EventIDs& event_ids,
+  const gsl::span<char const> bank_data)
+{
+
+  unsigned int* banks_offsets = nullptr;
+  // Number of offsets
+  size_t* n_banks_offsets = nullptr;
+
+  // Where to write the transposed bank data
+  uint32_t* banks_write = nullptr;
+
+  // Where should offsets to individual banks be written
+  uint32_t* banks_offsets_write = nullptr;
+
+  unsigned int bank_offset = 0;
+  unsigned int bank_counter = 1;
+  unsigned int bank_type_index = 0;
+
+  auto bank = bank_data.begin(), bank_end = bank_data.end();
+
+  // L0Calo doesn't exist in the upgrade
+  LHCb::RawBank::BankType prev_type = LHCb::RawBank::L0Calo;
+
+  // Loop over all bank data of this event
+  while (bank < bank_end) {
+    const auto* b = reinterpret_cast<const LHCb::RawBank*>(bank);
+
+    if (b->magic() != LHCb::RawBank::MagicPattern) {
+      error_cout << "Magic pattern failed: " << std::hex << b->magic() << std::dec << "\n";
+      return {false, false};
+      // Decode the odin bank
+    }
+
+    // Check what to do with this bank
+    auto bt = b->type();
+    if (bt == LHCb::RawBank::ODIN) {
+      // decode ODIN bank to obtain run and event numbers
+      auto odin = MDF::decode_odin(b);
+      event_ids.emplace_back(odin.run_number, odin.event_number);
+      bank += b->totalSize();
+      continue;
+    }
+    else if (bt >= LHCb::RawBank::LastType || bank_ids[bt] == -1) {
+      // This bank is not required: skip it
+      bank += b->totalSize();
+      continue;
+    }
+    else if (bt != prev_type) {
+      // Switch to new type of banks
+      bank_type_index = bank_ids[b->type()];
+      auto& slice = slices[bank_type_index][slice_index];
+      prev_type = bt;
+
+      bank_counter = 1;
+      banks_offsets = std::get<1>(slice).data();
+      n_banks_offsets = &std::get<2>(slice);
+
+      // Calculate the size taken by storing the number of banks
+      // and offsets to all banks within the event
+      auto preamble_words = 2 + banks_count[bank_type_index];
+
+      // Initialize offset to start of this set of banks from the
+      // previous one and increment with the preamble size
+      banks_offsets[*n_banks_offsets] = (banks_offsets[*n_banks_offsets - 1] + preamble_words * sizeof(uint32_t));
+
+      // Three things to write for a new set of banks:
+      // - number of banks/offsets
+      // - offsets to individual banks
+      // - bank data
+
+      // Initialize point to write from offset of previous set
+      banks_write = reinterpret_cast<uint32_t*>(std::get<0>(slice).data() + banks_offsets[*n_banks_offsets - 1]);
+
+      // New offset to increment
+      ++(*n_banks_offsets);
+
+      // Write the number of banks
+      banks_write[0] = banks_count[bank_type_index];
+
+      // All bank offsets are uit32_t so cast to that type
+      banks_offsets_write = banks_write + 1;
+      banks_offsets_write[0] = 0;
+
+      // Offset in number of uint32_t
+      bank_offset = 0;
+
+      // Start writing bank data after the preamble
+      banks_write += preamble_words;
+    }
+    else {
+      ++bank_counter;
+    }
+
+    // Write sourceID
+    banks_write[bank_offset] = b->sourceID();
+
+    // Write bank data
+    ::memcpy(banks_write + bank_offset + 1, b->data(), b->size());
+
+    auto n_word = b->size() / sizeof(uint32_t);
+    bank_offset += 1 + n_word;
+
+    // Write next offset in bytes
+    banks_offsets_write[bank_counter] = bank_offset * sizeof(uint32_t);
+
+    // Update "event" offset (in bytes)
+    banks_offsets[*n_banks_offsets - 1] += sizeof(uint32_t) * (1 + n_word);
+
+    // Increment overall bank pointer
+    bank += b->totalSize();
+  }
+
+  // Check if any of the per-bank-type slices potentially has too
+  // little space to fit the next event
+  for (auto bank_type : {Banks...}) {
+    auto ib = to_integral<BankTypes>(bank_type);
+    const auto& [slice, slice_offsets, offsets_size] = slices[ib][slice_index];
+    // Use the event size of the next event here instead of the
+    // per bank size because that's not yet known for the next
+    // event
+    if ((slice_offsets[offsets_size - 1] + bank_data.size()) > slice.size()) {
+      return {true, true};
+    }
+  }
+  return {true, false};
+}
+
+/**
+ * @brief      Reset a slice
+ *
+ * @param      slices
+ * @param      slice_index
+ * @param      event_ids
+ */
+template<BankTypes... Banks>
+void reset_slice(Slices& slices, int const slice_index, EventIDs& event_ids)
+{
+  // "Reset" the slice
+  for (auto bank_type : {Banks...}) {
+    auto ib = to_integral<BankTypes>(bank_type);
+    std::get<1>(slices[ib][slice_index])[0] = 0;
+    std::get<2>(slices[ib][slice_index]) = 1;
+  }
+  event_ids.clear();
+}
+
+/**
+ * @brief      Transpose events to Allen layout
+ *
+ * @param      ReadBuffer containing events to be transposed
+ * @param      slices to fill with transposed banks, slices are addressed by bank type
+ * @param      index of bank slices
+ * @param      event ids of banks in this slice
+ * @param      number of banks per event
+ * @param      number of events to transpose
+ *
+ * @return     (success, slice full for one of the bank types, number of events transposed)
+ */
+template<BankTypes... Banks>
+std::tuple<bool, bool, size_t> transpose_events(
+  const ReadBuffer& read_buffer,
+  Slices& slices,
+  int const slice_index,
+  std::vector<int> const& bank_ids,
+  std::array<unsigned int, NBankTypes> const& banks_count,
+  EventIDs& event_ids,
+  size_t n_events)
+{
+
+  bool full = false, success = true;
+  auto const& [n_filled, event_offsets, buffer] = read_buffer;
+
+  // Loop over events in the prefetch buffer
+  size_t i_event = 0;
+  for (; i_event < n_filled && i_event < n_events && success && !full; ++i_event) {
+    // Offsets are to the start of the event, which includes the header
+    auto const* bank = buffer.data() + event_offsets[i_event];
+    auto const* bank_end = buffer.data() + event_offsets[i_event + 1];
+    std::tie(success, full) =
+      transpose_event<Banks...>(slices, slice_index, bank_ids, banks_count, event_ids, {bank, bank_end});
+  }
+
+  return {success, full, i_event};
+}
diff --git a/main/src/Allen.cpp b/main/src/Allen.cpp
index 3078b175eb18b6ea83bbc758d3ff6759f03c6ed8..67551fbee68b954893d5ced71e27e49ad3599dae 100644
--- a/main/src/Allen.cpp
+++ b/main/src/Allen.cpp
@@ -78,8 +78,6 @@ void input_reader(const size_t io_id, IInputProvider* input_provider)
     // Check if there are messages
     zmq::poll(&items[0], 1, 0);
 
-    size_t idx = 0;
-    size_t fill = 0;
     if (items[0].revents & zmq::POLLIN) {
       auto msg = zmqSvc().receive<string>(control);
       if (msg == "DONE") {
diff --git a/main/src/InputTools.cpp b/main/src/InputTools.cpp
index d422d634a18947cc31b45e6d5c5cfa72e6a06cac..46dd865f837d41abb07bed176d3fb58dd41b46e5 100644
--- a/main/src/InputTools.cpp
+++ b/main/src/InputTools.cpp
@@ -49,7 +49,7 @@ namespace {
     return std::less<EventID> {}(name_to_number(lhs), name_to_number(rhs));
   };
 
-}; // namespace
+} // namespace
 
 /**
  * @brief      Convert "N.bin" to (0, N) and "N_M.bin" to (N, M)
@@ -70,7 +70,7 @@ EventID name_to_number(const std::string& arg)
   else {
     return {std::stoi(std::string {m[1].first, m[1].second}), std::stol(std::string {m[2].first, m[2].second})};
   }
-};
+}
 
 /**
  * @brief Test to check existence of filename.
diff --git a/main/src/MDFReader.cpp b/main/src/MDFReader.cpp
deleted file mode 100644
index 99e799124cd55d9aa4f0b6bde08b6fe6ad93c502..0000000000000000000000000000000000000000
--- a/main/src/MDFReader.cpp
+++ /dev/null
@@ -1,70 +0,0 @@
-#include <vector>
-#include <string>
-
-#include "MDFReader.h"
-
-#include "read_mdf.hpp"
-#include "odin.hpp"
-
-namespace {
-  using std::pair;
-  using std::string;
-  using std::vector;
-} // namespace
-
-std::vector<std::tuple<unsigned int, unsigned long>> MDFReader::read_events(
-  uint number_of_events_requested,
-  uint start_event_offset)
-{
-
-  size_t n_read = 0;
-  Allen::buffer_map buffers;
-  vector<LHCb::ODIN> odins;
-
-  auto bank_type = *begin(types());
-  auto foldername = folder(bank_type);
-  auto filenames = list_folder(foldername, "mdf");
-  vector<string> files;
-  files.reserve(filenames.size());
-  for (auto name : filenames) {
-    files.emplace_back(foldername + "/" + name);
-  }
-  std::tie(n_read, buffers, odins) = MDF::read_events(number_of_events_requested, files, types(), start_event_offset);
-  std::vector<std::tuple<unsigned int, unsigned long>> event_ids;
-  event_ids.reserve(odins.size());
-  for (auto& odin : odins) {
-    event_ids.emplace_back(odin.run_number, static_cast<unsigned long>(odin.event_number));
-  }
-  for (auto bank_type : types()) {
-    auto it = buffers.find(bank_type);
-    if (it == end(buffers)) {
-      throw StrException(string {"Cannot find buffer for bank type "} + bank_name(bank_type));
-    }
-    auto& entry = it->second;
-    ;
-    check_events(bank_type, entry.first, entry.second, number_of_events_requested);
-  }
-
-  // TODO Remove: Temporal check to understand if number_of_events_requested is the same as number_of_events
-  const uint number_of_events = begin(buffers)->second.second.size() - 1;
-  if (number_of_events_requested != number_of_events) {
-    throw StrException("Number of events requested differs from number of events read.");
-  }
-
-  for (auto bank_type : types()) {
-    auto it = buffers.find(bank_type);
-    const auto& ev_buf = it->second.first;
-    const auto& offsets_buf = it->second.second;
-
-    // Copy raw data to pinned host memory
-    char* events_mem = nullptr;
-    uint* offsets_mem = nullptr;
-    cudaCheck(cudaMallocHost((void**) &events_mem, ev_buf.size()));
-    cudaCheck(cudaMallocHost((void**) &offsets_mem, offsets_buf.size() * sizeof(uint)));
-    std::copy_n(std::begin(ev_buf), ev_buf.size(), events_mem);
-    std::copy_n(std::begin(offsets_buf), offsets_buf.size(), offsets_mem);
-
-    add_events(bank_type, {events_mem, ev_buf.size()}, {offsets_mem, offsets_buf.size()});
-  }
-  return event_ids;
-}
diff --git a/mdf/.gaudi_project_ignore b/mdf/.gaudi_project_ignore
new file mode 100644
index 0000000000000000000000000000000000000000..e69de29bb2d1d6434b8b29ae775ad8c2e48c5391
diff --git a/mdf/CMakeLists.txt b/mdf/CMakeLists.txt
index 5e91f11dfd02d749888532b125dad111ebc72f3f..44d1051b6e976b99908d35baed7950c914ee5a49 100644
--- a/mdf/CMakeLists.txt
+++ b/mdf/CMakeLists.txt
@@ -3,7 +3,7 @@ set(SOURCES
   src/raw_helpers.cpp
   src/read_mdf.cpp)
 
-add_library(mdf ${SOURCES})
+allen_add_library(mdf ${SOURCES})
 target_include_directories (mdf PUBLIC ${CMAKE_CURRENT_SOURCE_DIR}/include)
 
 target_include_directories(
@@ -72,21 +72,20 @@ function(test_program)
     target_compile_definitions(${PARSED_ARGS_NAME} PRIVATE NO_CUDA)
   endif()
   target_link_libraries(${PARSED_ARGS_NAME} PRIVATE Common mdf)
+  install(TARGETS ${PARSED_ARGS_NAME} RUNTIME DESTINATION bin OPTIONAL)
 endfunction()
 
-if(TARGET_DEVICE STREQUAL "CUDA")
-  test_program(NAME dump_banks)
+test_program(NAME dump_banks)
 
-  if (BUILD_TESTS)
-    include(Catch)
+if (BUILD_TESTS)
+  include(Catch)
 
-    test_program(NAME test_read)
-    test_program(NAME test_read_bin)
-    test_program(NAME bench_read NO_CUDA)
-    test_program(NAME bench_provider NO_CUDA)
-    test_program(NAME bench_transpose NO_CUDA)
+  test_program(NAME test_read)
+  test_program(NAME test_read_bin)
+  test_program(NAME bench_read NO_CUDA)
+  test_program(NAME bench_provider NO_CUDA)
+  test_program(NAME bench_transpose NO_CUDA)
 
-    test_program(NAME test_providers)
-    catch_discover_tests(test_providers EXTRA_ARGS --directory ${CMAKE_SOURCE_DIR}/input/minbias)
-  endif()
+  test_program(NAME test_providers)
+  catch_discover_tests(test_providers EXTRA_ARGS --directory ${CMAKE_SOURCE_DIR}/input/minbias)
 endif()
diff --git a/mdf/include/read_mdf.hpp b/mdf/include/read_mdf.hpp
index 7f337a729e905f4d4805f6d0b147cc7d87a44f9c..860942499208f328c71c2e3dd19366253b3d6b3c 100644
--- a/mdf/include/read_mdf.hpp
+++ b/mdf/include/read_mdf.hpp
@@ -50,13 +50,6 @@ namespace MDF {
     bool checkChecksum = true,
     size_t offset = 0);
 
-  void transpose_event(
-    std::vector<char> const& input_buffer,
-    size_t input_offset,
-    char* output_buffer,
-    size_t output_offset,
-    size_t output_size);
-
   LHCb::ODIN decode_odin(const LHCb::RawBank* bank);
 
 } // namespace MDF
diff --git a/mdf/src/read_mdf.cpp b/mdf/src/read_mdf.cpp
index ab0edd87a5a62aa2f9e06d0fe5e7787a90da41ca..5454d10558d35e14a81a33544e443235ba72a31c 100644
--- a/mdf/src/read_mdf.cpp
+++ b/mdf/src/read_mdf.cpp
@@ -394,11 +394,3 @@ void MDF::dump_hex(const char* start, int size)
   cout << std::dec << std::setfill(prev) << "\n";
   cout.setf(flags);
 }
-
-void MDF::transpose_event(
-  vector<char> const& input_buffer,
-  size_t input_offset,
-  char* output_buffer,
-  size_t output_offset,
-  size_t output_size)
-{}
diff --git a/mdf/test/bench_transpose.cpp b/mdf/test/bench_transpose.cpp
index 14670f4600d85e62cb4739ea62df5f2bbd3fbe41..f4285abb2d316790cdbcbd8df8a434681b284585 100644
--- a/mdf/test/bench_transpose.cpp
+++ b/mdf/test/bench_transpose.cpp
@@ -28,8 +28,6 @@ int main(int argc, char* argv[])
   }
 
   // Test parameters
-  size_t events_per_slice = 1000;
-  double n_filled = 0.;
   size_t n_events = 10000;
   size_t offsets_size = 10001;
   size_t n_reps = 50;
@@ -140,7 +138,8 @@ int main(int argc, char* argv[])
   cout << "read " << std::lround(n_read) << " events; " << n_read / t.get() << " events/s\n";
 
   // Count the number of banks of each type
-  auto [count_success, banks_count] = fill_counts(read_buffers[0]);
+  auto& [n_filled, event_offsets, read_buffer] = read_buffers[0];
+  auto [count_success, banks_count] = fill_counts({read_buffer.data(), event_offsets[1]});
 
   // Allocate space for event ids
   std::vector<EventIDs> event_ids(n_slices);
@@ -159,9 +158,14 @@ int main(int argc, char* argv[])
     threads.emplace_back(thread {[i, n_reps, n_events, &read_buffers, &slices, &bank_ids, &banks_count, &event_ids] {
       auto& read_buffer = read_buffers[i];
       for (size_t rep = 0; rep < n_reps; ++rep) {
+
+        // Reset the slice
+        reset_slice<BankTypes::VP, BankTypes::UT, BankTypes::FT, BankTypes::MUON>(slices, i, event_ids[i]);
+
+        // Transpose events
         auto [success, transpose_full, n_transposed] =
           transpose_events<BankTypes::VP, BankTypes::UT, BankTypes::FT, BankTypes::MUON>(
-            read_buffer, slices, i, event_ids[i], bank_ids, banks_count, n_events);
+            read_buffer, slices, i, bank_ids, banks_count, event_ids[i], n_events);
         info_cout << "thread " << i << " " << success << " " << transpose_full << " " << n_transposed << endl;
       }
     }});
diff --git a/mdf/test/dump_banks.cpp b/mdf/test/dump_banks.cpp
index 3de6fc27168124197c9368ee979112b71791f512..bc25908f36f7f758758a5fea353c739b5a28b312 100644
--- a/mdf/test/dump_banks.cpp
+++ b/mdf/test/dump_banks.cpp
@@ -60,7 +60,7 @@ int main(int argc, char* argv[])
     }
     else {
       const int status = mkdir(output_dir.c_str(), S_IRWXU | S_IRWXG | S_IROTH | S_IXOTH);
-      if (status != 0 & errno != EEXIST) {
+      if (status != 0 && errno != EEXIST) {
         cerr << "Error creating directory " << output_dir << endl;
         cerr << std::strerror(errno) << endl;
         return status;
@@ -81,7 +81,7 @@ int main(int argc, char* argv[])
   for (auto type : types) {
     auto dir = output_dir + "/" + bank_name(type);
     const int status = mkdir(dir.c_str(), S_IRWXU | S_IRWXG | S_IROTH | S_IXOTH);
-    if (status != 0 & errno != EEXIST) {
+    if (status != 0 && errno != EEXIST) {
       cerr << "Error creating directory " << output_dir << endl;
       cerr << std::strerror(errno) << endl;
       return status;
diff --git a/readme.md b/readme.md
index 9e110bc46bb59616ff0b9b876dd9cf60ae42c0f3..c2a8425eae2e8778ef7fee25f82046c9d3eb92bf 100644
--- a/readme.md
+++ b/readme.md
@@ -103,11 +103,11 @@ Here are some example run options:
 
     # Run one stream and print all memory allocations
     ./Allen -n 5000 -p
-    
-    
+
+
 How to profile it
 ------------------
-For profiling, Nvidia's nvprof can be used. 
+For profiling, Nvidia's nvprof can be used.
 Since CUDA version 10.1, profiling was limited to the root user by default for security reasons. However, the system administrator of a GPU server can add a kernel module option such that regular users can use the profiler by following these instructions:
 
 Add a file containing "option nvidia NVreg_RestrictProfilingToAdminUsers=0" to the `/etc/modprobe.d/` directory and reboot the machine. This will load the nvidia kernel module with "NVreg_RestrictProfilingToAdminUsers=0".
@@ -116,53 +116,43 @@ As a quick workaround one can also use the older version of nvprof:
 
     /usr/local/cuda-10.0/bin/nvprof ./Allen -c 0 -n 1000
 
-How to run build and run together with the LHCb stack
------------------------------------------------------
-
-Code is being developed in the LHCb stack that provides detector
-geometry data directly instead of reading it from binary files. To do
-this, the following is required:
- - Build with CUDA 10.1 (to allow gcc 8 as a host compiler)
- - Build the Rec project against the lhcb-gaudi-head nightly build
- - Build Allen using the same toolchain as Rec
- - Run Allen from a runtime environment provided by Rec
-
-### Building Rec
-Running Allen together with the LHCb stack requires some recent
-changes that have not yet been merged to master. Rec has to be built
-in its entirity against the lhcb-gaudi-head slot. First, decide on a
-directory where it will reside (`dev-dir` below) and then clone there:
- - `mkdir /path/to/dev-dir`
- - `cd /path/to/dev-dir`
-
-To setup the LHCb environment for building and running Rec, put the
-following in a script (e.g. `env.sh`) for easy access:
+Building as a Gaudi/LHCb project
+--------------------------------
+
+Allen can also be built as a Gaudi/LHCb cmake project; it then depends
+on Rec and Online. To build Allen like this, is the same as building
+any other Gaudi/LHCb project:
+
+    source /cvmfs/lhcb.cern.ch/lib/LbEnv
+    cd Allen
+    lb-project-init
+    make configure
+    make install
+
+### Build options
+By default the `DefaultSequence` is selected, Allen is built with
+CUDA, and the CUDA stack is searched for in `/usr/local/cuda`. These
+defaults (and other cmake variables) can be changed by adding the same
+flags that you would pass to a standalone build to the `CMAKEFLAGS`
+environment variable before calling `make configure`.
+
+For example, to specify another CUDA stack to be used set:
 ```console
-export CMTCONFIG=x86_64-centos7-gcc8-opt
-export CMTPROJECTPATH=/path/to/dev-dir:/cvmfs/lhcbdev.cern.ch/nightlies/lhcb-gaudi-head/Mon
-source /cvmfs/lhcb.cern.ch/lib/LbEnv
+$> export CMAKEFLAGS="-DCMAKE_CUDA_COMPILER=/path/to/alternative/nvcc"
 ```
 
-Then build Rec:
- - `source env.sh`
- - `git clone ssh://git@gitlab.cern.ch:7999/lhcb/Rec.git`
- - `cd Rec`
- - `lb-project-init`
- - `make install`
-
-### Building Allen with the toolchain used for Rec
-In the same environment, do the following
- - `cd /path/to/Allen`
- - `mkdir build-Rec`
- - `cd build-Rec`
- - `/path/to/dev-dir/Rec/build.${CMTCONFIG}/run bash --norc`
- - ```cmake -DCMAKE_C_COMPILER=`which gcc` -DCMAKE_CXX_COMPILER=`which g++` -DCMAKE_CUDA_COMPILER=/usr/local/cuda-10.1/bin/nvcc ..```
- - `make -j 10`
+### Runtime environment:
+To setup the runtime environment for Allen, the same tools as for
+other Gaudi/LHCb projects can be used:
+```console
+$> cd Allen
+$> ./build.${BINARY_TAG}/run Allen ...
+```
 
 ### Run Allen using the Python entry point:
 ```console
-$> /path/to/dev-dir/Rec/build.${CMTCONFIG}/run /path/to/Allen/bindings/Allen.py
+$> cd Allen
+$> ./build.${CMTCONFIG}/run bindings/Allen.py
 ```
 
-
 [This readme](contributing.md) explains how to add a new algorithm to the sequence and how to use the memory scheduler to define global memory variables for this sequence and pass on the dependencies. It also explains which checks to do before placing a merge request with your changes.
diff --git a/stream/.gaudi_project_ignore b/stream/.gaudi_project_ignore
new file mode 100644
index 0000000000000000000000000000000000000000..e69de29bb2d1d6434b8b29ae775ad8c2e48c5391
diff --git a/stream/CMakeLists.txt b/stream/CMakeLists.txt
index 1cd87b9ede26bbe246f53721f02dc911fe722142..079df287f60b82f6b7c37640f57f006055b6b8ac 100644
--- a/stream/CMakeLists.txt
+++ b/stream/CMakeLists.txt
@@ -109,7 +109,7 @@ add_custom_command(
 message(STATUS "Configured sequence: " ${SEQUENCE})
 
 if(TARGET_DEVICE STREQUAL "CPU")
-  foreach(source_file ${stream_gear} ${stream_memory_manager} ${stream_scheduler} 
+  foreach(source_file ${stream_gear} ${stream_memory_manager} ${stream_scheduler}
     ${stream_sequence} ${stream_sequence_cpp} ${stream_setup} ${stream_visitors_gec}
     ${stream_visitors_ip} ${stream_visitors_velo} ${stream_visitors_patPV}
     ${stream_visitors_beamlinePV} ${stream_visitors_assoc} ${stream_visitors_UT}
@@ -142,12 +142,11 @@ allen_add_library(Stream STATIC
 )
 
 target_link_libraries(Stream PRIVATE
-  Common
-  Utils
+  CudaCommon
   Associate
   GlobalEventCut
   Velo
-  PatPV
+  AllenPatPV
   PV_beamline
   x86Clustering
   x86beamlinePV
@@ -160,14 +159,18 @@ target_link_libraries(Stream PRIVATE
   x86LookingForward
   x86MomentumForward
   x86VeloUT
+  x86Clustering
   CpuGEC
   CpuUtils
-  Muon)
+  Muon
+  Utils)
+
+set_property(TARGET Stream PROPERTY CUDA_RESOLVE_DEVICE_SYMBOLS ON)
 
 if (ROOT_FOUND)
   target_compile_definitions(Stream PUBLIC ${ALLEN_ROOT_DEFINITIONS})
 
-  target_include_directories(Stream BEFORE PRIVATE
+  target_include_directories(Stream SYSTEM BEFORE PRIVATE
     ${ROOT_INCLUDE_DIRS}
   )
 endif()
diff --git a/tests/.gaudi_project_ignore b/tests/.gaudi_project_ignore
new file mode 100644
index 0000000000000000000000000000000000000000..e69de29bb2d1d6434b8b29ae775ad8c2e48c5391
diff --git a/x86/.gaudi_project_ignore b/x86/.gaudi_project_ignore
new file mode 100644
index 0000000000000000000000000000000000000000..e69de29bb2d1d6434b8b29ae775ad8c2e48c5391
diff --git a/x86/PV/beamlinePV/CMakeLists.txt b/x86/PV/beamlinePV/CMakeLists.txt
index 01c914f019360cdb6b10d5a6598a1ce2fad17950..368ce760d0dd06c86112dfd44f433da2e22f55b1 100644
--- a/x86/PV/beamlinePV/CMakeLists.txt
+++ b/x86/PV/beamlinePV/CMakeLists.txt
@@ -18,7 +18,7 @@ add_library(x86beamlinePV STATIC
 set_property(TARGET x86beamlinePV PROPERTY
              CUDA_RESOLVE_DEVICE_SYMBOLS ON)
 
-if ( ROOT_FOUND )
+if (USE_ROOT AND ROOT_FOUND)
   target_compile_definitions(x86beamlinePV PRIVATE WITH_ROOT)
   target_include_directories(x86beamlinePV PRIVATE ${ROOT_INCLUDE_DIRS})
 endif()
diff --git a/x86/SciFi/LookingForward/CMakeLists.txt b/x86/SciFi/LookingForward/CMakeLists.txt
index c09b981187b72d91737d51207f1627764bbaf6b0..99aab9f56bfb51a5ff4675aaa985b1f46552f816 100644
--- a/x86/SciFi/LookingForward/CMakeLists.txt
+++ b/x86/SciFi/LookingForward/CMakeLists.txt
@@ -32,7 +32,7 @@ set_property(TARGET x86LookingForward PROPERTY
 set_property(TARGET x86LookingForward PROPERTY
              CUDA_RESOLVE_DEVICE_SYMBOLS ON)
 
-if ( ROOT_FOUND )
+if (USE_ROOT AND ROOT_FOUND)
   target_compile_definitions(x86LookingForward PUBLIC WITH_ROOT)
   target_include_directories(x86LookingForward BEFORE PRIVATE ${ROOT_INCLUDE_DIRS})
 endif()
diff --git a/x86/SciFi/MomentumForward/CMakeLists.txt b/x86/SciFi/MomentumForward/CMakeLists.txt
index ed98cb7844efbbd96b138cc5e80a67f1bd9c0b21..72050c934647a14c1e2d72c4c83bd8948162951e 100644
--- a/x86/SciFi/MomentumForward/CMakeLists.txt
+++ b/x86/SciFi/MomentumForward/CMakeLists.txt
@@ -23,7 +23,7 @@ add_library(x86MomentumForward STATIC
 set_property(TARGET x86MomentumForward PROPERTY
              CUDA_RESOLVE_DEVICE_SYMBOLS ON)
 
-if ( ROOT_FOUND )
+if (USE_ROOT AND ROOT_FOUND)
   target_compile_definitions(x86MomentumForward PUBLIC WITH_ROOT)
-  target_include_directories(x86MomentumForward BEFORE PRIVATE ${ROOT_INCLUDE_DIRS})
+  target_include_directories(x86MomentumForward SYSTEM BEFORE PRIVATE ${ROOT_INCLUDE_DIRS})
 endif()
diff --git a/x86/UT/PrVeloUT/CMakeLists.txt b/x86/UT/PrVeloUT/CMakeLists.txt
index 838790b682258173feeca94fd515bd57ca2c0235..31799e86c5e48bde5f8d84562ec7288c7aaf7e1d 100644
--- a/x86/UT/PrVeloUT/CMakeLists.txt
+++ b/x86/UT/PrVeloUT/CMakeLists.txt
@@ -19,9 +19,7 @@ add_library(x86VeloUT STATIC
 set_property(TARGET x86VeloUT PROPERTY
              CUDA_SEPARABLE_COMPILATION ON)
 
-if ( ROOT_FOUND )
+if (USE_ROOT AND ROOT_FOUND)
   target_compile_definitions(x86VeloUT PUBLIC WITH_ROOT)
-  target_include_directories(x86VeloUT BEFORE PRIVATE
-    ${ROOT_INCLUDE_DIRS}
-  )
+  target_include_directories(x86VeloUT BEFORE PRIVATE ${ROOT_INCLUDE_DIRS})
 endif()
diff --git a/x86/muon/decoding/CMakeLists.txt b/x86/muon/decoding/CMakeLists.txt
index 9cfab8fbf66f4c2c52bb1d031e14b8f36da6f0c2..410fbbfaecceed9a63541be4966119e14d2b9651 100644
--- a/x86/muon/decoding/CMakeLists.txt
+++ b/x86/muon/decoding/CMakeLists.txt
@@ -14,7 +14,7 @@ add_library(x86MuonDecoding STATIC
 set_property(TARGET x86MuonDecoding PROPERTY
              CUDA_SEPARABLE_COMPILATION ON)
 
-if ( ROOT_FOUND )
+if (USE_ROOT AND ROOT_FOUND)
     target_compile_definitions(x86MuonDecoding PUBLIC WITH_ROOT)
-    target_include_directories(x86MuonDecoding BEFORE PRIVATE ${ROOT_INCLUDE_DIRS})
+    target_include_directories(x86MuonDecoding SYSTEM BEFORE PRIVATE ${ROOT_INCLUDE_DIRS})
 endif()
diff --git a/x86/muon/decoding/include/MuonGeometry.h b/x86/muon/decoding/include/MuonGeometry.h
index f9f88e7998ff08c008813927415a3255df44581e..3418d3c50b88bc88e71f228badac082e632a0619 100644
--- a/x86/muon/decoding/include/MuonGeometry.h
+++ b/x86/muon/decoding/include/MuonGeometry.h
@@ -15,4 +15,4 @@ namespace CPUMuon {
 
     unsigned int getADDInTell1(unsigned int Tell1_num, unsigned int ch) const;
   };
-}; // namespace CPUMuon
+} // namespace CPUMuon
diff --git a/x86/muon/decoding/include/MuonLayout.h b/x86/muon/decoding/include/MuonLayout.h
index ab06131ea112098c3cf3ba73e25246885c9f1dd2..c4d5792cbe39163e969f0ca52e543b5be0ddee1f 100644
--- a/x86/muon/decoding/include/MuonLayout.h
+++ b/x86/muon/decoding/include/MuonLayout.h
@@ -42,4 +42,4 @@ namespace CPUMuon {
   }
 
   inline bool operator!=(const MuonLayout& ml1, const MuonLayout& ml2) { return !(ml1 == ml2); }
-}; // namespace CPUMuon
+} // namespace CPUMuon
diff --git a/x86/muon/decoding/include/MuonRawToHits.h b/x86/muon/decoding/include/MuonRawToHits.h
index e56af7c800c6d2ddaaaf766698597825b3c3d57d..86a53dca2701bef24c4d2330d3259b2c1efc7d5e 100644
--- a/x86/muon/decoding/include/MuonRawToHits.h
+++ b/x86/muon/decoding/include/MuonRawToHits.h
@@ -51,4 +51,4 @@ namespace CPUMuon {
     MuonTable* stripY;
     MuonGeometry* muonGeometry;
   };
-}; // namespace CPUMuon
+} // namespace CPUMuon
diff --git a/x86/muon/decoding/include/MuonTable.h b/x86/muon/decoding/include/MuonTable.h
index 6a8b7fc75ee3fe59826b129ac1bb93a7b8a788e5..68d83934becd9749bab97dc70a6d64d1a43818af 100644
--- a/x86/muon/decoding/include/MuonTable.h
+++ b/x86/muon/decoding/include/MuonTable.h
@@ -43,4 +43,4 @@ namespace CPUMuon {
 
   void
   calcStripPos(MuonTable* strip, MuonTileID& tile, double& x, double& deltax, double& y, double& deltay, double& z);
-}; // namespace CPUMuon
+} // namespace CPUMuon
diff --git a/x86/muon/decoding/include/MuonTileID.h b/x86/muon/decoding/include/MuonTileID.h
index 67d286d4768aa3dc201ed15a8aba5e5c0eb50972..4393ba4a8f2d45a0f0def59578d1b356743e5815 100644
--- a/x86/muon/decoding/include/MuonTileID.h
+++ b/x86/muon/decoding/include/MuonTileID.h
@@ -68,4 +68,4 @@ namespace CPUMuon {
 
     unsigned int id() const { return m_muonid; }
   };
-}; // namespace CPUMuon
+} // namespace CPUMuon
diff --git a/x86/muon/decoding/src/MuonGeometry.cpp b/x86/muon/decoding/src/MuonGeometry.cpp
index 09ceac4f555fab56afb370fab12173d2688fe871..edc556c2a04db1c1bfeebd72b40fde30b3aa3b83 100644
--- a/x86/muon/decoding/src/MuonGeometry.cpp
+++ b/x86/muon/decoding/src/MuonGeometry.cpp
@@ -39,4 +39,4 @@ namespace CPUMuon {
       raw_input += sizeof(unsigned) * tilesSize;
     }
   }
-}; // namespace CPUMuon
+} // namespace CPUMuon
diff --git a/x86/muon/decoding/src/MuonRawToHits.cpp b/x86/muon/decoding/src/MuonRawToHits.cpp
index dc9afdc96432e435ff4a96560e537f8fa195900c..d1a6df4d107c7a2d0f909635b0647c92f25fbe42 100644
--- a/x86/muon/decoding/src/MuonRawToHits.cpp
+++ b/x86/muon/decoding/src/MuonRawToHits.cpp
@@ -33,23 +33,23 @@ namespace CPUMuon {
 
   void recalculateNumberOfHitsPerStationAndStationOffsets(::Muon::HitsSoA* hitsSoA, size_t totalNumberOfHits)
   {
-    int currentStation = MuonTileID::station(hitsSoA->tile[0]);
-    int initialCurrentStation = currentStation;
-    for (int i = 1; i < totalNumberOfHits; i++) {
+    auto currentStation = MuonTileID::station(hitsSoA->tile[0]);
+    auto initialCurrentStation = currentStation;
+    for (size_t i = 1; i < totalNumberOfHits; i++) {
       auto id = static_cast<unsigned int>(hitsSoA->tile[i]);
       auto currentTileStation = MuonTileID::station(id);
       if (currentTileStation != currentStation) {
-        for (int j = currentStation + 1; j <= currentTileStation; j++) {
+        for (unsigned int j = currentStation + 1; j <= currentTileStation; j++) {
           hitsSoA->station_offsets[j] = i;
         }
         currentStation = currentTileStation;
       }
     }
-    for (int j = currentStation; j + 1 < ::Muon::Constants::n_stations; j++) {
+    for (unsigned int j = currentStation; j + 1 < ::Muon::Constants::n_stations; j++) {
       hitsSoA->station_offsets[j + 1] = totalNumberOfHits;
     }
     if (initialCurrentStation == currentStation) {
-      for (int j = initialCurrentStation; j + 1 < ::Muon::Constants::n_stations; j++) {
+      for (unsigned int j = initialCurrentStation; j + 1 < ::Muon::Constants::n_stations; j++) {
         hitsSoA->station_offsets[j + 1] = totalNumberOfHits;
       }
     }
@@ -75,7 +75,6 @@ namespace CPUMuon {
 
     for (auto& decode : decoding) {
       std::vector<DigitsRange> perRegQua;
-      unsigned nReg = 0;
       auto it = decode.begin();
       for (auto jt = it; jt != decode.end(); ++jt) {
         if (regionAndQuarter(*jt) != regionAndQuarter(*it)) {
@@ -243,4 +242,4 @@ namespace CPUMuon {
       }
     }
   }
-}; // namespace CPUMuon
+} // namespace CPUMuon
diff --git a/x86/muon/decoding/src/MuonRawToHitsDecoding.cpp b/x86/muon/decoding/src/MuonRawToHitsDecoding.cpp
index 4d47510f25b67c4b994158582993b3080925fb42..6d86bbbdc1dd207e1070bcd24daf64a58f235b71 100644
--- a/x86/muon/decoding/src/MuonRawToHitsDecoding.cpp
+++ b/x86/muon/decoding/src/MuonRawToHitsDecoding.cpp
@@ -6,8 +6,8 @@ char muon_geometry_raw_input[100000];
 void muonRawToHitsDecode(
   char* events,
   unsigned int* offsets,
-  size_t events_size,
-  size_t offsets_size,
+  size_t,
+  size_t,
   std::vector<Muon::HitsSoA>& muon_hits_events,
   CPUMuon::MuonRawToHits* muonRawToHits)
 {
diff --git a/x86/muon/decoding/src/MuonTable.cpp b/x86/muon/decoding/src/MuonTable.cpp
index fd47103afe940738289603961b1356b3344b4fe1..f35c29d92d76e5c95b606a6d008acd499bf963db 100644
--- a/x86/muon/decoding/src/MuonTable.cpp
+++ b/x86/muon/decoding/src/MuonTable.cpp
@@ -125,12 +125,12 @@ namespace CPUMuon {
       std::copy_n((size_t*) raw_input, 1, &tableSize);
       (muonTable->points).resize(tableSize);
       raw_input += sizeof(size_t);
-      for (int i = 0; i < tableSize; i++) {
+      for (size_t i = 0; i < tableSize; i++) {
         size_t stationTableSize;
         std::copy_n((size_t*) raw_input, 1, &stationTableSize);
         raw_input += sizeof(size_t);
         (muonTable->points)[i].resize(stationTableSize);
-        for (int j = 0; j < stationTableSize; j++) {
+        for (size_t j = 0; j < stationTableSize; j++) {
           (muonTable->points)[i][j].insert(
             (muonTable->points)[i][j].end(), (float*) raw_input, ((float*) raw_input) + 3);
           raw_input += sizeof(float) * 3;
@@ -138,4 +138,4 @@ namespace CPUMuon {
       }
     }
   }
-}; // namespace CPUMuon
+} // namespace CPUMuon
diff --git a/zmq/CMakeLists.txt b/zmq/CMakeLists.txt
index 18a341c5a43c201bc771efe7bd23199dbf85ea8c..c313a360c72837cc1bf3df2f2083faf3276b5b0a 100644
--- a/zmq/CMakeLists.txt
+++ b/zmq/CMakeLists.txt
@@ -40,7 +40,7 @@ endif()
 find_package(Boost COMPONENTS serialization iostreams)
 
 file(GLOB SOURCES "src/*.cpp")
-add_library(AllenZMQ ${SOURCES})
+allen_add_library(AllenZMQ ${SOURCES})
 target_include_directories(AllenZMQ PUBLIC ${CMAKE_CURRENT_SOURCE_DIR}/include)
 target_include_directories(AllenZMQ PUBLIC ${CMAKE_SOURCE_DIR}/main/include)
 target_include_directories(AllenZMQ SYSTEM PUBLIC ${ZMQ_INCLUDE_DIRS} ${Boost_INCLUDE_DIR})
diff --git a/zmq/include/ZeroMQSvc.h b/zmq/include/ZeroMQSvc.h
index 948c3cc0bb83e872e7fbb0df32013a114f2c97e2..f7d87a800f50c98a6ac10ecaecbba32d79c279ec 100644
--- a/zmq/include/ZeroMQSvc.h
+++ b/zmq/include/ZeroMQSvc.h
@@ -117,6 +117,8 @@ class ZeroMQSvc {
 public:
   ZeroMQSvc() = default;
 
+  virtual ~ZeroMQSvc() = default;
+
   enum Encoding { Text = 0, Binary };
 
   Encoding encoding() const { return m_enc; };
diff --git a/zmq/include/zmq.hpp b/zmq/include/zmq.hpp
index 50aee3348c2ed38002f555050ad171f8e931e12c..77c31293a02423b4d076959d0b98ad67723a54b2 100644
--- a/zmq/include/zmq.hpp
+++ b/zmq/include/zmq.hpp
@@ -135,7 +135,7 @@ namespace zmq {
   public:
     error_t() : errnum(zmq_errno()) {}
 #ifdef ZMQ_CPP11
-    virtual const char* what() const noexcept { return zmq_strerror(errnum); }
+    virtual const char* what() const noexcept override { return zmq_strerror(errnum); }
 #else
     virtual const char* what() const throw() { return zmq_strerror(errnum); }
 #endif