Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Moved LST and Event classes to ALPAKA_ACCELERATOR_NAMESPACE #84

Merged
merged 3 commits into from
Aug 17, 2024
Merged

Moved LST and Event classes to ALPAKA_ACCELERATOR_NAMESPACE #84

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
9f6d612
Moved LST and Event classes to ALPAKA_ACCELERATOR_NAMESPACE
ariostas Aug 16, 2024
1a5e2c4
Merge branch 'CMSSW_14_1_0_pre3_LST_X_LSTCore_realfiles_batch5' into …
ariostas Aug 17, 2024
79ea879
Moved all alpaka code to ALPAKA_ACCELERATOR_NAMESPACE
ariostas Aug 17, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
6 changes: 3 additions & 3 deletions RecoTracker/LST/plugins/alpaka/LSTModulesDevESProducer.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9,7 +9,7 @@

// LST includes
#include "RecoTracker/LSTCore/interface/Module.h"
#include "RecoTracker/LSTCore/interface/LST.h"
#include "RecoTracker/LSTCore/interface/alpaka/LST.h"

namespace ALPAKA_ACCELERATOR_NAMESPACE {

Expand All @@ -22,8 +22,8 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
descriptions.addWithDefaultLabel(desc);
}

std::unique_ptr<lst::LSTESData<DevHost>> produce(TrackerRecoGeometryRecord const& iRecord) {
return lst::loadAndFillESHost();
std::unique_ptr<::lst::LSTESData<DevHost>> produce(TrackerRecoGeometryRecord const& iRecord) {
return ::lst::loadAndFillESHost();
}
};

Expand Down
6 changes: 3 additions & 3 deletions RecoTracker/LST/plugins/alpaka/LSTProducer.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -19,7 +19,7 @@

#include "RecoTracker/Record/interface/TrackerRecoGeometryRecord.h"

#include "RecoTracker/LSTCore/interface/LST.h"
#include "RecoTracker/LSTCore/interface/alpaka/LST.h"

namespace ALPAKA_ACCELERATOR_NAMESPACE {

Expand Down Expand Up @@ -87,11 +87,11 @@ namespace ALPAKA_ACCELERATOR_NAMESPACE {
private:
edm::EDGetTokenT<LSTPixelSeedInput> lstPixelSeedInputToken_;
edm::EDGetTokenT<LSTPhase2OTHitsInput> lstPhase2OTHitsInputToken_;
device::ESGetToken<lst::LSTESData<Device>, TrackerRecoGeometryRecord> lstESToken_;
device::ESGetToken<::lst::LSTESData<Device>, TrackerRecoGeometryRecord> lstESToken_;
const bool verbose_, nopLSDupClean_, tcpLSTriplets_;
edm::EDPutTokenT<LSTOutput> lstOutputToken_;

lst::LST<Acc3D> lst_;
lst::LST lst_;
};

} // namespace ALPAKA_ACCELERATOR_NAMESPACE
Expand Down
112 changes: 0 additions & 112 deletions RecoTracker/LSTCore/interface/LST.h
View file
Open in desktop

This file was deleted.

171 changes: 95 additions & 76 deletions RecoTracker/LSTCore/interface/alpaka/Constants.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9,99 +9,118 @@
#include <hip/hip_fp16.h>
#endif

