diff --git a/RecoTracker/LSTCore/BuildFile.xml b/RecoTracker/LSTCore/BuildFile.xml
new file mode 100644
index 0000000000000..69c03f883986e
--- /dev/null
+++ b/RecoTracker/LSTCore/BuildFile.xml
@@ -0,0 +1,9 @@
+
+
+
+
+
+
+
+
+
diff --git a/RecoTracker/LSTCore/README.md b/RecoTracker/LSTCore/README.md
new file mode 100644
index 0000000000000..cc313ba97ddf1
--- /dev/null
+++ b/RecoTracker/LSTCore/README.md
@@ -0,0 +1,42 @@
+# LSTCore proof of concept
+
+**This is a proof of concept for how I think we could continue working towards the CMSSW integration while keeping the standalone version alive.**
+
+This branch of CMSSW contains all of the relevant LST code and can be built entirely within CMSSW. The setup process is what you would expect.
+
+```bash
+export CMSSW_VERSION=CMSSW_14_1_0_pre3
+export CMSSW_BRANCH=${CMSSW_VERSION}_LST_X_LSTCore
+source /cvmfs/cms.cern.ch/cmsset_default.sh
+cmsrel $CMSSW_VERSION
+cd $CMSSW_VERSION/src
+cmsenv
+git cms-init
+git remote add SegLink https://github.com/SegmentLinking/cmssw.git
+git fetch SegLink ${CMSSW_BRANCH}:SegLink_cmssw
+git checkout SegLink_cmssw
+git cms-addpkg RecoTracker/LST RecoTracker/LSTCore Configuration/ProcessModifiers RecoTracker/ConversionSeedGenerators RecoTracker/FinalTrackSelectors RecoTracker/IterativeTracking
+git submodule update --init --recursive
+scram b -j 8
+```
+
+## How it works
+
+The [TrackLooper repository](https://github.com/SegmentLinking/TrackLooper) is included as a git submodule in `RecoTracker/LSTCore` and the rest of the structure is set up using symlinks. Since we have made a lot of progress getting the code ready for CMSSW, it was just a matter of writing a simple `BuildFile.xml` file.
+
+## Benefits
+
+- It would make it easier to work towards the full integration if we have a self-contained thing. It would probably be easier to slowly adapt more of the "proper" CMSSW conventions instead of having to switch them all at once.
+- We can keep the standalone version alive for as long as needed.
+- Our CI can start running the checks that are done by the `cms-bot` for CMSSW PRs.
+
+## Disadvantages
+
+- I might be better to work towards having a single CMSSW package instead of having them separated in `LST` and `LSTCore`. However, I think we could use a similar approach in that case.
+- I couldn't think of anything else, but there's likely other disadvantages.
+
+## Things to do
+
+- There are a few minor changes that need to be made to the current LST package to get it to work with LSTCore.
+- At some point we'll have to figure out how to properly integrate the `data` directory.
+
diff --git a/RecoTracker/LSTCore/interface/alpaka/Constants.h b/RecoTracker/LSTCore/interface/alpaka/Constants.h
new file mode 100644
index 0000000000000..d4f023631af1d
--- /dev/null
+++ b/RecoTracker/LSTCore/interface/alpaka/Constants.h
@@ -0,0 +1,157 @@
+#ifndef Constants_cuh
+#define Constants_cuh
+
+#include
+
+#include "HeterogeneousCore/AlpakaInterface/interface/config.h"
+
+#ifdef CACHE_ALLOC
+#include "HeterogeneousCore/AlpakaInterface/interface/CachedBufAlloc.h"
+#endif
+
+#ifdef ALPAKA_ACC_GPU_CUDA_ENABLED
+#include
+#endif
+
+namespace SDL {
+// Half precision wrapper functions.
+#if defined(FP16_Base)
+#define __F2H __float2half
+#define __H2F __half2float
+ typedef __half float FPX;
+#else
+#define __F2H
+#define __H2F
+ typedef float FPX;
+#endif
+
+ using Idx = alpaka_common::Idx;
+ using Dim = alpaka_common::Dim3D;
+ using Dim1d = alpaka_common::Dim1D;
+ using Vec = alpaka_common::Vec3D;
+ using Vec1d = alpaka_common::Vec1D;
+ using WorkDiv = alpaka_common::WorkDiv3D;
+
+ using Acc = ALPAKA_ACCELERATOR_NAMESPACE::Acc3D;
+ using Dev = ALPAKA_ACCELERATOR_NAMESPACE::Device;
+ using DevHost = ALPAKA_ACCELERATOR_NAMESPACE::DevHost;
+ using QueueAcc = ALPAKA_ACCELERATOR_NAMESPACE::Queue;
+
+ Vec const elementsPerThread(Vec::all(static_cast(1)));
+
+// Needed for files that are compiled by g++ to not throw an error.
+// uint4 is defined only for CUDA, so we will have to revisit this soon when running on other backends.
+#if !defined(ALPAKA_ACC_GPU_CUDA_ENABLED) && !defined(ALPAKA_ACC_GPU_HIP_ENABLED)
+ struct uint4 {
+ unsigned int x;
+ unsigned int y;
+ unsigned int z;
+ unsigned int w;
+ };
+#endif
+
+ // Buffer type for allocations where auto type can't be used.
+ template
+ using Buf = alpaka::Buf;
+
+ // Allocation wrapper function to make integration of the caching allocator easier and reduce code boilerplate.
+ template
+ ALPAKA_FN_HOST ALPAKA_FN_INLINE Buf, T> allocBufWrapper(TAcc const& devAccIn,
+ TSize nElements,
+ TQueue queue) {
+#ifdef CACHE_ALLOC
+ return cms::alpakatools::allocCachedBuf(devAccIn, queue, Vec1d(static_cast(nElements)));
+#else
+ return alpaka::allocBuf(devAccIn, Vec1d(static_cast(nElements)));
+#endif
+ }
+
+ // Second allocation wrapper function when queue is not given. Reduces code boilerplate.
+ template
+ ALPAKA_FN_HOST ALPAKA_FN_INLINE Buf, T> allocBufWrapper(TAcc const& devAccIn, TSize nElements) {
+ return alpaka::allocBuf(devAccIn, Vec1d(static_cast(nElements)));
+ }
+
+ // Wrapper function to reduce code boilerplate for defining grid/block sizes.
+ ALPAKA_FN_HOST ALPAKA_FN_INLINE Vec createVec(int x, int y, int z) {
+ return Vec(static_cast(x), static_cast(y), static_cast(z));
+ }
+
+ // Adjust grid and block sizes based on backend configuration
+ template
+ ALPAKA_FN_HOST ALPAKA_FN_INLINE WorkDiv createWorkDiv(const Vec& blocksPerGrid,
+ const Vec& threadsPerBlock,
+ const Vec& elementsPerThreadArg) {
+ Vec adjustedBlocks = blocksPerGrid;
+ Vec adjustedThreads = threadsPerBlock;
+
+ // Serial execution, so all launch parameters set to 1.
+#if defined(ALPAKA_ACC_CPU_B_SEQ_T_SEQ_ENABLED)
+ adjustedBlocks = Vec::all(static_cast(1));
+ adjustedThreads = Vec::all(static_cast(1));
+#endif
+
+ // Threads enabled, set number of blocks to 1.
+#if defined(ALPAKA_ACC_CPU_B_SEQ_T_THREADS_ENABLED)
+ adjustedBlocks = Vec::all(static_cast(1));
+#endif
+
+ return WorkDiv(adjustedBlocks, adjustedThreads, elementsPerThreadArg);
+ }
+
+// If a compile time flag does not define PT_CUT, default to 0.8 (GeV)
+#ifndef PT_CUT
+ constexpr float PT_CUT = 0.8f;
+#endif
+
+ const unsigned int MAX_BLOCKS = 80;
+ const unsigned int MAX_CONNECTED_MODULES = 40;
+
+ const unsigned int N_MAX_PIXEL_SEGMENTS_PER_MODULE = 50000;
+
+ const unsigned int N_MAX_PIXEL_MD_PER_MODULES = 2 * N_MAX_PIXEL_SEGMENTS_PER_MODULE;
+
+ const unsigned int N_MAX_PIXEL_TRIPLETS = 5000;
+ const unsigned int N_MAX_PIXEL_QUINTUPLETS = 15000;
+
+ const unsigned int N_MAX_PIXEL_TRACK_CANDIDATES = 30000;
+ const unsigned int N_MAX_NONPIXEL_TRACK_CANDIDATES = 1000;
+
+ const unsigned int size_superbins = 45000;
+
+ //defining the constant host device variables right up here
+ ALPAKA_STATIC_ACC_MEM_GLOBAL const float miniMulsPtScaleBarrel[6] = {0.0052, 0.0038, 0.0034, 0.0034, 0.0032, 0.0034};
+ ALPAKA_STATIC_ACC_MEM_GLOBAL const float miniMulsPtScaleEndcap[5] = {0.006, 0.006, 0.006, 0.006, 0.006};
+ ALPAKA_STATIC_ACC_MEM_GLOBAL const float miniRminMeanBarrel[6] = {
+ 25.007152356, 37.2186993757, 52.3104270826, 68.6658656666, 85.9770373007, 108.301772384};
+ ALPAKA_STATIC_ACC_MEM_GLOBAL const float miniRminMeanEndcap[5] = {
+ 130.992832231, 154.813883559, 185.352604327, 221.635123002, 265.022076742};
+ ALPAKA_STATIC_ACC_MEM_GLOBAL const float k2Rinv1GeVf = (2.99792458e-3 * 3.8) / 2;
+ ALPAKA_STATIC_ACC_MEM_GLOBAL const float kR1GeVf = 1. / (2.99792458e-3 * 3.8);
+ ALPAKA_STATIC_ACC_MEM_GLOBAL const float sinAlphaMax = 0.95;
+ ALPAKA_STATIC_ACC_MEM_GLOBAL const float ptCut = PT_CUT;
+ ALPAKA_STATIC_ACC_MEM_GLOBAL const float deltaZLum = 15.0;
+ ALPAKA_STATIC_ACC_MEM_GLOBAL const float pixelPSZpitch = 0.15;
+ ALPAKA_STATIC_ACC_MEM_GLOBAL const float strip2SZpitch = 5.0;
+ ALPAKA_STATIC_ACC_MEM_GLOBAL const float pt_betaMax = 7.0;
+ ALPAKA_STATIC_ACC_MEM_GLOBAL const float magnetic_field = 3.8112;
+ // Since C++ can't represent infinity, SDL_INF = 123456789 was used to represent infinity in the data table
+ ALPAKA_STATIC_ACC_MEM_GLOBAL const float SDL_INF = 123456789.0;
+} //namespace SDL
+
+namespace T5DNN {
+ // Working points matching LST fake rate (43.9%) or signal acceptance (82.0%)
+ ALPAKA_STATIC_ACC_MEM_GLOBAL const float LSTWP1 = 0.3418833f; // 94.0% TPR, 43.9% FPR
+ ALPAKA_STATIC_ACC_MEM_GLOBAL const float LSTWP2 = 0.6177366f; // 82.0% TPR, 20.0% FPR
+ // Other working points
+ ALPAKA_STATIC_ACC_MEM_GLOBAL const float WP70 = 0.7776195f; // 70.0% TPR, 10.0% FPR
+ ALPAKA_STATIC_ACC_MEM_GLOBAL const float WP75 = 0.7181118f; // 75.0% TPR, 13.5% FPR
+ ALPAKA_STATIC_ACC_MEM_GLOBAL const float WP80 = 0.6492643f; // 80.0% TPR, 17.9% FPR
+ ALPAKA_STATIC_ACC_MEM_GLOBAL const float WP85 = 0.5655319f; // 85.0% TPR, 23.8% FPR
+ ALPAKA_STATIC_ACC_MEM_GLOBAL const float WP90 = 0.4592205f; // 90.0% TPR, 32.6% FPR
+ ALPAKA_STATIC_ACC_MEM_GLOBAL const float WP95 = 0.3073708f; // 95.0% TPR, 47.7% FPR
+ ALPAKA_STATIC_ACC_MEM_GLOBAL const float WP97p5 = 0.2001348f; // 97.5% TPR, 61.2% FPR
+ ALPAKA_STATIC_ACC_MEM_GLOBAL const float WP99 = 0.1120605f; // 99.0% TPR, 75.9% FPR
+ ALPAKA_STATIC_ACC_MEM_GLOBAL const float WP99p9 = 0.0218196f; // 99.9% TPR, 95.4% FPR
+} // namespace T5DNN
+#endif
diff --git a/RecoTracker/LSTCore/interface/alpaka/LST.h b/RecoTracker/LSTCore/interface/alpaka/LST.h
new file mode 100644
index 0000000000000..7d9d11745ab3c
--- /dev/null
+++ b/RecoTracker/LSTCore/interface/alpaka/LST.h
@@ -0,0 +1,115 @@
+#ifndef LST_H
+#define LST_H
+
+#ifdef LST_IS_CMSSW_PACKAGE
+#include "RecoTracker/LSTCore/interface/alpaka/Constants.h"
+#include "RecoTracker/LSTCore/interface/alpaka/LSTESData.h"
+#else
+#include "Constants.h"
+#include "LSTESData.h"
+#endif
+
+#include
+#include
+#include
+#include
+
+namespace SDL {
+ template
+ class Event;
+
+ template
+ class LST;
+
+ template <>
+ class LST {
+ public:
+ LST() = default;
+
+ void run(QueueAcc& queue,
+ bool verbose,
+ const LSTESDeviceData* deviceESData,
+ const std::vector see_px,
+ const std::vector see_py,
+ const std::vector see_pz,
+ const std::vector see_dxy,
+ const std::vector see_dz,
+ const std::vector see_ptErr,
+ const std::vector see_etaErr,
+ const std::vector see_stateTrajGlbX,
+ const std::vector see_stateTrajGlbY,
+ const std::vector see_stateTrajGlbZ,
+ const std::vector see_stateTrajGlbPx,
+ const std::vector see_stateTrajGlbPy,
+ const std::vector see_stateTrajGlbPz,
+ const std::vector see_q,
+ const std::vector> see_hitIdx,
+ const std::vector ph2_detId,
+ const std::vector ph2_x,
+ const std::vector ph2_y,
+ const std::vector ph2_z);
+ std::vector> hits() { return out_tc_hitIdxs_; }
+ std::vector len() { return out_tc_len_; }
+ std::vector seedIdx() { return out_tc_seedIdx_; }
+ std::vector trackCandidateType() { return out_tc_trackCandidateType_; }
+
+ private:
+ void prepareInput(const std::vector see_px,
+ const std::vector see_py,
+ const std::vector see_pz,
+ const std::vector see_dxy,
+ const std::vector see_dz,
+ const std::vector see_ptErr,
+ const std::vector see_etaErr,
+ const std::vector see_stateTrajGlbX,
+ const std::vector see_stateTrajGlbY,
+ const std::vector see_stateTrajGlbZ,
+ const std::vector see_stateTrajGlbPx,
+ const std::vector see_stateTrajGlbPy,
+ const std::vector see_stateTrajGlbPz,
+ const std::vector see_q,
+ const std::vector> see_hitIdx,
+ const std::vector ph2_detId,
+ const std::vector ph2_x,
+ const std::vector ph2_y,
+ const std::vector ph2_z);
+
+ void getOutput(SDL::Event& event);
+ std::vector getHitIdxs(const short trackCandidateType,
+ const unsigned int TCIdx,
+ const unsigned int* TCHitIndices,
+ const unsigned int* hitIndices);
+
+ // Input and output vectors
+ std::vector in_trkX_;
+ std::vector in_trkY_;
+ std::vector in_trkZ_;
+ std::vector in_hitId_;
+ std::vector in_hitIdxs_;
+ std::vector in_hitIndices_vec0_;
+ std::vector in_hitIndices_vec1_;
+ std::vector in_hitIndices_vec2_;
+ std::vector in_hitIndices_vec3_;
+ std::vector in_deltaPhi_vec_;
+ std::vector in_ptIn_vec_;
+ std::vector in_ptErr_vec_;
+ std::vector in_px_vec_;
+ std::vector in_py_vec_;
+ std::vector in_pz_vec_;
+ std::vector in_eta_vec_;
+ std::vector in_etaErr_vec_;
+ std::vector in_phi_vec_;
+ std::vector in_charge_vec_;
+ std::vector in_seedIdx_vec_;
+ std::vector in_superbin_vec_;
+ std::vector in_pixelType_vec_;
+ std::vector in_isQuad_vec_;
+ std::vector> out_tc_hitIdxs_;
+ std::vector out_tc_len_;
+ std::vector out_tc_seedIdx_;
+ std::vector out_tc_trackCandidateType_;
+ };
+
+} // namespace SDL
+
+#endif
diff --git a/RecoTracker/LSTCore/interface/alpaka/LSTESData.h b/RecoTracker/LSTCore/interface/alpaka/LSTESData.h
new file mode 100644
index 0000000000000..29ae19a5484e8
--- /dev/null
+++ b/RecoTracker/LSTCore/interface/alpaka/LSTESData.h
@@ -0,0 +1,94 @@
+#ifndef LSTESData_H
+#define LSTESData_H
+
+#ifdef LST_IS_CMSSW_PACKAGE
+#include "RecoTracker/LSTCore/interface/alpaka/Constants.h"
+#else
+#include "Constants.h"
+#endif
+
+#include "HeterogeneousCore/AlpakaInterface/interface/CopyToDevice.h"
+
+#include
+#include
+
+namespace SDL {
+
+ struct pixelMap;
+
+ template
+ class TiltedGeometry;
+
+ template
+ class ModuleConnectionMap;
+ using MapPLStoLayer = std::array, 4>, 3>;
+
+ template
+ struct modulesBuffer;
+
+ template
+ class EndcapGeometryHost;
+
+ template
+ class EndcapGeometry;
+
+ template
+ struct LSTESHostData;
+
+ // FIXME: This shouldn't be a templated struct
+ template <>
+ struct LSTESHostData {
+ std::shared_ptr mapPLStoLayer;
+ std::shared_ptr> endcapGeometry;
+ std::shared_ptr> tiltedGeometry;
+ std::shared_ptr> moduleConnectionMap;
+
+ LSTESHostData(std::shared_ptr mapPLStoLayerIn,
+ std::shared_ptr> endcapGeometryIn,
+ std::shared_ptr> tiltedGeometryIn,
+ std::shared_ptr> moduleConnectionMapIn)
+ : mapPLStoLayer(mapPLStoLayerIn),
+ endcapGeometry(endcapGeometryIn),
+ tiltedGeometry(tiltedGeometryIn),
+ moduleConnectionMap(moduleConnectionMapIn) {}
+ };
+
+ template
+ struct LSTESDeviceData {
+ uint16_t nModules;
+ uint16_t nLowerModules;
+ unsigned int nPixels;
+ std::shared_ptr> modulesBuffers;
+ std::shared_ptr> endcapGeometry;
+ std::shared_ptr pixelMapping;
+
+ LSTESDeviceData(uint16_t nModulesIn,
+ uint16_t nLowerModulesIn,
+ unsigned int nPixelsIn,
+ std::shared_ptr> modulesBuffersIn,
+ std::shared_ptr> endcapGeometryIn,
+ std::shared_ptr pixelMappingIn)
+ : nModules(nModulesIn),
+ nLowerModules(nLowerModulesIn),
+ nPixels(nPixelsIn),
+ modulesBuffers(modulesBuffersIn),
+ endcapGeometry(endcapGeometryIn),
+ pixelMapping(pixelMappingIn) {}
+ };
+
+ std::unique_ptr> loadAndFillESHost();
+ std::unique_ptr> loadAndFillESDevice(SDL::QueueAcc& queue, const LSTESHostData* hostData);
+
+} // namespace SDL
+
+namespace cms::alpakatools {
+ template <>
+ struct CopyToDevice> {
+ template
+ static auto copyAsync(TQueue& queue, SDL::LSTESHostData const& hostData) {
+ return std::make_unique>(hostData);
+ }
+ };
+} // namespace cms::alpakatools
+
+#endif
diff --git a/RecoTracker/LSTCore/interface/alpaka/Module.h b/RecoTracker/LSTCore/interface/alpaka/Module.h
new file mode 100644
index 0000000000000..0a269eaa6e16a
--- /dev/null
+++ b/RecoTracker/LSTCore/interface/alpaka/Module.h
@@ -0,0 +1,372 @@
+#ifndef Module_cuh
+#define Module_cuh
+
+#include
+
+#ifdef LST_IS_CMSSW_PACKAGE
+#include "RecoTracker/LSTCore/interface/alpaka/Constants.h"
+#else
+#include "Constants.h"
+#endif
+
+namespace SDL {
+ enum SubDet { InnerPixel = 0, Barrel = 5, Endcap = 4 };
+
+ enum Side { NegZ = 1, PosZ = 2, Center = 3 };
+
+ enum ModuleType { PS, TwoS, PixelModule };
+
+ enum ModuleLayerType { Pixel, Strip, InnerPixelLayer };
+
+ struct objectRanges {
+ int* hitRanges;
+ int* hitRangesLower;
+ int* hitRangesUpper;
+ int8_t* hitRangesnLower;
+ int8_t* hitRangesnUpper;
+ int* mdRanges;
+ int* segmentRanges;
+ int* trackletRanges;
+ int* tripletRanges;
+ int* trackCandidateRanges;
+ // Others will be added later
+ int* quintupletRanges;
+
+ // This number is just nEligibleModules - 1, but still we want this to be independent of the TC kernel
+ uint16_t* nEligibleT5Modules;
+ // Will be allocated in createQuintuplets kernel!
+ uint16_t* indicesOfEligibleT5Modules;
+ // To store different starting points for variable occupancy stuff
+ int* quintupletModuleIndices;
+ int* quintupletModuleOccupancy;
+ int* miniDoubletModuleIndices;
+ int* miniDoubletModuleOccupancy;
+ int* segmentModuleIndices;
+ int* segmentModuleOccupancy;
+ int* tripletModuleIndices;
+ int* tripletModuleOccupancy;
+
+ unsigned int* device_nTotalMDs;
+ unsigned int* device_nTotalSegs;
+ unsigned int* device_nTotalTrips;
+ unsigned int* device_nTotalQuints;
+
+ template
+ void setData(TBuff& objectRangesbuf) {
+ hitRanges = alpaka::getPtrNative(objectRangesbuf.hitRanges_buf);
+ hitRangesLower = alpaka::getPtrNative(objectRangesbuf.hitRangesLower_buf);
+ hitRangesUpper = alpaka::getPtrNative(objectRangesbuf.hitRangesUpper_buf);
+ hitRangesnLower = alpaka::getPtrNative(objectRangesbuf.hitRangesnLower_buf);
+ hitRangesnUpper = alpaka::getPtrNative(objectRangesbuf.hitRangesnUpper_buf);
+ mdRanges = alpaka::getPtrNative(objectRangesbuf.mdRanges_buf);
+ segmentRanges = alpaka::getPtrNative(objectRangesbuf.segmentRanges_buf);
+ trackletRanges = alpaka::getPtrNative(objectRangesbuf.trackletRanges_buf);
+ tripletRanges = alpaka::getPtrNative(objectRangesbuf.tripletRanges_buf);
+ trackCandidateRanges = alpaka::getPtrNative(objectRangesbuf.trackCandidateRanges_buf);
+ quintupletRanges = alpaka::getPtrNative(objectRangesbuf.quintupletRanges_buf);
+
+ nEligibleT5Modules = alpaka::getPtrNative(objectRangesbuf.nEligibleT5Modules_buf);
+ indicesOfEligibleT5Modules = alpaka::getPtrNative(objectRangesbuf.indicesOfEligibleT5Modules_buf);
+
+ quintupletModuleIndices = alpaka::getPtrNative(objectRangesbuf.quintupletModuleIndices_buf);
+ quintupletModuleOccupancy = alpaka::getPtrNative(objectRangesbuf.quintupletModuleOccupancy_buf);
+ miniDoubletModuleIndices = alpaka::getPtrNative(objectRangesbuf.miniDoubletModuleIndices_buf);
+ miniDoubletModuleOccupancy = alpaka::getPtrNative(objectRangesbuf.miniDoubletModuleOccupancy_buf);
+ segmentModuleIndices = alpaka::getPtrNative(objectRangesbuf.segmentModuleIndices_buf);
+ segmentModuleOccupancy = alpaka::getPtrNative(objectRangesbuf.segmentModuleOccupancy_buf);
+ tripletModuleIndices = alpaka::getPtrNative(objectRangesbuf.tripletModuleIndices_buf);
+ tripletModuleOccupancy = alpaka::getPtrNative(objectRangesbuf.tripletModuleOccupancy_buf);
+
+ device_nTotalMDs = alpaka::getPtrNative(objectRangesbuf.device_nTotalMDs_buf);
+ device_nTotalSegs = alpaka::getPtrNative(objectRangesbuf.device_nTotalSegs_buf);
+ device_nTotalTrips = alpaka::getPtrNative(objectRangesbuf.device_nTotalTrips_buf);
+ device_nTotalQuints = alpaka::getPtrNative(objectRangesbuf.device_nTotalQuints_buf);
+ }
+ };
+
+ template
+ struct objectRangesBuffer : objectRanges {
+ Buf hitRanges_buf;
+ Buf hitRangesLower_buf;
+ Buf hitRangesUpper_buf;
+ Buf hitRangesnLower_buf;
+ Buf hitRangesnUpper_buf;
+ Buf mdRanges_buf;
+ Buf segmentRanges_buf;
+ Buf trackletRanges_buf;
+ Buf tripletRanges_buf;
+ Buf trackCandidateRanges_buf;
+ Buf quintupletRanges_buf;
+
+ Buf nEligibleT5Modules_buf;
+ Buf indicesOfEligibleT5Modules_buf;
+
+ Buf quintupletModuleIndices_buf;
+ Buf quintupletModuleOccupancy_buf;
+ Buf miniDoubletModuleIndices_buf;
+ Buf miniDoubletModuleOccupancy_buf;
+ Buf segmentModuleIndices_buf;
+ Buf segmentModuleOccupancy_buf;
+ Buf tripletModuleIndices_buf;
+ Buf tripletModuleOccupancy_buf;
+
+ Buf device_nTotalMDs_buf;
+ Buf device_nTotalSegs_buf;
+ Buf device_nTotalTrips_buf;
+ Buf device_nTotalQuints_buf;
+
+ template
+ objectRangesBuffer(unsigned int nMod, unsigned int nLowerMod, TDevAcc const& devAccIn, TQueue& queue)
+ : hitRanges_buf(allocBufWrapper(devAccIn, nMod * 2, queue)),
+ hitRangesLower_buf(allocBufWrapper(devAccIn, nMod, queue)),
+ hitRangesUpper_buf(allocBufWrapper(devAccIn, nMod, queue)),
+ hitRangesnLower_buf(allocBufWrapper(devAccIn, nMod, queue)),
+ hitRangesnUpper_buf(allocBufWrapper(devAccIn, nMod, queue)),
+ mdRanges_buf(allocBufWrapper(devAccIn, nMod * 2, queue)),
+ segmentRanges_buf(allocBufWrapper(devAccIn, nMod * 2, queue)),
+ trackletRanges_buf(allocBufWrapper(devAccIn, nMod * 2, queue)),
+ tripletRanges_buf(allocBufWrapper(devAccIn, nMod * 2, queue)),
+ trackCandidateRanges_buf(allocBufWrapper(devAccIn, nMod * 2, queue)),
+ quintupletRanges_buf(allocBufWrapper(devAccIn, nMod * 2, queue)),
+ nEligibleT5Modules_buf(allocBufWrapper(devAccIn, 1, queue)),
+ indicesOfEligibleT5Modules_buf(allocBufWrapper(devAccIn, nLowerMod, queue)),
+ quintupletModuleIndices_buf(allocBufWrapper(devAccIn, nLowerMod, queue)),
+ quintupletModuleOccupancy_buf(allocBufWrapper(devAccIn, nLowerMod, queue)),
+ miniDoubletModuleIndices_buf(allocBufWrapper(devAccIn, nLowerMod + 1, queue)),
+ miniDoubletModuleOccupancy_buf(allocBufWrapper(devAccIn, nLowerMod + 1, queue)),
+ segmentModuleIndices_buf(allocBufWrapper(devAccIn, nLowerMod + 1, queue)),
+ segmentModuleOccupancy_buf(allocBufWrapper(devAccIn, nLowerMod + 1, queue)),
+ tripletModuleIndices_buf(allocBufWrapper(devAccIn, nLowerMod, queue)),
+ tripletModuleOccupancy_buf(allocBufWrapper(devAccIn, nLowerMod, queue)),
+ device_nTotalMDs_buf(allocBufWrapper(devAccIn, 1, queue)),
+ device_nTotalSegs_buf(allocBufWrapper(devAccIn, 1, queue)),
+ device_nTotalTrips_buf(allocBufWrapper(devAccIn, 1, queue)),
+ device_nTotalQuints_buf(allocBufWrapper(devAccIn, 1, queue)) {
+ alpaka::memset(queue, hitRanges_buf, 0xff);
+ alpaka::memset(queue, hitRangesLower_buf, 0xff);
+ alpaka::memset(queue, hitRangesUpper_buf, 0xff);
+ alpaka::memset(queue, hitRangesnLower_buf, 0xff);
+ alpaka::memset(queue, hitRangesnUpper_buf, 0xff);
+ alpaka::memset(queue, mdRanges_buf, 0xff);
+ alpaka::memset(queue, segmentRanges_buf, 0xff);
+ alpaka::memset(queue, trackletRanges_buf, 0xff);
+ alpaka::memset(queue, tripletRanges_buf, 0xff);
+ alpaka::memset(queue, trackCandidateRanges_buf, 0xff);
+ alpaka::memset(queue, quintupletRanges_buf, 0xff);
+ alpaka::memset(queue, quintupletModuleIndices_buf, 0xff);
+ alpaka::wait(queue);
+ }
+ };
+
+ struct modules {
+ const unsigned int* detIds;
+ const uint16_t* moduleMap;
+ const unsigned int* mapdetId;
+ const uint16_t* mapIdx;
+ const uint16_t* nConnectedModules;
+ const float* drdzs;
+ const float* dxdys;
+ const uint16_t* nModules;
+ const uint16_t* nLowerModules;
+ const uint16_t* partnerModuleIndices;
+
+ const short* layers;
+ const short* rings;
+ const short* modules;
+ const short* rods;
+ const short* subdets;
+ const short* sides;
+ const float* eta;
+ const float* r;
+ const bool* isInverted;
+ const bool* isLower;
+ const bool* isAnchor;
+ const ModuleType* moduleType;
+ const ModuleLayerType* moduleLayerType;
+ const int* sdlLayers;
+ const unsigned int* connectedPixels;
+
+ static bool parseIsInverted(short subdet, short side, short module, short layer) {
+ if (subdet == Endcap) {
+ if (side == NegZ) {
+ return module % 2 == 1;
+ } else if (side == PosZ) {
+ return module % 2 == 0;
+ } else {
+ return false;
+ }
+ } else if (subdet == Barrel) {
+ if (side == Center) {
+ if (layer <= 3) {
+ return module % 2 == 1;
+ } else if (layer >= 4) {
+ return module % 2 == 0;
+ } else {
+ return false;
+ }
+ } else if (side == NegZ or side == PosZ) {
+ if (layer <= 2) {
+ return module % 2 == 1;
+ } else if (layer == 3) {
+ return module % 2 == 0;
+ } else {
+ return false;
+ }
+ } else {
+ return false;
+ }
+ } else {
+ return false;
+ }
+ };
+
+ static bool parseIsLower(bool isInvertedx, unsigned int detId) {
+ return (isInvertedx) ? !(detId & 1) : (detId & 1);
+ };
+
+ static unsigned int parsePartnerModuleId(unsigned int detId, bool isLowerx, bool isInvertedx) {
+ return isLowerx ? (isInvertedx ? detId - 1 : detId + 1) : (isInvertedx ? detId + 1 : detId - 1);
+ };
+
+ template
+ void setData(const TBuff& modulesbuf) {
+ detIds = alpaka::getPtrNative(modulesbuf.detIds_buf);
+ moduleMap = alpaka::getPtrNative(modulesbuf.moduleMap_buf);
+ mapdetId = alpaka::getPtrNative(modulesbuf.mapdetId_buf);
+ mapIdx = alpaka::getPtrNative(modulesbuf.mapIdx_buf);
+ nConnectedModules = alpaka::getPtrNative(modulesbuf.nConnectedModules_buf);
+ drdzs = alpaka::getPtrNative(modulesbuf.drdzs_buf);
+ dxdys = alpaka::getPtrNative(modulesbuf.dxdys_buf);
+ nModules = alpaka::getPtrNative(modulesbuf.nModules_buf);
+ nLowerModules = alpaka::getPtrNative(modulesbuf.nLowerModules_buf);
+ partnerModuleIndices = alpaka::getPtrNative(modulesbuf.partnerModuleIndices_buf);
+
+ layers = alpaka::getPtrNative(modulesbuf.layers_buf);
+ rings = alpaka::getPtrNative(modulesbuf.rings_buf);
+ modules = alpaka::getPtrNative(modulesbuf.modules_buf);
+ rods = alpaka::getPtrNative(modulesbuf.rods_buf);
+ subdets = alpaka::getPtrNative(modulesbuf.subdets_buf);
+ sides = alpaka::getPtrNative(modulesbuf.sides_buf);
+ eta = alpaka::getPtrNative(modulesbuf.eta_buf);
+ r = alpaka::getPtrNative(modulesbuf.r_buf);
+ isInverted = alpaka::getPtrNative(modulesbuf.isInverted_buf);
+ isLower = alpaka::getPtrNative(modulesbuf.isLower_buf);
+ isAnchor = alpaka::getPtrNative(modulesbuf.isAnchor_buf);
+ moduleType = alpaka::getPtrNative(modulesbuf.moduleType_buf);
+ moduleLayerType = alpaka::getPtrNative(modulesbuf.moduleLayerType_buf);
+ sdlLayers = alpaka::getPtrNative(modulesbuf.sdlLayers_buf);
+ connectedPixels = alpaka::getPtrNative(modulesbuf.connectedPixels_buf);
+ }
+ };
+
+ template
+ struct modulesBuffer : modules {
+ Buf detIds_buf;
+ Buf moduleMap_buf;
+ Buf mapdetId_buf;
+ Buf mapIdx_buf;
+ Buf nConnectedModules_buf;
+ Buf drdzs_buf;
+ Buf dxdys_buf;
+ Buf nModules_buf;
+ Buf nLowerModules_buf;
+ Buf partnerModuleIndices_buf;
+
+ Buf layers_buf;
+ Buf rings_buf;
+ Buf modules_buf;
+ Buf rods_buf;
+ Buf subdets_buf;
+ Buf sides_buf;
+ Buf eta_buf;
+ Buf r_buf;
+ Buf isInverted_buf;
+ Buf isLower_buf;
+ Buf isAnchor_buf;
+ Buf moduleType_buf;
+ Buf moduleLayerType_buf;
+ Buf sdlLayers_buf;
+ Buf connectedPixels_buf;
+
+ modulesBuffer(TDev const& dev, unsigned int nMod, unsigned int nPixs)
+ : detIds_buf(allocBufWrapper(dev, nMod)),
+ moduleMap_buf(allocBufWrapper(dev, nMod * MAX_CONNECTED_MODULES)),
+ mapdetId_buf(allocBufWrapper(dev, nMod)),
+ mapIdx_buf(allocBufWrapper(dev, nMod)),
+ nConnectedModules_buf(allocBufWrapper(dev, nMod)),
+ drdzs_buf(allocBufWrapper(dev, nMod)),
+ dxdys_buf(allocBufWrapper(dev, nMod)),
+ nModules_buf(allocBufWrapper(dev, 1)),
+ nLowerModules_buf(allocBufWrapper(dev, 1)),
+ partnerModuleIndices_buf(allocBufWrapper(dev, nMod)),
+
+ layers_buf(allocBufWrapper(dev, nMod)),
+ rings_buf(allocBufWrapper(dev, nMod)),
+ modules_buf(allocBufWrapper(dev, nMod)),
+ rods_buf(allocBufWrapper(dev, nMod)),
+ subdets_buf(allocBufWrapper(dev, nMod)),
+ sides_buf(allocBufWrapper(dev, nMod)),
+ eta_buf(allocBufWrapper(dev, nMod)),
+ r_buf(allocBufWrapper(dev, nMod)),
+ isInverted_buf(allocBufWrapper(dev, nMod)),
+ isLower_buf(allocBufWrapper(dev, nMod)),
+ isAnchor_buf(allocBufWrapper(dev, nMod)),
+ moduleType_buf(allocBufWrapper(dev, nMod)),
+ moduleLayerType_buf(allocBufWrapper(dev, nMod)),
+ sdlLayers_buf(allocBufWrapper(dev, nMod)),
+ connectedPixels_buf(allocBufWrapper(dev, nPixs)) {
+ setData(*this);
+ }
+
+ template
+ inline void copyFromSrc(TQueue queue, const modulesBuffer& src, bool isFull = true) {
+ alpaka::memcpy(queue, detIds_buf, src.detIds_buf);
+ if (isFull) {
+ alpaka::memcpy(queue, moduleMap_buf, src.moduleMap_buf);
+ alpaka::memcpy(queue, mapdetId_buf, src.mapdetId_buf);
+ alpaka::memcpy(queue, mapIdx_buf, src.mapIdx_buf);
+ alpaka::memcpy(queue, nConnectedModules_buf, src.nConnectedModules_buf);
+ alpaka::memcpy(queue, drdzs_buf, src.drdzs_buf);
+ alpaka::memcpy(queue, dxdys_buf, src.dxdys_buf);
+ }
+ alpaka::memcpy(queue, nModules_buf, src.nModules_buf);
+ alpaka::memcpy(queue, nLowerModules_buf, src.nLowerModules_buf);
+ if (isFull) {
+ alpaka::memcpy(queue, partnerModuleIndices_buf, src.partnerModuleIndices_buf);
+ }
+
+ alpaka::memcpy(queue, layers_buf, src.layers_buf);
+ alpaka::memcpy(queue, rings_buf, src.rings_buf);
+ alpaka::memcpy(queue, modules_buf, src.modules_buf);
+ alpaka::memcpy(queue, rods_buf, src.rods_buf);
+ alpaka::memcpy(queue, subdets_buf, src.subdets_buf);
+ alpaka::memcpy(queue, sides_buf, src.sides_buf);
+ alpaka::memcpy(queue, eta_buf, src.eta_buf);
+ alpaka::memcpy(queue, r_buf, src.r_buf);
+ if (isFull) {
+ alpaka::memcpy(queue, isInverted_buf, src.isInverted_buf);
+ }
+ alpaka::memcpy(queue, isLower_buf, src.isLower_buf);
+ if (isFull) {
+ alpaka::memcpy(queue, isAnchor_buf, src.isAnchor_buf);
+ }
+ alpaka::memcpy(queue, moduleType_buf, src.moduleType_buf);
+ if (isFull) {
+ alpaka::memcpy(queue, moduleLayerType_buf, src.moduleLayerType_buf);
+ alpaka::memcpy(queue, sdlLayers_buf, src.sdlLayers_buf);
+ alpaka::memcpy(queue, connectedPixels_buf, src.connectedPixels_buf);
+ }
+ alpaka::wait(queue);
+ }
+
+ template
+ modulesBuffer(TQueue queue, const modulesBuffer& src, unsigned int nMod, unsigned int nPixs)
+ : modulesBuffer(alpaka::getDev(queue), nMod, nPixs) {
+ copyFromSrc(queue, src);
+ }
+
+ inline SDL::modules const* data() const { return this; }
+ };
+
+} // namespace SDL
+#endif
diff --git a/RecoTracker/LSTCore/src/alpaka/EndcapGeometry.dev.cc b/RecoTracker/LSTCore/src/alpaka/EndcapGeometry.dev.cc
new file mode 100644
index 0000000000000..2b5be62ec94bc
--- /dev/null
+++ b/RecoTracker/LSTCore/src/alpaka/EndcapGeometry.dev.cc
@@ -0,0 +1,85 @@
+#include "EndcapGeometry.h"
+
+SDL::EndcapGeometry::EndcapGeometry(SDL::Dev const& devAccIn,
+ SDL::QueueAcc& queue,
+ SDL::EndcapGeometryHost const& endcapGeometryIn)
+ : geoMapDetId_buf(allocBufWrapper(devAccIn, endcapGeometryIn.centroid_phis_.size())),
+ geoMapPhi_buf(allocBufWrapper(devAccIn, endcapGeometryIn.centroid_phis_.size())) {
+ dxdy_slope_ = endcapGeometryIn.dxdy_slope_;
+ centroid_phis_ = endcapGeometryIn.centroid_phis_;
+ fillGeoMapArraysExplicit(queue);
+}
+
+void SDL::EndcapGeometryHost::load(std::string filename) {
+ dxdy_slope_.clear();
+ centroid_phis_.clear();
+
+ std::ifstream ifile(filename, std::ios::binary);
+ if (!ifile.is_open()) {
+ throw std::runtime_error("Unable to open file: " + filename);
+ }
+
+ while (!ifile.eof()) {
+ unsigned int detid;
+ float dxdy_slope, centroid_phi;
+
+ // Read the detid, dxdy_slope, and centroid_phi from binary file
+ ifile.read(reinterpret_cast(&detid), sizeof(detid));
+ ifile.read(reinterpret_cast(&dxdy_slope), sizeof(dxdy_slope));
+ ifile.read(reinterpret_cast(¢roid_phi), sizeof(centroid_phi));
+
+ if (ifile) {
+ dxdy_slope_[detid] = dxdy_slope;
+ centroid_phis_[detid] = centroid_phi;
+ } else {
+ // End of file or read failed
+ if (!ifile.eof()) {
+ throw std::runtime_error("Failed to read Endcap Geometry binary data.");
+ }
+ }
+ }
+}
+
+void SDL::EndcapGeometry::fillGeoMapArraysExplicit(SDL::QueueAcc& queue) {
+ unsigned int phi_size = centroid_phis_.size();
+
+ // Allocate buffers on host
+ SDL::DevHost const& devHost = cms::alpakatools::host();
+ auto mapPhi_host_buf = allocBufWrapper(devHost, phi_size);
+ auto mapDetId_host_buf = allocBufWrapper