diff --git a/cuda/veloUT/PrVeloUT/include/PrVeloUT.cuh b/cuda/veloUT/PrVeloUT/include/PrVeloUT.cuh
index 601175d74e6956d0817c9416ab40c32866bc85d6..69b84bd8ede79a5a3587398a6ef0790cb7a4b44a 100644
--- a/cuda/veloUT/PrVeloUT/include/PrVeloUT.cuh
+++ b/cuda/veloUT/PrVeloUT/include/PrVeloUT.cuh
@@ -58,7 +58,8 @@ __host__ __device__ bool getHits(
const int posLayers[4][85],
VeloUTTracking::HitsSoA *hits_layers,
const float* fudgeFactors,
- const VeloState& trState);
+ const VeloState& trState,
+ const float* ut_dxDy);
__host__ __device__ bool formClusters(
const int hitCandidatesInLayers[VeloUTTracking::n_layers][VeloUTTracking::max_hit_candidates_per_layer],
@@ -67,6 +68,7 @@ __host__ __device__ bool formClusters(
int hitCandidateIndices[VeloUTTracking::n_layers],
VeloUTTracking::HitsSoA *hits_layers,
TrackHelper& helper,
+ const float* ut_dxDy,
const bool forward);
__host__ __device__ void prepareOutputTrack(
@@ -105,6 +107,7 @@ __host__ __device__ void findHits(
VeloUTTracking::HitsSoA *hits_layers,
const int layer_offset,
const int i_layer,
+ const float* ut_dxDy,
const VeloState& myState,
const float xTolNormFact,
const float invNormFact,
@@ -125,16 +128,18 @@ __host__ __device__ void addHits(
const int hitCandidateIndices[VeloUTTracking::n_layers],
const int hitCandidatesInLayers[VeloUTTracking::n_layers][VeloUTTracking::max_hit_candidates_per_layer],
VeloUTTracking::HitsSoA *hits_layers,
- const int hitIndices[N] ) {
+ const int hitIndices[N],
+ const float* ut_dxDy) {
for ( int i_hit = 0; i_hit < N; ++i_hit ) {
const int hit_index = hitIndices[i_hit];
- const int planeCode = hits_layers->planeCode(hit_index);
+ const int planeCode = hits_layers->planeCode[hit_index];
const float ui = x_pos_layers[ planeCode ][ hitCandidateIndices[i_hit] ];
- const float ci = hits_layers->cosT(hit_index);
- const float z = hits_layers->zAtYEq0( hit_index );
+ const float dxDy = ut_dxDy[planeCode];
+ const float ci = hits_layers->cosT(hit_index, dxDy);
+ const float z = hits_layers->zAtYEq0[hit_index];
const float dz = 0.001*(z - PrVeloUTConst::zMidUT);
- const float wi = hits_layers->weight(hit_index);
+ const float wi = hits_layers->weight[hit_index];
mat[0] += wi * ci;
mat[1] += wi * ci * dz;
@@ -157,12 +162,12 @@ __host__ __device__ void addChi2s(
for ( int i_hit = 0; i_hit < N; ++i_hit ) {
const int hit_index = hitIndices[i_hit];
- const int planeCode = hits_layers->planeCode(hit_index);
- const float zd = hits_layers->zAtYEq0(hit_index);
+ const int planeCode = hits_layers->planeCode[hit_index];
+ const float zd = hits_layers->zAtYEq0[hit_index];
const float xd = xUTFit + xSlopeUTFit*(zd-PrVeloUTConst::zMidUT);
const float x = x_pos_layers[ planeCode ][ hitCandidateIndices[i_hit] ];
const float du = xd - x;
- chi2 += (du*du)*hits_layers->weight(hit_index);
+ chi2 += (du*du)*hits_layers->weight[hit_index];
}
}
@@ -174,7 +179,7 @@ __host__ __device__ int countHitsWithHighThreshold(
int nHighThres = 0;
for ( int i_hit = 0; i_hit < N; ++i_hit ) {
- nHighThres += hits_layers->highThreshold( hitIndices[i_hit] );
+ nHighThres += hits_layers->highThreshold[hitIndices[i_hit]];
}
return nHighThres;
}
@@ -187,7 +192,8 @@ __host__ __device__ void simpleFit(
const int hitCandidatesInLayers[VeloUTTracking::n_layers][VeloUTTracking::max_hit_candidates_per_layer],
VeloUTTracking::HitsSoA *hits_layers,
const int hitIndices[N],
- TrackHelper& helper ) {
+ TrackHelper& helper,
+ const float* ut_dxDy) {
assert( N==3||N==4 );
const int nHighThres = countHitsWithHighThreshold<N>(hitIndices, hits_layers);
@@ -209,7 +215,7 @@ __host__ __device__ void simpleFit(
float rhs[2] = { helper.wb* helper.xMidField, helper.wb*helper.xMidField*zDiff };
// then add to sum values from hits on track
- addHits<N>( mat, rhs, x_pos_layers, hitCandidateIndices, hitCandidatesInLayers, hits_layers, hitIndices );
+ addHits<N>( mat, rhs, x_pos_layers, hitCandidateIndices, hitCandidatesInLayers, hits_layers, hitIndices, ut_dxDy );
const float denom = 1. / (mat[0]*mat[2] - mat[1]*mat[1]);
const float xSlopeUTFit = 0.001*(mat[0]*rhs[1] - mat[1]*rhs[0]) * denom;
diff --git a/cuda/veloUT/PrVeloUT/include/VeloUT.cuh b/cuda/veloUT/PrVeloUT/include/VeloUT.cuh
index 6a001270872e5312a05c2bee4b461be116f29f97..76d43bf524f73401a15c14a6f354d3035a99b98c 100644
--- a/cuda/veloUT/PrVeloUT/include/VeloUT.cuh
+++ b/cuda/veloUT/PrVeloUT/include/VeloUT.cuh
@@ -13,5 +13,6 @@ __global__ void veloUT(
VeloState* dev_velo_states,
VeloUTTracking::TrackUT* dev_veloUT_tracks,
int* dev_atomics_veloUT,
- PrUTMagnetTool* dev_ut_magnet_tool
+ PrUTMagnetTool* dev_ut_magnet_tool,
+ float* dev_ut_dxDy
);
diff --git a/cuda/veloUT/PrVeloUT/src/PrVeloUT.cu b/cuda/veloUT/PrVeloUT/src/PrVeloUT.cu
index 197bb5851452ad448584f2232f454298f6976dd2..9e32530753b792ea812a4e6d93d2d4f59fd45f3e 100644
--- a/cuda/veloUT/PrVeloUT/src/PrVeloUT.cu
+++ b/cuda/veloUT/PrVeloUT/src/PrVeloUT.cu
@@ -50,7 +50,8 @@ __host__ __device__ bool getHits(
const int posLayers[4][85],
VeloUTTracking::HitsSoA *hits_layers,
const float* fudgeFactors,
- const VeloState& trState)
+ const VeloState& trState,
+ const float* ut_dxDy)
{
// -- This is hardcoded, so faster
// -- If you ever change the Table in the magnet tool, this will be wrong
@@ -86,8 +87,8 @@ __host__ __device__ bool getHits(
int layer_offset = hits_layers->layer_offset[layer];
if( hits_layers->n_hits_layers[layer] == 0 ) continue;
- const float dxDy = hits_layers->dxDy(layer_offset + 0);
- const float zLayer = hits_layers->zAtYEq0(layer_offset + 0);
+ const float dxDy = ut_dxDy[layer];
+ const float zLayer = hits_layers->zAtYEq0[layer_offset + 0];
const float yAtZ = trState.y + trState.ty*(zLayer - trState.z);
const float xLayer = trState.x + trState.tx*(zLayer - trState.z);
@@ -110,12 +111,12 @@ __host__ __device__ bool getHits(
size_t posBeg = posLayers[layer][ indexLow ];
size_t posEnd = posLayers[layer][ indexHi ];
- while ( (hits_layers->xAtYEq0(layer_offset + posBeg) < lowerBoundX) && (posBeg != hits_layers->n_hits_layers[layer] ) )
+ while ( (hits_layers->xAtYEq0[layer_offset + posBeg] < lowerBoundX) && (posBeg != hits_layers->n_hits_layers[layer] ) )
++posBeg;
if (posBeg == hits_layers->n_hits_layers[layer]) continue;
findHits(posBeg, posEnd,
- hits_layers, layer_offset, layer,
+ hits_layers, layer_offset, layer, ut_dxDy,
trState, xTol*invNormFact, invNormFact,
hitCandidatesInLayers[layer], n_hitCandidatesInLayers[layer],
x_pos_layers);
@@ -137,6 +138,7 @@ __host__ __device__ bool formClusters(
int bestHitCandidateIndices[VeloUTTracking::n_layers],
VeloUTTracking::HitsSoA *hits_layers,
TrackHelper& helper,
+ const float* ut_dxDy,
const bool forward)
{
@@ -158,7 +160,7 @@ __host__ __device__ bool formClusters(
const int layer_offset0 = hits_layers->layer_offset[ layers[0] ];
const int hit_index0 = layer_offset0 + hitCandidatesInLayers[ layers[0] ][i_hit0];
const float xhitLayer0 = x_pos_layers[layers[0]][i_hit0];
- const float zhitLayer0 = hits_layers->zAtYEq0( hit_index0 );
+ const float zhitLayer0 = hits_layers->zAtYEq0[hit_index0];
hitCandidateIndices[0] = i_hit0;
for ( int i_hit2 = 0; i_hit2 < n_hitCandidatesInLayers[ layers[2] ]; ++i_hit2 ) {
@@ -166,7 +168,7 @@ __host__ __device__ bool formClusters(
const int layer_offset2 = hits_layers->layer_offset[ layers[2] ];
const int hit_index2 = layer_offset2 + hitCandidatesInLayers[ layers[2] ][i_hit2];
const float xhitLayer2 = x_pos_layers[layers[2]][i_hit2];
- const float zhitLayer2 = hits_layers->zAtYEq0( hit_index2 );
+ const float zhitLayer2 = hits_layers->zAtYEq0[hit_index2];
hitCandidateIndices[2] = i_hit2;
const float tx = (xhitLayer2 - xhitLayer0)/(zhitLayer2 - zhitLayer0);
@@ -179,7 +181,7 @@ __host__ __device__ bool formClusters(
const int layer_offset1 = hits_layers->layer_offset[ layers[1] ];
const int hit_index1 = layer_offset1 + hitCandidatesInLayers[ layers[1] ][i_hit1];
const float xhitLayer1 = x_pos_layers[layers[1]][i_hit1];
- const float zhitLayer1 = hits_layers->zAtYEq0( hit_index1 );
+ const float zhitLayer1 = hits_layers->zAtYEq0[hit_index1];
const float xextrapLayer1 = xhitLayer0 + tx*(zhitLayer1-zhitLayer0);
if(std::abs(xhitLayer1 - xextrapLayer1) < hitTol){
@@ -198,7 +200,7 @@ __host__ __device__ bool formClusters(
const int layer_offset3 = hits_layers->layer_offset[ layers[3] ];
const int hit_index3 = layer_offset3 + hitCandidatesInLayers[ layers[3] ][i_hit3];
const float xhitLayer3 = x_pos_layers[layers[3]][i_hit3];
- const float zhitLayer3 = hits_layers->zAtYEq0( hit_index3 );
+ const float zhitLayer3 = hits_layers->zAtYEq0[hit_index3];
const float xextrapLayer3 = xhitLayer2 + tx*(zhitLayer3-zhitLayer2);
if(std::abs(xhitLayer3 - xextrapLayer3) < hitTol){
@@ -211,7 +213,7 @@ __host__ __device__ bool formClusters(
// -- All hits found
if ( IndexBestHit1 > 0 && IndexBestHit3 > 0 ) {
const int hitIndices[4] = {hit_index0, IndexBestHit1, hit_index2, IndexBestHit3};
- simpleFit<4>(x_pos_layers, hitCandidateIndices, bestHitCandidateIndices, hitCandidatesInLayers, hits_layers, hitIndices, helper);
+ simpleFit<4>(x_pos_layers, hitCandidateIndices, bestHitCandidateIndices, hitCandidatesInLayers, hits_layers, hitIndices, helper, ut_dxDy);
if(!fourLayerSolution && helper.n_hits > 0){
fourLayerSolution = true;
@@ -222,14 +224,14 @@ __host__ __device__ bool formClusters(
// -- Nothing found in layer 3
if( !fourLayerSolution && IndexBestHit1 > 0 ){
const int hitIndices[3] = {hit_index0, IndexBestHit1, hit_index2};
- simpleFit<3>(x_pos_layers, hitCandidateIndices, bestHitCandidateIndices, hitCandidatesInLayers, hits_layers, hitIndices, helper);
+ simpleFit<3>(x_pos_layers, hitCandidateIndices, bestHitCandidateIndices, hitCandidatesInLayers, hits_layers, hitIndices, helper, ut_dxDy);
continue;
}
// -- Nothing found in layer 1
if( !fourLayerSolution && IndexBestHit3 > 0 ){
hitCandidateIndices[1] = hitCandidateIndices[3]; // hit3 saved in second position of hits4fit
const int hitIndices[3] = {hit_index0, IndexBestHit3, hit_index2};
- simpleFit<3>(x_pos_layers, hitCandidateIndices, bestHitCandidateIndices, hitCandidatesInLayers, hits_layers, hitIndices, helper);
+ simpleFit<3>(x_pos_layers, hitCandidateIndices, bestHitCandidateIndices, hitCandidatesInLayers, hits_layers, hitIndices, helper, ut_dxDy);
continue;
}
@@ -327,17 +329,17 @@ __host__ __device__ void prepareOutputTrack(
for ( int i_hit = 0; i_hit < helper.n_hits; ++i_hit ) {
const int hit_index = helper.bestHitIndices[i_hit];
- track.addLHCbID( hits_layers->LHCbID(hit_index) );
+ track.addLHCbID( hits_layers->LHCbID[hit_index] );
assert( track.hitsNum < VeloUTTracking::max_track_size);
- const int planeCode = hits_layers->planeCode(hit_index);
+ const int planeCode = hits_layers->planeCode[hit_index];
const float xhit = x_pos_layers[ planeCode ][ hitCandidateIndices[i_hit] ];
- const float zhit = hits_layers->zAtYEq0( hit_index );
+ const float zhit = hits_layers->zAtYEq0[hit_index];
const int layer_offset = hits_layers->layer_offset[ planeCode ];
for ( int i_ohit = 0; i_ohit < n_hitCandidatesInLayers[planeCode]; ++i_ohit ) {
const int ohit_index = hitCandidatesInLayers[planeCode][i_ohit];
- const float zohit = hits_layers->zAtYEq0( layer_offset + ohit_index );
+ const float zohit = hits_layers->zAtYEq0[layer_offset + ohit_index];
if(zohit==zhit) continue;
@@ -346,7 +348,7 @@ __host__ __device__ void prepareOutputTrack(
if( xohit-xextrap < -PrVeloUTConst::overlapTol) continue;
if( xohit-xextrap > PrVeloUTConst::overlapTol) break;
- track.addLHCbID( hits_layers->LHCbID(layer_offset + ohit_index) );
+ track.addLHCbID( hits_layers->LHCbID[layer_offset + ohit_index] );
assert( track.hitsNum < VeloUTTracking::max_track_size);
// -- only one overlap hit
@@ -414,7 +416,7 @@ __host__ __device__ void fillIterators(
// TODO add bounds checking
for ( ; pos != hits_layers->n_hits_layers[layer]; ++pos) {
- while( hits_layers->xAtYEq0( layer_offset + pos ) > val){
+ while( hits_layers->xAtYEq0[layer_offset + pos] > val){
posLayers[layer][bound+42] = pos;
++bound;
val = std::copysign(float(bound*bound)/2.0, bound);
@@ -440,8 +442,8 @@ __host__ __device__ void findHits(
const size_t posEnd,
VeloUTTracking::HitsSoA *hits_layers,
const int layer_offset,
- // to do: pass array for this layer -> get rid of i_layer index
const int i_layer,
+ const float* ut_dxDy,
const VeloState& myState,
const float xTolNormFact,
const float invNormFact,
@@ -449,7 +451,7 @@ __host__ __device__ void findHits(
int &n_hitCandidatesInLayer,
float x_pos_layers[VeloUTTracking::n_layers][VeloUTTracking::max_hit_candidates_per_layer])
{
- const auto zInit = hits_layers->zAtYEq0( layer_offset + posBeg );
+ const auto zInit = hits_layers->zAtYEq0[layer_offset + posBeg];
const auto yApprox = myState.y + myState.ty * (zInit - myState.z);
size_t pos = posBeg;
@@ -462,8 +464,8 @@ __host__ __device__ void findHits(
const auto yyProto = myState.y - myState.ty*myState.z;
for (int i=pos; i<posEnd; ++i) {
-
- const auto xx = hits_layers->xAt( layer_offset + i, yApprox );
+ const float dxDy = ut_dxDy[i_layer];
+ const auto xx = hits_layers->xAt( layer_offset + i, yApprox, dxDy );
const auto dx = xx - xOnTrackProto;
if( dx < -xTolNormFact ) continue;
@@ -472,9 +474,9 @@ __host__ __device__ void findHits(
// -- Now refine the tolerance in Y
if ( hits_layers->isNotYCompatible( layer_offset + i, yApprox, PrVeloUTConst::yTol + PrVeloUTConst::yTolSlope * std::abs(dx*invNormFact)) ) continue;
- const auto zz = hits_layers->zAtYEq0( layer_offset + i );
+ const auto zz = hits_layers->zAtYEq0[layer_offset + i];
const auto yy = yyProto + myState.ty*zz;
- const auto xx2 = hits_layers->xAt( layer_offset + i, yy );
+ const auto xx2 = hits_layers->xAt( layer_offset + i, yy, dxDy );
hitCandidatesInLayer[n_hitCandidatesInLayer] = i;
x_pos_layers[i_layer][n_hitCandidatesInLayer] = xx2;
diff --git a/cuda/veloUT/PrVeloUT/src/VeloUT.cu b/cuda/veloUT/PrVeloUT/src/VeloUT.cu
index 65ec38a325e825af67a536ca82b2910f9e7fc318..bfadd1720122dc37d0bedccdfe1fa89759ea33e2 100644
--- a/cuda/veloUT/PrVeloUT/src/VeloUT.cu
+++ b/cuda/veloUT/PrVeloUT/src/VeloUT.cu
@@ -8,7 +8,8 @@ __global__ void veloUT(
VeloState* dev_velo_states,
VeloUTTracking::TrackUT* dev_veloUT_tracks,
int* dev_atomics_veloUT,
- PrUTMagnetTool* dev_ut_magnet_tool
+ PrUTMagnetTool* dev_ut_magnet_tool,
+ float* dev_ut_dxDy
) {
const int number_of_events = gridDim.x;
@@ -77,7 +78,8 @@ __global__ void veloUT(
posLayers,
hits_layers_event,
fudgeFactors,
- velo_states_event[i_track] )
+ velo_states_event[i_track],
+ dev_ut_dxDy )
) continue;
TrackHelper helper(velo_states_event[i_track]);
@@ -93,6 +95,7 @@ __global__ void veloUT(
hitCandidateIndices,
hits_layers_event,
helper,
+ dev_ut_dxDy,
true )){
// go through UT layers in backward direction
@@ -103,6 +106,7 @@ __global__ void veloUT(
hitCandidateIndices,
hits_layers_event,
helper,
+ dev_ut_dxDy,
false);
}
diff --git a/cuda/veloUT/common/include/VeloUTDefinitions.cuh b/cuda/veloUT/common/include/VeloUTDefinitions.cuh
index 322b39c24277bf917cec82b4c33c6ccbfa539f72..76465722fc8cfa2a5f2bd1c5022972b7bc0edde4 100644
--- a/cuda/veloUT/common/include/VeloUTDefinitions.cuh
+++ b/cuda/veloUT/common/include/VeloUTDefinitions.cuh
@@ -17,12 +17,7 @@ namespace VeloUTTracking {
*/
static constexpr uint n_layers = 4;
static constexpr uint n_ut_hit_variables = 8;
- /* For now, the planeCode is an attribute of every hit,
- -> check which information
- needs to be saved for the forward tracking
- */
- static constexpr int planeCode[n_layers] = {0, 1, 2, 3};
-
+
/* Cut-offs */
static constexpr uint max_numhits_per_layer = 500;
static constexpr uint max_numhits_per_event = 6000;
@@ -39,49 +34,21 @@ namespace VeloUTTracking {
int layer_offset[n_layers];
int n_hits_layers[n_layers];
- float m_cos[max_numhits_per_event];
- float m_yBegin[max_numhits_per_event];
- float m_yEnd[max_numhits_per_event];
- float m_zAtYEq0[max_numhits_per_event];
- float m_xAtYEq0[max_numhits_per_event];
- float m_weight[max_numhits_per_event];
- int m_highThreshold[max_numhits_per_event];
- unsigned int m_LHCbID[max_numhits_per_event];
- int m_planeCode[max_numhits_per_event];
+ float yBegin[max_numhits_per_event];
+ float yEnd[max_numhits_per_event];
+ float zAtYEq0[max_numhits_per_event];
+ float xAtYEq0[max_numhits_per_event];
+ float weight[max_numhits_per_event];
+ int highThreshold[max_numhits_per_event];
+ unsigned int LHCbID[max_numhits_per_event];
+ int planeCode[max_numhits_per_event];
- __host__ __device__ inline float cos(const int i_hit) const { return m_cos[i_hit]; }
- __host__ __device__ inline int planeCode(const int i_hit) const { return m_planeCode[i_hit]; }
- // layer configuration: XUVX, U and V layers tilted by +/- 5 degrees = 0.087 radians
- __host__ __device__ inline float dxDy(const int i_hit) const {
- const int i_plane = m_planeCode[i_hit];
- if ( i_plane == 0 || i_plane == 3 )
- return 0.;
- else if ( i_plane == 1 )
- return 0.08748867;
- else if ( i_plane == 2 )
- return -0.08748867;
- else return -1;
- }
- __host__ __device__ inline float cosT(const int i_hit) const { return ( std::fabs( m_xAtYEq0[i_hit] ) < 1.0E-9 ) ? 1. / std::sqrt( 1 + dxDy(i_hit) * dxDy(i_hit) ) : cos(i_hit); }
- __host__ __device__ inline bool highThreshold(const int i_hit) const { return m_highThreshold[i_hit]; }
- __host__ __device__ inline bool isYCompatible( const int i_hit, const float y, const float tol ) const { return yMin(i_hit) - tol <= y && y <= yMax(i_hit) + tol; }
+ __host__ __device__ inline float cosT(const int i_hit, const float dxDy) const { return ( std::fabs( xAtYEq0[i_hit] ) < 1.0E-9 ) ? 1. / std::sqrt( 1 + dxDy * dxDy ) : std::cos(dxDy); }
__host__ __device__ inline bool isNotYCompatible( const int i_hit, const float y, const float tol ) const { return yMin(i_hit) - tol > y || y > yMax(i_hit) + tol; }
- __host__ __device__ inline int LHCbID(const int i_hit) const { return m_LHCbID[i_hit]; }
- __host__ __device__ inline float sinT(const int i_hit) const { return tanT(i_hit) * cosT(i_hit); }
- __host__ __device__ inline float tanT(const int i_hit) const { return -1 * dxDy(i_hit); }
- __host__ __device__ inline float weight(const int i_hit) const { return m_weight[i_hit]; }
- __host__ __device__ inline float xAt( const int i_hit, const float globalY ) const { return m_xAtYEq0[i_hit] + globalY * dxDy(i_hit); }
- __host__ __device__ inline float xAtYEq0(const int i_hit) const { return m_xAtYEq0[i_hit]; }
- __host__ __device__ inline float xMax(const int i_hit) const { return std::max( xAt( i_hit, yBegin(i_hit) ), xAt( i_hit, yEnd(i_hit) ) ); }
- __host__ __device__ inline float xMin(const int i_hit) const { return std::min( xAt( i_hit, yBegin(i_hit) ), xAt( i_hit, yEnd(i_hit) ) ); }
- __host__ __device__ inline float xT(const int i_hit) const { return cos(i_hit); }
- __host__ __device__ inline float yBegin(const int i_hit) const { return m_yBegin[i_hit]; }
- __host__ __device__ inline float yEnd(const int i_hit) const { return m_yEnd[i_hit]; }
- __host__ __device__ inline float yMax(const int i_hit) const { return std::max( yBegin(i_hit), yEnd(i_hit) ); }
- __host__ __device__ inline float yMid(const int i_hit) const { return 0.5 * ( yBegin(i_hit) + yEnd(i_hit) ); }
- __host__ __device__ inline float yMin(const int i_hit) const { return std::min( yBegin(i_hit), yEnd(i_hit) ); }
- __host__ __device__ inline float zAtYEq0(const int i_hit) const { return m_zAtYEq0[i_hit]; }
+ __host__ __device__ inline float xAt( const int i_hit, const float globalY, const float dxDy ) const { return xAtYEq0[i_hit] + globalY * dxDy; }
+ __host__ __device__ inline float yMax(const int i_hit) const { return std::max( yBegin[i_hit], yEnd[i_hit] ); }
+ __host__ __device__ inline float yMin(const int i_hit) const { return std::min( yBegin[i_hit], yEnd[i_hit] ); }
};
diff --git a/cuda/veloUT/sorting/src/SortByX.cu b/cuda/veloUT/sorting/src/SortByX.cu
index afe37f8cfa24b8a64354e3ac624f10f429a210fa..6f36c7676b3163186309d894675a97075f7b9e55 100644
--- a/cuda/veloUT/sorting/src/SortByX.cu
+++ b/cuda/veloUT/sorting/src/SortByX.cu
@@ -32,7 +32,7 @@ __global__ void sort_by_x(
}
find_permutation<float>(
- hits_layers->m_xAtYEq0,
+ hits_layers->xAtYEq0,
layer_offset,
hit_permutations,
n_hits
@@ -40,15 +40,14 @@ __global__ void sort_by_x(
__syncthreads();
- apply_permutation<float>( hit_permutations, layer_offset, n_hits, hits_layers->m_cos, hits_layers_sorted->m_cos );
- apply_permutation<float>( hit_permutations, layer_offset, n_hits, hits_layers->m_weight, hits_layers_sorted->m_weight );
- apply_permutation<float>( hit_permutations, layer_offset, n_hits, hits_layers->m_xAtYEq0, hits_layers_sorted->m_xAtYEq0 );
- apply_permutation<float>( hit_permutations, layer_offset, n_hits, hits_layers->m_yBegin, hits_layers_sorted->m_yBegin );
- apply_permutation<float>( hit_permutations, layer_offset, n_hits, hits_layers->m_yEnd, hits_layers_sorted->m_yEnd );
- apply_permutation<float>( hit_permutations, layer_offset, n_hits, hits_layers->m_zAtYEq0, hits_layers_sorted->m_zAtYEq0 );
- apply_permutation<unsigned int>( hit_permutations, layer_offset, n_hits, hits_layers->m_LHCbID, hits_layers_sorted->m_LHCbID );
- apply_permutation<int>( hit_permutations, layer_offset, n_hits, hits_layers->m_planeCode, hits_layers_sorted->m_planeCode );
- apply_permutation<int>( hit_permutations, layer_offset, n_hits, hits_layers->m_highThreshold, hits_layers_sorted->m_highThreshold );
+ apply_permutation<float>( hit_permutations, layer_offset, n_hits, hits_layers->weight, hits_layers_sorted->weight );
+ apply_permutation<float>( hit_permutations, layer_offset, n_hits, hits_layers->xAtYEq0, hits_layers_sorted->xAtYEq0 );
+ apply_permutation<float>( hit_permutations, layer_offset, n_hits, hits_layers->yBegin, hits_layers_sorted->yBegin );
+ apply_permutation<float>( hit_permutations, layer_offset, n_hits, hits_layers->yEnd, hits_layers_sorted->yEnd );
+ apply_permutation<float>( hit_permutations, layer_offset, n_hits, hits_layers->zAtYEq0, hits_layers_sorted->zAtYEq0 );
+ apply_permutation<unsigned int>( hit_permutations, layer_offset, n_hits, hits_layers->LHCbID, hits_layers_sorted->LHCbID );
+ apply_permutation<int>( hit_permutations, layer_offset, n_hits, hits_layers->planeCode, hits_layers_sorted->planeCode );
+ apply_permutation<int>( hit_permutations, layer_offset, n_hits, hits_layers->highThreshold, hits_layers_sorted->highThreshold );
}
}
diff --git a/main/src/Tools.cpp b/main/src/Tools.cpp
index 080bcd2003225daec40d470089890250c89d41b1..af8309477413e9732eb61a7b5693de2c38c8b18f 100644
--- a/main/src/Tools.cpp
+++ b/main/src/Tools.cpp
@@ -84,28 +84,28 @@ void read_ut_events_into_arrays( VeloUTTracking::HitsSoA *hits_layers_events,
// then the hit variables, sorted by layer
for ( int i_layer = 0; i_layer < VeloUTTracking::n_layers; ++i_layer ) {
int layer_offset = hits_layers->layer_offset[i_layer];
- std::copy_n((float*) raw_input, hits_layers->n_hits_layers[i_layer], &( hits_layers->m_cos[ layer_offset ]) );
+ // to do: change tracker dumper to not dump cos
+ //std::copy_n((float*) raw_input, hits_layers->n_hits_layers[i_layer], &( hits_layers->cos[ layer_offset ]) );
raw_input += sizeof(float) * hits_layers->n_hits_layers[i_layer];
- std::copy_n((float*) raw_input, hits_layers->n_hits_layers[i_layer], &(hits_layers->m_yBegin[ layer_offset ]) );
+ std::copy_n((float*) raw_input, hits_layers->n_hits_layers[i_layer], &(hits_layers->yBegin[ layer_offset ]) );
raw_input += sizeof(float) * hits_layers->n_hits_layers[i_layer];
- std::copy_n((float*) raw_input, hits_layers->n_hits_layers[i_layer], &(hits_layers->m_yEnd[ layer_offset ]) );
+ std::copy_n((float*) raw_input, hits_layers->n_hits_layers[i_layer], &(hits_layers->yEnd[ layer_offset ]) );
raw_input += sizeof(float) * hits_layers->n_hits_layers[i_layer];
// to do: change tracker dumper to not dump dxDy
- //std::copy_n((float*) raw_input, hits_layers->n_hits_layers[i_layer], &(hits_layers->m_dxDy[ layer_offset ]) );
raw_input += sizeof(float) * hits_layers->n_hits_layers[i_layer];
- std::copy_n((float*) raw_input, hits_layers->n_hits_layers[i_layer], &(hits_layers->m_zAtYEq0[ layer_offset ]) );
+ std::copy_n((float*) raw_input, hits_layers->n_hits_layers[i_layer], &(hits_layers->zAtYEq0[ layer_offset ]) );
raw_input += sizeof(float) * hits_layers->n_hits_layers[i_layer];
- std::copy_n((float*) raw_input, hits_layers->n_hits_layers[i_layer], &(hits_layers->m_xAtYEq0[ layer_offset ]) );
+ std::copy_n((float*) raw_input, hits_layers->n_hits_layers[i_layer], &(hits_layers->xAtYEq0[ layer_offset ]) );
raw_input += sizeof(float) * hits_layers->n_hits_layers[i_layer];
- std::copy_n((float*) raw_input, hits_layers->n_hits_layers[i_layer], &(hits_layers->m_weight[ layer_offset ]) );
+ std::copy_n((float*) raw_input, hits_layers->n_hits_layers[i_layer], &(hits_layers->weight[ layer_offset ]) );
raw_input += sizeof(float) * hits_layers->n_hits_layers[i_layer];
- std::copy_n((int*) raw_input, hits_layers->n_hits_layers[i_layer], &(hits_layers->m_highThreshold[ layer_offset ]) );
+ std::copy_n((int*) raw_input, hits_layers->n_hits_layers[i_layer], &(hits_layers->highThreshold[ layer_offset ]) );
raw_input += sizeof(int) * hits_layers->n_hits_layers[i_layer];
- std::copy_n((unsigned int*) raw_input, hits_layers->n_hits_layers[i_layer], &(hits_layers->m_LHCbID[ layer_offset ]) );
+ std::copy_n((unsigned int*) raw_input, hits_layers->n_hits_layers[i_layer], &(hits_layers->LHCbID[ layer_offset ]) );
raw_input += sizeof(unsigned int) * hits_layers->n_hits_layers[i_layer];
for ( int i_hit = 0; i_hit < hits_layers->n_hits_layers[i_layer]; ++i_hit ) {
- hits_layers->m_planeCode[ layer_offset + i_hit ] = i_layer;
+ hits_layers->planeCode[ layer_offset + i_hit ] = i_layer;
}
}
@@ -128,17 +128,15 @@ void check_ut_events( const VeloUTTracking::HitsSoA *hits_layers_events,
number_of_hits += hits_layers.n_hits_layers[i_layer];
int layer_offset = hits_layers.layer_offset[i_layer];
for ( int i_hit = 0; i_hit < 3; ++i_hit ) {
- printf("\t at hit %u, cos = %f, yBegin = %f, yEnd = %f, zAtyEq0 = %f, xAtyEq0 = %f, weight = %f, highThreshold = %u, LHCbID = %u, dxDy = %f \n",
+ printf("\t at hit %u, yBegin = %f, yEnd = %f, zAtyEq0 = %f, xAtyEq0 = %f, weight = %f, highThreshold = %u, LHCbID = %u \n",
i_hit,
- hits_layers.m_cos[ layer_offset + i_hit ],
- hits_layers.m_yBegin[ layer_offset + i_hit ],
- hits_layers.m_yEnd[ layer_offset + i_hit ],
- hits_layers.m_zAtYEq0[ layer_offset + i_hit ],
- hits_layers.m_xAtYEq0[ layer_offset + i_hit ],
- hits_layers.m_weight[ layer_offset + i_hit ],
- hits_layers.m_highThreshold[ layer_offset + i_hit ],
- hits_layers.m_LHCbID[ layer_offset + i_hit ],
- hits_layers.dxDy( layer_offset + i_hit ) );
+ hits_layers.yBegin[ layer_offset + i_hit ],
+ hits_layers.yEnd[ layer_offset + i_hit ],
+ hits_layers.zAtYEq0[ layer_offset + i_hit ],
+ hits_layers.xAtYEq0[ layer_offset + i_hit ],
+ hits_layers.weight[ layer_offset + i_hit ],
+ hits_layers.highThreshold[ layer_offset + i_hit ],
+ hits_layers.LHCbID[ layer_offset + i_hit ]);
}
}
diff --git a/main/src/main.cpp b/main/src/main.cpp
index 936ba7e295b737c1fbf994edf14747ace91d6b08..36622641c4ef2bd1ac262011885ccd0acbf0a37c 100644
--- a/main/src/main.cpp
+++ b/main/src/main.cpp
@@ -248,8 +248,8 @@ int main(int argc, char *argv[])
// Initialize detector constants on GPU
- GpuConstants gpu_constants;
- gpu_constants.reserve_and_initialize();
+ Constants constants;
+ constants.reserve_and_initialize();
// Create streams
StreamWrapper stream_wrapper;
@@ -266,7 +266,7 @@ int main(int argc, char *argv[])
folder_name_MC,
start_event_offset,
reserve_mb,
- gpu_constants
+ constants
);
// Attempt to execute all in one go
diff --git a/stream/CMakeLists.txt b/stream/CMakeLists.txt
index a57db20e1b6450c22e631872d6fbb088ffb1cc0e..36876dde8ad3921bedae5e710819ca843bde3d95 100644
--- a/stream/CMakeLists.txt
+++ b/stream/CMakeLists.txt
@@ -41,6 +41,7 @@ cuda_add_library(Stream STATIC
if ( ROOT_FOUND )
target_compile_definitions(Stream PUBLIC WITH_ROOT)
+ message("-- Found ROOT, setting WITH_ROOT to true")
target_link_libraries(Stream
Velo
VeloUT
diff --git a/stream/sequence/include/Constants.cuh b/stream/sequence/include/Constants.cuh
index 54f469dc122dd2f7f0690b781ac92c4f45532d6c..bd3d08107704436495c981928cab69e581a2cc91 100644
--- a/stream/sequence/include/Constants.cuh
+++ b/stream/sequence/include/Constants.cuh
@@ -6,6 +6,7 @@
#include "VeloDefinitions.cuh"
#include "ClusteringDefinitions.cuh"
#include "ClusteringCommon.h"
+#include "VeloUTDefinitions.cuh"
/**
* @brief Struct intended as a singleton with constants defined on GPU.
@@ -16,13 +17,15 @@
*
* The pointers are hard-coded. Feel free to write more as needed.
*/
-struct GpuConstants {
+struct Constants {
float* dev_velo_module_zs;
uint8_t* dev_velo_candidate_ks;
uint8_t* dev_velo_sp_patterns;
float* dev_velo_sp_fx;
float* dev_velo_sp_fy;
-
+ float* dev_ut_dxDy;
+ float host_ut_dxDy[VeloUTTracking::n_layers];
+
void reserve_and_initialize() {
reserve_constants();
initialize_constants();
diff --git a/stream/sequence/include/Stream.cuh b/stream/sequence/include/Stream.cuh
index d8b7d57857f88cbc82602dc90f9f601a8bb6ca27..1d50d4082935b21de716cc73151e20fdea3aac32 100644
--- a/stream/sequence/include/Stream.cuh
+++ b/stream/sequence/include/Stream.cuh
@@ -69,8 +69,8 @@ struct Stream {
std::string folder_name_MC;
uint start_event_offset;
- // Gpu constants
- GpuConstants gpu_constants;
+ // Constants
+ Constants constants;
cudaError_t initialize(
const std::vector<char>& velopix_geometry,
@@ -85,7 +85,7 @@ struct Stream {
const uint param_start_event_offset,
const size_t param_reserve_mb,
const uint param_stream_number,
- const GpuConstants& param_gpu_constants
+ const Constants& param_constants
);
cudaError_t run_sequence(
diff --git a/stream/sequence/include/StreamWrapper.cuh b/stream/sequence/include/StreamWrapper.cuh
index b3c15092bb49ba9c1fb4a089a0350b75733f28a8..a7c928afc139dafbab5d835783c4e941c9f51fcb 100644
--- a/stream/sequence/include/StreamWrapper.cuh
+++ b/stream/sequence/include/StreamWrapper.cuh
@@ -43,7 +43,7 @@ struct StreamWrapper {
const std::string& folder_name_MC,
const uint start_event_offset,
const size_t reserve_mb,
- const GpuConstants& gpu_constants
+ const Constants& constants
);
/**
diff --git a/stream/sequence/src/Constants.cu b/stream/sequence/src/Constants.cu
index 4bf13ecb4a7aed43b8f8e4bd63bf12459a2a7218..279fef65c0d59d7483f3772539b1ef31b40ceedb 100644
--- a/stream/sequence/src/Constants.cu
+++ b/stream/sequence/src/Constants.cu
@@ -1,14 +1,15 @@
#include "Constants.cuh"
-void GpuConstants::reserve_constants() {
+void Constants::reserve_constants() {
cudaCheck(cudaMalloc((void**)&dev_velo_module_zs, VeloTracking::n_modules * sizeof(float)));
cudaCheck(cudaMalloc((void**)&dev_velo_candidate_ks, 9 * sizeof(uint8_t)));
cudaCheck(cudaMalloc((void**)&dev_velo_sp_patterns, 256 * sizeof(uint8_t)));
cudaCheck(cudaMalloc((void**)&dev_velo_sp_fx, 512 * sizeof(float)));
cudaCheck(cudaMalloc((void**)&dev_velo_sp_fy, 512 * sizeof(float)));
+ cudaCheck(cudaMalloc((void**)&dev_ut_dxDy, VeloUTTracking::n_layers * sizeof(float)));
}
-void GpuConstants::initialize_constants() {
+void Constants::initialize_constants() {
// Velo module constants
const std::array<float, VeloTracking::n_modules> velo_module_zs = {-287.5, -275, -262.5, -250, -237.5, -225, -212.5, \
-200, -137.5, -125, -62.5, -50, -37.5, -25, -12.5, 0, 12.5, 25, 37.5, 50, 62.5, 75, 87.5, 100, \
@@ -31,4 +32,13 @@ void GpuConstants::initialize_constants() {
cudaCheck(cudaMemcpy(dev_velo_sp_patterns, sp_patterns.data(), sp_patterns.size(), cudaMemcpyHostToDevice));
cudaCheck(cudaMemcpy(dev_velo_sp_fx, sp_fx.data(), sp_fx.size() * sizeof(float), cudaMemcpyHostToDevice));
cudaCheck(cudaMemcpy(dev_velo_sp_fy, sp_fy.data(), sp_fy.size() * sizeof(float), cudaMemcpyHostToDevice));
+
+ // UT geometry constants
+ // layer configuration: XUVX, U and V layers tilted by +/- 5 degrees = 0.087 radians
+ host_ut_dxDy[0] = 0.;
+ host_ut_dxDy[1] = 0.08748867;
+ host_ut_dxDy[2] = -0.0874886;
+ host_ut_dxDy[3] = 0.;
+
+ cudaCheck(cudaMemcpy(dev_ut_dxDy, host_ut_dxDy, VeloUTTracking::n_layers * sizeof(float), cudaMemcpyHostToDevice));
}
diff --git a/stream/sequence/src/Stream.cu b/stream/sequence/src/Stream.cu
index 508d9466e398889b288fc01291a340fbcf4dae20..a848da0f5fa8bd1d7d629729e6c455f0d878201d 100644
--- a/stream/sequence/src/Stream.cu
+++ b/stream/sequence/src/Stream.cu
@@ -16,7 +16,7 @@ cudaError_t Stream::initialize(
const uint param_start_event_offset,
const size_t reserve_mb,
const uint param_stream_number,
- const GpuConstants& param_gpu_constants
+ const Constants& param_constants
) {
// Set stream and events
cudaCheck(cudaStreamCreate(&stream));
@@ -33,7 +33,7 @@ cudaError_t Stream::initialize(
run_on_x86 = param_run_on_x86;
folder_name_MC = param_folder_name_MC;
start_event_offset = param_start_event_offset;
- gpu_constants = param_gpu_constants;
+ constants = param_constants;
// Special case
// Populate velo geometry
diff --git a/stream/sequence/src/StreamSequence.cu b/stream/sequence/src/StreamSequence.cu
index 281f37bdc0f0318b9d90a841bab03221567f604e..fbe095a59768aebe90ca1b665e885a8ef64d13a8 100644
--- a/stream/sequence/src/StreamSequence.cu
+++ b/stream/sequence/src/StreamSequence.cu
@@ -45,7 +45,7 @@ cudaError_t Stream::run_sequence(
argen.generate<arg::dev_module_cluster_num>(argument_offsets),
argen.generate<arg::dev_module_candidate_num>(argument_offsets),
argen.generate<arg::dev_cluster_candidates>(argument_offsets),
- gpu_constants.dev_velo_candidate_ks
+ constants.dev_velo_candidate_ks
);
cudaCheck(cudaMemcpyAsync(argen.generate<arg::dev_raw_input>(argument_offsets), host_velopix_events, host_velopix_events_size, cudaMemcpyHostToDevice, stream));
cudaCheck(cudaMemcpyAsync(argen.generate<arg::dev_raw_input_offsets>(argument_offsets), host_velopix_event_offsets, host_velopix_event_offsets_size * sizeof(uint), cudaMemcpyHostToDevice, stream));
@@ -108,9 +108,9 @@ cudaError_t Stream::run_sequence(
argen.generate<arg::dev_cluster_candidates>(argument_offsets),
argen.generate<arg::dev_velo_cluster_container>(argument_offsets),
dev_velo_geometry,
- gpu_constants.dev_velo_sp_patterns,
- gpu_constants.dev_velo_sp_fx,
- gpu_constants.dev_velo_sp_fy
+ constants.dev_velo_sp_patterns,
+ constants.dev_velo_sp_fx,
+ constants.dev_velo_sp_fy
);
sequence.item<seq::masked_velo_clustering>().invoke();
@@ -307,7 +307,8 @@ cudaError_t Stream::run_sequence(
argen.generate<arg::dev_velo_states>(argument_offsets),
argen.generate<arg::dev_veloUT_tracks>(argument_offsets),
argen.generate<arg::dev_atomics_veloUT>(argument_offsets),
- dev_ut_magnet_tool );
+ dev_ut_magnet_tool,
+ constants.dev_ut_dxDy );
sequence.item<seq::veloUT>().invoke();
// Transmission device to host
@@ -382,6 +383,7 @@ cudaError_t Stream::run_sequence(
ut_tracks_events,
host_ut_hits_events,
host_ut_magnet_tool,
+ constants.host_ut_dxDy,
host_velo_states,
host_accumulated_tracks,
host_velo_track_hit_number,
diff --git a/stream/sequence/src/StreamWrapper.cu b/stream/sequence/src/StreamWrapper.cu
index 36d14bcc41da5bb137b38dd695c55db500aed007..006cc0a05ec5b01198a33dba8f1d7fb93d7ec3c6 100644
--- a/stream/sequence/src/StreamWrapper.cu
+++ b/stream/sequence/src/StreamWrapper.cu
@@ -14,7 +14,7 @@ void StreamWrapper::initialize_streams(
const std::string& folder_name_MC,
const uint start_event_offset,
const size_t reserve_mb,
- const GpuConstants& gpu_constants
+ const Constants& constants
) {
for (uint i=0; i<n; ++i) {
streams.push_back(new Stream());
@@ -34,7 +34,7 @@ void StreamWrapper::initialize_streams(
start_event_offset,
reserve_mb,
i,
- gpu_constants
+ constants
);
// Memory consumption
diff --git a/x86/veloUT/PrVeloUT/CMakeLists.txt b/x86/veloUT/PrVeloUT/CMakeLists.txt
index 362c7bc48c344b4a903d1364959d4be2023c776f..861e99bea4259ba44f5ad608ead86dbd0b4c3534 100644
--- a/x86/veloUT/PrVeloUT/CMakeLists.txt
+++ b/x86/veloUT/PrVeloUT/CMakeLists.txt
@@ -11,3 +11,10 @@ file(GLOB x86VeloUT_cpp "src/*cpp")
cuda_add_library(x86VeloUT STATIC
${x86VeloUT_cpp}
)
+if ( ROOT_FOUND )
+ target_compile_definitions(x86VeloUT PUBLIC WITH_ROOT)
+ message("Found Root, compiling WITH_ROOT")
+ target_link_libraries(x86VeloUT
+ ${ROOT_LIBRARIES}
+ )
+endif()
diff --git a/x86/veloUT/PrVeloUT/include/PrVeloUTWrapper.h b/x86/veloUT/PrVeloUT/include/PrVeloUTWrapper.h
index 7dbb1279458c7c0b5d3218e74be47f6a649a8bac..05915aca645cb14d2c833d31236d1e1275d2d09f 100644
--- a/x86/veloUT/PrVeloUT/include/PrVeloUTWrapper.h
+++ b/x86/veloUT/PrVeloUT/include/PrVeloUTWrapper.h
@@ -9,6 +9,7 @@ void call_PrVeloUT(
const VeloState* velo_states_event,
VeloUTTracking::HitsSoA *hits_layers_events,
const PrUTMagnetTool *magnet_tool,
+ const float* ut_dxDy,
VeloUTTracking::TrackUT VeloUT_tracks[VeloUTTracking::max_num_tracks],
int &n_velo_tracks_in_UT,
int &n_veloUT_tracks
diff --git a/x86/veloUT/PrVeloUT/include/run_VeloUT_CPU.h b/x86/veloUT/PrVeloUT/include/run_VeloUT_CPU.h
index 95d6293c2cdf8854b373bb892e4cf2d1584e0432..b001dafa12e973cf43e970dd73b89c098bdb6c00 100644
--- a/x86/veloUT/PrVeloUT/include/run_VeloUT_CPU.h
+++ b/x86/veloUT/PrVeloUT/include/run_VeloUT_CPU.h
@@ -5,11 +5,13 @@
#include "PrVeloUTWrapper.h"
#include "Tools.h"
#include "Sorting.cuh"
+#include "VeloUTDefinitions.cuh"
int run_veloUT_on_CPU (
std::vector<trackChecker::Tracks>& ut_tracks_events,
VeloUTTracking::HitsSoA * hits_layers_events,
const PrUTMagnetTool* host_ut_magnet_tool,
+ const float host_ut_dxDy[VeloUTTracking::n_layers],
const VeloState * host_velo_states,
const int * host_accumulated_tracks,
const uint* host_velo_track_hit_number_pinned,
diff --git a/x86/veloUT/PrVeloUT/src/PrVeloUTWrapper.cpp b/x86/veloUT/PrVeloUT/src/PrVeloUTWrapper.cpp
index fe053bff059a91e379e2b6ce0de00d1a219b2583..85b47634cf6771ac58521f4411fa8f935dd891ad 100644
--- a/x86/veloUT/PrVeloUT/src/PrVeloUTWrapper.cpp
+++ b/x86/veloUT/PrVeloUT/src/PrVeloUTWrapper.cpp
@@ -21,6 +21,7 @@ void call_PrVeloUT (
const VeloState* velo_states_event,
VeloUTTracking::HitsSoA *hits_layers,
const PrUTMagnetTool *magnet_tool,
+ const float* ut_dxDy,
VeloUTTracking::TrackUT VeloUT_tracks[VeloUTTracking::max_num_tracks],
int &n_velo_tracks_in_UT,
int &n_veloUT_tracks )
@@ -55,7 +56,8 @@ void call_PrVeloUT (
posLayers,
hits_layers,
fudgeFactors,
- velo_states_event[i_track] ) ) continue;
+ velo_states_event[i_track],
+ ut_dxDy ) ) continue;
TrackHelper helper(velo_states_event[i_track]);
@@ -70,6 +72,7 @@ void call_PrVeloUT (
hitCandidateIndices,
hits_layers,
helper,
+ ut_dxDy,
true) ){
// go through UT layers in backward direction
@@ -80,6 +83,7 @@ void call_PrVeloUT (
hitCandidateIndices,
hits_layers,
helper,
+ ut_dxDy,
false);
}
diff --git a/x86/veloUT/PrVeloUT/src/run_VeloUT_CPU.cpp b/x86/veloUT/PrVeloUT/src/run_VeloUT_CPU.cpp
index 3ef52889f722be97d73bc6bd8a7e6c8215bc8c42..b0ca52d917506ef8f7481f10e27f64823884dace 100644
--- a/x86/veloUT/PrVeloUT/src/run_VeloUT_CPU.cpp
+++ b/x86/veloUT/PrVeloUT/src/run_VeloUT_CPU.cpp
@@ -10,6 +10,7 @@ int run_veloUT_on_CPU (
std::vector<trackChecker::Tracks>& ut_tracks_events,
VeloUTTracking::HitsSoA* hits_layers_events,
const PrUTMagnetTool* host_ut_magnet_tool,
+ const float * host_ut_dxDy,
const VeloState* host_velo_states,
const int* host_accumulated_tracks,
const uint* host_velo_track_hit_number_pinned,
@@ -26,7 +27,7 @@ int run_veloUT_on_CPU (
TTree *t_velo_states = new TTree("velo_states", "velo_states");
TTree *t_track_hits = new TTree("track_hits", "track_hits");
TTree *t_veloUT_tracks = new TTree("veloUT_tracks", "veloUT_tracks");
- float cos, yBegin, yEnd, dxDy, zAtYEq0, xAtYEq0, weight;
+ float yBegin, yEnd, dxDy, zAtYEq0, xAtYEq0, weight;
float x, y, tx, ty, chi2, z, drdz;
unsigned int LHCbID;
int highThreshold, layer;
@@ -36,7 +37,6 @@ int run_veloUT_on_CPU (
float qop;
- t_ut_hits->Branch("cos", &cos);
t_ut_hits->Branch("yBegin", &yBegin);
t_ut_hits->Branch("yEnd", &yEnd);
t_ut_hits->Branch("dxDy", &dxDy);
@@ -88,35 +88,33 @@ int run_veloUT_on_CPU (
// sort according to xAtyEq0
find_permutation<float>(
- hits_layers.m_xAtYEq0,
+ hits_layers.xAtYEq0,
layer_offset,
hit_permutations,
n_hits
);
- apply_permutation<float>( hit_permutations, layer_offset, n_hits, hits_layers.m_cos, hits_layers_sorted.m_cos );
- apply_permutation<float>( hit_permutations, layer_offset, n_hits, hits_layers.m_weight, hits_layers_sorted.m_weight );
- apply_permutation<float>( hit_permutations, layer_offset, n_hits, hits_layers.m_xAtYEq0, hits_layers_sorted.m_xAtYEq0 );
- apply_permutation<float>( hit_permutations, layer_offset, n_hits, hits_layers.m_yBegin, hits_layers_sorted.m_yBegin );
- apply_permutation<float>( hit_permutations, layer_offset, n_hits, hits_layers.m_yEnd, hits_layers_sorted.m_yEnd );
- apply_permutation<float>( hit_permutations, layer_offset, n_hits, hits_layers.m_zAtYEq0, hits_layers_sorted.m_zAtYEq0 );
- apply_permutation<unsigned int>( hit_permutations, layer_offset, n_hits, hits_layers.m_LHCbID, hits_layers_sorted.m_LHCbID );
- apply_permutation<int>( hit_permutations, layer_offset, n_hits, hits_layers.m_planeCode, hits_layers_sorted.m_planeCode );
- apply_permutation<int>( hit_permutations, layer_offset, n_hits, hits_layers.m_highThreshold, hits_layers_sorted.m_highThreshold );
+ apply_permutation<float>( hit_permutations, layer_offset, n_hits, hits_layers.weight, hits_layers_sorted.weight );
+ apply_permutation<float>( hit_permutations, layer_offset, n_hits, hits_layers.xAtYEq0, hits_layers_sorted.xAtYEq0 );
+ apply_permutation<float>( hit_permutations, layer_offset, n_hits, hits_layers.yBegin, hits_layers_sorted.yBegin );
+ apply_permutation<float>( hit_permutations, layer_offset, n_hits, hits_layers.yEnd, hits_layers_sorted.yEnd );
+ apply_permutation<float>( hit_permutations, layer_offset, n_hits, hits_layers.zAtYEq0, hits_layers_sorted.zAtYEq0 );
+ apply_permutation<unsigned int>( hit_permutations, layer_offset, n_hits, hits_layers.LHCbID, hits_layers_sorted.LHCbID );
+ apply_permutation<int>( hit_permutations, layer_offset, n_hits, hits_layers.planeCode, hits_layers_sorted.planeCode );
+ apply_permutation<int>( hit_permutations, layer_offset, n_hits, hits_layers.highThreshold, hits_layers_sorted.highThreshold );
#ifdef WITH_ROOT
for ( int i_hit = 0; i_hit < n_hits; ++i_hit ) {
- cos = hits_layers.m_cos[layer_offset + i_hit];
- weight = hits_layers.m_weight[layer_offset + i_hit];
- xAtYEq0 = hits_layers.m_xAtYEq0[layer_offset + i_hit];
- yBegin = hits_layers.m_yBegin[layer_offset + i_hit];
- yEnd = hits_layers.m_yEnd[layer_offset + i_hit];
- zAtYEq0 = hits_layers.m_zAtYEq0[layer_offset + i_hit];
- LHCbID = hits_layers.m_LHCbID[layer_offset + i_hit];
+ weight = hits_layers.weight[layer_offset + i_hit];
+ xAtYEq0 = hits_layers.xAtYEq0[layer_offset + i_hit];
+ yBegin = hits_layers.yBegin[layer_offset + i_hit];
+ yEnd = hits_layers.yEnd[layer_offset + i_hit];
+ zAtYEq0 = hits_layers.zAtYEq0[layer_offset + i_hit];
+ LHCbID = hits_layers.LHCbID[layer_offset + i_hit];
layer = i_layer;
- highThreshold = hits_layers.m_highThreshold[layer_offset + i_hit];
- dxDy = hits_layers.dxDy(layer_offset + i_hit);
+ highThreshold = hits_layers.highThreshold[layer_offset + i_hit];
+ dxDy = host_ut_dxDy[layer];
t_ut_hits->Fill();
}
@@ -175,12 +173,13 @@ int run_veloUT_on_CPU (
VeloUTTracking::TrackUT veloUT_tracks[VeloUTTracking::max_num_tracks];
call_PrVeloUT(
host_velo_track_hit_number_pinned,
- host_velo_track_hits_pinned,
+ host_velo_track_hits_pinned,
host_number_of_tracks_pinned[i_event],
host_accumulated_tracks[i_event],
host_velo_states_event,
&(hits_layers_sorted),
host_ut_magnet_tool,
+ host_ut_dxDy,
veloUT_tracks,
n_velo_tracks_in_UT,
n_veloUT_tracks_event