namespace lst {

using namespace ALPAKA_ACCELERATOR_NAMESPACE;
namespace ALPAKA_ACCELERATOR_NAMESPACE {
namespace lst {

// Re-export some useful things from the main namespace
using ::lst::allocBufWrapper;
using ::lst::Buf;
using ::lst::max_blocks;
using ::lst::max_connected_modules;
using ::lst::n_max_nonpixel_track_candidates;
using ::lst::n_max_pixel_md_per_modules;
using ::lst::n_max_pixel_quintuplets;
using ::lst::n_max_pixel_segments_per_module;
using ::lst::n_max_pixel_track_candidates;
using ::lst::n_max_pixel_triplets;
using ::lst::Params_LS;
using ::lst::Params_pLS;
using ::lst::Params_pT3;
using ::lst::Params_pT5;
using ::lst::Params_T3;
using ::lst::Params_T5;
using ::lst::size_superbins;

// Half precision wrapper functions.
#if defined(FP16_Base)
#define __F2H __float2half
#define __H2F __half2float
typedef __half float FPX;
typedef __half float FPX;
#else
#define __F2H
#define __H2F
typedef float FPX;
typedef float FPX;
#endif

Vec3D constexpr elementsPerThread(Vec3D::all(static_cast<Idx>(1)));
Vec3D constexpr elementsPerThread(Vec3D::all(static_cast<Idx>(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;
};
struct uint4 {
unsigned int x;
unsigned int y;
unsigned int z;
unsigned int w;
};
#endif

// Adjust grid and block sizes based on backend configuration
template <typename Vec, typename TAcc = ALPAKA_ACCELERATOR_NAMESPACE::Acc<typename Vec::Dim>>
ALPAKA_FN_HOST ALPAKA_FN_INLINE WorkDiv<typename Vec::Dim> createWorkDiv(const Vec& blocksPerGrid,
const Vec& threadsPerBlock,
const Vec& elementsPerThreadArg) {
Vec adjustedBlocks = blocksPerGrid;
Vec adjustedThreads = threadsPerBlock;

// special overrides for CPU/host cases
if constexpr (std::is_same_v<Platform, alpaka::PlatformCpu>) {
adjustedBlocks = Vec::all(static_cast<Idx>(1));

if constexpr (alpaka::accMatchesTags<TAcc, alpaka::TagCpuSerial>) {
// Serial execution, set threads to 1 as well
adjustedThreads = Vec::all(static_cast<Idx>(1)); // probably redundant
// Adjust grid and block sizes based on backend configuration
template <typename Vec, typename TAcc = Acc<typename Vec::Dim>>
ALPAKA_FN_HOST ALPAKA_FN_INLINE WorkDiv<typename Vec::Dim> createWorkDiv(const Vec& blocksPerGrid,
const Vec& threadsPerBlock,
const Vec& elementsPerThreadArg) {
Vec adjustedBlocks = blocksPerGrid;
Vec adjustedThreads = threadsPerBlock;

// special overrides for CPU/host cases
if constexpr (std::is_same_v<Platform, alpaka::PlatformCpu>) {
adjustedBlocks = Vec::all(static_cast<Idx>(1));

if constexpr (alpaka::accMatchesTags<TAcc, alpaka::TagCpuSerial>) {
// Serial execution, set threads to 1 as well
adjustedThreads = Vec::all(static_cast<Idx>(1)); // probably redundant
}
}

return WorkDiv<typename Vec::Dim>(adjustedBlocks, adjustedThreads, elementsPerThreadArg);
}

return WorkDiv<typename Vec::Dim>(adjustedBlocks, adjustedThreads, elementsPerThreadArg);
}

// The constants below are usually used in functions like alpaka::math::min(),
// expecting a reference (T const&) in the arguments. Hence,
// ALPAKA_STATIC_ACC_MEM_GLOBAL needs to be used in addition to constexpr.

// 15 MeV constant from the approximate Bethe-Bloch formula
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kMulsInGeV = 0.015;
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kMiniMulsPtScaleBarrel[6] = {
0.0052, 0.0038, 0.0034, 0.0034, 0.0032, 0.0034};
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kMiniMulsPtScaleEndcap[5] = {0.006, 0.006, 0.006, 0.006, 0.006};
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kMiniRminMeanBarrel[6] = {
25.007152356, 37.2186993757, 52.3104270826, 68.6658656666, 85.9770373007, 108.301772384};
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kMiniRminMeanEndcap[5] = {
130.992832231, 154.813883559, 185.352604327, 221.635123002, 265.022076742};
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float k2Rinv1GeVf = (2.99792458e-3 * 3.8) / 2;
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kR1GeVf = 1. / (2.99792458e-3 * 3.8);
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kSinAlphaMax = 0.95;
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float ptCut = PT_CUT;
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kDeltaZLum = 15.0;
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kPixelPSZpitch = 0.15;
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kStripPSZpitch = 2.4;
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kStrip2SZpitch = 5.0;
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWidth2S = 0.009;
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWidthPS = 0.01;
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kPt_betaMax = 7.0;
// Since C++ can't represent infinity, lst_INF = 123456789 was used to represent infinity in the data table
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float lst_INF = 123456789.0;

namespace t5dnn {

// Working points matching LST fake rate (43.9%) or signal acceptance (82.0%)
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kLSTWp1 = 0.3418833f; // 94.0% TPR, 43.9% FPR
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kLSTWp2 = 0.6177366f; // 82.0% TPR, 20.0% FPR
// Other working points
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWp70 = 0.7776195f; // 70.0% TPR, 10.0% FPR
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWp75 = 0.7181118f; // 75.0% TPR, 13.5% FPR
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWp80 = 0.6492643f; // 80.0% TPR, 17.9% FPR
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWp85 = 0.5655319f; // 85.0% TPR, 23.8% FPR
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWp90 = 0.4592205f; // 90.0% TPR, 32.6% FPR
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWp95 = 0.3073708f; // 95.0% TPR, 47.7% FPR
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWp97p5 = 0.2001348f; // 97.5% TPR, 61.2% FPR
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWp99 = 0.1120605f; // 99.0% TPR, 75.9% FPR
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWp99p9 = 0.0218196f; // 99.9% TPR, 95.4% FPR

} // namespace t5dnn

} //namespace lst
// The constants below are usually used in functions like alpaka::math::min(),
// expecting a reference (T const&) in the arguments. Hence,
// ALPAKA_STATIC_ACC_MEM_GLOBAL needs to be used in addition to constexpr.

// 15 MeV constant from the approximate Bethe-Bloch formula
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kMulsInGeV = 0.015;
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kMiniMulsPtScaleBarrel[6] = {
0.0052, 0.0038, 0.0034, 0.0034, 0.0032, 0.0034};
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kMiniMulsPtScaleEndcap[5] = {0.006, 0.006, 0.006, 0.006, 0.006};
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kMiniRminMeanBarrel[6] = {
25.007152356, 37.2186993757, 52.3104270826, 68.6658656666, 85.9770373007, 108.301772384};
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kMiniRminMeanEndcap[5] = {
130.992832231, 154.813883559, 185.352604327, 221.635123002, 265.022076742};
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float k2Rinv1GeVf = (2.99792458e-3 * 3.8) / 2;
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kR1GeVf = 1. / (2.99792458e-3 * 3.8);
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kSinAlphaMax = 0.95;
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float ptCut = PT_CUT;
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kDeltaZLum = 15.0;
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kPixelPSZpitch = 0.15;
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kStripPSZpitch = 2.4;
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kStrip2SZpitch = 5.0;
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWidth2S = 0.009;
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWidthPS = 0.01;
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kPt_betaMax = 7.0;
// Since C++ can't represent infinity, lst_INF = 123456789 was used to represent infinity in the data table
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float lst_INF = 123456789.0;

namespace t5dnn {

// Working points matching LST fake rate (43.9%) or signal acceptance (82.0%)
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kLSTWp1 = 0.3418833f; // 94.0% TPR, 43.9% FPR
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kLSTWp2 = 0.6177366f; // 82.0% TPR, 20.0% FPR
// Other working points
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWp70 = 0.7776195f; // 70.0% TPR, 10.0% FPR
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWp75 = 0.7181118f; // 75.0% TPR, 13.5% FPR
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWp80 = 0.6492643f; // 80.0% TPR, 17.9% FPR
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWp85 = 0.5655319f; // 85.0% TPR, 23.8% FPR
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWp90 = 0.4592205f; // 90.0% TPR, 32.6% FPR
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWp95 = 0.3073708f; // 95.0% TPR, 47.7% FPR
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWp97p5 = 0.2001348f; // 97.5% TPR, 61.2% FPR
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWp99 = 0.1120605f; // 99.0% TPR, 75.9% FPR
ALPAKA_STATIC_ACC_MEM_GLOBAL constexpr float kWp99p9 = 0.0218196f; // 99.9% TPR, 95.4% FPR

} // namespace t5dnn

} //namespace lst
} //namespace ALPAKA_ACCELERATOR_NAMESPACE
#endif
Loading