Commit 5c1b1a74 authored by Daniel Charles Craik's avatar Daniel Charles Craik Committed by Daniel Hugo Campora Perez
Browse files

started to build constant configuration into algorithms; constants defined but...

started to build constant configuration into algorithms; constants defined but neither configured nor used
parent 5f13004f
{
"compass_ut_t": {"sigma_velo_slope" : "0.00010", "min_momentum_final" : "2500.0", "min_pt_final" : "425.0", "hit_tol_2" : "0.8f", "delta_tx_2" : "0.018f"},
"consolidate_ut_tracks_t": {"block_dim" : "[256, 1, 1]"},
"consolidate_scifi_tracks_t": {"block_dim" : "[256, 1, 1]"},
"consolidate_velo_tracks_t": {"block_dim" : "[256, 1, 1]"},
"copy_and_prefix_sum_single_block_sv_t": {"grid_dim" : "[1, 1, 1]", "block_dim" : "[1024, 1, 1]"},
"copy_and_prefix_sum_single_block_scifi_t": {"grid_dim" : "[1, 1, 1]", "block_dim" : "[1024, 1, 1]"},
"copy_and_prefix_sum_single_block_velo_t": {"grid_dim" : "[1, 1, 1]", "block_dim" : "[1024, 1, 1]"},
"copy_and_prefix_sum_single_block_ut_t": {"grid_dim" : "[1, 1, 1]", "block_dim" : "[1024, 1, 1]"},
"copy_scifi_track_hit_number_t": {"block_dim" : "[512, 1, 1]"},
"copy_ut_track_hit_number_t": {"block_dim" : "[512, 1, 1]"},
"copy_velo_track_hit_number_t": {"block_dim" : "[512, 1, 1]"},
"fit_secondary_vertices_t": {"block_dim" : "[16, 16, 1]", "track_min_pt" : "200.0", "track_min_ipchi2" : "9.0", "track_muon_min_ipchi2" : "4.0", "max_assoc_ipchi2" : "16.0"},
"global_event_cut_t": {"block_dim" : "[32, 1, 1]"},
"kalman_filter_t": {"block_dim" : "[256, 1, 1]"},
"kalman_pv_ipchi2_t": {"block_dim" : "[32, 1, 1]"},
"kalman_velo_only_t": {"block_dim" : "[256, 1, 1]"},
"lf_calculate_candidate_extrapolation_window_t": {"block_dim" : "[256, 1, 1]"},
"lf_calculate_first_layer_window_t": {"block_dim" : "[64, 1, 1]"},
"lf_calculate_second_layer_window_t": {"block_dim" : "[2, 16, 1]"},
"lf_calculate_track_extrapolation_window_t": {"block_dim" : "[256, 1, 1]"},
"lf_collect_candidates_t": {"block_dim" : "[64, 6, 1]"},
"lf_extend_missing_x_t": {"block_dim" : "[64, 1, 1]"},
"lf_extend_tracks_t": {"block_dim" : "[256, 1, 1]"},
"lf_fit_t": {"block_dim" : "[512, 1, 1]"},
"lf_form_seeds_from_candidates_t": {"block_dim" : "[1024, 1, 1]"},
"lf_promote_candidates_t": {"block_dim" : "[256, 1, 1]"},
"lf_quality_filter_length_t": {"block_dim" : "[256, 1, 1]"},
"lf_quality_filter_t": {"block_dim" : "[256, 1, 1]"},
"lf_quality_filter_x_t": {"block_dim" : "[32, 1, 1]"},
"lf_search_initial_windows_t": {"block_dim" : "[256, 1, 1]"},
"muon_add_coords_crossing_maps_t": {"block_dim" : "[256, 1, 1]"},
"muon_catboost_evaluator_t": {"block_dim" : "[32, 1, 1]"},
"muon_catboost_features_extraction_t": {"block_dim" : "[32, 1, 1]"},
"muon_sort_by_station_t": {"block_dim" : "[256, 1, 1]"},
"muon_sort_station_region_quarter_t": {"block_dim" : "[256, 1, 1]"},
"pv_beamline_cleanup_t": {"block_dim" : "[32, 1, 1]"},
"pv_beamline_extrapolate_t": {"block_dim" : "[128, 1, 1]"},
"pv_beamline_histo_t": {"block_dim" : "[128, 1, 1]"},
"pv_beamline_multi_fitter_t": {"block_dim" : "[32, 1, 1]"},
"pv_get_seeds_t": {"block_dim" : "[1, 1, 1]", "max_chi2_merge" : "25.0", "factor_to_increase_errors" : "15.0", "min_cluster_mult" : "4", "min_close_tracks_in_cluster" : "3", "dz_close_tracks_in_cluster" : "5.0", "high_mult" : "10", "ratio_sig2_high_mult" : "1.0", "ratio_sig2_low_mult" : "0.9"},
"pv_fit_seeds_t": {"block_dim" : "[1, 1, 1]"},
"run_hlt1_t": {"block_dim" : "[256, 1, 1]"},
"scifi_calculate_cluster_count_v4_t": {"block_dim" : "[240, 1, 1]"},
"scifi_direct_decoder_v4_t": {"block_dim" : "[2, 16, 1]"},
"scifi_raw_bank_decoder_v4_t": {"block_dim" : "[256, 1, 1]"},
"scifi_raw_bank_decoder_v5_t": {"block_dim" : "[256, 1, 1]"},
"scifi_raw_bank_decoder_v6_t": {"block_dim" : "[256, 1, 1]"},
"ut_calculate_number_of_hits_t": {"block_dim" : "[64, 4, 1]"},
"ut_decode_raw_banks_in_order_t": {"block_dim" : "[64, 1, 1]"},
"ut_find_permutation_t": {"block_dim" : "[16, 1, 1]"},
"ut_pre_decode_t": {"block_dim" : "[64, 4, 1]"},
"ut_search_windows_t": {"min_momentum" : "1500.0", "min_pt" : "300.0", "y_tol" : "0.5", "y_tol_slope" : "0.08"},
"velo_calculate_phi_and_sort_t": {"block_dim" : "[64, 1, 1]"},
"velo_estimate_input_size_t": {"block_dim" : "[16, 16, 1]"},
"velo_fill_candidates_t": {"block_dim" : "[128, 1, 1]"},
"velo_kalman_fit_t": {"block_dim" : "[256, 1, 1]"},
"velo_masked_clustering_t": {"block_dim" : "[256, 1, 1]"},
"velo_pv_ip_t": {"block_dim" : "[32, 1, 1]"},
"velo_search_by_triplet_t": {"block_dim" : "[32, 1, 1]", "forward_phi_tolerance" : "0.052", "max_chi2" : "20.0", "max_scatter_forwarding" : "0.1", "max_scatter_seeding" : "0.1", "max_skipped_modules" : "1", "max_weak_tracks" : "500", "phi_extrapolation_base" : "0.03", "phi_extrapolation_coef" : "0.0002", "ttf_modulo" : "2048", "ttf_modulo_mask" : "0x7FF" },
"velo_weak_tracks_adder_t": {"block_dim" : "[256, 1, 1]"}
}
......@@ -342,3 +342,148 @@ dev_velo_track_hit_number (0.01), unused (0.05), dev_atomics_storage (0.00), unu
Max memory required: 9.61 MiB
```
Adding configurable parameters
==============================
To allow a parameter to be configurable via the JSON configuration interface, a `Property` must be
added to the corresponding `ALGORITHM` call. This makes uses of variadic macros so multiple `Property`
objects can be included and will be appended verbatim to the class definition written by the `ALGORITHM` macro.
For example, the following code will add two properties to the `search_by_triplet` algorithm:
```
ALGORITHM(search_by_triplet,
velo_search_by_triplet_t,
ARGUMENTS(
dev_velo_cluster_container,
...
dev_rel_indices),
Property<float> m_tol {this,
"forward_phi_tolerance",
Configuration::velo_search_by_triplet_t::forward_phi_tolerance,
0.052f,
"tolerance"};
Property<float> m_scat {this,
"max_scatter_forwarding",
Configuration::velo_search_by_triplet_t::max_scatter_forwarding,
0.1f,
"scatter forwarding"};
)
```
The arguments passed to the `Property` constructor are
* the `Algorithm` that "owns" it;
* the name of the property in the JSON configuration;
* the underlying variable - this must be in `__constant__` memory for regular properties (see below);
* the default value of the property;
* a description of the property.
As the underlying parameters make use of GPU constant memory, they may not be defined within the
algorithm's class. They should instead be placed inside of namespace of the same name within the
`Configuration` namespace. For the example above, the following needs to be added to the header file:
```
namespace Configuration {
namespace velo_search_by_triplet_t {
// Forward tolerance in phi
extern __constant__ float forward_phi_tolerance;
// Max scatter for forming triplets (seeding) and forwarding
extern __constant__ float max_scatter_forwarding;
} // namespace velo_search_by_triplet_t
} // namespace Configuration
```
and the following to the code file:
```
__constant__ float Configuration::velo_search_by_triplet_t::forward_phi_tolerance;
__constant__ float Configuration::velo_search_by_triplet_t::max_scatter_forwarding;
```
Finally, the following can be added to the configuration file (default: `configuration/constants/default.json`)
to configure the values of these parameters at runtime:
```
"velo_search_by_triplet_t": {"forward_phi_tolerance" : "0.052", "max_scatter_forwarding" : "0.1"}
```
Derived properties
------------------
For properties derived from other configurable properties, the `DerivedProperty` class may be used:
```
Property<float> m_slope {this,
"sigma_velo_slope",
Configuration::compass_ut_t::sigma_velo_slope,
0.010f * Gaudi::Units::mrad,
"sigma velo slope [radians]"};
DerivedProperty<float> m_inv_slope {this,
"inv_sigma_velo_slope",
Configuration::compass_ut_t::inv_sigma_velo_slope,
Configuration::Relations::inverse,
std::vector<Property<float>*> {&this->m_slope},
"inv sigma velo slope"};
```
Here, the value of the `m_inv_slope` property is determined by the function and the
vector of properties given in the third and fourth arguments. Additional functions
may be added to the `Configuration::Relations` and defined in `stream/gear/src/Configuration.cu`.
All functions take a vector of properties as an argument, to allow for functions of an
arbitrary number of properties.
CPU properties
--------------
Regular properties are designed to be used in GPU algorithms and are stored
in GPU constant memory with a cached copy within the `Property` class.
For properties that are only needed on the CPU, e.g. grid and block dimensions,
a `CPUProperty` can be used, which only stores the configured value internally.
This is also useful for properties tht are only needed when first configuring the
algorithm, such as properties only used in the visitor class.
Note that regular properties may also be used in this case
(e.g. `../stream/visitors/velo/src/SearchByTripletVisitor.cu` accesses non-CPU properties)
but if a property is *only* needed on the CPU then there is a reduced overhead in using a `CPUProperty`.
These are defined in the same way as a `Property` but take one fewer argument as there is no underlying
constant memory object to reference.
```
CPUProperty<std::array<int, 3>> m_block_dim {this, "block_dim", {32, 1, 1}, "block dimensions"};
CPUProperty<std::array<int, 3>> m_grid_dim {this, "grid_dim", {1, 1, 1}, "grid dimensions"};
```
Shared properties
-----------------
For properties that are shared between multiple top-level algorithms, it may be preferred
to keep the properties in a neutral location. This ensures that properties are configured
regardless of which algorithms are used in the configured sequence and can be achieved by
using a `SharedProperty`.
Shared properties are owned by a `SharedPropertySet` rather than an `Algorithm`
and example of which is given below.
```
#include "Configuration.cuh"
namespace Configuration {
namespace example_common {
extern __constant__ float param;
}
}
struct ExampleConfiguration : public SharedPropertySet {
ExampleConfiguration() = default;
constexpr static auto name{ "example_common" };
private:
Property<float> m_par{this, "param", Configuration::example_common::param, 0., "an example parameter"};
};
```
This may be used by any algorithm by including the header and adding the following line
to the end of the arguments of the `ALGORITHM` call.
```
SharedProperty<float> m_shared{this, "example_common", "param"};
```
These must also be plumbed in to `Configuration::getSharedPropertySet` in `stream/gear/src/Configuration.cu`
to allow the property set to be found by algorithms.
......@@ -25,4 +25,4 @@ ALGORITHM(
dev_multi_fit_vertices,
dev_number_of_multi_fit_vertices,
dev_multi_final_vertices,
dev_number_of_multi_final_vertices))
\ No newline at end of file
dev_number_of_multi_final_vertices))
......@@ -9,28 +9,6 @@ namespace PatPV {
// auxiliary class for searching of clusters of tracks
// configuration for seeding
// steering parameters for merging procedure
static constexpr float mcu_maxChi2Merge = 25.;
static constexpr float mcu_factorToIncreaseErrors = 15.;
// try parameters from RecoUpgradeTracking.py
static constexpr int mcu_minClusterMult = 4;
static constexpr int mcu_minCloseTracksInCluster = 3;
// steering parameters for final cluster selection
// int m_minClusterMult = 3;
static constexpr float mcu_dzCloseTracksInCluster = 5.; // unit: mm
// int m_minCloseTracksInCluster = 3;
static constexpr int mcu_highMult = 10;
static constexpr float mcu_ratioSig2HighMult = 1.0;
static constexpr float mcu_ratioSig2LowMult = 0.9;
static constexpr int mcu_max_clusters = 1200; // maximmum number of clusters
static constexpr float mcu_x0MS = 0.01; // X0 (tunable) of MS to add for extrapolation of
// track parameters to PV
// don't forget to actually calculate this!!
// double m_scatCons = 0; // calculated from m_x0MS
static constexpr float X0cu = 0.01;
......
......@@ -19,7 +19,57 @@ __global__ void get_seeds(
__device__ int find_clusters(PatPV::vtxCluster* vclus, float* zclusters, int number_of_clusters);
namespace Configuration {
namespace pv_get_seeds_t {
// configuration for seeding
// steering parameters for merging procedure
extern __constant__ float max_chi2_merge;
extern __constant__ float factor_to_increase_errors;
// try parameters from RecoUpgradeTracking.py
extern __constant__ int min_cluster_mult;
extern __constant__ int min_close_tracks_in_cluster;
// steering parameters for final cluster selection
// unit: mm
extern __constant__ float dz_close_tracks_in_cluster;
extern __constant__ int high_mult;
extern __constant__ float ratio_sig2_high_mult;
extern __constant__ float ratio_sig2_low_mult;
} // namespace pv_get_seeds_t
} // namespace Configuration
ALGORITHM(
get_seeds,
pv_get_seeds_t,
ARGUMENTS(dev_velo_kalman_beamline_states, dev_atomics_velo, dev_velo_track_hit_number, dev_seeds, dev_number_seeds))
ARGUMENTS(dev_velo_kalman_beamline_states, dev_atomics_velo, dev_velo_track_hit_number, dev_seeds, dev_number_seeds),
Property<float>
m_chi2 {this, "max_chi2_merge", Configuration::pv_get_seeds_t::max_chi2_merge, 25.f, "max chi2 merge"};
Property<float> m_ferr {this,
"factor_to_increase_errors",
Configuration::pv_get_seeds_t::factor_to_increase_errors,
15.f,
"factor to increase errors"};
Property<int>
m_mult {this, "min_cluster_mult", Configuration::pv_get_seeds_t::min_cluster_mult, 4, "min cluster mult"};
Property<int> m_close {this,
"min_close_tracks_in_cluster",
Configuration::pv_get_seeds_t::min_close_tracks_in_cluster,
3,
"min close tracks in cluster"};
Property<float> m_dz {this,
"dz_close_tracks_in_cluster",
Configuration::pv_get_seeds_t::dz_close_tracks_in_cluster,
5.f,
"dz close tracks in cluster [mm]"};
Property<int> m_himult {this, "high_mult", Configuration::pv_get_seeds_t::high_mult, 10, "high mult"};
Property<float> m_ratiohi {this,
"ratio_sig2_high_mult",
Configuration::pv_get_seeds_t::ratio_sig2_high_mult,
1.0f,
"ratio sig2 high mult"};
Property<float> m_ratiolo {this,
"ratio_sig2_low_mult",
Configuration::pv_get_seeds_t::ratio_sig2_low_mult,
0.9f,
"ratio sig2 low mult"};)
#include "GetSeeds.cuh"
// simplficiations: no tracks2disable
__constant__ float Configuration::pv_get_seeds_t::max_chi2_merge;
__constant__ float Configuration::pv_get_seeds_t::factor_to_increase_errors;
__constant__ int Configuration::pv_get_seeds_t::min_cluster_mult;
__constant__ int Configuration::pv_get_seeds_t::min_close_tracks_in_cluster;
__constant__ float Configuration::pv_get_seeds_t::dz_close_tracks_in_cluster;
__constant__ int Configuration::pv_get_seeds_t::high_mult;
__constant__ float Configuration::pv_get_seeds_t::ratio_sig2_high_mult;
__constant__ float Configuration::pv_get_seeds_t::ratio_sig2_low_mult;
__device__ float zCloseBeam(KalmanVeloState track, const PatPV::XYZPoint& beamspot)
{
......@@ -101,7 +112,8 @@ __device__ int find_clusters(PatPV::vtxCluster* vclus, float* zclusters, int num
{
for (int i = 0; i < number_of_clusters; i++) {
vclus[i].sigsq *= PatPV::mcu_factorToIncreaseErrors * PatPV::mcu_factorToIncreaseErrors; // blow up errors
vclus[i].sigsq *= Configuration::pv_get_seeds_t::factor_to_increase_errors *
Configuration::pv_get_seeds_t::factor_to_increase_errors; // blow up errors
vclus[i].sigsqmin = vclus[i].sigsq;
}
......@@ -138,7 +150,7 @@ __device__ int find_clusters(PatPV::vtxCluster* vclus, float* zclusters, int num
float zdist = z1 - z2;
float chi2dist = zdist * zdist / (s1 + s2);
// merge if chi2dist is smaller than max
if (chi2dist < PatPV::mcu_maxChi2Merge) {
if (chi2dist < Configuration::pv_get_seeds_t::max_chi2_merge) {
no_merges = false;
float w_inv = (s1 * s2 / (s1 + s2));
float zmerge = w_inv * (z1 / s1 + z2 / s2);
......@@ -177,7 +189,8 @@ __device__ int find_clusters(PatPV::vtxCluster* vclus, float* zclusters, int num
int n_tracks_close = 0;
for (int i = 0; i < number_of_clusters; i++)
if (fabsf(vclus[i].z - pvclus[index].z) < PatPV::mcu_dzCloseTracksInCluster) n_tracks_close++;
if (fabsf(vclus[i].z - pvclus[index].z) < Configuration::pv_get_seeds_t::dz_close_tracks_in_cluster)
n_tracks_close++;
float dist_to_closest = 1000000.;
if (return_number_of_clusters > 1) {
......@@ -191,13 +204,18 @@ __device__ int find_clusters(PatPV::vtxCluster* vclus, float* zclusters, int num
float rat = pvclus[index].sigsq / pvclus[index].sigsqmin;
bool igood = false;
int ntracks = pvclus[index].ntracks;
if (ntracks >= PatPV::mcu_minClusterMult) {
if (ntracks >= Configuration::pv_get_seeds_t::min_cluster_mult) {
if (dist_to_closest > 10.f && rat < 0.95f) igood = true;
if (ntracks >= PatPV::mcu_highMult && rat < PatPV::mcu_ratioSig2HighMult) igood = true;
if (ntracks < PatPV::mcu_highMult && rat < PatPV::mcu_ratioSig2LowMult) igood = true;
if (
ntracks >= Configuration::pv_get_seeds_t::high_mult &&
rat < Configuration::pv_get_seeds_t::ratio_sig2_high_mult)
igood = true;
if (
ntracks < Configuration::pv_get_seeds_t::high_mult && rat < Configuration::pv_get_seeds_t::ratio_sig2_low_mult)
igood = true;
}
// veto
if (n_tracks_close < PatPV::mcu_minCloseTracksInCluster) igood = false;
if (n_tracks_close < Configuration::pv_get_seeds_t::min_close_tracks_in_cluster) igood = false;
if (igood) {
zclusters[number_good_clusters] = pvclus[index].z;
number_good_clusters++;
......
......@@ -7,7 +7,7 @@
#include "LFExtendTracksUV.cuh"
#include "LFSearchUVWindows.cuh"
struct lf_composite_extend_tracks_uv_t {
struct lf_composite_extend_tracks_uv_t : public Algorithm {
constexpr static auto name {"lf_composite_extend_tracks_uv_t"};
using Arguments = std::tuple<
dev_scifi_hits,
......
......@@ -16,7 +16,7 @@
// decltype(composite_handler(FUNCTIONS)) handlers {composite_handler(FUNCTIONS)};
// };
struct lf_composite_track_seeding_t {
struct lf_composite_track_seeding_t : public Algorithm {
constexpr static auto name {"lf_composite_track_seeding_t"};
using Arguments = std::tuple<
dev_scifi_hits,
......
......@@ -53,21 +53,9 @@ namespace UT {
static constexpr float zMidUT = 2484.6f;
// distToMomentum is properly recalculated in UTMagnetTool when B field changes
static constexpr float distToMomentum = 4.0212e-05f;
static constexpr float sigmaVeloSlope = 0.10f * Gaudi::Units::mrad;
static constexpr float invSigmaVeloSlope = 1.0f / sigmaVeloSlope;
static constexpr float zKink = 1780.0f;
static constexpr float minMomentum = 1.5f * Gaudi::Units::GeV;
static constexpr float minMomentumFinal = 2.5f * Gaudi::Units::GeV;
static constexpr float minPT = 0.3f * Gaudi::Units::GeV;
static constexpr float minPTFinal = 0.425f * Gaudi::Units::GeV;
static constexpr float maxPseudoChi2 = 1280.0f;
static constexpr float yTol = 0.5f * Gaudi::Units::mm;
static constexpr float yTolSlope = 0.08f;
static constexpr float hitTol1 = 6.0f * Gaudi::Units::mm;
static constexpr float hitTol2 = 0.8f * Gaudi::Units::mm;
static constexpr float deltaTx1 = 0.035f;
static constexpr float deltaTx2 = 0.018f;
static constexpr float maxXSlope = 0.350f;
static constexpr float maxYSlope = 0.300f;
static constexpr float centralHoleSize = 33.0f * Gaudi::Units::mm;
......@@ -86,7 +74,6 @@ namespace UT {
//
static constexpr float LD3Hits = -0.5f;
static constexpr float LD4Hits = -0.5f;
} // namespace Constants
} // namespace UT
......
......@@ -33,7 +33,7 @@ __device__ std::tuple<int, int> find_candidates_in_sector_group(
const float xTol,
const int sector_group);
__host__ __device__ void tol_refine(
__device__ void tol_refine(
int& first_candidate,
int& last_candidate,
const UT::Hits& ut_hits,
......
......@@ -70,17 +70,59 @@ __device__ void save_track(
UT::TrackHits* VeloUT_tracks,
const int event_hit_offset);
ALGORITHM(
compass_ut,
compass_ut_t,
ARGUMENTS(
dev_ut_hits,
dev_ut_hit_offsets,
dev_atomics_velo,
dev_velo_track_hit_number,
dev_velo_states,
dev_ut_tracks,
dev_atomics_ut,
dev_ut_active_tracks,
dev_ut_windows_layers,
dev_accepted_velo_tracks))
namespace Configuration {
namespace compass_ut_t {
extern __constant__ float sigma_velo_slope;
extern __constant__ float inv_sigma_velo_slope;
extern __constant__ float min_momentum_final;
extern __constant__ float min_pt_final;
extern __constant__ float hit_tol_2;
extern __constant__ float delta_tx_2;
} // namespace compass_ut_t
} // namespace Configuration
ALGORITHM(compass_ut,
compass_ut_t,
ARGUMENTS(
dev_ut_hits,
dev_ut_hit_offsets,
dev_atomics_velo,
dev_velo_track_hit_number,
dev_velo_states,
dev_ut_tracks,
dev_atomics_ut,
dev_ut_active_tracks,
dev_ut_windows_layers,
dev_accepted_velo_tracks),
Property<float> m_slope {this,
"sigma_velo_slope",
Configuration::compass_ut_t::sigma_velo_slope,
0.010f * Gaudi::Units::mrad,
"sigma velo slope [radians]"};
DerivedProperty<float> m_inv_slope {this,
"inv_sigma_velo_slope",
Configuration::compass_ut_t::inv_sigma_velo_slope,
Configuration::Relations::inverse,
std::vector<Property<float>*> {&this->m_slope},
"inv sigma velo slope"};
Property<float> m_mom_fin {this,
"min_momentum_final",
Configuration::compass_ut_t::min_momentum_final,
2.5f * Gaudi::Units::GeV,
"final min momentum cut [MeV/c]"};
Property<float> m_pt_fin {this,
"min_pt_final",
Configuration::compass_ut_t::min_pt_final,
0.425f * Gaudi::Units::GeV,
"final min pT cut [MeV/c]"};
Property<float> m_hit_tol_2 {this,
"hit_tol_2",
Configuration::compass_ut_t::hit_tol_2,
0.8f * Gaudi::Units::mm,
"hit_tol_2 [mm]"};
Property<float> m_delta_tx_2 {this,
"delta_tx_2",
Configuration::compass_ut_t::delta_tx_2,
0.018f,
"delta_tx_2"};
)
......@@ -34,11 +34,3 @@ __device__ __inline__ int calc_index(
const int layer0,
const int layer2,
const UT::HitOffsets& ut_hit_offsets);
__device__ __inline__ bool check_tol_refine(
const int hit_index,
const UT::Hits& ut_hits,
const MiniState& velo_state,
const float normFactNum,
const float xTol,
const float dxDy);
......@@ -22,16 +22,38 @@ __global__ void ut_search_windows(
uint* dev_active_tracks,
bool* dev_accepted_velo_tracks);
ALGORITHM(
ut_search_windows,
ut_search_windows_t,
ARGUMENTS(
dev_ut_hits,
dev_ut_hit_offsets,
dev_atomics_velo,
dev_velo_track_hit_number,
dev_velo_track_hits,
dev_velo_states,
dev_ut_windows_layers,
dev_accepted_velo_tracks,
dev_ut_active_tracks))
namespace Configuration {
namespace ut_search_windows_t {
extern __constant__ float min_momentum;
extern __constant__ float min_pt;
extern __constant__ float y_tol;
extern __constant__ float y_tol_slope;
} // namespace ut_search_windows_t
} // namespace Configuration
ALGORITHM(ut_search_windows,
ut_search_windows_t,
ARGUMENTS(
dev_ut_hits,
dev_ut_hit_offsets,
dev_atomics_velo,
dev_velo_track_hit_number,
dev_velo_track_hits,
dev_velo_states,
dev_ut_windows_layers,
dev_accepted_velo_tracks,
dev_ut_active_tracks),
Property<float> m_mom {this,
"min_momentum",
Configuration::ut_search_windows_t::min_momentum,
1.5f * Gaudi::Units::GeV,
"min momentum cut [MeV/c]"};
Property<float> m_pt {this,
"min_pt",
Configuration::ut_search_windows_t::min_pt,
0.3f * Gaudi::Units::GeV,
"min pT cut [MeV/c]"};
Property<float>
m_ytol {this, "y_tol", Configuration::ut_search_windows_t::y_tol, 0.5f * Gaudi::Units::mm, "y tol [mm]"};
Property<float>
m_yslope {this, "y_tol_slope", Configuration::ut_search_windows_t::y_tol_slope, 0.08f, "y tol slope [mm]"};)
#include "BinarySearch.cuh"
#include "VeloTools.cuh"
#include "CalculateWindows.cuh"
#include "SearchWindows.cuh"
//=============================================================================
// Reject tracks outside of acceptance or pointing to the beam pipe
......@@ -25,7 +26,7 @@ __device__ bool velo_track_in_UTA_acceptance(const MiniState& state)
//=========================================================================
// Check if hit is inside tolerance and refine by Y
//=========================================================================
__host__ __device__ void tol_refine(
__device__ void tol_refine(
int& first_candidate,
int& last_candidate,
const UT::Hits& ut_hits,
......@@ -45,7 +46,11 @@ __host__ __device__ void tol_refine(