Skip to content
Snippets Groups Projects
Commit 0f39ca10 authored by Daniel Campora's avatar Daniel Campora
Browse files

Make all lines homogeneous.

parent 4786e703
No related branches found
No related tags found
1 merge request!846Delayed selections
Hide whitespace changes
Inline Side-by-side
Showing
with 59 additions and 115 deletions
configuration/python/AllenConf/hlt1_monitoring_lines.py
+ 2
− 0
View file @ 0f39ca10
......@@ -63,6 +63,8 @@ def make_odin_event_type_line(
return make_algorithm(
odin_event_type_line_t,
name=name_map[odin_event_type],
dev_odin_raw_input_t=odin["dev_odin_raw_input"],
dev_odin_raw_input_offsets_t=odin["dev_odin_raw_input_offsets"],
odin_event_type=odin_event_type,
host_number_of_events_t=number_of_events["host_number_of_events"],
pre_scaler_hash_string=pre_scaler_hash_string,
......
device/selections/line_types/include/EventLine.cuh
+ 9
− 4
View file @ 0f39ca10
......@@ -10,10 +10,15 @@
*/
template<typename Derived, typename Parameters>
struct EventLine : public Line<Derived, Parameters> {
/**
* @brief Set the tag to event_iteration_tag, to mark how to iterate events.
*/
using iteration_t = LineIteration::event_iteration_tag;
__device__ static unsigned offset(const Parameters&, const unsigned event_number)
{
return event_number;
}
__device__ static unsigned input_size(const Parameters&, const unsigned)
{
return 1;
}
/**
* @brief Decision size is the number of events.
......
device/selections/line_types/include/Line.cuh
+ 34
− 102
View file @ 0f39ca10
......@@ -23,14 +23,6 @@
template __device__ void process_line<LINE, PARAMETERS>(char*, unsigned, unsigned, unsigned, unsigned, unsigned); \
INSTANTIATE_ALGORITHM(LINE)
// "Enum of types" to determine dispatch to global_function
namespace LineIteration {
struct default_iteration_tag {
};
struct event_iteration_tag {
};
} // namespace LineIteration
// Type-erased line function type
using line_fn_t = void (*)(char*, unsigned, unsigned, unsigned, unsigned, unsigned);
......@@ -78,8 +70,6 @@ private:
uint32_t m_post_scaler_hash;
public:
using iteration_t = LineIteration::default_iteration_tag;
void init()
{
auto derived_instance = static_cast<const Derived*>(this);
......@@ -139,105 +129,47 @@ template<typename Derived, typename Parameters>
__device__ void
process_line(char* input, unsigned run_no, unsigned evt_hi, unsigned evt_lo, unsigned gps_hi, unsigned gps_lo)
{
if constexpr (!std::is_same_v<typename Derived::iteration_t, LineIteration::event_iteration_tag>) {
const auto& type_casted_input = *reinterpret_cast<const std::tuple<Parameters, size_t, unsigned, unsigned>*>(input);
const auto& parameters = std::get<0>(type_casted_input);
const auto number_of_events = std::get<2>(type_casted_input);
// Check if blockIdx.x (event_number) is in dev_event_list
unsigned mask = 0;
for (unsigned i = 0; i < (number_of_events + warp_size - 1) / warp_size; ++i) {
const auto index = i * warp_size + threadIdx.x;
mask |= __ballot_sync(0xFFFFFFFF, index < number_of_events ? threadIdx.x == parameters.dev_event_list[i] : false);
}
// Do initialization for all events, regardless of mask
// * Populate offsets in first block
if (blockIdx.x == 0) {
for (unsigned i = threadIdx.x * blockIdx.y + threadIdx.y; i < number_of_events; i += blockDim.x * blockDim.y) {
parameters.dev_decisions_offsets[i] = mask ? Derived::offset(parameters, i) : 0;
}
}
const auto& type_casted_input = *reinterpret_cast<const std::tuple<Parameters, size_t, unsigned, unsigned>*>(input);
const auto& parameters = std::get<0>(type_casted_input);
const auto event_list_size = std::get<1>(type_casted_input);
const auto number_of_events = std::get<2>(type_casted_input);
// Check if blockIdx.x (event_number) is in dev_event_list
unsigned mask = 0;
for (unsigned i = 0; i < (event_list_size + warp_size - 1) / warp_size; ++i) {
const auto index = i * warp_size + threadIdx.x;
mask |= __ballot_sync(0xFFFFFFFF, index < event_list_size ? threadIdx.x == parameters.dev_event_list[i] : false);
}
// * Populate IMultiEventContainer* if relevant
if (blockIdx.x == 0 && threadIdx.x == 0 && threadIdx.y == 0) {
if constexpr (Allen::has_dev_particle_container<Derived, device_datatype, input_datatype>::value) {
const auto particle_container_ptr =
static_cast<const Allen::IMultiEventContainer*>(parameters.dev_particle_container);
parameters.dev_particle_container_ptr[0] = const_cast<Allen::IMultiEventContainer*>(particle_container_ptr);
} else {
parameters.dev_particle_container_ptr[0] = nullptr;
}
// Do initialization for all events, regardless of mask
// * Populate offsets in first block
if (blockIdx.x == 0) {
for (unsigned i = threadIdx.x * blockIdx.y + threadIdx.y; i < number_of_events; i += blockDim.x * blockDim.y) {
parameters.dev_decisions_offsets[i] = mask ? Derived::offset(parameters, i) : 0;
}
}
// * Populate decisions
const auto pre_scaler_hash = std::get<3>(type_casted_input);
const bool pre_scaler_result = deterministic_scaler(pre_scaler_hash, parameters.pre_scaler, run_no, evt_hi, evt_lo, gps_hi, gps_lo);
const unsigned input_size = Derived::input_size(parameters, blockIdx.x);
for (unsigned i = threadIdx.x; i < input_size; i += blockDim.x) {
const bool sel = mask && pre_scaler_result && Derived::select(parameters, Derived::get_input(parameters, blockIdx.x, i));
unsigned index = Derived::offset(parameters, blockIdx.x) + i;
parameters.dev_decisions[index] = sel;
// * Populate IMultiEventContainer* if relevant
if (blockIdx.x == 0 && threadIdx.x == 0 && threadIdx.y == 0) {
if constexpr (Allen::has_dev_particle_container<Derived, device_datatype, input_datatype>::value) {
const auto particle_container_ptr =
static_cast<const Allen::IMultiEventContainer*>(parameters.dev_particle_container);
parameters.dev_particle_container_ptr[0] = const_cast<Allen::IMultiEventContainer*>(particle_container_ptr);
} else {
parameters.dev_particle_container_ptr[0] = nullptr;
}
}
// if constexpr (std::is_same_v<typename Derived::iteration_t, LineIteration::event_iteration_tag>) {
// if (blockIdx.x == 0) {
// // Iterates over events and processes the line
// const auto& type_casted_input =
// *reinterpret_cast<const std::tuple<Parameters, size_t, unsigned, unsigned>*>(input);
// const auto& [parameters, number_of_events_in_event_list, number_of_events, pre_scaler_hash] =
// type_casted_input;
// * Populate decisions
const auto pre_scaler_hash = std::get<3>(type_casted_input);
const bool pre_scaler_result = deterministic_scaler(pre_scaler_hash, parameters.pre_scaler, run_no, evt_hi, evt_lo, gps_hi, gps_lo);
const unsigned input_size = Derived::input_size(parameters, blockIdx.x);
// // Populate IMultiEventContainer* if relevant
// if constexpr (Allen::has_dev_particle_container<Derived, device_datatype, input_datatype>::value) {
// if (blockIdx.x == 0 && threadIdx.x == 0) {
// const auto particle_container_ptr =
// static_cast<const Allen::IMultiEventContainer*>(parameters.dev_particle_container);
// parameters.dev_particle_container_ptr[0] =
// const_cast<Allen::IMultiEventContainer*>(particle_container_ptr);
// }
// }
// // Do selection
// for (unsigned i = threadIdx.x; i < number_of_events_in_event_list; i += blockDim.x) {
// const auto event_number = parameters.dev_event_list[i];
// // ODIN data
// const LHCb::ODIN odin {
// {*parameters.dev_mep_layout ?
// odin_data_mep_t::data(parameters.dev_odin_raw_input, parameters.dev_odin_raw_input_offsets,
// event_number) : odin_data_t::data(parameters.dev_odin_raw_input, parameters.dev_odin_raw_input_offsets,
// event_number),
// 10}};
// const uint32_t run_no = odin.runNumber();
// const uint32_t evt_hi = static_cast<uint32_t>(odin.eventNumber() >> 32);
// const uint32_t evt_lo = static_cast<uint32_t>(odin.eventNumber() & 0xffffffff);
// const uint32_t gps_hi = static_cast<uint32_t>(odin.gpsTime() >> 32);
// const uint32_t gps_lo = static_cast<uint32_t>(odin.gpsTime() & 0xffffffff);
// bool decision = false;
// if (deterministic_scaler(pre_scaler_hash, parameters.pre_scaler, run_no, evt_hi, evt_lo, gps_hi, gps_lo)) {
// auto input = Derived::get_input(parameters, event_number);
// decision = Derived::select(parameters, input);
// parameters.dev_decisions[event_number] = decision;
// Derived::monitor(parameters, input, event_number, decision);
// }
// }
// // Populate offsets
// for (unsigned event_number = threadIdx.x; event_number < number_of_events; event_number += blockDim.x) {
// parameters.dev_decisions_offsets[event_number] = event_number;
// }
// }
// }
// else {
// Processes a line by iterating over all events and all "input sizes" (ie. tracks, vertices, etc.).
// }
for (unsigned i = threadIdx.x; i < input_size; i += blockDim.x) {
const bool sel = mask && pre_scaler_result && Derived::select(parameters, Derived::get_input(parameters, blockIdx.x, i));
unsigned index = Derived::offset(parameters, blockIdx.x) + i;
parameters.dev_decisions[index] = sel;
}
}
template<typename Derived, typename Parameters>
......
device/selections/line_types/include/ODINLine.cuh
+ 1
− 1
View file @ 0f39ca10
......@@ -15,7 +15,7 @@
*/
template<typename Derived, typename Parameters>
struct ODINLine : public EventLine<Derived, Parameters> {
__device__ static std::tuple<const unsigned*> get_input(const Parameters& parameters, const unsigned event_number)
__device__ static std::tuple<const unsigned*> get_input(const Parameters& parameters, const unsigned event_number, const unsigned)
{
const unsigned* event_odin_data = nullptr;
if (parameters.dev_mep_layout[0]) {
......
device/selections/lines/calibration/include/PassthroughLine.cuh
+ 1
− 1
View file @ 0f39ca10
......@@ -27,7 +27,7 @@ namespace passthrough_line {
};
struct passthrough_line_t : public SelectionAlgorithm, Parameters, EventLine<passthrough_line_t, Parameters> {
__device__ static std::tuple<const bool> get_input(const Parameters& parameters, const unsigned event_number);
__device__ static std::tuple<const bool> get_input(const Parameters& parameters, const unsigned event_number, const unsigned);
__device__ static bool select(const Parameters& parameters, std::tuple<const bool> input);
......
device/selections/lines/calibration/src/PassthroughLine.cu
+ 1
− 1
View file @ 0f39ca10
......@@ -6,7 +6,7 @@
// Explicit instantiation
INSTANTIATE_LINE(passthrough_line::passthrough_line_t, passthrough_line::Parameters)
__device__ std::tuple<const bool> passthrough_line::passthrough_line_t::get_input(const Parameters&, const unsigned)
__device__ std::tuple<const bool> passthrough_line::passthrough_line_t::get_input(const Parameters&, const unsigned, const unsigned)
{
return std::forward_as_tuple(true);
}
......
device/selections/lines/electron/include/DisplacedLeptonsLine.cuh
+ 1
− 1
View file @ 0f39ca10
......@@ -38,7 +38,7 @@ namespace displaced_leptons_line {
Parameters,
EventLine<displaced_leptons_line_t, Parameters> {
__device__ static std::tuple<const Allen::Views::Physics::BasicParticles, const unsigned, const bool*, const float*>
get_input(const Parameters& parameters, const unsigned event_number);
get_input(const Parameters& parameters, const unsigned event_number, const unsigned);
__device__ static bool select(
const Parameters& parameters,
......
device/selections/lines/electron/src/DisplacedLeptonsLine.cu
+ 1
− 1
View file @ 0f39ca10
......@@ -7,7 +7,7 @@
INSTANTIATE_LINE(displaced_leptons_line::displaced_leptons_line_t, displaced_leptons_line::Parameters)
__device__ std::tuple<const Allen::Views::Physics::BasicParticles, const unsigned, const bool*, const float*>
displaced_leptons_line::displaced_leptons_line_t::get_input(const Parameters& parameters, const unsigned event_number)
displaced_leptons_line::displaced_leptons_line_t::get_input(const Parameters& parameters, const unsigned event_number, const unsigned)
{
const auto event_tracks = parameters.dev_track_container->container(event_number);
const unsigned N_tracks = event_tracks.size();
......
device/selections/lines/monitoring/include/BeamCrossingLine.cuh
+ 3
− 1
View file @ 0f39ca10
......@@ -17,7 +17,9 @@ namespace beam_crossing_line {
HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
HOST_OUTPUT(host_fn_parameters_t, char) host_fn_parameters;
DEVICE_OUTPUT(dev_fn_parameters_t, char) dev_fn_parameters;
DEVICE_INPUT(dev_odin_raw_input_t, char) dev_odin_raw_input;
DEVICE_INPUT(dev_odin_raw_input_offsets_t, unsigned) dev_odin_raw_input_offsets;
DEVICE_INPUT(dev_mep_layout_t, unsigned) dev_mep_layout;
DEVICE_OUTPUT(dev_particle_container_ptr_t, Allen::IMultiEventContainer*)
dev_particle_container_ptr;
......
device/selections/lines/monitoring/include/ODINEventTypeLine.cuh
+ 3
− 1
View file @ 0f39ca10
......@@ -18,7 +18,9 @@ namespace odin_event_type_line {
HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
HOST_OUTPUT(host_fn_parameters_t, char) host_fn_parameters;
DEVICE_OUTPUT(dev_fn_parameters_t, char) dev_fn_parameters;
DEVICE_INPUT(dev_odin_raw_input_t, char) dev_odin_raw_input;
DEVICE_INPUT(dev_odin_raw_input_offsets_t, unsigned) dev_odin_raw_input_offsets;
DEVICE_INPUT(dev_mep_layout_t, unsigned) dev_mep_layout;
DEVICE_OUTPUT(dev_particle_container_ptr_t, Allen::IMultiEventContainer*)
dev_particle_container_ptr;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
......
device/selections/lines/monitoring/include/VeloMicroBiasLine.cuh
+ 1
− 1
View file @ 0f39ca10
......@@ -32,7 +32,7 @@ namespace velo_micro_bias_line {
};
struct velo_micro_bias_line_t : public SelectionAlgorithm, Parameters, EventLine<velo_micro_bias_line_t, Parameters> {
__device__ static std::tuple<const unsigned> get_input(const Parameters& parameters, const unsigned event_number);
__device__ static std::tuple<const unsigned> get_input(const Parameters& parameters, const unsigned event_number, const unsigned);
__device__ static bool select(const Parameters& parameters, std::tuple<const unsigned> input);
......
device/selections/lines/monitoring/src/VeloMicroBiasLine.cu
+ 2
− 1
View file @ 0f39ca10
......@@ -8,7 +8,8 @@ INSTANTIATE_LINE(velo_micro_bias_line::velo_micro_bias_line_t, velo_micro_bias_l
__device__ std::tuple<const unsigned> velo_micro_bias_line::velo_micro_bias_line_t::get_input(
const Parameters& parameters,
const unsigned event_number)
const unsigned event_number,
const unsigned)
{
Velo::Consolidated::ConstTracks velo_tracks {parameters.dev_offsets_velo_tracks,
parameters.dev_offsets_velo_track_hit_number,
......
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