Skip to content
Snippets Groups Projects

Compare revisions

Changes are shown as if the source revision was being merged into the target revision. Learn more about comparing revisions.

Source

Select target project
No results found

Target

Select target project
  • pvfinder/inference-engine
  • rmatev/Allen
  • suali/Allen
  • mstahl/Allen
  • roneil/Allen
  • graemes/Allen
  • cburr/Allen
  • jonrob/Allen
  • bjashal/Allen-HIP
  • dcampora/MiniAllen
  • brij/Allen
  • raaij/cuda_hlt
  • bsm-fleet/cuda_hlt
  • abrearod/cuda_hlt
  • aalvesju/cuda_hlt
  • lhcb/Allen
16 results
Show changes
Commits on Source (55)
......@@ -230,6 +230,21 @@ namespace Allen {
struct has_monitoring_types<T, std::void_t<typename T::monitoring_types>> : std::true_type {
};
template<typename T1, typename = void>
struct monitoring_has_evtNo : std::false_type {
};
template<typename T1>
struct monitoring_has_evtNo<T1, std::void_t<typename T1::evtNo_t>> : std::true_type {
};
template<typename T1, typename = void>
struct monitoring_has_runNo : std::false_type {
};
template<typename T1>
struct monitoring_has_runNo<T1, std::void_t<typename T1::runNo_t>> : std::true_type {
};
template<typename T>
void initialize_algorithm(T& alg)
{
......
......@@ -28,7 +28,7 @@ file(MAKE_DIRECTORY ${ALLEN_ALGORITHMS_DIR})
if(CMAKE_TOOLCHAIN_FILE)
set(PARSER_ENV PYTHONPATH=$ENV{PYTHONPATH}:${PROJECT_SOURCE_DIR}/scripts LD_LIBRARY_PATH=${LIBCLANG_LIBDIR}:$ENV{LD_LIBRARY_PATH})
else()
set(PARSER_ENV LD_LIBRARY_PATH=${LIBCLANG_LIBDIR}:$ENV{LD_LIBRARY_PATH})
set(PARSER_ENV PYTHONPATH=$ENV{PYTHONPATH}:${LIBCLANG_LIBDIR}/python${Python_VERSION_MAJOR}.${Python_VERSION_MINOR}/site-packages LD_LIBRARY_PATH=${LIBCLANG_LIBDIR}:$ENV{LD_LIBRARY_PATH})
endif()
# Parse Allen algorithms
......
......@@ -312,6 +312,12 @@ __device__ void process_line(
decisions[index] = decision;
if constexpr (Allen::has_enable_monitoring<Parameters>::value) {
if (parameters.enable_monitoring) {
if constexpr (Allen::monitoring_has_evtNo<Parameters>::value) {
parameters.evtNo[index] = (uint64_t(evt_hi) << 32) + evt_lo;
}
if constexpr (Allen::monitoring_has_runNo<Parameters>::value) {
parameters.runNo[index] = run_no;
}
Derived::monitor(parameters, input, index, decision);
}
}
......
/*****************************************************************************\
* (c) Copyright 2018-2020 CERN for the benefit of the LHCb Collaboration *
\*****************************************************************************/
* (c) Copyright 2018-2020 CERN for the benefit of the LHCb Collaboration *
\*****************************************************************************/
#pragma once
#include "AlgorithmTypes.cuh"
......@@ -31,10 +31,24 @@ namespace track_mva_line {
PROPERTY(param3_t, "param3", "param3 description", float) param3;
PROPERTY(alpha_t, "alpha", "alpha description", float) alpha;
PROPERTY(minBPVz_t, "minBPVz", "minimum z for the best associated primary vertex", float) minBPVz;
DEVICE_OUTPUT(pt_t, float) pt;
DEVICE_OUTPUT(ipchi2_t, float) ipchi2;
DEVICE_OUTPUT(evtNo_t, uint64_t) evtNo;
DEVICE_OUTPUT(runNo_t, unsigned) runNo;
PROPERTY(enable_monitoring_t, "enable_monitoring", "Enable line monitoring", bool) enable_monitoring;
};
struct track_mva_line_t : public SelectionAlgorithm, Parameters, OneTrackLine<track_mva_line_t, Parameters> {
__device__ static bool select(const Parameters& ps, std::tuple<const Allen::Views::Physics::BasicParticle> input);
__device__ static void monitor(
const Parameters& parameters,
std::tuple<const Allen::Views::Physics::BasicParticle> input,
unsigned index,
bool sel);
using monitoring_types = std::tuple<pt_t, ipchi2_t, evtNo_t, runNo_t>;
private:
Property<pre_scaler_t> m_pre_scaler {this, 1.f};
......@@ -50,5 +64,7 @@ namespace track_mva_line {
Property<param3_t> m_param3 {this, 1.248f};
Property<alpha_t> m_alpha {this, 296.f * Gaudi::Units::MeV}; // tuned to about 330 kHz (modulo GEC)
Property<minBPVz_t> m_minBPVz {this, -341.f * Gaudi::Units::mm};
Property<enable_monitoring_t> m_enableMonitoring {this, false};
};
} // namespace track_mva_line
......@@ -23,3 +23,16 @@ __device__ bool track_mva_line::track_mva_line_t::select(
return decision;
}
__device__ void track_mva_line::track_mva_line_t::monitor(
const Parameters& parameters,
std::tuple<const Allen::Views::Physics::BasicParticle> input,
unsigned index,
bool sel)
{
if (sel) {
const auto track = std::get<0>(input);
parameters.ipchi2[index] = track.ip_chi2();
parameters.pt[index] = track.state().pt();
}
}
......@@ -820,7 +820,20 @@ Your selection algorithm should define an additional tuple type that contains th
using monitoring_types = std::tuple<pt_t, sv_masses_t>
The tuple type is a convenient way to store a parameter pack, which is then used internally (see `Line.cuh <https://gitlab.cern.ch/lhcb/Allen/-/blob/master/device/selections/line_types/include/Line.cuh>`_ for details) to handle the initialisation of the containers and the transport of the arrays to the nTuple.
Then we set up the `monitor` function that will be handled by the kernel in order to retrieve the information that we want to monitor:
Two useful additional pieces of information to include in the nTuple are the event and run number.
These are automatically filled per candidate if the tuple of monitoring types includes the types runNo_t and evtNo_t, which should be declared as DEVICE_OUTPUTs of type uint64_t and unsigned for runNo and evtNo, respectively.
So in this example, we would add
.. code-blocks:: c++
struct Parameters {
(...)
DEVICE_OUTPUT(evtNo_t, uint64_t) evtNo;
DEVICE_OUTPUT(runNo_t, unsigned) runNo;
}
and the monitoring types becomes ::
using monitoring_types = std::tuple<pt_t, sv_masses_t, evtNo_t, runNo_t>
After all of the values we wish to monitor have been declared,
then we set up the `monitor` function that will be handled by the kernel in order to retrieve the information that we want to monitor:
.. code-block:: c++
......