Commit 9d954227 authored by Thomas Boettcher's avatar Thomas Boettcher
Browse files

Further simplify line creation and add documentation.

parent dabac3d4
Pipeline #1349080 passed with stages
in 20 minutes and 4 seconds
......@@ -458,6 +458,8 @@ include_directories(cuda/event_model/velo/include)
include_directories(cuda/event_model/UT/include)
include_directories(cuda/event_model/SciFi/include)
include_directories(cuda/event_model/common/include)
include_directories(cuda/vertex_fit/common/include)
include_directories(cuda/selections/Hlt1/include)
include_directories(cuda/raw_banks/include)
include_directories(checker/tracking/include)
include_directories(checker/pv/include)
......
......@@ -14,6 +14,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/vertex_fit/common/include)
include_directories(${CMAKE_SOURCE_DIR}/cuda/raw_banks/include)
include_directories(${CMAKE_SOURCE_DIR}/cuda/selections/Hlt1/include)
include_directories(${CMAKE_SOURCE_DIR}/checker/tracking/include)
......
......@@ -4,6 +4,7 @@
#include <CheckerTypes.h>
#include <CheckerInvoker.h>
#include "RawBanksDefinitions.cuh"
#include "LineInfo.cuh"
void checkHlt1Rate(
const bool* decisions,
......
......@@ -12,6 +12,7 @@
#include "ParKalmanDefinitions.cuh"
#include "VertexDefinitions.cuh"
#include "RawBanksDefinitions.cuh"
#include "LineInfo.cuh"
#include "ROOTHeaders.h"
#include <algorithm>
......
......@@ -549,3 +549,51 @@ Allen has finished executing. In principle, this could be performed on a regular
ideally would require monitoring threads to be paused for thread safety.
Histograms are currently written to `monitoringHists.root`.
Adding selection lines
======================
This will cover how to add trigger lines to Allen that select events
based on reconstructed trigger candidates. Special lines (e.g. NoBias
or pass-through lines) should be handled on a case-by-case basis.
Writing the selection
---------------------
Trigger selections should be `__device__` functions that take either a
`const ParKalmanFilter::FittedTrack&` or a `const
VertexFit::TrackMVAVertex&` as an argument and return a `bool`. For
example, a line selecting high-pT tracks might look like:
'''
__device__ bool HighPtTrack(const ParKalmanFilter::FittedTrack& track)
{
return track.pt() > 10.0 / Gaudi::Units::GeV
}
'''
The header file for the selection should be placed in
`cuda/selections/Hlt1/include` and the implementation should be placed
in `cuda/selections/Hlt1/src`.
Adding the line to the Allen sequence
-------------------------------------
Bookkeeping information for the Hlt1 lines is found in
`cuda/selections/Hlt1/include/LineInfo.cuh`. In order for a line
to run, it must be added to `Hlt1::Hlt1Lines` and a name must be added
to `Hlt1::Hlt1LineNames`. This will ensure that space is allocated to
store the selection decision for each candidate.
Special lines are listed first, followed by 1-track lines, then 2-,
3-, and finally 4-track lines. The new line should be added to the
appropriate place in the list. In addition, the number of lines of
that type should be incremented by 1. For example, the above
`HighPtTrack` line should be added after `// Begin 1-track lines.` and
before `Begin 2-track lines.` The line name should be added at the
same position in `Hlt1::Hlt1LineNames`.
Finally, add the selection function to the relevant array of pointers
to selections (e.g. `Hlt1::OneTrackSelections` or
`Hlt1::TwoTrackSelections`). These must be in the same order as in
`Hlt1::Hlt1LineNames` and `Hlt1::Hlt1Lines`.
\ No newline at end of file
......@@ -5,8 +5,11 @@ include_directories(${CMAKE_SOURCE_DIR}/cuda/event_model/SciFi/include)
include_directories(${CMAKE_SOURCE_DIR}/cuda/event_model/common/include)
include_directories(${CMAKE_SOURCE_DIR}/cuda/SciFi/common/include)
include_directories(${CMAKE_SOURCE_DIR}/cuda/kalman/ParKalman/include)
include_directories(${CMAKE_SOURCE_DIR}/cuda/selections/Hlt1/include)
include_directories(${CMAKE_SOURCE_DIR}/stream/setup/include)
include_directories(${CMAKE_SOURCE_DIR}/stream/gear/include)
include_directories(${CMAKE_SOURCE_DIR}/cuda/PV/beamlinePV/include)
include_directories(${CMAKE_SOURCE_DIR}/cuda/PV/common/include)
file(GLOB raw_banks_src "src/*cu")
......
......@@ -2,6 +2,7 @@
#include "HltDecReport.cuh"
#include "RawBanksDefinitions.cuh"
#include "LineInfo.cuh"
#include "Handler.cuh"
#include "ArgumentsSelections.cuh"
......
......@@ -8,39 +8,5 @@ namespace Hlt1 {
// Task ID.
const unsigned int taskID = 1;
// Number of special lines
const unsigned int nSpecialLines = 1;
const unsigned int nOneTrackLines = 2;
const unsigned int nTwoTrackLines = 4;
const unsigned int nThreeTrackLines = 0;
const unsigned int nFourTrackLines = 0;
const unsigned int startOneTrackLines = nSpecialLines;
const unsigned int startTwoTrackLines = startOneTrackLines + nOneTrackLines;
const unsigned int startThreeTrackLines = startTwoTrackLines + nTwoTrackLines;
const unsigned int startFourTrackLines = startThreeTrackLines + nThreeTrackLines;
// Hlt1 lines.
enum Hlt1Lines {
PassThrough,
OneTrackMVA,
SingleMuon,
TwoTrackMVA,
DisplacedDiMuon,
HighMassDiMuon,
SoftDiMuon,
End
};
// Hlt1 line names.
static const std::string Hlt1LineNames[] = {
"PassThrough",
"OneTrackMVA",
"SingleMuon",
"TwoTrackMVA",
"DisplacedDiMuon",
"HighMassDiMuon",
"SoftDiMuon"
};
}
#pragma once
#include "TrackMVALines.cuh"
#include "MuonLines.cuh"
#include <string>
namespace Hlt1 {
// Number of special lines
const unsigned int nSpecialLines = 1;
const unsigned int nOneTrackLines = 2;
const unsigned int nTwoTrackLines = 4;
const unsigned int nThreeTrackLines = 0;
const unsigned int nFourTrackLines = 0;
const unsigned int startOneTrackLines = nSpecialLines;
const unsigned int startTwoTrackLines = startOneTrackLines + nOneTrackLines;
const unsigned int startThreeTrackLines = startTwoTrackLines + nTwoTrackLines;
const unsigned int startFourTrackLines = startThreeTrackLines + nThreeTrackLines;
// Hlt1 lines.
enum Hlt1Lines {
PassThrough,
// Begin 1-track lines.
OneTrackMVA,
SingleMuon,
// Begin 2-track lines.
TwoTrackMVA,
DisplacedDiMuon,
HighMassDiMuon,
SoftDiMuon,
// Begin 3-track lines.
// Begin 4-track lines.
End
};
// Hlt1 line names.
static const std::string Hlt1LineNames[] = {
"PassThrough",
// Begin 1-track lines.
"OneTrackMVA",
"SingleMuon",
// Begin 2-track lines.
"TwoTrackMVA",
"DisplacedDiMuon",
"HighMassDiMuon",
"SoftDiMuon"
// Begin 3-track lines.
// Begin 4-track lines.
};
// See if this works...
static __device__ bool (*OneTrackSelections[])(const ParKalmanFilter::FittedTrack&) = {
TrackMVALines::OneTrackMVA,
MuonLines::SingleMuon
};
static __device__ bool (*TwoTrackSelections[])(const VertexFit::TrackMVAVertex&) = {
TrackMVALines::TwoTrackMVA,
MuonLines::DisplacedDiMuon,
MuonLines::HighMassDiMuon,
MuonLines::DiMuonSoft
};
}
\ No newline at end of file
#include "RunHlt1.cuh"
#include "RawBanksDefinitions.cuh"
#include "TrackMVALines.cuh"
#include "MuonLines.cuh"
#include "LineInfo.cuh"
#include "Handler.cuh"
#include "ArgumentsSciFi.cuh"
......@@ -20,9 +20,6 @@ __global__ void run_hlt1(
{
const uint number_of_events = gridDim.x;
const uint event_number = blockIdx.x;
const uint total_tracks = dev_atomics_scifi[2 * number_of_events];
const uint total_svs = dev_sv_atomics[2 * number_of_events];
const uint* dev_sel_result_offsets = dev_sel_results_atomics + Hlt1::Hlt1Lines::End;
......@@ -30,38 +27,24 @@ __global__ void run_hlt1(
const auto* event_tracks_offsets = dev_atomics_scifi + number_of_events;
const auto* event_svs_offsets = dev_sv_atomics + number_of_events;
const ParKalmanFilter::FittedTrack* event_tracks = dev_kf_tracks + event_tracks_offsets[event_number];
bool* event_one_track_results = dev_sel_results +
dev_sel_result_offsets[Hlt1::Hlt1Lines::OneTrackMVA] + event_tracks_offsets[event_number];
bool* event_single_muon_results = dev_sel_results +
dev_sel_result_offsets[Hlt1::Hlt1Lines::SingleMuon] + event_tracks_offsets[event_number];
const auto n_tracks_event = dev_atomics_scifi[event_number];
// Vertices.
const VertexFit::TrackMVAVertex* event_vertices = dev_consolidated_svs + event_svs_offsets[event_number];
bool* event_two_track_results = dev_sel_results +
dev_sel_result_offsets[Hlt1::Hlt1Lines::TwoTrackMVA] + event_svs_offsets[event_number];
bool* event_disp_dimuon_results = dev_sel_results +
dev_sel_result_offsets[Hlt1::Hlt1Lines::DisplacedDiMuon] + event_svs_offsets[event_number];
bool* event_high_mass_dimuon_results = dev_sel_results +
dev_sel_result_offsets[Hlt1::Hlt1Lines::HighMassDiMuon] + event_svs_offsets[event_number];
bool* event_dimuon_soft_results = dev_sel_results +
dev_sel_result_offsets[Hlt1::Hlt1Lines::SoftDiMuon] + event_svs_offsets[event_number];
const auto n_vertices_event = dev_sv_atomics[event_number];
LineHandler<ParKalmanFilter::FittedTrack> oneTrackHandler {TrackMVALines::OneTrackMVA};
LineHandler<ParKalmanFilter::FittedTrack> singleMuonHandler {MuonLines::SingleMuon};
LineHandler<VertexFit::TrackMVAVertex> twoTrackHandler {TrackMVALines::TwoTrackMVA};
LineHandler<VertexFit::TrackMVAVertex> dispDiMuonHandler {MuonLines::DisplacedDiMuon};
LineHandler<VertexFit::TrackMVAVertex> highMassDiMuonHandler {MuonLines::HighMassDiMuon};
LineHandler<VertexFit::TrackMVAVertex> diMuonSoftHandler {MuonLines::DiMuonSoft};
// One track lines.
oneTrackHandler(event_tracks, n_tracks_event, event_one_track_results);
singleMuonHandler(event_tracks, n_tracks_event, event_single_muon_results);
// Process 1-track lines.
for (uint i_line = Hlt1::startOneTrackLines; i_line < Hlt1::startTwoTrackLines; i_line++) {
bool* decs = dev_sel_results + dev_sel_result_offsets[i_line] + event_tracks_offsets[event_number];
LineHandler<ParKalmanFilter::FittedTrack> handler {Hlt1::OneTrackSelections[i_line - Hlt1::startOneTrackLines]};
handler(event_tracks, n_tracks_event, decs);
}
// Two track lines.
twoTrackHandler(event_vertices, n_vertices_event, event_two_track_results);
dispDiMuonHandler(event_vertices, n_vertices_event, event_disp_dimuon_results);
highMassDiMuonHandler(event_vertices, n_vertices_event, event_high_mass_dimuon_results);
diMuonSoftHandler(event_vertices, n_vertices_event, event_dimuon_soft_results);
// Process 2-track lines.
for (uint i_line = Hlt1::startTwoTrackLines; i_line < Hlt1::startThreeTrackLines; i_line++) {
bool* decs = dev_sel_results + dev_sel_result_offsets[i_line] + event_svs_offsets[event_number];
LineHandler<VertexFit::TrackMVAVertex> handler {Hlt1::TwoTrackSelections[i_line - Hlt1::startTwoTrackLines]};
handler(event_vertices, n_vertices_event, decs);
}
}
......@@ -15,6 +15,7 @@ include_directories(${CMAKE_SOURCE_DIR}/cuda/SciFi/common/include)
include_directories(${CMAKE_SOURCE_DIR}/cuda/UT/common/include)
include_directories(${CMAKE_SOURCE_DIR}/cuda/velo/common/include)
include_directories(${CMAKE_SOURCE_DIR}/cuda/vertex_fit/common/include)
include_directories(${CMAKE_SOURCE_DIR}/cuda/selections/Hlt1/include)
include_directories(${CMAKE_SOURCE_DIR}/stream/sequence/include)
include_directories(${ROOT_INCLUDE_DIRS})
......
......@@ -5,6 +5,7 @@
#include "HltDecReport.cuh"
#include "RawBanksDefinitions.cuh"
#include "LineInfo.cuh"
#ifdef WITH_ROOT
void RateMonitor::fill(uint i_buf, bool useWallTime)
......
......@@ -5,6 +5,7 @@
#include "HltDecReport.cuh"
#include "RawBanksDefinitions.cuh"
#include "LineInfo.cuh"
#include <cmath>
......
......@@ -15,6 +15,7 @@
#include <InputProvider.h>
#include <OutputHandler.h>
#include <RawBanksDefinitions.cuh>
#include <LineInfo.cuh>
bool OutputHandler::output_selected_events(
size_t const slice_index,
......
......@@ -3,6 +3,7 @@
#include "BeamlinePVConstants.cuh"
#include "LookingForwardConstants.cuh"
#include "RawBanksDefinitions.cuh"
#include "LineInfo.cuh"
void HostBuffers::reserve(const uint max_number_of_events, const bool do_check)
{
......
......@@ -3,6 +3,7 @@
#include "HltDecReport.cuh"
#include "RawBanksDefinitions.cuh"
#include "LineInfo.cuh"
#include "Logger.h"
......
#include "RunHlt1.cuh"
#include "SequenceVisitor.cuh"
#include "RawBanksDefinitions.cuh"
#include "LineInfo.cuh"
template<>
void SequenceVisitor::set_arguments_size<run_hlt1_t>(
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment