Skip to content

Commit

Permalink
Merge pull request #89 from SegmentLinking/more_review_fixes
Browse files Browse the repository at this point in the history
More review fixes: smart pointers, cleaner namespaces, naming fixes
  • Loading branch information
slava77 authored Aug 30, 2024
2 parents f245be0 + f0e560b commit d2b9e89
Show file tree
Hide file tree
Showing 20 changed files with 1,410 additions and 1,555 deletions.
12 changes: 6 additions & 6 deletions RecoTracker/LST/plugins/alpaka/LSTModulesDevESProducer.cc
Original file line number Diff line number Diff line change
@@ -1,3 +1,7 @@
// LST includes
#include "RecoTracker/LSTCore/interface/Module.h"
#include "RecoTracker/LSTCore/interface/alpaka/LST.h"

#include "FWCore/ParameterSet/interface/ParameterSet.h"

#include "HeterogeneousCore/AlpakaCore/interface/alpaka/ESProducer.h"
Expand All @@ -7,10 +11,6 @@

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

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

namespace ALPAKA_ACCELERATOR_NAMESPACE {

class LSTModulesDevESProducer : public ESProducer {
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
Original file line number Diff line number Diff line change
@@ -1,5 +1,7 @@
#include <alpaka/alpaka.hpp>

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

#include "FWCore/MessageLogger/interface/MessageLogger.h"
#include "FWCore/ParameterSet/interface/ConfigurationDescriptions.h"
#include "FWCore/ParameterSet/interface/ParameterSet.h"
Expand All @@ -19,8 +21,6 @@

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

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

namespace ALPAKA_ACCELERATOR_NAMESPACE {

class LSTProducer : public stream::SynchronizingEDProducer<> {
Expand Down Expand Up @@ -86,7 +86,7 @@ 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_;

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

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;
namespace ALPAKA_ACCELERATOR_NAMESPACE::lst {

using namespace ::lst;

// 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 = 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);
}

// 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
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 ALPAKA_ACCELERATOR_NAMESPACE::lst
#endif
Loading

0 comments on commit d2b9e89

Please sign in to comment.