diff --git a/Rec/Allen/src/GaudiAllenCountAndDumpLineDecisions.cpp b/Rec/Allen/src/GaudiAllenCountAndDumpLineDecisions.cpp
index fb14f03d1a005e45cb99de3ce5a62f2a0474b1eb..0f59241626757226542b38cf6207bea46ea1d033 100644
--- a/Rec/Allen/src/GaudiAllenCountAndDumpLineDecisions.cpp
+++ b/Rec/Allen/src/GaudiAllenCountAndDumpLineDecisions.cpp
@@ -28,7 +28,7 @@
struct GaudiAllenCountAndDumpLineDecisions final : public Gaudi::Functional::Consumer<void(
const std::vector<unsigned>&,
const std::vector<char>&,
- const std::vector<char>&,
+ const std::vector<unsigned>&,
const std::vector<unsigned>&)> {
// Standard constructor
GaudiAllenCountAndDumpLineDecisions(const std::string& name, ISvcLocator* pSvcLocator);
@@ -36,7 +36,7 @@ struct GaudiAllenCountAndDumpLineDecisions final : public Gaudi::Functional::Con
void operator()(
const std::vector<unsigned>& allen_number_of_active_lines,
const std::vector<char>& allen_names_of_active_lines,
- const std::vector<char>& allen_selections,
+ const std::vector<unsigned>& allen_selections,
const std::vector<unsigned>& allen_selections_offsets) const override;
private:
@@ -84,7 +84,7 @@ GaudiAllenCountAndDumpLineDecisions::GaudiAllenCountAndDumpLineDecisions(
void GaudiAllenCountAndDumpLineDecisions::operator()(
const std::vector<unsigned>& allen_number_of_active_lines,
const std::vector<char>& allen_names_of_active_lines,
- const std::vector<char>& allen_selections,
+ const std::vector<unsigned>& allen_selections,
const std::vector<unsigned>& allen_selections_offsets) const
{
assert(allen_number_of_active_lines[0] == m_hlt1_line_rates.size());
@@ -98,17 +98,13 @@ void GaudiAllenCountAndDumpLineDecisions::operator()(
const unsigned number_of_events = 1;
// Selections view
- const Selections::Selections_t<const char> selections {
+ const Selections::ConstSelections selections {
allen_selections.data(), allen_selections_offsets.data(), number_of_events};
// Increment counters
bool global_dec = false;
for (unsigned line_index = 0; line_index < allen_number_of_active_lines[0]; line_index++) {
- bool line_dec = false;
- auto decs = selections.get_span(line_index, i_event);
- for (unsigned idec = 0; idec < decs.size(); idec++)
- line_dec |= decs[idec];
-
+ bool line_dec = !selections.is_span_empty(line_index, i_event);
m_hlt1_line_rates[line_index] += line_dec;
global_dec |= line_dec;
}
diff --git a/configuration/parser/ParseAlgorithms.py b/configuration/parser/ParseAlgorithms.py
index 5cf964345342d3b7af0f5cba215a628603857323..505d982245cfbfee52cd65b5c521ef752abec406 100755
--- a/configuration/parser/ParseAlgorithms.py
+++ b/configuration/parser/ParseAlgorithms.py
@@ -722,7 +722,7 @@ class AllenCore():
code += "\n"
if separable_compilation:
for alg in selection_algorithms:
- code += f"extern template __device__ void process_line<{alg.namespace}::{alg.name}, {alg.namespace}::Parameters>(char*, bool*, unsigned*, Allen::IMultiEventContainer**, unsigned, unsigned, unsigned, unsigned, unsigned, unsigned, const unsigned);\n"
+ code += f"extern template __device__ void process_line<{alg.namespace}::{alg.name}, {alg.namespace}::Parameters>(char*, uint32_t*, unsigned*, const LineData*, const ODINData*, const unsigned*, const unsigned, const unsigned, const unsigned, const unsigned);\n"
code += "\n"
for alg in selection_algorithms:
code += f"extern template void line_output_monitor<{alg.namespace}::{alg.name}, {alg.namespace}::Parameters>(char*, const RuntimeOptions&, const Allen::Context&);\n"
@@ -732,11 +732,11 @@ class AllenCore():
if i != len(selection_algorithms) - 1:
code += ",\n"
code += "\n};\n\n"
- code += "__device__ inline void invoke_line_functions(unsigned index, char* a, bool* b, unsigned* c, Allen::IMultiEventContainer** d, unsigned e, unsigned f, unsigned g, unsigned h, unsigned i, unsigned j, const unsigned k) {\n"
+ code += "__device__ inline void invoke_line_functions(unsigned index, char* a, uint32_t* b, unsigned* c, const LineData* d, const ODINData* e, const unsigned* f, const unsigned g, const unsigned h, const unsigned i, const unsigned j) {\n"
code += f" assert(index < {len(selection_algorithms)});\n"
code += " switch (index) {\n"
for i, alg in enumerate(selection_algorithms):
- code += f" case {i}: process_line<{alg.namespace}::{alg.name}, {alg.namespace}::Parameters>(a, b, c, d, e, f, g, h, i, j, k); break;\n"
+ code += f" case {i}: process_line<{alg.namespace}::{alg.name}, {alg.namespace}::Parameters>(a, b, c, d, e, f, g, h, i, j); break;\n"
code += " }\n}\n\n"
code += f"constexpr std::array<void(*)(char*, const RuntimeOptions&, const Allen::Context&), {len(selection_algorithms)}> line_output_monitor_functions = {{\n"
for i, alg in enumerate(selection_algorithms):
diff --git a/configuration/python/AllenConf/persistency.py b/configuration/python/AllenConf/persistency.py
index f742cca43d351d4bc4946bae883ac94e7db44399..290b2da8990ec12a330f24f137a157e1de5648bd 100644
--- a/configuration/python/AllenConf/persistency.py
+++ b/configuration/python/AllenConf/persistency.py
@@ -220,13 +220,7 @@ def make_gather_selections(lines):
gather_selections_t,
name="gather_selections",
host_number_of_events_t=number_of_events["host_number_of_events"],
- host_decisions_sizes_t=[line.host_decisions_size_t for line in lines],
- host_input_post_scale_factors_t=[
- line.host_post_scaler_t for line in lines
- ],
- host_input_post_scale_hashes_t=[
- line.host_post_scaler_hash_t for line in lines
- ],
+ host_input_line_data_t=[line.host_line_data_t for line in lines],
dev_odin_data_t=odin["dev_odin_data"],
names_of_active_lines=",".join([line.name for line in lines]),
names_of_active_line_algorithms=",".join(
@@ -353,8 +347,6 @@ def make_sel_report_writer(lines):
host_total_sum_holder_t,
dev_number_of_active_lines_t=gather_selections.
dev_number_of_active_lines_t,
- dev_selections_t=gather_selections.dev_selections_t,
- dev_selections_offsets_t=gather_selections.dev_selections_offsets_t,
dev_max_objects_offsets_t=prefix_sum_max_objects.dev_output_buffer_t,
dev_sel_count_t=make_selected_object_lists.dev_sel_count_t,
dev_sel_list_t=make_selected_object_lists.dev_sel_list_t,
diff --git a/device/event_model/selections/include/SelectionsEventModel.cuh b/device/event_model/selections/include/SelectionsEventModel.cuh
index 04d47aab83af2c776e940d50b603b2c40b6f198d..2a4dba3d013d59f298d52219cbfa9964c228d8e7 100644
--- a/device/event_model/selections/include/SelectionsEventModel.cuh
+++ b/device/event_model/selections/include/SelectionsEventModel.cuh
@@ -14,48 +14,92 @@
#include "BackendCommon.h"
namespace Selections {
+ template<typename T>
+ struct SelectionSpan {
+ const T* m_data;
+ const unsigned m_size;
+ constexpr static unsigned store_size = sizeof(T) * 8;
+ __host__ __device__ SelectionSpan(T* data, unsigned size) : m_data(data), m_size(size) {}
+ __host__ __device__ bool operator[](unsigned index) const
+ {
+ return (m_data[index / store_size] >> (index % store_size)) & 1;
+ }
+ __host__ __device__ unsigned size() const { return m_size; }
+ __host__ __device__ bool empty() const { return m_size == 0; }
+ };
+
template<typename T>
struct Selections_t {
protected:
- typename ForwardType<T, bool>::t* m_base_pointer;
+ T* m_base_pointer;
const unsigned* m_offsets;
const unsigned m_number_of_events;
public:
- constexpr static unsigned element_size = sizeof(bool);
+ using store_t = T;
+ constexpr static unsigned store_size = sizeof(store_t) * 8;
- __host__ __device__ Selections_t(T* base_pointer, const unsigned* offsets, const unsigned number_of_events) :
- m_base_pointer(reinterpret_cast<typename ForwardType<T, bool>::t*>(base_pointer)), m_offsets(offsets),
- m_number_of_events(number_of_events)
+ __host__ __device__ Selections_t(store_t* base_pointer, const unsigned* offsets, const unsigned number_of_events) :
+ m_base_pointer(base_pointer), m_offsets(offsets), m_number_of_events(number_of_events)
{}
- __host__ __device__ Selections_t(const Selections_t<T>& selections) :
+ __host__ __device__ Selections_t(const Selections_t<store_t>& selections) :
m_base_pointer(selections.m_base_pointer), m_offsets(selections.m_offsets),
m_number_of_events(selections.m_number_of_events)
{}
- __host__ __device__ bool selection(const unsigned line, const unsigned event, const unsigned index = 0) const
+ __host__ __device__ unsigned get_bin_index(unsigned span_index) const
+ {
+ return m_offsets[span_index] / store_size + span_index; // assume 1 padding word per span_index
+ }
+
+ __host__ __device__ void fill_span(const unsigned line, const unsigned event, const bool value) const
+ {
+ assert(event < m_number_of_events);
+ store_t fill_value = value ? -1 : 0;
+ auto span_index = line * m_number_of_events + event;
+ unsigned start = get_bin_index(span_index);
+ unsigned end = get_bin_index(span_index + 1);
+ for (unsigned i = start; i < end; i++) {
+ m_base_pointer[i] = fill_value;
+ }
+ }
+
+ __host__ __device__ bool is_span_empty(const unsigned line, const unsigned event) const
{
assert(event < m_number_of_events);
- return m_base_pointer[m_offsets[line * m_number_of_events + event] + index];
+ auto span_index = line * m_number_of_events + event;
+ unsigned start = get_bin_index(span_index);
+ unsigned end = get_bin_index(span_index + 1);
+ for (unsigned i = start; i < end; i++) {
+ if (m_base_pointer[i]) return false;
+ }
+ return true;
}
- __host__ __device__ bool& selection(const unsigned line, const unsigned event, const unsigned index = 0)
+ __host__ __device__ unsigned count_span_population(const unsigned line, const unsigned event) const
{
assert(event < m_number_of_events);
- return m_base_pointer[m_offsets[line * m_number_of_events + event] + index];
+ auto span_index = line * m_number_of_events + event;
+ unsigned start = get_bin_index(span_index);
+ unsigned end = get_bin_index(span_index + 1);
+ int sum = 0;
+ for (unsigned i = start; i < end; i++) {
+ sum += __popc(m_base_pointer[i]);
+ }
+ return sum;
}
- __host__ __device__ Allen::device::span<typename ForwardType<T, bool>::t> get_span(
- const unsigned line,
- const unsigned event) const
+ __host__ __device__ SelectionSpan<store_t> get_span(const unsigned line, const unsigned event) const
{
assert(event < m_number_of_events);
- return {m_base_pointer + m_offsets[line * m_number_of_events + event],
+ auto span_index = line * m_number_of_events + event;
+ auto offset = get_bin_index(span_index);
+ return {m_base_pointer + offset,
m_offsets[line * m_number_of_events + event + 1] - m_offsets[line * m_number_of_events + event]};
}
};
- using ConstSelections = const Selections_t<const bool>;
- using Selections = Selections_t<bool>;
+ using ConstSelections = const Selections_t<const uint32_t>;
+ using Selections = Selections_t<uint32_t>;
} // namespace Selections
diff --git a/device/lumi/include/CalcLumiSumSize.cuh b/device/lumi/include/CalcLumiSumSize.cuh
index 5fa26bf181ccbfb7d0b1830988c90fb24b238742..595504caa9ee3d57d28562efec7867f0ca8f9ed0 100644
--- a/device/lumi/include/CalcLumiSumSize.cuh
+++ b/device/lumi/include/CalcLumiSumSize.cuh
@@ -15,7 +15,7 @@
namespace calc_lumi_sum_size {
struct Parameters {
HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events;
- DEVICE_INPUT(dev_selections_t, bool) dev_selections;
+ DEVICE_INPUT(dev_selections_t, uint32_t) dev_selections;
DEVICE_INPUT(dev_selections_offsets_t, unsigned) dev_selections_offsets;
DEVICE_OUTPUT(dev_lumi_sum_sizes_t, unsigned) dev_lumi_sum_sizes;
DEVICE_OUTPUT(dev_lumi_sum_present_t, unsigned) dev_lumi_sum_present;
diff --git a/device/selections/Hlt1/include/DecReporter.cuh b/device/selections/Hlt1/include/DecReporter.cuh
index d922948eea29c35820e2ace901dd0fa7d0149bcb..3131a0142270ae4d08e57932e46ccf21f248ff62 100644
--- a/device/selections/Hlt1/include/DecReporter.cuh
+++ b/device/selections/Hlt1/include/DecReporter.cuh
@@ -17,7 +17,7 @@ namespace dec_reporter {
HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events;
HOST_INPUT(host_number_of_active_lines_t, unsigned) host_number_of_active_lines;
DEVICE_INPUT(dev_number_of_active_lines_t, unsigned) dev_number_of_active_lines;
- DEVICE_INPUT(dev_selections_t, bool) dev_selections;
+ DEVICE_INPUT(dev_selections_t, uint32_t) dev_selections;
DEVICE_INPUT(dev_selections_offsets_t, unsigned) dev_selections_offsets;
DEVICE_OUTPUT(dev_selected_candidates_counts_t, unsigned) dev_selected_candidates_counts;
DEVICE_OUTPUT(dev_dec_reports_t, unsigned) dev_dec_reports;
diff --git a/device/selections/Hlt1/include/GatherSelections.cuh b/device/selections/Hlt1/include/GatherSelections.cuh
index 2fa0fd5e0eadd4276d81ff31611296e977bc4988..71ea5e833cfc2c5b98fa16eaaabb6166a70a7f1c 100644
--- a/device/selections/Hlt1/include/GatherSelections.cuh
+++ b/device/selections/Hlt1/include/GatherSelections.cuh
@@ -30,9 +30,7 @@ namespace gather_selections {
HOST_OUTPUT(host_selections_offsets_t, unsigned) host_selections_offsets;
HOST_OUTPUT(host_number_of_active_lines_t, unsigned) host_number_of_active_lines;
HOST_OUTPUT(host_names_of_active_lines_t, char) host_names_of_active_lines;
- HOST_INPUT_AGGREGATE(host_decisions_sizes_t, unsigned) host_decisions_sizes;
- HOST_INPUT_AGGREGATE(host_input_post_scale_factors_t, float) host_input_post_scale_factors;
- HOST_INPUT_AGGREGATE(host_input_post_scale_hashes_t, uint32_t) host_input_post_scale_hashes;
+ HOST_INPUT_AGGREGATE(host_input_line_data_t, LineData) host_input_line_data;
HOST_INPUT_AGGREGATE(host_fn_parameters_agg_t, char) host_fn_parameters_agg;
DEVICE_OUTPUT(dev_fn_parameters_t, char) dev_fn_parameters;
HOST_OUTPUT(host_fn_parameter_pointers_t, char*) host_fn_parameter_pointers;
@@ -41,14 +39,14 @@ namespace gather_selections {
DEVICE_OUTPUT(dev_fn_indices_t, unsigned) dev_fn_indices;
HOST_OUTPUT(host_fn_parameters_t, char) host_fn_parameters;
DEVICE_INPUT(dev_odin_data_t, ODINData) dev_odin_data;
- DEVICE_OUTPUT(dev_selections_t, bool) dev_selections;
+ DEVICE_OUTPUT(dev_selections_t, uint32_t) dev_selections;
DEVICE_OUTPUT(dev_selections_lines_offsets_t, unsigned) dev_selections_lines_offsets;
DEVICE_OUTPUT(dev_selections_offsets_t, unsigned) dev_selections_offsets;
DEVICE_OUTPUT(dev_number_of_active_lines_t, unsigned) dev_number_of_active_lines;
- HOST_OUTPUT(host_post_scale_factors_t, float) host_post_scale_factors;
- HOST_OUTPUT(host_post_scale_hashes_t, uint32_t) host_post_scale_hashes;
- DEVICE_OUTPUT(dev_post_scale_factors_t, float) dev_post_scale_factors;
- DEVICE_OUTPUT(dev_post_scale_hashes_t, uint32_t) dev_post_scale_hashes;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
+ DEVICE_OUTPUT(dev_line_data_t, LineData) dev_line_data;
+ DEVICE_OUTPUT(dev_pre_scale_event_lists_t, unsigned) dev_pre_scale_event_lists;
+ DEVICE_OUTPUT(dev_pre_scale_event_lists_size_t, unsigned) dev_pre_scale_event_lists_size;
DEVICE_OUTPUT_WITH_DEPENDENCIES(
dev_particle_containers_t,
DEPENDENCIES(host_fn_parameters_agg_t),
diff --git a/device/selections/Hlt1/include/MakeSelectedObjectLists.cuh b/device/selections/Hlt1/include/MakeSelectedObjectLists.cuh
index 132019e94afaf2d8d50e4a9c4fe62c5958bdca8b..62fbaed0f8ad19510d5912f859d8f5fea19bbd20 100644
--- a/device/selections/Hlt1/include/MakeSelectedObjectLists.cuh
+++ b/device/selections/Hlt1/include/MakeSelectedObjectLists.cuh
@@ -23,7 +23,7 @@ namespace make_selected_object_lists {
DEVICE_INPUT(dev_number_of_events_t, unsigned) dev_number_of_events;
DEVICE_INPUT(dev_multi_event_particle_containers_t, Allen::IMultiEventContainer*)
dev_multi_event_particle_containers;
- DEVICE_INPUT(dev_selections_t, bool) dev_selections;
+ DEVICE_INPUT(dev_selections_t, uint32_t) dev_selections;
DEVICE_INPUT(dev_selections_offsets_t, unsigned) dev_selections_offsets;
DEVICE_INPUT(dev_max_objects_offsets_t, unsigned) dev_max_objects_offsets;
DEVICE_OUTPUT(dev_candidate_count_t, unsigned) dev_candidate_count;
diff --git a/device/selections/Hlt1/include/MakeSubBanks.cuh b/device/selections/Hlt1/include/MakeSubBanks.cuh
index 7f1ff6af5d39659b2d338020416fc60ce576742b..0807b6e9b18a23520976b0b08ee381d4708b7d6d 100644
--- a/device/selections/Hlt1/include/MakeSubBanks.cuh
+++ b/device/selections/Hlt1/include/MakeSubBanks.cuh
@@ -21,8 +21,6 @@ namespace make_subbanks {
HOST_INPUT(host_objtyp_bank_size_t, unsigned) host_objtyp_bank_size;
HOST_INPUT(host_stdinfo_bank_size_t, unsigned) host_stdinfo_bank_size;
DEVICE_INPUT(dev_number_of_active_lines_t, unsigned) dev_number_of_active_lines;
- DEVICE_INPUT(dev_selections_t, bool) dev_selections;
- DEVICE_INPUT(dev_selections_offsets_t, unsigned) dev_selections_offsets;
DEVICE_INPUT(dev_max_objects_offsets_t, unsigned) dev_max_objects_offsets;
DEVICE_INPUT(dev_sel_count_t, unsigned) dev_sel_count;
DEVICE_INPUT(dev_sel_list_t, unsigned) dev_sel_list;
diff --git a/device/selections/Hlt1/src/DecReporter.cu b/device/selections/Hlt1/src/DecReporter.cu
index 0c21ce18065ef463aa2d7d0c7ce2707d28dd537a..5cf9a8db617dda406c68e95539833f6a55b23d47 100644
--- a/device/selections/Hlt1/src/DecReporter.cu
+++ b/device/selections/Hlt1/src/DecReporter.cu
@@ -68,14 +68,9 @@ __global__ void dec_reporter::dec_reporter(dec_reporter::Parameters parameters)
for (unsigned line_index = threadIdx.x; line_index < reports.number_of_lines(); line_index += blockDim.x) {
// Iterate all elements and get a decision for the current {event, line}
- bool final_decision = false;
- auto decs = selections.get_span(line_index, event_index);
- for (unsigned i = 0; i < decs.size(); ++i) {
- final_decision |= decs[i];
- if (decs[i]) {
- event_selected_candidates_counts[line_index]++;
- }
- }
+ auto span_popcount = selections.count_span_population(line_index, event_index);
+ bool final_decision = span_popcount > 0;
+ event_selected_candidates_counts[line_index] = span_popcount;
reports.set_dec_report(
line_index,
diff --git a/device/selections/Hlt1/src/GatherSelections.cu b/device/selections/Hlt1/src/GatherSelections.cu
index d27e3318d809868783c9b1e812dd8990c77e0fbc..788f5fc94edfb3dd4593b3190f9733e8cb348f24 100644
--- a/device/selections/Hlt1/src/GatherSelections.cu
+++ b/device/selections/Hlt1/src/GatherSelections.cu
@@ -39,35 +39,75 @@ std::vector<std::string> split(const std::string& s, char delim)
namespace gather_selections {
__global__ void
- run_lines(gather_selections::Parameters params, const unsigned number_of_events, const unsigned number_of_lines)
+ prescaler(gather_selections::Parameters params, const unsigned number_of_events, const unsigned number_of_lines)
{
- // Process each event with a different block
- // ODIN data
- LHCb::ODIN odin {params.dev_odin_data[blockIdx.x]};
+ unsigned line = blockIdx.x;
+ auto pre_scaler_hash = params.dev_line_data[line].pre_scaler_hash;
+ auto pre_scaler = params.dev_line_data[line].pre_scaler;
+ auto event_list = params.dev_line_data[line].event_list;
+ auto event_list_size = params.dev_line_data[line].event_list_size;
+
+ if (pre_scaler >= 1.f) { // No prescaler, just copy the event list
+ for (unsigned i = threadIdx.x; i < event_list_size; i += blockDim.x) {
+ auto event_number = event_list[i];
+ params.dev_pre_scale_event_lists[line * number_of_events + i] = event_number;
+ }
+ if (threadIdx.x == 0) params.dev_pre_scale_event_lists_size[line] = event_list_size;
+ }
+ else {
+ for (unsigned i = threadIdx.x; i < event_list_size; i += blockDim.x) {
+ auto event_number = event_list[i];
+ LHCb::ODIN odin {params.dev_odin_data[event_number]};
+
+ 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);
+
+ if (deterministic_scaler(pre_scaler_hash, pre_scaler, run_no, evt_hi, evt_lo, gps_hi, gps_lo)) {
+ auto index = atomicAdd(¶ms.dev_pre_scale_event_lists_size[line], 1);
+ params.dev_pre_scale_event_lists[line * number_of_events + index] = event_number;
+ }
+ }
+ }
- 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);
+ if (blockIdx.x == 0) {
+ for (unsigned i = threadIdx.x; i < number_of_lines; i += blockDim.x) {
+ params.dev_particle_containers[i] = params.dev_line_data[i].particle_container_ptr;
+ }
+ }
+ }
- for (unsigned i = threadIdx.y; i < number_of_lines; i += blockDim.y) {
+ __global__ void
+ run_lines(gather_selections::Parameters params, const unsigned number_of_events, const unsigned number_of_lines)
+ {
+ // Process each line with a different block
+ unsigned line = blockIdx.x;
+ char* input = params.dev_fn_parameter_pointers[line];
+ auto line_offset = params.dev_selections_lines_offsets[line];
+
+ if (input == nullptr) {
+ for (unsigned i = threadIdx.x; i < number_of_events; i += blockDim.x) {
+ params.dev_selections_offsets[line * number_of_events + i] = line_offset;
+ }
+ }
+ else {
invoke_line_functions(
- params.dev_fn_indices[i],
- params.dev_fn_parameter_pointers[i],
- params.dev_selections + params.dev_selections_lines_offsets[i],
- params.dev_selections_offsets + i * number_of_events,
- params.dev_particle_containers + i,
- run_no,
- evt_hi,
- evt_lo,
- gps_hi,
- gps_lo,
- params.dev_selections_lines_offsets[i],
+ params.dev_fn_indices[line],
+ input,
+ params.dev_selections,
+ params.dev_selections_offsets + line * number_of_events,
+ params.dev_line_data + line,
+ params.dev_odin_data,
+ params.dev_pre_scale_event_lists + line * number_of_events,
+ params.dev_pre_scale_event_lists_size[line],
+ line_offset,
+ line,
number_of_events);
}
- if (blockIdx.x == 0 && threadIdx.x == 0 && threadIdx.y == 0) {
+ if (blockIdx.x == 0 && threadIdx.x == 0) {
params.dev_selections_offsets[number_of_lines * number_of_events] =
params.dev_selections_lines_offsets[number_of_lines];
}
@@ -104,34 +144,26 @@ namespace gather_selections {
const uint32_t gps_lo = static_cast<uint32_t>(odin.gpsTime() & 0xffffffff);
for (unsigned i = threadIdx.x; i < number_of_lines; i += blockDim.x) {
- auto span = sels.get_span(i, event_number);
+ if (!sels.is_span_empty(i, event_number)) {
+ dev_decisions_per_event_line[event_number * number_of_lines + i] = true;
+ atomicAdd(&dev_histo_line_passes[i], 1);
+ }
- for (unsigned j = 0; j < span.size(); ++j) {
- if (span[j]) {
- dev_decisions_per_event_line[event_number * number_of_lines + i] = true;
- atomicAdd(&dev_histo_line_passes[i], 1);
- break;
- }
+ if (!deterministic_scaler(
+ params.dev_line_data[i].post_scaler_hash,
+ params.dev_line_data[i].post_scaler,
+ run_no,
+ evt_hi,
+ evt_lo,
+ gps_hi,
+ gps_lo)) {
+ sels.fill_span(i, event_number, false);
}
- deterministic_post_scaler(
- params.dev_post_scale_hashes[i],
- params.dev_post_scale_factors[i],
- span.size(),
- span.data(),
- run_no,
- evt_hi,
- evt_lo,
- gps_hi,
- gps_lo);
-
- for (unsigned j = 0; j < span.size(); ++j) {
- if (span[j]) {
- dev_postscaled_decisions_per_event_line[event_number * number_of_lines + i] = true;
- atomicAdd(&dev_histo_line_rates[i], 1);
- event_decision = true;
- break;
- }
+ if (!sels.is_span_empty(i, event_number)) {
+ dev_postscaled_decisions_per_event_line[event_number * number_of_lines + i] = true;
+ atomicAdd(&dev_histo_line_rates[i], 1);
+ event_decision = true;
}
}
@@ -187,19 +219,17 @@ void gather_selections::gather_selections_t::set_arguments_size(
return total_size;
}();
- const auto host_decisions_sizes = input_aggregate<host_decisions_sizes_t>(arguments);
- const auto total_size_host_decisions_sizes = [&host_decisions_sizes]() {
+ const auto size_of_aggregates = input_aggregate<host_input_line_data_t>(arguments).size_of_aggregate();
+
+ const auto host_line_data = input_aggregate<host_input_line_data_t>(arguments);
+ const auto total_size_host_decisions_sizes = [&host_line_data]() {
unsigned sum = 0;
- for (unsigned i = 0; i < host_decisions_sizes.size_of_aggregate(); ++i) {
- sum += (host_decisions_sizes.size(i) > 0) ? host_decisions_sizes.first(i) : 0;
+ for (unsigned i = 0; i < host_line_data.size_of_aggregate(); ++i) {
+ sum += (host_line_data.size(i) > 0) ? host_line_data.first(i).decisions_size : 0;
}
return sum;
- }();
-
- const auto size_of_aggregates = input_aggregate<host_decisions_sizes_t>(arguments).size_of_aggregate();
+ }() / 32 + size_of_aggregates * first<host_number_of_events_t>(arguments);
- assert(input_aggregate<host_input_post_scale_factors_t>(arguments).size_of_aggregate() == size_of_aggregates);
- assert(input_aggregate<host_input_post_scale_hashes_t>(arguments).size_of_aggregate() == size_of_aggregates);
assert(m_indices_active_line_algorithms.size() == size_of_aggregates);
set_size<host_number_of_active_lines_t>(arguments, 1);
@@ -210,10 +240,10 @@ void gather_selections::gather_selections_t::set_arguments_size(
set_size<host_selections_offsets_t>(arguments, first<host_number_of_events_t>(arguments) * size_of_aggregates + 1);
set_size<dev_selections_offsets_t>(arguments, first<host_number_of_events_t>(arguments) * size_of_aggregates + 1);
set_size<dev_selections_t>(arguments, total_size_host_decisions_sizes);
- set_size<host_post_scale_factors_t>(arguments, size_of_aggregates);
- set_size<host_post_scale_hashes_t>(arguments, size_of_aggregates);
- set_size<dev_post_scale_factors_t>(arguments, size_of_aggregates);
- set_size<dev_post_scale_hashes_t>(arguments, size_of_aggregates);
+ set_size<host_line_data_t>(arguments, size_of_aggregates);
+ set_size<dev_line_data_t>(arguments, size_of_aggregates);
+ set_size<dev_pre_scale_event_lists_t>(arguments, first<host_number_of_events_t>(arguments) * size_of_aggregates);
+ set_size<dev_pre_scale_event_lists_size_t>(arguments, size_of_aggregates);
set_size<dev_particle_containers_t>(arguments, size_of_aggregates);
set_size<host_fn_parameters_t>(arguments, total_size_host_fn_parameters_agg);
set_size<dev_fn_parameters_t>(arguments, total_size_host_fn_parameters_agg);
@@ -238,7 +268,21 @@ void gather_selections::gather_selections_t::operator()(
[[maybe_unused]] const Constants& constants,
const Allen::Context& context) const
{
- // Run the selection algorithms
+ // * Pass the number of lines for posterior algorithms
+ const auto host_line_data = input_aggregate<host_input_line_data_t>(arguments);
+ data<host_number_of_active_lines_t>(arguments)[0] = host_line_data.size_of_aggregate();
+ Allen::copy_async<dev_number_of_active_lines_t, host_number_of_active_lines_t>(arguments, context);
+
+ // === Run the prescalers
+ Allen::aggregate::store_contiguous_async<host_line_data_t, host_input_line_data_t>(arguments, context, true);
+ Allen::copy_async<dev_line_data_t, host_line_data_t>(arguments, context);
+
+ Allen::memset_async<dev_pre_scale_event_lists_size_t>(arguments, 0, context);
+
+ global_function(prescaler)(host_line_data.size_of_aggregate(), property<block_dim_x_t>().get(), context)(
+ arguments, first<host_number_of_events_t>(arguments), first<host_number_of_active_lines_t>(arguments));
+
+ // === Run the selection algorithms
// * Aggregate parameter fns
Allen::aggregate::store_contiguous_async<host_fn_parameters_t, host_fn_parameters_agg_t>(arguments, context);
Allen::copy_async<dev_fn_parameters_t, host_fn_parameters_t>(arguments, context);
@@ -258,16 +302,11 @@ void gather_selections::gather_selections_t::operator()(
}
Allen::copy_async<dev_fn_parameter_pointers_t, host_fn_parameter_pointers_t>(arguments, context);
- // * Pass the number of lines for posterior algorithms
- const auto host_decisions_sizes = input_aggregate<host_decisions_sizes_t>(arguments);
- data<host_number_of_active_lines_t>(arguments)[0] = host_decisions_sizes.size_of_aggregate();
- Allen::copy_async<dev_number_of_active_lines_t, host_number_of_active_lines_t>(arguments, context);
-
// * Calculate prefix sum of host_decisions_sizes_t sizes into host_selections_lines_offsets_t
auto* container = data<host_selections_lines_offsets_t>(arguments);
container[0] = 0;
- for (size_t i = 0; i < host_decisions_sizes.size_of_aggregate(); ++i) {
- container[i + 1] = container[i] + (host_decisions_sizes.size(i) ? host_decisions_sizes.first(i) : 0);
+ for (size_t i = 0; i < host_line_data.size_of_aggregate(); ++i) {
+ container[i + 1] = container[i] + (host_line_data.size(i) ? host_line_data.first(i).decisions_size : 0);
}
Allen::copy_async<dev_selections_lines_offsets_t, host_selections_lines_offsets_t>(arguments, context);
@@ -277,9 +316,11 @@ void gather_selections::gather_selections_t::operator()(
}
Allen::copy_async<dev_fn_indices_t, host_fn_indices_t>(arguments, context);
+ Allen::memset_async<dev_selections_t>(arguments, 0, context); // <<===
+
// * Run all selections in one go
global_function(gather_selections::run_lines)(
- first<host_number_of_events_t>(arguments), dim3(warp_size, 256 / warp_size), context)(
+ host_line_data.size_of_aggregate(), property<block_dim_x_t>().get(), context)(
arguments, first<host_number_of_events_t>(arguments), first<host_number_of_active_lines_t>(arguments));
// Run monitoring if configured
@@ -293,18 +334,7 @@ void gather_selections::gather_selections_t::operator()(
const auto line_names = std::string(property<names_of_active_lines_t>());
line_names.copy(data<host_names_of_active_lines_t>(arguments), line_names.size());
- // Populate host_post_scale_factors_t
- Allen::aggregate::store_contiguous_async<host_post_scale_factors_t, host_input_post_scale_factors_t>(
- arguments, context, true);
-
- // Populate host_post_scale_hashes_t
- Allen::aggregate::store_contiguous_async<host_post_scale_hashes_t, host_input_post_scale_hashes_t>(
- arguments, context, true);
-
- // Copy host_post_scale_factors_t to dev_post_scale_factors_t,
- // and host_post_scale_hashes_t to dev_post_scale_hashes_t
- Allen::copy_async<dev_post_scale_factors_t, host_post_scale_factors_t>(arguments, context);
- Allen::copy_async<dev_post_scale_hashes_t, host_post_scale_hashes_t>(arguments, context);
+ // === Run the postscalers
// Initialize output mask size
Allen::memset_async<dev_event_list_output_size_t>(arguments, 0, context);
@@ -322,17 +352,6 @@ void gather_selections::gather_selections_t::operator()(
Allen::memset_async(dev_histo_line_passes.data(), 0, dev_histo_line_passes.size() * sizeof(unsigned), context);
Allen::memset_async(dev_histo_line_rates.data(), 0, dev_histo_line_rates.size() * sizeof(unsigned), context);
- const auto host_input_post_scale_hashes = input_aggregate<host_input_post_scale_hashes_t>(arguments);
- for (unsigned i = 0; i < host_input_post_scale_hashes.size_of_aggregate(); ++i) {
- if (host_input_post_scale_hashes.size(i) > 0 && host_input_post_scale_hashes.first(i) == 0) {
- std::vector<std::string> names = split(line_names, ',');
- throw std::runtime_error(
- "Postscaler hash was not properly initialized for " + names[i] +
- ". Did you forget to call the line's init() ?");
- }
- }
-
- // Run the postscaler
global_function(postscaler)(first<host_number_of_events_t>(arguments), property<block_dim_x_t>().get(), context)(
arguments,
first<host_number_of_active_lines_t>(arguments),
@@ -368,20 +387,16 @@ void gather_selections::gather_selections_t::operator()(
const auto host_selections = make_host_buffer<dev_selections_t>(arguments, context);
Allen::copy<host_selections_offsets_t, dev_selections_offsets_t>(arguments, context);
- Selections::ConstSelections sels {reinterpret_cast<const bool*>(host_selections.data()),
- data<host_selections_offsets_t>(arguments),
- first<host_number_of_events_t>(arguments)};
+ Selections::ConstSelections sels {
+ host_selections.data(), data<host_selections_offsets_t>(arguments), first<host_number_of_events_t>(arguments)};
std::vector<uint8_t> event_decisions {};
for (auto i = 0u; i < first<host_number_of_events_t>(arguments); ++i) {
bool dec = false;
for (auto j = 0u; j < first<host_number_of_active_lines_t>(arguments); ++j) {
auto decs = sels.get_span(j, i);
- bool span_decision = false;
- for (auto k = 0u; k < decs.size(); ++k) {
- dec |= decs[k];
- span_decision |= decs[k];
- }
+ bool span_decision = !sels.is_span_empty(j, i);
+ dec |= span_decision;
std::cout << "Span (event " << i << ", line " << j << "), size " << decs.size()
<< ", decision: " << span_decision << "\n";
}
diff --git a/device/selections/line_types/include/Line.cuh b/device/selections/line_types/include/Line.cuh
index f342dc70579ca2bda2ce41faefbf976f6c4bc5d1..ea12bc3a76e17921ea7681f038d01f11fff94363 100755
--- a/device/selections/line_types/include/Line.cuh
+++ b/device/selections/line_types/include/Line.cuh
@@ -18,6 +18,7 @@
#include <tuple>
#include <ROOTHeaders.h>
#include "ROOTService.h"
+#include "ODINBank.cuh"
// Helper macro to explicitly instantiate lines
#define INSTANTIATE_LINE(DERIVED, PARAMETERS) \
@@ -25,35 +26,35 @@
const ArgumentReferences<PARAMETERS>&, const RuntimeOptions&, const Constants&, const Allen::Context&) const; \
template __device__ void process_line<DERIVED, PARAMETERS>( \
char*, \
- bool*, \
+ uint32_t*, \
unsigned*, \
- Allen::IMultiEventContainer**, \
- unsigned, \
- unsigned, \
- unsigned, \
- unsigned, \
- unsigned, \
- unsigned, \
+ const LineData*, \
+ const ODINData*, \
+ const unsigned*, \
+ const unsigned, \
+ const unsigned, \
+ const unsigned, \
const unsigned); \
template void line_output_monitor<DERIVED, PARAMETERS>(char*, const RuntimeOptions&, const Allen::Context&); \
INSTANTIATE_ALGORITHM(DERIVED)
-// Type-erased line function type
-using line_fn_t = void (*)(
- char*,
- bool*,
- unsigned*,
- Allen::IMultiEventContainer**,
- unsigned,
- unsigned,
- unsigned,
- unsigned,
- unsigned,
- unsigned,
- const unsigned);
-
template<typename Derived, typename Parameters>
-using type_erased_tuple_t = std::tuple<Parameters, size_t, unsigned, ArgumentReferences<Parameters>, const Derived*>;
+using type_erased_tuple_t = std::tuple<Parameters, ArgumentReferences<Parameters>, const Derived*>;
+
+struct LineData {
+ float pre_scaler {0};
+ uint32_t pre_scaler_hash {0};
+
+ float post_scaler {0};
+ uint32_t post_scaler_hash {0};
+
+ unsigned decisions_size {0};
+
+ const mask_t* event_list {nullptr};
+ unsigned event_list_size {0};
+
+ Allen::IMultiEventContainer* particle_container_ptr {nullptr};
+};
template<typename Derived, typename Parameters>
struct Line {
@@ -152,9 +153,7 @@ public:
void set_arguments_size(ArgumentReferences<Parameters> arguments, const RuntimeOptions&, const Constants&) const
{
- Allen::ArgumentOperations::set_size<typename Parameters::host_decisions_size_t>(arguments, 1);
- Allen::ArgumentOperations::set_size<typename Parameters::host_post_scaler_t>(arguments, 1);
- Allen::ArgumentOperations::set_size<typename Parameters::host_post_scaler_hash_t>(arguments, 1);
+ Allen::ArgumentOperations::set_size<typename Parameters::host_line_data_t>(arguments, 1);
// Set the size of the type-erased fn parameters
Allen::ArgumentOperations::set_size<typename Parameters::host_fn_parameters_t>(
@@ -244,18 +243,18 @@ void line_output_monitor(char* input, const RuntimeOptions& runtime_options, con
if constexpr (Allen::has_enable_monitoring<Parameters>::value) {
if (input != nullptr) {
const auto& type_casted_input = *reinterpret_cast<type_erased_tuple_t<Derived, Parameters>*>(input);
- auto derived_instance = std::get<4>(type_casted_input);
+ auto derived_instance = std::get<2>(type_casted_input);
if (derived_instance->template property<typename Parameters::enable_monitoring_t>()) {
- derived_instance->output_monitor(std::get<3>(type_casted_input), runtime_options, context);
+ derived_instance->output_monitor(std::get<1>(type_casted_input), runtime_options, context);
}
}
}
if constexpr (Allen::has_enable_tupling<Parameters>::value) {
if (input != nullptr) {
const auto& type_casted_input = *reinterpret_cast<type_erased_tuple_t<Derived, Parameters>*>(input);
- auto derived_instance = std::get<4>(type_casted_input);
+ auto derived_instance = std::get<2>(type_casted_input);
if (derived_instance->template property<typename Parameters::enable_tupling_t>()) {
- derived_instance->output_tuples(std::get<3>(type_casted_input), runtime_options, context);
+ derived_instance->output_tuples(std::get<1>(type_casted_input), runtime_options, context);
}
}
}
@@ -265,87 +264,60 @@ void line_output_monitor(char* input, const RuntimeOptions& runtime_options, con
template<typename Derived, typename Parameters>
__device__ void process_line(
char* input,
- bool* decisions,
+ uint32_t* decisions,
unsigned* decisions_offsets,
- Allen::IMultiEventContainer** particle_container_ptr,
- unsigned run_no,
- unsigned evt_hi,
- unsigned evt_lo,
- unsigned gps_hi,
- unsigned gps_lo,
- unsigned line_offset,
+ [[maybe_unused]] const LineData* dev_line_data,
+ const ODINData* dev_odin_data,
+ const unsigned* pre_scale_event_list,
+ const unsigned pre_scale_event_list_size,
+ const unsigned line_offset,
+ const unsigned line_index,
const unsigned number_of_events)
{
- if (input == nullptr) {
- if (blockIdx.x == 0 && threadIdx.x == 0) {
- *particle_container_ptr = nullptr;
- }
+ const auto& type_casted_input = *reinterpret_cast<type_erased_tuple_t<Derived, Parameters>*>(input);
+ const auto& parameters = std::get<0>(type_casted_input);
- if (blockIdx.x == 0) {
- for (unsigned i = threadIdx.x; i < number_of_events; i += blockDim.x) {
- decisions_offsets[i] = line_offset;
- }
- }
+ // Do initialization for all events, regardless of mask
+ for (unsigned i = threadIdx.x; i < number_of_events; i += blockDim.x) {
+ decisions_offsets[i] = line_offset + Derived::offset(parameters, i);
}
- else {
- const auto& type_casted_input = *reinterpret_cast<type_erased_tuple_t<Derived, Parameters>*>(input);
- const auto& parameters = std::get<0>(type_casted_input);
- const auto event_list_size = std::get<1>(type_casted_input);
- const auto event_number = blockIdx.x;
-
- // 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 ? event_number == parameters.dev_event_list[index] : false);
- }
- // Do initialization for all events, regardless of mask
- // * Populate offsets in first block
- if (blockIdx.x == 0) {
- for (unsigned i = threadIdx.x; i < number_of_events; i += blockDim.x) {
- decisions_offsets[i] = line_offset + Derived::offset(parameters, i);
- }
- }
+ // Switch between 3 settings:
+ // * 1 thread per event
+ // * 1 thread per object
+ // * balanced: (32 threads per objects, 8 event per block)
+ const auto threads_per_event =
+ (dev_line_data->decisions_size == number_of_events) ? 1 : (pre_scale_event_list_size == 1) ? blockDim.x : warp_size;
+ const auto events_per_block = blockDim.x / threads_per_event;
- // * Populate IMultiEventContainer* if relevant
- if (blockIdx.x == 0 && threadIdx.x == 0) {
- if constexpr (Allen::has_dev_particle_container<
- Derived,
- Allen::Store::device_datatype,
- Allen::Store::input_datatype>::value) {
- const auto ptr = static_cast<const Allen::IMultiEventContainer*>(parameters.dev_particle_container);
- *particle_container_ptr = const_cast<Allen::IMultiEventContainer*>(ptr);
- }
- else {
- *particle_container_ptr = nullptr;
- }
- }
+ for (unsigned j = threadIdx.x / threads_per_event; j < pre_scale_event_list_size; j += events_per_block) {
+ const auto event_number = pre_scale_event_list[j];
// * Populate decisions
- const auto pre_scaler_hash = std::get<2>(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, event_number);
-
- for (unsigned i = threadIdx.x; i < input_size; i += blockDim.x) {
+ for (unsigned i = threadIdx.x % threads_per_event; i < input_size; i += threads_per_event) {
const auto input = Derived::get_input(parameters, event_number, i);
- const bool decision = mask > 0 && pre_scaler_result && Derived::select(parameters, input);
- unsigned index = Derived::offset(parameters, event_number) + i;
- decisions[index] = decision;
+ bool decision = Derived::select(parameters, input);
+
+ unsigned span_index = line_index * number_of_events + event_number;
+ unsigned index = (line_offset + Derived::offset(parameters, event_number)) / 32 + span_index + i / 32;
+ if (decision) atomicOr(&decisions[index], 1 << (i % 32));
+
if constexpr (Allen::has_enable_monitoring<Parameters>::value) {
if (parameters.enable_monitoring) {
+ unsigned index = Derived::offset(parameters, event_number) + i;
Derived::monitor(parameters, input, index, decision);
}
}
if constexpr (Allen::has_enable_tupling<Parameters>::value) {
if (parameters.enable_tupling) {
+ LHCb::ODIN odin {dev_odin_data[event_number]};
+ unsigned index = Derived::offset(parameters, event_number) + i;
if constexpr (Allen::monitoring_has_evtNo<Parameters>::value) {
- parameters.evtNo[index] = (uint64_t(evt_hi) << 32) + evt_lo;
+ parameters.evtNo[index] = odin.eventNumber();
}
if constexpr (Allen::monitoring_has_runNo<Parameters>::value) {
- parameters.runNo[index] = run_no;
+ parameters.runNo[index] = odin.runNumber();
}
Derived::fill_tuples(parameters, input, index, decision);
}
@@ -363,21 +335,29 @@ void Line<Derived, Parameters>::operator()(
{
const auto* derived_instance = static_cast<const Derived*>(this);
- // Copy post scaler and hash to an output, such that GatherSelections can later
- // perform the postscaling
- Allen::ArgumentOperations::data<typename Parameters::host_post_scaler_t>(arguments)[0] =
- derived_instance->template property<typename Parameters::post_scaler_t>();
- Allen::ArgumentOperations::data<typename Parameters::host_post_scaler_hash_t>(arguments)[0] = m_post_scaler_hash;
- Allen::ArgumentOperations::data<typename Parameters::host_decisions_size_t>(arguments)[0] =
- Derived::get_decisions_size(arguments);
+ // Copy infos needed by GatherSelections to an output:
+ auto& line_data = Allen::ArgumentOperations::data<typename Parameters::host_line_data_t>(arguments)[0];
+ line_data.pre_scaler = derived_instance->template property<typename Parameters::pre_scaler_t>();
+ line_data.pre_scaler_hash = m_pre_scaler_hash;
+ line_data.post_scaler = derived_instance->template property<typename Parameters::post_scaler_t>();
+ line_data.post_scaler_hash = m_post_scaler_hash;
+ line_data.decisions_size = Derived::get_decisions_size(arguments);
+ line_data.event_list = Allen::ArgumentOperations::data<typename Parameters::dev_event_list_t>(arguments);
+ line_data.event_list_size = Allen::ArgumentOperations::size<typename Parameters::dev_event_list_t>(arguments);
+ if constexpr (Allen::has_dev_particle_container<
+ Derived,
+ Allen::Store::device_datatype,
+ Allen::Store::input_datatype>::value) {
+ const auto ptr = static_cast<const Allen::IMultiEventContainer*>(
+ Allen::ArgumentOperations::data<typename Parameters::dev_particle_container_t>(arguments));
+ line_data.particle_container_ptr = const_cast<Allen::IMultiEventContainer*>(ptr);
+ }
+ else {
+ line_data.particle_container_ptr = nullptr;
+ }
// Delay the execution of the line: Pass the parameters
- auto parameters = std::make_tuple(
- derived_instance->make_parameters(1, 1, 0, arguments),
- Allen::ArgumentOperations::size<typename Parameters::dev_event_list_t>(arguments),
- m_pre_scaler_hash,
- arguments,
- derived_instance);
+ auto parameters = std::make_tuple(derived_instance->make_parameters(1, 1, 0, arguments), arguments, derived_instance);
assert(sizeof(type_erased_tuple_t<Derived, Parameters>) == sizeof(parameters));
std::memcpy(
diff --git a/device/selections/lines/SMOG2/include/SMOG2_DiMuonHighMassLine.cuh b/device/selections/lines/SMOG2/include/SMOG2_DiMuonHighMassLine.cuh
index b3beda742159f90e58464b10885e257ffe76daa6..a29a183da795cf399cef65435cb485dc9df98d54 100644
--- a/device/selections/lines/SMOG2/include/SMOG2_DiMuonHighMassLine.cuh
+++ b/device/selections/lines/SMOG2/include/SMOG2_DiMuonHighMassLine.cuh
@@ -32,9 +32,7 @@ namespace SMOG2_dimuon_highmass_line {
HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events;
HOST_INPUT(host_number_of_svs_t, unsigned) host_number_of_svs;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_particle_container_t), char)
host_fn_parameters;
diff --git a/device/selections/lines/SMOG2/include/SMOG2_DiTrack.cuh b/device/selections/lines/SMOG2/include/SMOG2_DiTrack.cuh
index d9eb0646351a727886b1e364fba7f5be9019ebc4..66197b85a2bb24e136df6c6dc12f6f9c36587ffd 100644
--- a/device/selections/lines/SMOG2/include/SMOG2_DiTrack.cuh
+++ b/device/selections/lines/SMOG2/include/SMOG2_DiTrack.cuh
@@ -21,9 +21,7 @@ namespace SMOG2_ditrack_line {
HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events;
HOST_INPUT(host_number_of_svs_t, unsigned) host_number_of_svs;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_particle_container_t), char)
host_fn_parameters;
diff --git a/device/selections/lines/SMOG2/include/SMOG2_KsToPiPi.cuh b/device/selections/lines/SMOG2/include/SMOG2_KsToPiPi.cuh
index a7c3df38469a9f8b937474bcd34d17c8fef353a2..a22e3830ba41bdc1ef2ef8bc1f2ad3097e74bd6c 100644
--- a/device/selections/lines/SMOG2/include/SMOG2_KsToPiPi.cuh
+++ b/device/selections/lines/SMOG2/include/SMOG2_KsToPiPi.cuh
@@ -26,10 +26,7 @@ namespace SMOG2_kstopipi_line {
HOST_INPUT(host_number_of_svs_t, unsigned) host_number_of_svs;
DEVICE_INPUT(dev_particle_container_t, Allen::Views::Physics::MultiEventCompositeParticles) dev_particle_container;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
-
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_particle_container_t), char)
host_fn_parameters;
diff --git a/device/selections/lines/SMOG2/include/SMOG2_MinimumBiasLine.cuh b/device/selections/lines/SMOG2/include/SMOG2_MinimumBiasLine.cuh
index 032166a8a655bde06a66fc2fcaaedeceaf8abc98..e37f53f8e92aa0a1d4b937d0ac8cf1c85bc9c9c9 100644
--- a/device/selections/lines/SMOG2/include/SMOG2_MinimumBiasLine.cuh
+++ b/device/selections/lines/SMOG2/include/SMOG2_MinimumBiasLine.cuh
@@ -23,9 +23,7 @@ namespace SMOG2_minimum_bias_line {
HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events;
HOST_INPUT(host_number_of_reconstructed_velo_tracks_t, unsigned) host_number_of_reconstructed_velo_tracks;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_tracks_container_t), char) host_fn_parameters;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
diff --git a/device/selections/lines/SMOG2/include/SMOG2_SingleMuon.cuh b/device/selections/lines/SMOG2/include/SMOG2_SingleMuon.cuh
index 58afe8f473ba7667ba5ea7ee0a02ca6abffa9236..d88015d770454a41f0d3e1f31f3be8b8db44908a 100644
--- a/device/selections/lines/SMOG2/include/SMOG2_SingleMuon.cuh
+++ b/device/selections/lines/SMOG2/include/SMOG2_SingleMuon.cuh
@@ -19,10 +19,7 @@ namespace SMOG2_single_muon_line {
HOST_INPUT(host_number_of_reconstructed_scifi_tracks_t, unsigned) host_number_of_reconstructed_scifi_tracks;
DEVICE_INPUT(dev_particle_container_t, Allen::Views::Physics::MultiEventBasicParticles) dev_particle_container;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
-
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_particle_container_t), char)
host_fn_parameters;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
diff --git a/device/selections/lines/SMOG2/include/SMOG2_SingleTrack.cuh b/device/selections/lines/SMOG2/include/SMOG2_SingleTrack.cuh
index 5c2947d681f773464f200c3d8cacd927292bf8d3..101dd6042ae95114a4deff729f42cb090e0944d8 100644
--- a/device/selections/lines/SMOG2/include/SMOG2_SingleTrack.cuh
+++ b/device/selections/lines/SMOG2/include/SMOG2_SingleTrack.cuh
@@ -21,9 +21,7 @@ namespace SMOG2_singletrack_line {
HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events;
HOST_INPUT(host_number_of_reconstructed_scifi_tracks_t, unsigned) host_number_of_reconstructed_scifi_tracks;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_particle_container_t), char)
host_fn_parameters;
diff --git a/device/selections/lines/calibration/include/D2KPiLine.cuh b/device/selections/lines/calibration/include/D2KPiLine.cuh
index 9d201bd350f14d0c1f72b80f0a490664acf681f0..08b1cb28c6fb543d9bd4ad5bafb158a9edd01338 100644
--- a/device/selections/lines/calibration/include/D2KPiLine.cuh
+++ b/device/selections/lines/calibration/include/D2KPiLine.cuh
@@ -26,10 +26,7 @@ namespace d2kpi_line {
HOST_INPUT(host_number_of_svs_t, unsigned) host_number_of_svs;
DEVICE_INPUT(dev_particle_container_t, Allen::Views::Physics::MultiEventCompositeParticles) dev_particle_container;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
-
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_particle_container_t), char)
host_fn_parameters;
diff --git a/device/selections/lines/calibration/include/DiMuonMassAlignmentLine.cuh b/device/selections/lines/calibration/include/DiMuonMassAlignmentLine.cuh
index e40339e989c6c09da061150073519196780b8ae0..2151592ee221d9ac87c4becb080f920ce15b6c13 100644
--- a/device/selections/lines/calibration/include/DiMuonMassAlignmentLine.cuh
+++ b/device/selections/lines/calibration/include/DiMuonMassAlignmentLine.cuh
@@ -19,10 +19,7 @@ namespace di_muon_mass_alignment_line {
HOST_INPUT(host_number_of_svs_t, unsigned) host_number_of_svs;
DEVICE_INPUT(dev_particle_container_t, Allen::Views::Physics::MultiEventCompositeParticles) dev_particle_container;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
-
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_particle_container_t), char)
host_fn_parameters;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
diff --git a/device/selections/lines/calibration/include/DisplacedDiMuonMassLine.cuh b/device/selections/lines/calibration/include/DisplacedDiMuonMassLine.cuh
index 8be054a94fef9a152d198aa2dad5fe0bee2282d7..35e95b694572929518388a66acb31eb47417c024 100644
--- a/device/selections/lines/calibration/include/DisplacedDiMuonMassLine.cuh
+++ b/device/selections/lines/calibration/include/DisplacedDiMuonMassLine.cuh
@@ -19,10 +19,7 @@ namespace displaced_di_muon_mass_line {
HOST_INPUT(host_number_of_svs_t, unsigned) host_number_of_svs;
DEVICE_INPUT(dev_particle_container_t, Allen::Views::Physics::MultiEventCompositeParticles) dev_particle_container;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
-
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_particle_container_t), char)
host_fn_parameters;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
diff --git a/device/selections/lines/calibration/include/ODINCalibTrigger.cuh b/device/selections/lines/calibration/include/ODINCalibTrigger.cuh
index 3e39cf0fddb79e0a8bd745d249834194c124e389..02d1a36147d9bb2c45c9978c066f59a070cf29d4 100644
--- a/device/selections/lines/calibration/include/ODINCalibTrigger.cuh
+++ b/device/selections/lines/calibration/include/ODINCalibTrigger.cuh
@@ -18,10 +18,8 @@ namespace odin_calib_line {
struct Parameters {
HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
DEVICE_INPUT(dev_odin_data_t, ODINData) dev_odin_data;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT(host_fn_parameters_t, char) host_fn_parameters;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
diff --git a/device/selections/lines/calibration/include/PassthroughLine.cuh b/device/selections/lines/calibration/include/PassthroughLine.cuh
index 5fb93cd119d76494334c544aea5b49d234f44338..0e9ec6b9ac1b2a60aeb614f80aa566ecd3c92c52 100644
--- a/device/selections/lines/calibration/include/PassthroughLine.cuh
+++ b/device/selections/lines/calibration/include/PassthroughLine.cuh
@@ -19,9 +19,7 @@ namespace passthrough_line {
HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events;
DEVICE_INPUT(dev_number_of_events_t, unsigned) dev_number_of_events;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT(host_fn_parameters_t, char) host_fn_parameters;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
PROPERTY(post_scaler_t, "post_scaler", "Post-scaling factor", float) post_scaler;
diff --git a/device/selections/lines/calibration/include/RICH1Line.cuh b/device/selections/lines/calibration/include/RICH1Line.cuh
index 35c9257955d7f6851d7d98099b6d121e37da944a..6562f00804655f742ad535d387a7514b0707b8ec 100644
--- a/device/selections/lines/calibration/include/RICH1Line.cuh
+++ b/device/selections/lines/calibration/include/RICH1Line.cuh
@@ -19,9 +19,7 @@ namespace rich_1_line {
// Commonly required inputs, outputs and properties
HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
PROPERTY(post_scaler_t, "post_scaler", "Post-scaling factor", float) post_scaler;
PROPERTY(pre_scaler_hash_string_t, "pre_scaler_hash_string", "Pre-scaling hash string", std::string);
diff --git a/device/selections/lines/calibration/include/RICH2Line.cuh b/device/selections/lines/calibration/include/RICH2Line.cuh
index 527f2fb3d29421cb4da65a7a59ef5616e2e3318a..e44068c542bb56f1662570fcb2df53c10140d439 100644
--- a/device/selections/lines/calibration/include/RICH2Line.cuh
+++ b/device/selections/lines/calibration/include/RICH2Line.cuh
@@ -19,9 +19,7 @@ namespace rich_2_line {
// Commonly required inputs, outputs and properties
HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
PROPERTY(post_scaler_t, "post_scaler", "Post-scaling factor", float) post_scaler;
diff --git a/device/selections/lines/charm/include/D2KKLine.cuh b/device/selections/lines/charm/include/D2KKLine.cuh
index f4b0753f5e59cb463692700ff6011a66f7591c3d..b3b7140dd3a14c53c64258681f6bd832692ed688 100644
--- a/device/selections/lines/charm/include/D2KKLine.cuh
+++ b/device/selections/lines/charm/include/D2KKLine.cuh
@@ -26,10 +26,7 @@ namespace d2kk_line {
HOST_INPUT(host_number_of_svs_t, unsigned) host_number_of_svs;
DEVICE_INPUT(dev_particle_container_t, Allen::Views::Physics::MultiEventCompositeParticles) dev_particle_container;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
-
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_particle_container_t), char)
host_fn_parameters;
diff --git a/device/selections/lines/charm/include/D2PiPiLine.cuh b/device/selections/lines/charm/include/D2PiPiLine.cuh
index e5e58ee315094cd1e067d66634962d20aaf6126c..1f54b8cb9d1fafbd1fe9cc133011c60204808a6c 100644
--- a/device/selections/lines/charm/include/D2PiPiLine.cuh
+++ b/device/selections/lines/charm/include/D2PiPiLine.cuh
@@ -26,9 +26,7 @@ namespace d2pipi_line {
HOST_INPUT(host_number_of_svs_t, unsigned) host_number_of_svs;
DEVICE_INPUT(dev_particle_container_t, Allen::Views::Physics::MultiEventCompositeParticles) dev_particle_container;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_particle_container_t), char)
host_fn_parameters;
diff --git a/device/selections/lines/charm/include/TwoKsLine.cuh b/device/selections/lines/charm/include/TwoKsLine.cuh
index a32634f8cd9298f5567084932977e4809e9a499a..e394ab23770f2aa45f49e06ec72cbe0972c97a8d 100644
--- a/device/selections/lines/charm/include/TwoKsLine.cuh
+++ b/device/selections/lines/charm/include/TwoKsLine.cuh
@@ -20,9 +20,7 @@ namespace two_ks_line {
struct Parameters {
HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
PROPERTY(post_scaler_t, "post_scaler", "Post-scaling factor", float) post_scaler;
diff --git a/device/selections/lines/charm/include/TwoTrackMVACharmXSec.cuh b/device/selections/lines/charm/include/TwoTrackMVACharmXSec.cuh
index f415233e590950e813f325dc3ee6ff411c57d8d3..fc9dfe686bbd6e2a05eb2a40a11f3a141210e4ff 100644
--- a/device/selections/lines/charm/include/TwoTrackMVACharmXSec.cuh
+++ b/device/selections/lines/charm/include/TwoTrackMVACharmXSec.cuh
@@ -27,10 +27,7 @@ namespace two_track_mva_charm_xsec_line {
DEVICE_INPUT(dev_particle_container_t, Allen::Views::Physics::MultiEventCompositeParticles) dev_particle_container;
DEVICE_INPUT(dev_two_track_mva_evaluation_t, float) dev_two_track_mva_evaluation;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
-
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_particle_container_t), char)
host_fn_parameters;
diff --git a/device/selections/lines/downstream/include/DownstreamTrackPassThroughLine.cuh b/device/selections/lines/downstream/include/DownstreamTrackPassThroughLine.cuh
index 4b69b0dc60d3121184d6416a44162fbd6239deea..98e1d059a1c5fb6dd4eeb1d07802e77d625bc483 100644
--- a/device/selections/lines/downstream/include/DownstreamTrackPassThroughLine.cuh
+++ b/device/selections/lines/downstream/include/DownstreamTrackPassThroughLine.cuh
@@ -19,9 +19,7 @@ namespace downstream_track_line {
HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
PROPERTY(post_scaler_t, "post_scaler", "Post-scaling factor", float) post_scaler;
diff --git a/device/selections/lines/electron/include/DisplacedDielectronLine.cuh b/device/selections/lines/electron/include/DisplacedDielectronLine.cuh
index ef7e60593bd807b72477ae3c987bb793137f196c..9e74cefac08b281ec9469e0bc1fbcf70aa78b826 100644
--- a/device/selections/lines/electron/include/DisplacedDielectronLine.cuh
+++ b/device/selections/lines/electron/include/DisplacedDielectronLine.cuh
@@ -25,10 +25,7 @@ namespace displaced_dielectron_line {
DEVICE_INPUT(dev_track_isElectron_t, bool) dev_track_isElectron;
DEVICE_INPUT(dev_brem_corrected_pt_t, float) dev_brem_corrected_pt;
// Outputs
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
-
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_particle_container_t), char)
host_fn_parameters;
// Properties
diff --git a/device/selections/lines/electron/include/DisplacedLeptonsLine.cuh b/device/selections/lines/electron/include/DisplacedLeptonsLine.cuh
index 925dd7e21f999a46d10ea5ebde5dbc9a35f72008..8eb8507f80b433a9ba95a77fd75967000f3e8894 100644
--- a/device/selections/lines/electron/include/DisplacedLeptonsLine.cuh
+++ b/device/selections/lines/electron/include/DisplacedLeptonsLine.cuh
@@ -26,9 +26,7 @@ namespace displaced_leptons_line {
DEVICE_INPUT(dev_track_container_t, Allen::Views::Physics::MultiEventBasicParticles) dev_track_container;
DEVICE_INPUT(dev_track_isElectron_t, bool) dev_track_isElectron;
DEVICE_INPUT(dev_brem_corrected_pt_t, float) dev_brem_corrected_pt;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT(host_fn_parameters_t, char) host_fn_parameters;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
PROPERTY(post_scaler_t, "post_scaler", "Post-scaling factor", float) post_scaler;
diff --git a/device/selections/lines/electron/include/HighMassDielectronLine.cuh b/device/selections/lines/electron/include/HighMassDielectronLine.cuh
index 3d63d83d5f46ca42bcb12c3fdf6134c25b4753f4..02895e910d23822644b7fbd52e581cbb8b609254 100644
--- a/device/selections/lines/electron/include/HighMassDielectronLine.cuh
+++ b/device/selections/lines/electron/include/HighMassDielectronLine.cuh
@@ -29,9 +29,7 @@ namespace highmass_dielectron_line {
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_particle_container_t), char)
host_fn_parameters;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
// Properties
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
diff --git a/device/selections/lines/electron/include/LowMassDielectronLine.cuh b/device/selections/lines/electron/include/LowMassDielectronLine.cuh
index c71b3fc1738ed203c641519c3a978e1c7db97eca..549c7df102be6c7373424d38968b5d176949ba90 100644
--- a/device/selections/lines/electron/include/LowMassDielectronLine.cuh
+++ b/device/selections/lines/electron/include/LowMassDielectronLine.cuh
@@ -33,9 +33,7 @@ namespace lowmass_dielectron_line {
// Outputs
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_particle_container_t), char)
host_fn_parameters;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
// Outputs for ROOT tupling
DEVICE_OUTPUT(dev_die_masses_raw_t, float) dev_die_masses_raw;
DEVICE_OUTPUT(dev_die_masses_bremcorr_t, float) dev_die_masses_bremcorr;
diff --git a/device/selections/lines/electron/include/SingleHighETLine.cuh b/device/selections/lines/electron/include/SingleHighETLine.cuh
index 81efa1366ced1342ff4b6a14aff0023a126170d8..cd1f451d18e4d37184bc2a99fbddb60b55b696c2 100644
--- a/device/selections/lines/electron/include/SingleHighETLine.cuh
+++ b/device/selections/lines/electron/include/SingleHighETLine.cuh
@@ -25,9 +25,7 @@ namespace single_high_et_line {
// ECAL
DEVICE_INPUT(dev_brem_ET_t, float) dev_brem_ET;
// Outputs
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT(host_fn_parameters_t, char) host_fn_parameters;
diff --git a/device/selections/lines/electron/include/SingleHighPtElectronLine.cuh b/device/selections/lines/electron/include/SingleHighPtElectronLine.cuh
index 8b5473e2a626f0aef3c16bd32b69aa9441cbdecc..94377f1dfa166598a9006f30215ed814894c8a06 100644
--- a/device/selections/lines/electron/include/SingleHighPtElectronLine.cuh
+++ b/device/selections/lines/electron/include/SingleHighPtElectronLine.cuh
@@ -22,10 +22,7 @@ namespace single_high_pt_electron_line {
MASK_INPUT(dev_event_list_t) dev_event_list;
DEVICE_INPUT(dev_track_isElectron_t, bool) dev_track_isElectron;
DEVICE_INPUT(dev_brem_corrected_pt_t, float) dev_brem_corrected_pt;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
-
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_particle_container_t), char)
host_fn_parameters;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
diff --git a/device/selections/lines/electron/include/TrackElectronMVALine.cuh b/device/selections/lines/electron/include/TrackElectronMVALine.cuh
index 5951dd67d80c719c0b8b3c06fd34b392ba105ae9..338063e8542e6387de3833aeb2b8cd5f48558eef 100644
--- a/device/selections/lines/electron/include/TrackElectronMVALine.cuh
+++ b/device/selections/lines/electron/include/TrackElectronMVALine.cuh
@@ -21,10 +21,7 @@ namespace track_electron_mva_line {
MASK_INPUT(dev_event_list_t) dev_event_list;
DEVICE_INPUT(dev_track_isElectron_t, bool) dev_track_isElectron;
DEVICE_INPUT(dev_brem_corrected_pt_t, float) dev_brem_corrected_pt;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
-
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_particle_container_t), char)
host_fn_parameters;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
diff --git a/device/selections/lines/heavy_ions/include/HeavyIonEventLine.cuh b/device/selections/lines/heavy_ions/include/HeavyIonEventLine.cuh
index 3547227f67adb359c31afdde497eaf798fbd5a45..acf2ce4d9eb798b9f7b83a9db20eec261b86bcfe 100644
--- a/device/selections/lines/heavy_ions/include/HeavyIonEventLine.cuh
+++ b/device/selections/lines/heavy_ions/include/HeavyIonEventLine.cuh
@@ -29,9 +29,7 @@ namespace heavy_ion_event_line {
DEVICE_INPUT(dev_total_ecal_e_t, float) dev_total_ecal_e;
DEVICE_INPUT(dev_pvs_t, PV::Vertex) dev_pvs;
DEVICE_INPUT(dev_number_of_pvs_t, unsigned) dev_number_of_pvs;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT(host_fn_parameters_t, char) host_fn_parameters;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
PROPERTY(post_scaler_t, "post_scaler", "Post-scaling factor", float) post_scaler;
diff --git a/device/selections/lines/inclusive_hadron/include/KsToPiPiLine.cuh b/device/selections/lines/inclusive_hadron/include/KsToPiPiLine.cuh
index 5e4d43fabc085c698593b12294a3df9a8323c79d..9f9286ae5d12b853ff0b9862809caf41238bc92e 100644
--- a/device/selections/lines/inclusive_hadron/include/KsToPiPiLine.cuh
+++ b/device/selections/lines/inclusive_hadron/include/KsToPiPiLine.cuh
@@ -26,10 +26,7 @@ namespace kstopipi_line {
HOST_INPUT(host_number_of_svs_t, unsigned) host_number_of_svs;
DEVICE_INPUT(dev_particle_container_t, Allen::Views::Physics::MultiEventCompositeParticles) dev_particle_container;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
-
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_particle_container_t), char)
host_fn_parameters;
diff --git a/device/selections/lines/inclusive_hadron/include/Lambda2PPiLine.cuh b/device/selections/lines/inclusive_hadron/include/Lambda2PPiLine.cuh
index 3423b43dcf8a732837cf1895e5ffc8a28edaf91e..29650d336aedb2c3e9c524822bcffeba6cbf2e27 100644
--- a/device/selections/lines/inclusive_hadron/include/Lambda2PPiLine.cuh
+++ b/device/selections/lines/inclusive_hadron/include/Lambda2PPiLine.cuh
@@ -20,9 +20,7 @@ namespace lambda2ppi_line {
HOST_INPUT(host_number_of_svs_t, unsigned) host_number_of_svs;
DEVICE_INPUT(dev_particle_container_t, Allen::Views::Physics::MultiEventCompositeParticles) dev_particle_container;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_particle_container_t), char)
host_fn_parameters;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
diff --git a/device/selections/lines/inclusive_hadron/include/LambdaLLDetachedTrackLine.cuh b/device/selections/lines/inclusive_hadron/include/LambdaLLDetachedTrackLine.cuh
index 61053ec06a18c3258d9380dbdb0faced7f78f842..de0c686e38a0cc315ee8e45beec85811cc6f0dae 100644
--- a/device/selections/lines/inclusive_hadron/include/LambdaLLDetachedTrackLine.cuh
+++ b/device/selections/lines/inclusive_hadron/include/LambdaLLDetachedTrackLine.cuh
@@ -25,9 +25,7 @@ namespace lambda_ll_detached_track_line {
HOST_INPUT(host_number_of_svs_t, unsigned) host_number_of_svs;
DEVICE_INPUT(dev_particle_container_t, Allen::Views::Physics::MultiEventCompositeParticles) dev_particle_container;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_particle_container_t), char)
host_fn_parameters;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
diff --git a/device/selections/lines/inclusive_hadron/include/TrackMVALine.cuh b/device/selections/lines/inclusive_hadron/include/TrackMVALine.cuh
index f6b3fd65be10404a6889dfd542cf304088f0907b..5468e46e72a3d151d9daabb7ed9775fb70e91115 100644
--- a/device/selections/lines/inclusive_hadron/include/TrackMVALine.cuh
+++ b/device/selections/lines/inclusive_hadron/include/TrackMVALine.cuh
@@ -19,10 +19,7 @@ namespace track_mva_line {
HOST_INPUT(host_number_of_reconstructed_scifi_tracks_t, unsigned) host_number_of_reconstructed_scifi_tracks;
DEVICE_INPUT(dev_particle_container_t, Allen::Views::Physics::MultiEventBasicParticles) dev_particle_container;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
-
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_particle_container_t), char)
host_fn_parameters;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
diff --git a/device/selections/lines/inclusive_hadron/include/TwoTrackKsLine.cuh b/device/selections/lines/inclusive_hadron/include/TwoTrackKsLine.cuh
index 80f7b19cdacb6213a6cc03e0e338b9f334d0d6d6..651f51d51faf2b9e3519962613b0b25ef3aab436 100644
--- a/device/selections/lines/inclusive_hadron/include/TwoTrackKsLine.cuh
+++ b/device/selections/lines/inclusive_hadron/include/TwoTrackKsLine.cuh
@@ -18,10 +18,7 @@ namespace two_track_line_ks {
struct Parameters {
HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
-
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
PROPERTY(post_scaler_t, "post_scaler", "Post-scaling factor", float) post_scaler;
PROPERTY(pre_scaler_hash_string_t, "pre_scaler_hash_string", "Pre-scaling hash string", std::string);
diff --git a/device/selections/lines/inclusive_hadron/include/TwoTrackMVALine.cuh b/device/selections/lines/inclusive_hadron/include/TwoTrackMVALine.cuh
index fc5333069130eb01eb37583c08f836b5ced953b5..91ba1129858d265399a1e307bd0a5e28a4faed80 100644
--- a/device/selections/lines/inclusive_hadron/include/TwoTrackMVALine.cuh
+++ b/device/selections/lines/inclusive_hadron/include/TwoTrackMVALine.cuh
@@ -21,10 +21,7 @@ namespace two_track_mva_line {
DEVICE_INPUT(dev_particle_container_t, Allen::Views::Physics::MultiEventCompositeParticles) dev_particle_container;
DEVICE_INPUT(dev_two_track_mva_evaluation_t, float) dev_two_track_mva_evaluation;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
-
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_particle_container_t), char)
host_fn_parameters;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
diff --git a/device/selections/lines/inclusive_hadron/include/XiOmegaLLLLine.cuh b/device/selections/lines/inclusive_hadron/include/XiOmegaLLLLine.cuh
index fa3fcb1804444c643e661521b916ae435f79d4c1..b2dc6a7f5ae44e040e411b7d638b9ee878904338 100644
--- a/device/selections/lines/inclusive_hadron/include/XiOmegaLLLLine.cuh
+++ b/device/selections/lines/inclusive_hadron/include/XiOmegaLLLLine.cuh
@@ -25,9 +25,7 @@ namespace xi_omega_lll_line {
HOST_INPUT(host_number_of_svs_t, unsigned) host_number_of_svs;
DEVICE_INPUT(dev_particle_container_t, Allen::Views::Physics::MultiEventCompositeParticles) dev_particle_container;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_particle_container_t), char)
host_fn_parameters;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
diff --git a/device/selections/lines/monitoring/include/BeamCrossingLine.cuh b/device/selections/lines/monitoring/include/BeamCrossingLine.cuh
index 1b8938aae1b5630f4aba76e8d33ccf2e62ed50a5..fcbfcfb2d4402dec921fe9a4d0fc61000c8885c1 100644
--- a/device/selections/lines/monitoring/include/BeamCrossingLine.cuh
+++ b/device/selections/lines/monitoring/include/BeamCrossingLine.cuh
@@ -18,9 +18,7 @@ namespace beam_crossing_line {
struct Parameters {
HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
DEVICE_INPUT(dev_odin_data_t, ODINData) dev_odin_data;
HOST_OUTPUT(host_fn_parameters_t, char) host_fn_parameters;
diff --git a/device/selections/lines/monitoring/include/BeamGasLine.cuh b/device/selections/lines/monitoring/include/BeamGasLine.cuh
index e3086fe7a8511b712261d9e32ea49cb13dee105b..d092640f99de0ba46588e22396c7715358e22559 100644
--- a/device/selections/lines/monitoring/include/BeamGasLine.cuh
+++ b/device/selections/lines/monitoring/include/BeamGasLine.cuh
@@ -26,9 +26,7 @@ namespace beam_gas_line {
DEVICE_INPUT(dev_offsets_velo_tracks_t, unsigned) dev_offsets_velo_tracks;
DEVICE_INPUT(dev_offsets_velo_track_hit_number_t, unsigned) dev_offsets_velo_track_hit_number;
DEVICE_INPUT(dev_odin_data_t, ODINData) dev_odin_data;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT(host_fn_parameters_t, char) host_fn_parameters;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
PROPERTY(post_scaler_t, "post_scaler", "Post-scaling factor", float) post_scaler;
diff --git a/device/selections/lines/monitoring/include/CaloDigitsMinADCLine.cuh b/device/selections/lines/monitoring/include/CaloDigitsMinADCLine.cuh
index 99968e830293aecb9a2f1d82de2aa436e0358547..267942dcb0c043e41d4b33e8b6e9c80d493c32d7 100644
--- a/device/selections/lines/monitoring/include/CaloDigitsMinADCLine.cuh
+++ b/device/selections/lines/monitoring/include/CaloDigitsMinADCLine.cuh
@@ -25,9 +25,7 @@ namespace calo_digits_minADC {
DEVICE_INPUT(dev_ecal_digits_t, CaloDigit) dev_ecal_digits;
DEVICE_INPUT(dev_ecal_digits_offsets_t, unsigned) dev_ecal_digits_offsets;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT(host_fn_parameters_t, char) host_fn_parameters;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
diff --git a/device/selections/lines/monitoring/include/ODINEventAndOrbitLine.cuh b/device/selections/lines/monitoring/include/ODINEventAndOrbitLine.cuh
index cdbc05861c17c8b9b3a058b0e55176ff24370b78..5741855a1e9eab17af63602f3eecc718d718bada 100644
--- a/device/selections/lines/monitoring/include/ODINEventAndOrbitLine.cuh
+++ b/device/selections/lines/monitoring/include/ODINEventAndOrbitLine.cuh
@@ -19,9 +19,7 @@ namespace odin_event_and_orbit_line {
struct Parameters {
HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
DEVICE_INPUT(dev_odin_data_t, ODINData) dev_odin_data;
HOST_OUTPUT(host_fn_parameters_t, char) host_fn_parameters;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
diff --git a/device/selections/lines/monitoring/include/ODINEventTypeLine.cuh b/device/selections/lines/monitoring/include/ODINEventTypeLine.cuh
index 54585c870631b48257bcf195742c7c63e6271151..f02ba05e35a3a0dc301c421d0bb424456bb600af 100644
--- a/device/selections/lines/monitoring/include/ODINEventTypeLine.cuh
+++ b/device/selections/lines/monitoring/include/ODINEventTypeLine.cuh
@@ -19,9 +19,7 @@ namespace odin_event_type_line {
struct Parameters {
HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
DEVICE_INPUT(dev_odin_data_t, ODINData) dev_odin_data;
HOST_OUTPUT(host_fn_parameters_t, char) host_fn_parameters;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
diff --git a/device/selections/lines/monitoring/include/ODINEventTypeWithDecodingLine.cuh b/device/selections/lines/monitoring/include/ODINEventTypeWithDecodingLine.cuh
index e5c9c5d52ead63574702d984381e6b01f03b73df..91971639a876f3434f6cf4216bb3191449d5cfaa 100644
--- a/device/selections/lines/monitoring/include/ODINEventTypeWithDecodingLine.cuh
+++ b/device/selections/lines/monitoring/include/ODINEventTypeWithDecodingLine.cuh
@@ -21,9 +21,7 @@ namespace odin_event_type_with_decoding_line {
struct Parameters {
HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
DEVICE_INPUT(dev_odin_data_t, ODINData) dev_odin_data;
// dummy inputs to establish a dependence on the decoding algorithms
DEVICE_INPUT(dev_sorted_velo_cluster_container_t, char) dev_sorted_velo_cluster_container;
diff --git a/device/selections/lines/monitoring/include/PlumeActivityLine.cuh b/device/selections/lines/monitoring/include/PlumeActivityLine.cuh
index 9a15fa4458830c8798f3b3b12f505a25165b0413..57a7955d739428a8b4cd2395aff111d4f6eebcca 100644
--- a/device/selections/lines/monitoring/include/PlumeActivityLine.cuh
+++ b/device/selections/lines/monitoring/include/PlumeActivityLine.cuh
@@ -22,10 +22,7 @@ namespace plume_activity_line {
DEVICE_INPUT(dev_plume_t, Plume_) dev_plume;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
-
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT(host_fn_parameters_t, char) host_fn_parameters;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
diff --git a/device/selections/lines/monitoring/include/TTrackCosmicLine.cuh b/device/selections/lines/monitoring/include/TTrackCosmicLine.cuh
index 456265b07da7eb0bb5207ff6bcb74a68c02b860b..ec9954026931ed79ce5d4bc4398fdc01e34a5f9d 100644
--- a/device/selections/lines/monitoring/include/TTrackCosmicLine.cuh
+++ b/device/selections/lines/monitoring/include/TTrackCosmicLine.cuh
@@ -21,9 +21,7 @@ namespace t_track_cosmic_line {
DEVICE_INPUT(dev_seeding_offsets_t, unsigned) dev_seeding_offsets;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT(host_fn_parameters_t, char) host_fn_parameters;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
diff --git a/device/selections/lines/monitoring/include/VeloClustersMicroBiasLine.cuh b/device/selections/lines/monitoring/include/VeloClustersMicroBiasLine.cuh
index f409321d062dedf01f31137b2b195d39b53c77ca..48609600e23e38e4d81f81a65dc0a26d326f1cf1 100644
--- a/device/selections/lines/monitoring/include/VeloClustersMicroBiasLine.cuh
+++ b/device/selections/lines/monitoring/include/VeloClustersMicroBiasLine.cuh
@@ -22,10 +22,7 @@ namespace velo_clusters_micro_bias_line {
DEVICE_INPUT(dev_offsets_estimated_input_size_t, unsigned) dev_offsets_estimated_input_size;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
-
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT(host_fn_parameters_t, char) host_fn_parameters;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
diff --git a/device/selections/lines/monitoring/include/VeloMicroBiasLine.cuh b/device/selections/lines/monitoring/include/VeloMicroBiasLine.cuh
index 3075fead85de95798afb387b1c03d7918c91fcf8..7b99acb28bbe26ed8573a48b4a96f8cdaacd228c 100644
--- a/device/selections/lines/monitoring/include/VeloMicroBiasLine.cuh
+++ b/device/selections/lines/monitoring/include/VeloMicroBiasLine.cuh
@@ -22,10 +22,7 @@ namespace velo_micro_bias_line {
MASK_INPUT(dev_event_list_t) dev_event_list;
DEVICE_INPUT(dev_offsets_velo_tracks_t, unsigned) dev_offsets_velo_tracks;
DEVICE_INPUT(dev_offsets_velo_track_hit_number_t, unsigned) dev_offsets_velo_track_hit_number;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
-
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT(host_fn_parameters_t, char) host_fn_parameters;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
PROPERTY(post_scaler_t, "post_scaler", "Post-scaling factor", float) post_scaler;
diff --git a/device/selections/lines/monitoring/include/ZRange_MaterialVertexSeedLine.cuh b/device/selections/lines/monitoring/include/ZRange_MaterialVertexSeedLine.cuh
index 016fc3550af422f3960efa6e00b35a6106208aa9..e3a53788aa695dc39ce4de48c7415259a642e755 100644
--- a/device/selections/lines/monitoring/include/ZRange_MaterialVertexSeedLine.cuh
+++ b/device/selections/lines/monitoring/include/ZRange_MaterialVertexSeedLine.cuh
@@ -11,9 +11,7 @@ namespace z_range_materialvertex_seed_line {
MASK_INPUT(dev_event_list_t) dev_event_list;
HOST_INPUT(host_total_number_of_interaction_seeds_t, unsigned) host_total_number_of_interactions_seeds;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT(host_fn_parameters_t, char) host_fn_parameters;
DEVICE_INPUT(dev_consolidated_interaction_seeds_t, float3) dev_consolidated_interaction_seeds;
diff --git a/device/selections/lines/muon/include/DiMuonDrellYanLine.cuh b/device/selections/lines/muon/include/DiMuonDrellYanLine.cuh
index d90e8a3a86418f6cfb5e8dd7f0fdb6605ba6722b..4f82922ea1bd7910e8afdb85c335b49e2e182a5f 100644
--- a/device/selections/lines/muon/include/DiMuonDrellYanLine.cuh
+++ b/device/selections/lines/muon/include/DiMuonDrellYanLine.cuh
@@ -25,9 +25,7 @@ namespace di_muon_drell_yan_line {
HOST_INPUT(host_number_of_svs_t, unsigned) host_number_of_svs;
DEVICE_INPUT(dev_particle_container_t, Allen::Views::Physics::MultiEventCompositeParticles) dev_particle_container;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_particle_container_t), char)
host_fn_parameters;
diff --git a/device/selections/lines/muon/include/DiMuonMassLine.cuh b/device/selections/lines/muon/include/DiMuonMassLine.cuh
index 2b3d845a8142b06a3cd305f9752a60fb64fd3fe8..39e16c38fe1d7afd6c0049f1d886acb0ccad4d96 100644
--- a/device/selections/lines/muon/include/DiMuonMassLine.cuh
+++ b/device/selections/lines/muon/include/DiMuonMassLine.cuh
@@ -24,9 +24,7 @@ namespace di_muon_mass_line {
HOST_INPUT(host_number_of_svs_t, unsigned) host_number_of_svs;
DEVICE_INPUT(dev_particle_container_t, Allen::Views::Physics::MultiEventCompositeParticles) dev_particle_container;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_particle_container_t), char)
host_fn_parameters;
diff --git a/device/selections/lines/muon/include/DiMuonNoIPLine.cuh b/device/selections/lines/muon/include/DiMuonNoIPLine.cuh
index bd7e66f48b1f2e430c3f15fa9cad46833c81e8a7..4042c1ba880fb1806312c9494a2c658a984e41a8 100644
--- a/device/selections/lines/muon/include/DiMuonNoIPLine.cuh
+++ b/device/selections/lines/muon/include/DiMuonNoIPLine.cuh
@@ -45,9 +45,7 @@ namespace di_muon_no_ip_line {
HOST_INPUT(host_number_of_svs_t, unsigned) host_number_of_svs;
DEVICE_INPUT(dev_particle_container_t, Allen::Views::Physics::MultiEventCompositeParticles) dev_particle_container;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_particle_container_t), char)
host_fn_parameters;
diff --git a/device/selections/lines/muon/include/DiMuonSoftLine.cuh b/device/selections/lines/muon/include/DiMuonSoftLine.cuh
index 04ed89925dc18c3d49c018a3ced2a9b111367ea7..e6b3e3cfff97824c2520aeeb6ab36a289061d8a7 100644
--- a/device/selections/lines/muon/include/DiMuonSoftLine.cuh
+++ b/device/selections/lines/muon/include/DiMuonSoftLine.cuh
@@ -19,10 +19,7 @@ namespace di_muon_soft_line {
HOST_INPUT(host_number_of_svs_t, unsigned) host_number_of_svs;
DEVICE_INPUT(dev_particle_container_t, Allen::Views::Physics::MultiEventCompositeParticles) dev_particle_container;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
-
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_particle_container_t), char)
host_fn_parameters;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
diff --git a/device/selections/lines/muon/include/DiMuonTrackEffLine.cuh b/device/selections/lines/muon/include/DiMuonTrackEffLine.cuh
index 3864c9801da9da99c69208b95a0c8249a3a9782e..0139c8a18b8bd68b18ea954a390091f47f8410f6 100644
--- a/device/selections/lines/muon/include/DiMuonTrackEffLine.cuh
+++ b/device/selections/lines/muon/include/DiMuonTrackEffLine.cuh
@@ -19,10 +19,7 @@ namespace di_muon_track_eff_line {
HOST_INPUT(host_number_of_svs_t, unsigned) host_number_of_svs;
DEVICE_INPUT(dev_particle_container_t, Allen::Views::Physics::MultiEventCompositeParticles) dev_particle_container;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
-
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_particle_container_t), char)
host_fn_parameters;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
diff --git a/device/selections/lines/muon/include/DisplacedDiMuonLine.cuh b/device/selections/lines/muon/include/DisplacedDiMuonLine.cuh
index 4b1c96d64910dd7c12aee4fe23b50963dcd0a81d..c5b6e386a4e4242f2d6ceea66badea81db31f540 100644
--- a/device/selections/lines/muon/include/DisplacedDiMuonLine.cuh
+++ b/device/selections/lines/muon/include/DisplacedDiMuonLine.cuh
@@ -24,10 +24,7 @@ namespace displaced_di_muon_line {
HOST_INPUT(host_number_of_svs_t, unsigned) host_number_of_svs;
DEVICE_INPUT(dev_particle_container_t, Allen::Views::Physics::MultiEventCompositeParticles) dev_particle_container;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
-
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_particle_container_t), char)
host_fn_parameters;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
diff --git a/device/selections/lines/muon/include/LowPtDiMuonLine.cuh b/device/selections/lines/muon/include/LowPtDiMuonLine.cuh
index 98f6a5ddc2a7128d13261eae8525ec8bb51ec6dd..8c1dcfd5254ef8a2ec5b8b1a8fdf1e42717a971f 100644
--- a/device/selections/lines/muon/include/LowPtDiMuonLine.cuh
+++ b/device/selections/lines/muon/include/LowPtDiMuonLine.cuh
@@ -19,10 +19,7 @@ namespace low_pt_di_muon_line {
HOST_INPUT(host_number_of_svs_t, unsigned) host_number_of_svs;
DEVICE_INPUT(dev_particle_container_t, Allen::Views::Physics::MultiEventCompositeParticles) dev_particle_container;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
-
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_particle_container_t), char)
host_fn_parameters;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
diff --git a/device/selections/lines/muon/include/LowPtMuonLine.cuh b/device/selections/lines/muon/include/LowPtMuonLine.cuh
index dca1db2c3bf92a0a27bc37386502c1f17627b873..65dca070d084368a01f683d08b3a556065b427bb 100644
--- a/device/selections/lines/muon/include/LowPtMuonLine.cuh
+++ b/device/selections/lines/muon/include/LowPtMuonLine.cuh
@@ -19,10 +19,7 @@ namespace low_pt_muon_line {
HOST_INPUT(host_number_of_reconstructed_scifi_tracks_t, unsigned) host_number_of_reconstructed_scifi_tracks;
DEVICE_INPUT(dev_particle_container_t, Allen::Views::Physics::MultiEventBasicParticles) dev_particle_container;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
-
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_particle_container_t), char)
host_fn_parameters;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
diff --git a/device/selections/lines/muon/include/OneMuonTrackLine.cuh b/device/selections/lines/muon/include/OneMuonTrackLine.cuh
index 20be2b7f6edf7b06022c1b36509ecc17ce55c8d1..40c8555bbf9f85dc06628996f359c38a75b2bd11 100644
--- a/device/selections/lines/muon/include/OneMuonTrackLine.cuh
+++ b/device/selections/lines/muon/include/OneMuonTrackLine.cuh
@@ -20,9 +20,7 @@ namespace one_muon_track_line {
HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events;
HOST_INPUT(host_muon_total_number_of_tracks_t, unsigned) host_muon_total_number_of_tracks;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT(host_fn_parameters_t, char) host_fn_parameters;
DEVICE_INPUT(dev_number_of_events_t, unsigned) dev_number_of_events;
diff --git a/device/selections/lines/muon/include/SingleHighPtMuonLine.cuh b/device/selections/lines/muon/include/SingleHighPtMuonLine.cuh
index 1e855ff5b568e02e00ad24184962e63fe1781063..05caf072670d55a21062767cbdad9379fdbfd637 100644
--- a/device/selections/lines/muon/include/SingleHighPtMuonLine.cuh
+++ b/device/selections/lines/muon/include/SingleHighPtMuonLine.cuh
@@ -19,10 +19,7 @@ namespace single_high_pt_muon_line {
HOST_INPUT(host_number_of_reconstructed_scifi_tracks_t, unsigned) host_number_of_reconstructed_scifi_tracks;
DEVICE_INPUT(dev_particle_container_t, Allen::Views::Physics::MultiEventBasicParticles) dev_particle_container;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
-
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_particle_container_t), char)
host_fn_parameters;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
diff --git a/device/selections/lines/muon/include/SingleHighPtMuonLineNoMuID.cuh b/device/selections/lines/muon/include/SingleHighPtMuonLineNoMuID.cuh
index 9338f41bcfa29639c5eb681d7163c52b074d7b0d..d534dfa2f096be25762bb24dd68131259c0a1a17 100644
--- a/device/selections/lines/muon/include/SingleHighPtMuonLineNoMuID.cuh
+++ b/device/selections/lines/muon/include/SingleHighPtMuonLineNoMuID.cuh
@@ -19,10 +19,7 @@ namespace single_high_pt_muon_no_muid_line {
HOST_INPUT(host_number_of_reconstructed_scifi_tracks_t, unsigned) host_number_of_reconstructed_scifi_tracks;
DEVICE_INPUT(dev_particle_container_t, Allen::Views::Physics::MultiEventBasicParticles) dev_particle_container;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
-
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_particle_container_t), char)
host_fn_parameters;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
diff --git a/device/selections/lines/muon/include/TrackMuonMVALine.cuh b/device/selections/lines/muon/include/TrackMuonMVALine.cuh
index 400b6b73db07ff7edc44aa6d4b9726ed7265ee7a..73c68d5500edd43b1645f3ce565ea3273f5c9a2f 100644
--- a/device/selections/lines/muon/include/TrackMuonMVALine.cuh
+++ b/device/selections/lines/muon/include/TrackMuonMVALine.cuh
@@ -19,10 +19,7 @@ namespace track_muon_mva_line {
HOST_INPUT(host_number_of_reconstructed_scifi_tracks_t, unsigned) host_number_of_reconstructed_scifi_tracks;
DEVICE_INPUT(dev_particle_container_t, Allen::Views::Physics::MultiEventBasicParticles) dev_particle_container;
MASK_INPUT(dev_event_list_t) dev_event_list;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
-
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_particle_container_t), char)
host_fn_parameters;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
diff --git a/device/selections/lines/photon/include/SingleCaloCluster.cuh b/device/selections/lines/photon/include/SingleCaloCluster.cuh
index af2d32cd3bec325301cb7322d9888e70e7e8d59b..b36ae70e7214b00fe5169a99ff762572554a1813 100755
--- a/device/selections/lines/photon/include/SingleCaloCluster.cuh
+++ b/device/selections/lines/photon/include/SingleCaloCluster.cuh
@@ -26,9 +26,7 @@ namespace single_calo_cluster_line {
DEVICE_INPUT(dev_particle_container_t, Allen::Views::Physics::MultiEventNeutralBasicParticles)
dev_particle_container;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_particle_container_t), char)
host_fn_parameters;
diff --git a/device/selections/lines/photon/include/TwoCaloClusters.cuh b/device/selections/lines/photon/include/TwoCaloClusters.cuh
index 85c34463b3acb21e8ec678b68b88de5c2f6de649..088057543f99b957b570183fcd85d50948d656a6 100644
--- a/device/selections/lines/photon/include/TwoCaloClusters.cuh
+++ b/device/selections/lines/photon/include/TwoCaloClusters.cuh
@@ -36,11 +36,7 @@ namespace two_calo_clusters_line {
dev_cluster_particle_container;
DEVICE_INPUT(dev_number_of_pvs_t, unsigned) dev_number_of_pvs;
- HOST_OUTPUT(host_decisions_size_t, unsigned) host_decisions_size;
-
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
-
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_particle_container_t), char)
host_fn_parameters;
diff --git a/device/utils/scaler/include/DeterministicScaler.cuh b/device/utils/scaler/include/DeterministicScaler.cuh
index 5a30760c687416ddc74b69c58c57f9a1ea188c60..985c3fa11196b55e5d9cad56fcd5c02d83d3229d 100644
--- a/device/utils/scaler/include/DeterministicScaler.cuh
+++ b/device/utils/scaler/include/DeterministicScaler.cuh
@@ -95,23 +95,4 @@ namespace {
auto x = mix64(mix32(mix64(initial_value, gps_time_hi, gps_time_lo), run_number), evt_number_hi, evt_number_lo);
return x < accept_threshold;
}
-
- __device__ inline void deterministic_post_scaler(
- const unsigned initial_value,
- const float scale_factor,
- const int n_candidates,
- bool* results,
- const uint32_t run_number,
- const uint32_t evt_number_hi,
- const uint32_t evt_number_lo,
- const uint32_t gps_time_hi,
- const uint32_t gps_time_lo)
- {
- if (!deterministic_scaler(
- initial_value, scale_factor, run_number, evt_number_hi, evt_number_lo, gps_time_hi, gps_time_lo)) {
- for (auto i_cand = 0; i_cand < n_candidates; ++i_cand) {
- results[i_cand] = 0;
- }
- }
- }
} // namespace
diff --git a/doc/develop/selections.rst b/doc/develop/selections.rst
index d40e194d639658ac34eed64709c13fc049c5155e..cd7db52c610984a7667e41867cf544e11eb05cf6 100644
--- a/doc/develop/selections.rst
+++ b/doc/develop/selections.rst
@@ -5,6 +5,16 @@ Writing selections
This tutorial will cover adding trigger selections to Allen using the
main reconstruction sequence.
+Line execution
+^^^^^^^^^^^^^^^^^^^^^^^
+
+A line is a function that take some object list as input and returns a decision (boolean) for each object.
+The decision can be prescaled and postscaled per event. The prescaler runs before the line, if the decision
+of the prescaler is false, the line function will not be executed for that event. The postscaler runs after the line,
+it will not affect the line function itself but only the output decision. This distinction is important if the line is
+used for monitoring or filling tuples and for performances. If a line need to run over a lot of objects, tuning
+the prescaler can make a difference in the throughput of the line.
+
Types of selections
^^^^^^^^^^^^^^^^^^^^^^^
Selections are fully configurable algorithms in Allen. Lines that select events
@@ -80,10 +90,6 @@ Event list to which the selection is applied::
MASK_INPUT(dev_event_list_t);
-Size of the decision object::
-
- HOST_OUTPUT(host_decisions_size_t, unsigned), host_decisions_size;
-
Type-erased parameters to be passed to the line functions for delayed line processing::
HOST_OUTPUT(host_fn_parameters_t, char) host_fn_parameters;
@@ -92,13 +98,9 @@ In case that the selection algorithm requires a `dev_particle_container_t`, then
HOST_OUTPUT_WITH_DEPENDENCIES(host_fn_parameters_t, DEPENDENCIES(dev_particle_container_t), char) host_fn_parameters;
-Post-scaler factor, such that an upcoming algorithm (usually `gather_selections_t`) can do the post-scaling::
-
- HOST_OUTPUT(host_post_scaler_t, float), host_post_scaler;
-
-Hash resulting from applying the hash function to the property "post_scaler_hash_string". Needed such that an upcoming algorithm can do the post-scaling::
+Line data that will be passed to an upcoming algorithm (usually `gather_selections_t`). It contains copies of various properties and inputs, such as the pre and post scaler::
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t), host_post_scaler_hash;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
Pre-scaling factor::
@@ -195,9 +197,7 @@ header.
// Commonly required inputs, outputs and properties
HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events;
MASK_INPUT(dev_event_list_t);
- HOST_OUTPUT(host_decisions_size_t, unsigned), host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
PROPERTY(post_scaler_t, "post_scaler", "Post-scaling factor", float) post_scaler;
PROPERTY(pre_scaler_hash_string_t, "pre_scaler_hash_string", "Pre-scaling hash string", std::string)
@@ -267,11 +267,7 @@ secondary vertices with no postscale. This line inherits from `CompositeParticle
// Commonly required inputs, outputs and properties
HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events;
MASK_INPUT(dev_event_list_t);
- HOST_OUTPUT(host_decisions_size_t, unsigned), host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
PROPERTY(post_scaler_t, "post_scaler", "Post-scaling factor", float) post_scaler;
PROPERTY(pre_scaler_hash_string_t, "pre_scaler_hash_string", "Pre-scaling hash string", std::string)
@@ -351,12 +347,8 @@ The header `monitoring/include/VeloMicroBiasLine.cuh <https://gitlab.cern.ch/lhc
// Commonly required inputs, outputs and properties
HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events;
MASK_INPUT(dev_event_list_t);
- HOST_OUTPUT(host_decisions_size_t, unsigned), host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT(host_fn_parameters_t, char) host_fn_parameters;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
PROPERTY(post_scaler_t, "post_scaler", "Post-scaling factor", float) post_scaler;
PROPERTY(pre_scaler_hash_string_t, "pre_scaler_hash_string", "Pre-scaling hash string", std::string)
@@ -437,12 +429,8 @@ The header `ExampleOneVeloTrackLine.cuh` is as follows:
// Commonly required inputs, outputs and properties
HOST_INPUT(host_number_of_events_t, unsigned) host_number_of_events;
MASK_INPUT(dev_event_list_t);
- HOST_OUTPUT(host_decisions_size_t, unsigned), host_decisions_size;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
+ HOST_OUTPUT(host_line_data_t, LineData) host_line_data;
HOST_OUTPUT(host_fn_parameters_t, char) host_fn_parameters;
- HOST_OUTPUT(host_post_scaler_t, float) host_post_scaler;
- HOST_OUTPUT(host_post_scaler_hash_t, uint32_t) host_post_scaler_hash;
PROPERTY(pre_scaler_t, "pre_scaler", "Pre-scaling factor", float) pre_scaler;
PROPERTY(post_scaler_t, "post_scaler", "Post-scaling factor", float) post_scaler;
PROPERTY(pre_scaler_hash_string_t, "pre_scaler_hash_string", "Pre-scaling hash string", std::string)
diff --git a/stream/store/include/ArgumentOps.cuh b/stream/store/include/ArgumentOps.cuh
index e886984c43dc8f6d90c6b726aa5c3bf22709053e..754af988ff7898a2ea61fffea076d8f7fa415b06 100644
--- a/stream/store/include/ArgumentOps.cuh
+++ b/stream/store/include/ArgumentOps.cuh
@@ -249,9 +249,8 @@ namespace Allen {
void store_contiguous_async(
const Args& arguments,
const Allen::Context& context,
- bool fill_if_empty_container = false,
- int fill_value = 0,
- int fill_count = 1)
+ bool skip_if_empty_container = false,
+ int skip_count = 1)
{
auto container = arguments.template get<A>();
auto aggregate = arguments.template input_aggregate<B>();
@@ -275,9 +274,8 @@ namespace Allen {
Allen::copy_async(container, aggregate.get(i), context, kind, aggregate.size(i), container_offset);
container_offset += aggregate.size(i);
}
- else if (fill_if_empty_container) {
- Allen::memset_async<A>(arguments, fill_value, context, fill_count, container_offset);
- container_offset += fill_count;
+ else if (skip_if_empty_container) {
+ container_offset += skip_count;
}
}
}