Commit 739d9aae authored by Dorothea Vom Bruch's avatar Dorothea Vom Bruch
Browse files

Merge branch 'vava_cleanup_prforward'

parents 5552b7cd dd087268
Pipeline #938763 passed with stages
in 13 minutes and 33 seconds
......@@ -215,7 +215,7 @@ include_directories(cuda/SciFi/looking_forward_sbt/extend_tracks_uv/include)
include_directories(cuda/SciFi/looking_forward_sbt/quality_filter/include)
include_directories(cuda/SciFi/looking_forward_sbt/quality_filter_x/include)
include_directories(cuda/SciFi/looking_forward_sbt/search_uv_windows/include)
include_directories(cuda/SciFi/PrForward/include)
include_directories(cuda/SciFi/classifiers/include)
include_directories(cuda/SciFi/consolidate/include)
include_directories(cuda/muon/common/include)
include_directories(cuda/utils/prefix_sum/include)
......@@ -227,7 +227,6 @@ include_directories(checker/tracking/include)
include_directories(checker/pv/include)
include_directories(checker/selections/include)
include_directories(stream/sequence/include)
include_directories(x86/SciFi/PrForward/include)
include_directories(x86/SciFi/LookingForward/include)
include_directories(x86/SciFi/MomentumForward/include)
include_directories(cuda/UT/UTDecoding/include)
......
......@@ -9,7 +9,6 @@ include_directories(${CMAKE_SOURCE_DIR}/cuda/event_model/common/include)
include_directories(${CMAKE_SOURCE_DIR}/cuda/event_model/UT/include)
include_directories(${CMAKE_SOURCE_DIR}/cuda/event_model/SciFi/include)
include_directories(${CMAKE_SOURCE_DIR}/cuda/SciFi/common/include)
include_directories(${CMAKE_SOURCE_DIR}/cuda/SciFi/PrForward/include)
include_directories(${CMAKE_SOURCE_DIR}/cuda/kalman/ParKalman/include)
include_directories(${CMAKE_SOURCE_DIR}/cuda/PV/common/include)
include_directories(${CMAKE_SOURCE_DIR}/cuda/PV/beamlinePV/include)
......
......@@ -19,7 +19,8 @@ file(GLOB scifi_lf_quality_filter "looking_forward_sbt/quality_filter/src/*cu")
file(GLOB scifi_lf_quality_filter_x "looking_forward_sbt/quality_filter_x/src/*cu")
file(GLOB scifi_lf_search_uv_windows "looking_forward_sbt/search_uv_windows/src/*cu")
file(GLOB scifi_lf_fit "looking_forward_sbt/fit/src/*cu")
file(GLOB scifi_prforward "PrForward/src/*cu")
file(GLOB scifi_utils "utils/src/*cu")
file(GLOB scifi_classifiers "classifiers/src/*cu")
file(GLOB scifi_consolidate "consolidate/src/*cu")
include_directories(${CMAKE_CUDA_TOOLKIT_INCLUDE_DIRECTORIES})
......@@ -61,12 +62,14 @@ include_directories(looking_forward_sbt/quality_filter/include)
include_directories(looking_forward_sbt/quality_filter_x/include)
include_directories(looking_forward_sbt/search_uv_windows/include)
include_directories(looking_forward_sbt/fit/include)
include_directories(PrForward/include)
include_directories(utils/include)
include_directories(classifiers/include)
include_directories(consolidate/include)
add_library(SciFi STATIC
${scifi_common}
${scifi_preprocessing}
${scifi_utils}
${scifi_lf_common}
${scifi_lf_calculate_first_layer_window}
${scifi_lf_calculate_second_layer_window}
......@@ -86,7 +89,7 @@ add_library(SciFi STATIC
${scifi_lf_quality_filter_x}
${scifi_lf_search_uv_windows}
${scifi_lf_fit}
${scifi_prforward}
${scifi_classifiers}
${scifi_consolidate}
)
......
# Implementation of the forward based on Rec v23r3
#pragma once
#include <cmath>
#include <array>
#include <vector>
#include <algorithm>
#include <fstream>
#include "SciFiDefinitions.cuh"
#include "SciFiEventModel.cuh"
#include "TrackUtils.cuh"
#include "HitUtils.cuh"
/**
Functions related to selecting hits on the uv planes,
which match to the VeloUT input track
*/
__host__ __device__ void collectStereoHits(
const SciFi::Hits& scifi_hits,
const SciFi::HitCount& scifi_hit_count,
SciFi::Tracking::Track& track,
const MiniState& velo_state,
SciFi::Tracking::HitSearchCuts& pars,
const SciFi::Tracking::Arrays* constArrays,
float stereoCoords[SciFi::Tracking::max_stereo_hits],
int stereoHits[SciFi::Tracking::max_stereo_hits],
int& n_stereoHits);
__host__ __device__ bool selectStereoHits(
const SciFi::Hits& scifi_hits,
const SciFi::HitCount& scifi_hit_count,
SciFi::Tracking::Track& track,
const SciFi::Tracking::Arrays* constArrays,
float stereoCoords[SciFi::Tracking::max_stereo_hits],
int stereoHits[SciFi::Tracking::max_stereo_hits],
int& n_stereoHits,
const MiniState& velo_state,
SciFi::Tracking::HitSearchCuts& pars_cur);
__host__ __device__ bool addHitsOnEmptyStereoLayers(
const SciFi::Hits& scifi_hits,
const SciFi::HitCount& scifi_hit_count,
SciFi::Tracking::Track& track,
int stereoHits[SciFi::Tracking::max_stereo_hits],
int& n_stereoHits,
const SciFi::Tracking::Arrays* constArrays,
PlaneCounter& planeCounter,
const MiniState& velo_state,
SciFi::Tracking::HitSearchCuts& pars_cur);
#pragma once
#include <cmath>
#include <array>
#include <vector>
#include <algorithm>
#include <fstream>
#include "SciFiDefinitions.cuh"
#include "PrForwardConstants.cuh"
#include "UTDefinitions.cuh"
#include "TrackUtils.cuh"
#include "HitUtils.cuh"
#include "LinearFitting.cuh"
#include "ReferencePlaneProjection.cuh"
#include "SciFiEventModel.cuh"
#include "LookingForwardUtils.h"
/**
Functions related to selecting hits on the x planes,
which match to the VeloUT input track
*/
__host__ void collectAllXHits_proto_p(
const SciFi::Hits& scifi_hits,
const SciFi::HitCount& scifi_hit_count,
const SciFi::Tracking::Arrays* constArrays,
const float magnet_polarity,
const MiniState& velo_state,
const MiniState& UT_state,
const float qOverP,
int side,
std::array<int, 2 * 6>& windows_x,
std::array<int, 2 * 6>& windows_uv,
std::array<float, 4 * 6>& parameters_uv,
const SciFiWindowsParams& window_params,
const std::array<int, 12> true_scifi_indices_per_layer);
__host__ void x_limits_from_dxRef(
const SciFi::Tracking::Arrays* constArrays,
const MiniState& velo_state,
const float InvPz,
const float p,
const float tx2,
const float ty2,
const bool wSignTreatment,
float& xBoundOnRef,
float& xBoundOnRefWS);
__host__ void collectAllXHits_proto(
const SciFi::Hits& scifi_hits,
const SciFi::HitCount& scifi_hit_count,
const float xParams_seed[4],
const float yParams_seed[4],
const SciFi::Tracking::Arrays* constArrays,
const float magnet_polarity,
const MiniState& UT_state,
const float qOverP,
int side,
std::array<int, 2 * 6>& windows_x,
std::array<int, 2 * 6>& windows_uv,
std::array<float, 4 * 6>& parameters_uv);
__host__ __device__ void collectAllXHits(
const SciFi::Hits& scifi_hits,
const SciFi::HitCount& scifi_hit_count,
int allXHits[SciFi::Tracking::max_x_hits],
int& n_x_hits,
float coordX[SciFi::Tracking::max_x_hits],
const float xParams_seed[4],
const float yParams_seed[4],
const SciFi::Tracking::Arrays* constArrays,
const float magnet_polarity,
const MiniState& velo_state,
const float qop,
int side);
__host__ __device__ void improveXCluster(
int& it2,
const int it1,
const int itEnd,
const int n_x_hits,
const bool usedHits[SciFi::Tracking::max_x_hits],
const float coordX[SciFi::Tracking::max_x_hits],
const float xWindow,
const SciFi::Tracking::HitSearchCuts& pars,
PlaneCounter& planeCounter,
const int allXHits[SciFi::Tracking::max_x_hits],
const SciFi::Hits& scifi_hits);
__host__ __device__ void selectXCandidates(
const SciFi::Hits& scifi_hits,
const SciFi::HitCount& scifi_hit_count,
int allXHits[SciFi::Tracking::max_x_hits],
int& n_x_hits,
bool usedHits[SciFi::Tracking::max_x_hits],
float coordX[SciFi::Tracking::max_x_hits],
SciFi::Tracking::Track candidate_tracks[SciFi::Constants::max_tracks],
int& n_candidate_tracks,
const float zRef_track,
const float xParams_seed[4],
const float yParams_seed[4],
const MiniState& velo_state,
SciFi::Tracking::HitSearchCuts& pars,
const SciFi::Tracking::Arrays* constArrays,
int side,
const bool secondLoop);
__host__ __device__ bool addHitsOnEmptyXLayers(
const SciFi::Hits& scifi_hits,
const SciFi::HitCount& scifi_hit_count,
float trackParameters[SciFi::Tracking::nTrackParams],
const float xParams_seed[4],
const float yParams_seed[4],
bool fullFit,
int coordToFit[SciFi::Tracking::max_coordToFit],
int& n_coordToFit,
const SciFi::Tracking::Arrays* constArrays,
PlaneCounter& planeCounter,
SciFi::Tracking::HitSearchCuts& pars_cur,
int side);
#pragma once
#include "SciFiDefinitions.cuh"
#include "PrForwardConstants.cuh"
#include "SciFiEventModel.cuh"
/**
Helper functions related to properties of hits on planes
*/
// Helper used to keep track of how many x / stereo hits per lane have
// been added to a candidate track
struct PlaneCounter {
int planeList[SciFi::Constants::n_layers] = {0};
unsigned int nbDifferent = 0;
__host__ __device__ inline void addHit(int plane)
{
assert(plane < SciFi::Constants::n_layers);
nbDifferent += (int) ((planeList[plane] += 1) == 1);
}
__host__ __device__ inline void removeHit(int plane)
{
assert(plane < SciFi::Constants::n_layers);
nbDifferent -= ((int) ((planeList[plane] -= 1) == 0));
}
__host__ __device__ inline int nbInPlane(int plane) const
{
assert(plane < SciFi::Constants::n_layers);
return planeList[plane];
}
__host__ __device__ inline int nbSingle() const
{
int single = 0;
for (int i = 0; i < SciFi::Constants::n_layers; ++i) {
single += planeList[i] == 1 ? 1 : 0;
}
return single;
}
__host__ __device__ inline void clear()
{
nbDifferent = 0;
for (int i = 0; i < SciFi::Constants::n_layers; ++i) {
planeList[i] = 0;
}
}
};
__host__ __device__ void countPlanesOfXHits(
PlaneCounter& planeCounter,
const int it1,
const int it2,
const int n_x_hits,
const int allXHits[SciFi::Tracking::max_x_hits],
const bool usedHits[SciFi::Tracking::max_x_hits],
const SciFi::Hits& scifi_hits);
__host__ __device__ void countUnusedXHitsOnPlanes(
PlaneCounter& lplaneCounter,
const int itWindowStart,
const int itWindowEnd,
const int n_x_hits,
const int allXHits[SciFi::Tracking::max_x_hits],
const bool usedHits[SciFi::Tracking::max_x_hits],
const SciFi::Hits& scifi_hits);
__host__ __device__ void addXHitsForCandidateWithTooFewPlanes(
int& itWindowStart,
int& itWindowEnd,
const int it2,
const int itEnd,
float& minInterval,
PlaneCounter& lplaneCounter,
const int nPlanes,
const float coordX[SciFi::Tracking::max_x_hits],
int& best,
int& bestEnd,
const bool usedHits[SciFi::Tracking::max_x_hits],
const int n_x_hits,
const int allXHits[SciFi::Tracking::max_x_hits],
const SciFi::Hits& scifi_hits);
__host__ __device__ void collectXHitsToFit(
const int it1,
const int it2,
const int n_x_hits,
const int allXHits[SciFi::Tracking::max_x_hits],
bool usedHits[SciFi::Tracking::max_x_hits],
int coordToFit[SciFi::Tracking::max_x_hits],
int& n_coordToFit,
const float coordX[SciFi::Tracking::max_x_hits],
float& xAtRef);
__host__ __device__ int findBestXHitOnEmptyLayer(
const int itEnd,
const int itH,
const SciFi::Hits& scifi_hits,
const float maxX,
const float xPred);
template<int N>
__host__ __device__ void sortHitsByKey(float* keys, int n, int* hits)
{
// find permutations
uint permutations[N];
assert(n <= N);
for (int i = 0; i < n; ++i) {
uint position = 0;
for (int j = 0; j < n; ++j) {
// sort keys in ascending order
int sort_result = -1;
if (keys[i] > keys[j]) sort_result = 1;
if (keys[i] == keys[j]) sort_result = 0;
position += sort_result > 0 || (sort_result == 0 && i > j);
}
permutations[position] = i;
}
// apply permutations, store hits in temporary container
int hits_tmp[N];
float keys_tmp[N];
for (int i = 0; i < n; ++i) {
const int index = permutations[i];
hits_tmp[i] = hits[index];
keys_tmp[i] = keys[index];
}
// copy hits back to original container
for (int i = 0; i < n; ++i) {
hits[i] = hits_tmp[i];
keys[i] = keys_tmp[i];
}
}
// check that val is within [min, max]
__host__ __device__ inline bool isInside(float val, const float min, const float max)
{
return (val > min) && (val < max);
}
// get lowest index where range[index] > value, within [start,end] of range
__host__ __device__ inline int getLowerBound(float range[], float value, int start, int end)
{
int i = start;
for (; i < end; i++) {
if (range[i] > value) break;
}
return i;
}
// match stereo hits to x hits
__host__ __device__ bool matchStereoHit(
const int itUV1,
const int uv_zone_offset_end,
const SciFi::Hits& scifi_hits,
const float xMinUV,
const float xMaxUV);
__host__ __device__ bool matchStereoHitWithTriangle(
const int itUV2,
const int triangle_zone_offset_end,
const float yInZone,
const SciFi::Hits& scifi_hits,
const float xMinUV,
const float xMaxUV,
const int side);
__host__ __device__ void removeOutlier(
const SciFi::Hits& scifi_hits,
PlaneCounter& planeCounter,
int* coordToFit,
int& n_coordToFit,
const int worst);
__host__ __device__ void findStereoHitsWithinXTol(
const int itBegin,
const int itEnd,
const SciFi::Hits& scifi_hits,
const float yZone,
const float xPred,
const float dxTol,
const bool triangleSearch,
const float dxDySign,
int& n_stereoHits,
float stereoCoords[SciFi::Tracking::max_stereo_hits],
int stereoHits[SciFi::Tracking::max_stereo_hits]);
__host__ __device__ void findStereoHitClusterByDx(
PlaneCounter& planeCounter,
int& endRange,
const SciFi::Tracking::HitSearchCuts& pars,
float stereoCoords[SciFi::Tracking::max_stereo_hits],
int stereoHits[SciFi::Tracking::max_stereo_hits],
const int n_stereoHits,
const SciFi::Hits& scifi_hits,
float& sumCoord,
int& first_hit);
__host__ __device__ void cleanStereoHitCluster(
int& beginRange,
int& endRange,
const int n_stereoHits,
const int stereoHits[SciFi::Tracking::max_stereo_hits],
const float stereoCoords[SciFi::Tracking::max_stereo_hits],
float& sumCoord,
PlaneCounter& planeCounter,
const SciFi::Hits& scifi_hits);
__host__ __device__ int findBestStereoHitOnEmptyLayer(
const int itBegin,
const int itEnd,
const SciFi::Hits& scifi_hits,
const float yZone,
const float xPred,
const float dxTol,
const bool triangleSearch);
#pragma once
#include "SciFiDefinitions.cuh"
#include "TrackUtils.cuh"
#include "HitUtils.cuh"
#include "SciFiEventModel.cuh"
#include <cmath>
namespace SciFi {
namespace Tracking {
struct LineFitterPars {
float m_z0 = 0.;
float m_c0 = 0.;
float m_tc = 0.;
float m_s0 = 0.;
float m_sz = 0.;
float m_sz2 = 0.;
float m_sc = 0.;
float m_scz = 0.;
};
} // namespace Tracking
} // namespace SciFi
__host__ __device__ void incrementLineFitParameters(
SciFi::Tracking::LineFitterPars& parameters,
const SciFi::Hits& scifi_hits,
const float coordX[SciFi::Tracking::max_x_hits],
const int allXHits[SciFi::Tracking::max_x_hits],
const int it);
__host__ __device__ void fitHitsFromSingleHitPlanes(
const int it1,
const int it2,
const bool usedHits[SciFi::Tracking::max_x_hits],
const SciFi::Hits& scifi_hits,
const int allXHits[SciFi::Tracking::max_x_hits],
const int n_x_hits,
const PlaneCounter planeCounter,
SciFi::Tracking::LineFitterPars& lineFitParameters,
const float coordX[SciFi::Tracking::max_x_hits],
int otherHits[SciFi::Constants::n_layers][SciFi::Tracking::max_other_hits],
int nOtherHits[SciFi::Constants::n_layers]);
__host__ __device__ void addAndFitHitsFromMultipleHitPlanes(
const int nOtherHits[SciFi::Constants::n_layers],
SciFi::Tracking::LineFitterPars& lineFitParameters,
const SciFi::Hits& scifi_hits,
const float coordX[SciFi::Tracking::max_x_hits],
const int allXHits[SciFi::Tracking::max_x_hits],
const int otherHits[SciFi::Constants::n_layers][SciFi::Tracking::max_other_hits]);
__host__ __device__ float getLineFitDistance(
SciFi::Tracking::LineFitterPars& parameters,
const SciFi::Hits& scifi_hits,
float coordX[SciFi::Tracking::max_x_hits],
int allXHits[SciFi::Tracking::max_x_hits],
int it);
__host__ __device__ float getLineFitChi2(
SciFi::Tracking::LineFitterPars& parameters,
const SciFi::Hits& scifi_hits,
float coordX[SciFi::Tracking::max_x_hits],
int allXHits[SciFi::Tracking::max_x_hits],
int it);
__host__ __device__ void solveLineFit(SciFi::Tracking::LineFitterPars& parameters);
__host__ __device__ void fastLinearFit(
const SciFi::Hits& scifi_hits,
float trackParameters[SciFi::Tracking::nTrackParams],
int coordToFit[SciFi::Tracking::max_coordToFit],
int& n_coordToFit,
PlaneCounter planeCounter,
SciFi::Tracking::HitSearchCuts& pars_cur);
#pragma once
#include <cmath>
#include <array>
#include <vector>
#include <algorithm>
#include <fstream>
#include "SciFiDefinitions.cuh"
#include "TrackUtils.cuh"
#include "SciFiEventModel.cuh"
__host__ __device__ int fitParabola(
int* coordToFit,
const int n_coordToFit,
const SciFi::Hits& scifi_hits,
float trackParameters[SciFi::Tracking::nTrackParams],
const bool xFit);
#pragma once
#include "PrForwardTools.cuh"
#include "Handler.cuh"
// #include "ArgumentsCommon.cuh"
#include "ArgumentsVelo.cuh"
#include "ArgumentsUT.cuh"
#include "ArgumentsSciFi.cuh"
/** @class PrForward PrForward.h
*
* - InputTracksName: Input location for VeloUT tracks
* - OutputTracksName: Output location for Forward tracks
* Based on code written by
* 2012-03-20 : Olivier Callot
* 2013-03-15 : Thomas Nikodem
* 2015-02-13 : Sevda Esen [additional search in the triangles by Marian Stahl]
* 2016-03-09 : Thomas Nikodem [complete restructuring]
* 2018-08 : Vava Gligorov [extract code from Rec, make compile within GPU framework
* 2018-09 : Dorothea vom Bruch [convert to CUDA, runs on GPU]
*/
__global__ void scifi_pr_forward(
uint32_t* dev_scifi_hits,
const uint32_t* dev_scifi_hit_count,
const int* dev_atomics_velo,
const uint* dev_velo_track_hit_number,
const char* dev_velo_states,
const int* dev_atomics_ut,
const char* dev_ut_track_hits,
const uint* dev_ut_track_hit_number,
const float* dev_ut_qop,
const uint* dev_ut_track_velo_indices,
SciFi::TrackHits* dev_scifi_tracks,
int* dev_atomics_scifi,
const SciFi::Tracking::TMVA