]> git.uio.no Git - u/mrichter/AliRoot.git/blame - HLT/TPCLib/tracking-ca/AliHLTTPCCATrackletConstructor.cxx
git://git.uio.no / u / mrichter / AliRoot.git / blame
? search:
summary | shortlog | log | commit | commitdiff | tree
blob | blame (incremental) | history | HEAD
cosmetic changes
[u/mrichter/AliRoot.git] / HLT / TPCLib / tracking-ca / AliHLTTPCCATrackletConstructor.cxx
CommitLineData
00d07bcd 1// @(#) $Id: AliHLTTPCCATrackletConstructor.cxx 27042 2008-07-02 12:06:02Z richterm $
ce565086 2// **************************************************************************
fbb9b71b 3// This file is property of and copyright by the ALICE HLT Project *
00d07bcd 4// ALICE Experiment at CERN, All rights reserved. *
5// *
6// Primary Authors: Sergey Gorbunov <sergey.gorbunov@kip.uni-heidelberg.de> *
7// Ivan Kisel <kisel@kip.uni-heidelberg.de> *
8// for The ALICE HLT Project. *
9// *
10// Permission to use, copy, modify and distribute this software and its *
11// documentation strictly for non-commercial purposes is hereby granted *
12// without fee, provided that the above copyright notice appears in all *
13// copies and that both the copyright notice and this permission notice *
14// appear in the supporting documentation. The authors make no claims *
15// about the suitability of this software for any purpose. It is *
16// provided "as is" without express or implied warranty. *
ce565086 17// *
00d07bcd 18//***************************************************************************
19
20#include "AliHLTTPCCATracker.h"
21#include "AliHLTTPCCATrackParam.h"
4687b8fc 22#include "AliHLTTPCCATrackParam.h"
00d07bcd 23#include "AliHLTTPCCAGrid.h"
00d07bcd 24#include "AliHLTTPCCAMath.h"
25#include "AliHLTTPCCADef.h"
ce565086 26#include "AliHLTTPCCATracklet.h"
00d07bcd 27#include "AliHLTTPCCATrackletConstructor.h"
b22af1bf 28#include "MemoryAssignmentHelpers.h"
29
693d2443 30//#include "AliHLTTPCCAPerformance.h"
31//#include "TH1D.h"
32
33//#define DRAW
34
35#ifdef DRAW
36#include "AliHLTTPCCADisplay.h"
37#endif
00d07bcd 38
b22af1bf 39#define kMaxRowGap 4
00d07bcd 40
b22af1bf 41GPUd() void AliHLTTPCCATrackletConstructor::InitTracklet( AliHLTTPCCATrackParam &tParam )
00d07bcd 42{
b22af1bf 43 //Initialize Tracklet Parameters using default values
fbb9b71b 44 tParam.SetSinPhi( 0 );
45 tParam.SetDzDs( 0 );
46 tParam.SetQPt( 0 );
47 tParam.SetSignCosPhi( 1 );
48 tParam.SetChi2( 0 );
49 tParam.SetNDF( -3 );
50 tParam.SetCov( 0, 1 );
51 tParam.SetCov( 1, 0 );
52 tParam.SetCov( 2, 1 );
53 tParam.SetCov( 3, 0 );
54 tParam.SetCov( 4, 0 );
55 tParam.SetCov( 5, 1 );
56 tParam.SetCov( 6, 0 );
57 tParam.SetCov( 7, 0 );
58 tParam.SetCov( 8, 0 );
59 tParam.SetCov( 9, 1 );
60 tParam.SetCov( 10, 0 );
61 tParam.SetCov( 11, 0 );
62 tParam.SetCov( 12, 0 );
63 tParam.SetCov( 13, 0 );
64 tParam.SetCov( 14, 10. );
00d07bcd 65}
66
fbb9b71b 67GPUd() void AliHLTTPCCATrackletConstructor::ReadData
b22af1bf 68#ifndef HLTCA_GPU_PREFETCHDATA
69( int /*iThread*/, AliHLTTPCCASharedMemory& /*s*/, AliHLTTPCCAThreadMemory& /*r*/, AliHLTTPCCATracker& /*tracker*/, int /*iRow*/ )
70{
71 //Prefetch Data to shared memory
72#else
73( int iThread, AliHLTTPCCASharedMemory& s, AliHLTTPCCAThreadMemory& r, AliHLTTPCCATracker& tracker, int iRow )
fbb9b71b 74{
00d07bcd 75 // reconstruction of tracklets, read data step
fbb9b71b 76 const AliHLTTPCCARow &row = tracker.Row( iRow );
b22af1bf 77 //bool jr = !r.fCurrentData;
4acc2401 78
79 // copy hits, grid content and links
80
81 // FIXME: inefficient copy
b22af1bf 82 //const int numberOfHitsAligned = NextMultipleOf<sizeof(HLTCA_GPU_ROWALIGNMENT) / sizeof(ushort_v)>(row.NHits());
83
84/*
85#ifdef HLTCA_GPU_REORDERHITDATA
4acc2401 86 ushort2 *sMem1 = reinterpret_cast<ushort2 *>( s.fData[jr] );
b22af1bf 87 for ( int i = iThread; i < numberOfHitsAligned; i += TRACKLET_CONSTRUCTOR_NMEMTHREDS ) {
4acc2401 88 sMem1[i].x = tracker.HitDataY( row, i );
89 sMem1[i].y = tracker.HitDataZ( row, i );
90 }
b22af1bf 91#else
92 ushort_v *sMem1 = reinterpret_cast<ushort_v *>( s.fData[jr] );
93 for ( int i = iThread; i < numberOfHitsAligned; i += TRACKLET_CONSTRUCTOR_NMEMTHREDS ) {
94 sMem1[i] = tracker.HitDataY( row, i );
95 }
96
97 ushort_v *sMem1a = reinterpret_cast<ushort_v *>( s.fData[jr] ) + numberOfHitsAligned;
98 for ( int i = iThread; i < numberOfHitsAligned; i += TRACKLET_CONSTRUCTOR_NMEMTHREDS ) {
99 sMem1a[i] = tracker.HitDataZ( row, i );
4acc2401 100 }
b22af1bf 101#endif
4acc2401 102
b22af1bf 103 short *sMem2 = reinterpret_cast<short *>( s.fData[jr] ) + 2 * numberOfHitsAligned;
104 for ( int i = iThread; i < numberOfHitsAligned; i += TRACKLET_CONSTRUCTOR_NMEMTHREDS ) {
105 sMem2[i] = tracker.HitLinkUpData( row, i );
106 }
107
108 unsigned short *sMem3 = reinterpret_cast<unsigned short *>( s.fData[jr] ) + 3 * numberOfHitsAligned;
4acc2401 109 const int n = row.FullSize(); // + grid content size
7be9b0d7 110 for ( int i = iThread; i < n; i += TRACKLET_CONSTRUCTOR_NMEMTHREDS ) {
4acc2401 111 sMem3[i] = tracker.FirstHitInBin( row, i );
b22af1bf 112 }*/
113
114 /*for (int k = 0;k < 2;k++)
115 {
116 HLTCA_GPU_ROWCOPY* sharedMem;
117 const HLTCA_GPU_ROWCOPY* sourceMem;
118 int copyCount;
119 switch (k)
120 {
121 case 0:
122 sourceMem = reinterpret_cast<const HLTCA_GPU_ROWCOPY *>( tracker.HitDataY(row) );
123 sharedMem = reinterpret_cast<HLTCA_GPU_ROWCOPY *> (reinterpret_cast<ushort_v *>( s.fData[jr] ) + k * numberOfHitsAligned);
124 copyCount = numberOfHitsAligned * sizeof(ushort_v) / sizeof(HLTCA_GPU_ROWCOPY);
125 break;
126 case 1:
127 sourceMem = reinterpret_cast<const HLTCA_GPU_ROWCOPY *>( tracker.HitDataZ(row) );
128 sharedMem = reinterpret_cast<HLTCA_GPU_ROWCOPY *> (reinterpret_cast<ushort_v *>( s.fData[jr] ) + k * numberOfHitsAligned);
129 copyCount = numberOfHitsAligned * sizeof(ushort_v) / sizeof(HLTCA_GPU_ROWCOPY);
130 break;
131 case 2:
132 sourceMem = reinterpret_cast<const HLTCA_GPU_ROWCOPY *>( tracker.HitLinkUpData(row) );
133 sharedMem = reinterpret_cast<HLTCA_GPU_ROWCOPY *> (reinterpret_cast<ushort_v *>( s.fData[jr] ) + k * numberOfHitsAligned);
134 copyCount = numberOfHitsAligned * sizeof(ushort_v) / sizeof(HLTCA_GPU_ROWCOPY);
135 break;
136 case 1:
137 sourceMem = reinterpret_cast<const HLTCA_GPU_ROWCOPY *>( tracker.FirstHitInBin(row) );
138 sharedMem = reinterpret_cast<HLTCA_GPU_ROWCOPY *> (reinterpret_cast<ushort_v *>( s.fData[jr] ) + k * numberOfHitsAligned);
139 copyCount = NextMultipleOf<sizeof(HLTCA_GPU_ROWALIGNMENT) / sizeof(ushort_v)>(row.FullSize()) * sizeof(ushort_v) / sizeof(HLTCA_GPU_ROWCOPY);
140 break;
141 }
142 for (int i = iThread;i < copyCount;i += TRACKLET_CONSTRUCTOR_NMEMTHREDS)
143 {
144 sharedMem[i] = sourceMem[i];
145 }
146 }*/
147
148 for (unsigned int i = iThread;i < row.FullSize() * sizeof(ushort_v) / sizeof(HLTCA_GPU_ROWCOPY);i += TRACKLET_CONSTRUCTOR_NMEMTHREDS)
149 {
150 reinterpret_cast<HLTCA_GPU_ROWCOPY *> (reinterpret_cast<ushort_v *>( s.fData[!r.fCurrentData] ))[i] = reinterpret_cast<const HLTCA_GPU_ROWCOPY *>( tracker.FirstHitInBin(row) )[i];
151 }
152
153 const HLTCA_GPU_ROWCOPY* const sourceMem = (const HLTCA_GPU_ROWCOPY *) &row;
154 HLTCA_GPU_ROWCOPY* const sharedMem = reinterpret_cast<HLTCA_GPU_ROWCOPY *> ( &s.fRow[!r.fCurrentData] );
155 for (unsigned int i = iThread;i < sizeof(AliHLTTPCCARow) / sizeof(HLTCA_GPU_ROWCOPY);i += TRACKLET_CONSTRUCTOR_NMEMTHREDS)
156 {
157 sharedMem[i] = sourceMem[i];
158 }
159#endif
00d07bcd 160}
161
162
fbb9b71b 163GPUd() void AliHLTTPCCATrackletConstructor::StoreTracklet
164( int /*nBlocks*/, int /*nThreads*/, int /*iBlock*/, int /*iThread*/,
a59a784e 165 AliHLTTPCCASharedMemory &/*s*/, AliHLTTPCCAThreadMemory &r, AliHLTTPCCATracker &tracker, AliHLTTPCCATrackParam &tParam )
fbb9b71b 166{
00d07bcd 167 // reconstruction of tracklets, tracklet store step
fbb9b71b 168
693d2443 169 //AliHLTTPCCAPerformance::Instance().HNHitsPerTrackCand()->Fill(r.fNHits);
170
fbb9b71b 171 do {
693d2443 172 {
7be9b0d7 173 //std::cout<<"tracklet to store: "<<r.fItr<<", nhits = "<<r.fNHits<<std::endl;
693d2443 174 }
175
b22af1bf 176 if ( r.fNHits < TRACKLET_SELECTOR_MIN_HITS ) {
00d07bcd 177 r.fNHits = 0;
178 break;
179 }
ce565086 180
fbb9b71b 181 if ( 0 ) {
182 if ( 1. / .5 < CAMath::Abs( tParam.QPt() ) ) { //SG!!!
183 r.fNHits = 0;
184 break;
ce565086 185 }
186 }
fbb9b71b 187
188 {
189 bool ok = 1;
190 const float *c = tParam.Cov();
191 for ( int i = 0; i < 15; i++ ) ok = ok && CAMath::Finite( c[i] );
192 for ( int i = 0; i < 5; i++ ) ok = ok && CAMath::Finite( tParam.Par()[i] );
193 ok = ok && ( tParam.X() > 50 );
194
195 if ( c[0] <= 0 || c[2] <= 0 || c[5] <= 0 || c[9] <= 0 || c[14] <= 0 ) ok = 0;
196
197 if ( !ok ) {
198 r.fNHits = 0;
199 break;
00d07bcd 200 }
fbb9b71b 201 }
202 } while ( 0 );
203
204 if ( !SAVE() ) return;
205
ce565086 206 AliHLTTPCCATracklet &tracklet = tracker.Tracklets()[r.fItr];
207
fbb9b71b 208 tracklet.SetNHits( r.fNHits );
209
210 if ( r.fNHits > 0 ) {
693d2443 211#ifdef DRAW
fbb9b71b 212 if ( 0 ) {
213 std::cout << "store tracklet " << r.fItr << ", nhits = " << r.fNHits << std::endl;
214 if ( AliHLTTPCCADisplay::Instance().DrawTracklet( tParam, hitstore, kBlue, 1. ) ) {
215 AliHLTTPCCADisplay::Instance().Ask();
693d2443 216 }
217 }
218#endif
91794c67 219 if ( CAMath::Abs( tParam.Par()[4] ) < 1.e-4 ) tParam.SetPar( 4, 1.e-4 );
b22af1bf 220 if (r.fStartRow < r.fFirstRow) r.fFirstRow = r.fStartRow;
221 tracklet.SetFirstRow( r.fFirstRow );
ce565086 222 tracklet.SetLastRow( r.fLastRow );
223 tracklet.SetParam( tParam );
fbb9b71b 224 int w = ( r.fNHits << 16 ) + r.fItr;
b22af1bf 225 for ( int iRow = r.fFirstRow; iRow <= r.fLastRow; iRow++ ) {
226#ifdef EXTERN_ROW_HITS
227 int ih = tracker.TrackletRowHits()[iRow * s.fNTracklets + r.fItr];
228#else
229 int ih = tracklet.RowHit( iRow );
230#endif
fbb9b71b 231 if ( ih >= 0 ) {
b22af1bf 232#if defined(HLTCA_GPUCODE) & !defined(HLTCA_GPU_PREFETCHDATA) & !defined(HLTCA_GPU_PREFETCH_ROWBLOCK_ONLY)
233 tracker.MaximizeHitWeight( s.fRows[ iRow ], ih, w );
234#else
235 tracker.MaximizeHitWeight( tracker.Row( iRow ), ih, w );
236#endif
00d07bcd 237 }
238 }
fbb9b71b 239 }
b22af1bf 240
00d07bcd 241}
242
243GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
fbb9b71b 244( int /*nBlocks*/, int /*nThreads*/, int /*iBlock*/, int /*iThread*/,
a59a784e 245 AliHLTTPCCASharedMemory &/*s*/, AliHLTTPCCAThreadMemory &r, AliHLTTPCCATracker &tracker, AliHLTTPCCATrackParam &tParam, int iRow )
fbb9b71b 246{
00d07bcd 247 // reconstruction of tracklets, tracklets update step
248
693d2443 249 //std::cout<<"Update tracklet: "<<r.fItr<<" "<<r.fGo<<" "<<r.fStage<<" "<<iRow<<std::endl;
fbb9b71b 250 bool drawSearch = 0;//r.fItr==2;
251 bool drawFit = 0;//r.fItr==2;
252 bool drawFitted = drawFit ;//|| 1;//r.fItr==16;
693d2443 253
fbb9b71b 254 if ( !r.fGo ) return;
255
b22af1bf 256#ifndef EXTERN_ROW_HITS
ce565086 257 AliHLTTPCCATracklet &tracklet = tracker.Tracklets()[r.fItr];
b22af1bf 258#endif
ce565086 259
b22af1bf 260#ifdef HLTCA_GPU_PREFETCHDATA
261 const AliHLTTPCCARow &row = s.fRow[r.fCurrentData];
262#elif defined(HLTCA_GPUCODE)
263 const AliHLTTPCCARow &row = s.fRows[iRow];
264#else
fbb9b71b 265 const AliHLTTPCCARow &row = tracker.Row( iRow );
b22af1bf 266#endif
fbb9b71b 267
268 float y0 = row.Grid().YMin();
269 float stepY = row.HstepY();
270 float z0 = row.Grid().ZMin();
271 float stepZ = row.HstepZ();
272 float stepYi = row.HstepYi();
273 float stepZi = row.HstepZi();
274
275 if ( r.fStage == 0 ) { // fitting part
276 do {
277
278 if ( iRow < r.fStartRow || r.fCurrIH < 0 ) break;
279
b22af1bf 280 if ( ( iRow - r.fStartRow ) % 2 != 0 )
281 {
282#ifndef EXTERN_ROW_HITS
283 tracklet.SetRowHit(iRow, -1);
284#else
285 tracker.TrackletRowHits()[iRow * s.fNTracklets + r.fItr] = -1;
286#endif
287 break; // SG!!! - jump over the row
288 }
289
290//#ifdef HLTCA_GPU_PREFETCHDATA
291// uint4 *tmpint4 = s.fData[r.fCurrentData];
292//#endif
293 ushort2 hh;
294//#ifdef HLTCA_GPU_REORDERHITDATA
295// hh = reinterpret_cast<ushort2*>( tmpint4 )[r.fCurrIH];
296//#else
297//#ifdef HLTCA_GPU_PREFETCHDATA
298// hh.x = reinterpret_cast<ushort_v*>( tmpint4 )[r.fCurrIH];
299// hh.y = reinterpret_cast<ushort_v*>( tmpint4 )[NextMultipleOf<sizeof(HLTCA_GPU_ROWALIGNMENT) / sizeof(ushort_v)>(row.NHits()) + r.fCurrIH];
300//#else
301#if defined(HLTCA_GPU_TEXTURE_FETCH)
302 hh = tex1Dfetch(gAliTexRefu2, ((char*) tracker.Data().HitData() - tracker.pData()->GPUTextureBase()) / sizeof(ushort2) + row.HitNumberOffset() + r.fCurrIH);
303#else
304 hh = tracker.HitData(row)[r.fCurrIH];
305#endif
306//#endif
307//#endif
fbb9b71b 308
309 int oldIH = r.fCurrIH;
b22af1bf 310//#ifdef HLTCA_GPU_PREFETCHDATA
311// r.fCurrIH = reinterpret_cast<short*>( tmpint4 )[2 * NextMultipleOf<sizeof(HLTCA_GPU_ROWALIGNMENT) / sizeof(ushort_v)>(row.NHits()) + r.fCurrIH]; // read from linkup data
312//#else
313#if defined(HLTCA_GPU_TEXTURE_FETCH)
314 r.fCurrIH = tex1Dfetch(gAliTexRefs, ((char*) tracker.Data().HitLinkUpData(row) - tracker.pData()->GPUTextureBase()) / sizeof(unsigned short) + r.fCurrIH);
315#else
316 r.fCurrIH = tracker.HitLinkUpData(row)[r.fCurrIH]; // read from linkup data
317#endif
318//#endif
fbb9b71b 319
320 float x = row.X();
321 float y = y0 + hh.x * stepY;
322 float z = z0 + hh.y * stepZ;
7be9b0d7 323#ifdef DRAW
fbb9b71b 324 if ( drawFit ) std::cout << " fit tracklet: new hit " << oldIH << ", xyz=" << x << " " << y << " " << z << std::endl;
7be9b0d7 325#endif
fbb9b71b 326
327 if ( iRow == r.fStartRow ) {
328 tParam.SetX( x );
329 tParam.SetY( y );
330 tParam.SetZ( z );
331 r.fLastY = y;
332 r.fLastZ = z;
7be9b0d7 333 #ifdef DRAW
fbb9b71b 334 if ( drawFit ) std::cout << " fit tracklet " << r.fItr << ", row " << iRow << " first row" << std::endl;
7be9b0d7 335 #endif
fbb9b71b 336 } else {
337
338 float err2Y, err2Z;
339 float dx = x - tParam.X();
340 float dy = y - r.fLastY;//tParam.Y();
341 float dz = z - r.fLastZ;//tParam.Z();
342 r.fLastY = y;
343 r.fLastZ = z;
344
345 float ri = 1. / CAMath::Sqrt( dx * dx + dy * dy );
5abba7a4 346 if ( iRow == r.fStartRow + 2 ) { //SG!!! important - thanks to Matthias
fbb9b71b 347 tParam.SetSinPhi( dy*ri );
348 tParam.SetSignCosPhi( dx );
349 tParam.SetDzDs( dz*ri );
4acc2401 350 //std::cout << "Init. errors... " << r.fItr << std::endl;
fbb9b71b 351 tracker.GetErrors2( iRow, tParam, err2Y, err2Z );
4acc2401 352 //std::cout << "Init. errors = " << err2Y << " " << err2Z << std::endl;
fbb9b71b 353 tParam.SetCov( 0, err2Y );
354 tParam.SetCov( 2, err2Z );
355 }
356 if ( drawFit ) {
7be9b0d7 357 #ifdef DRAW
fbb9b71b 358 std::cout << " fit tracklet " << r.fItr << ", row " << iRow << " transporting.." << std::endl;
359 std::cout << " params before transport=" << std::endl;
360 tParam.Print();
7be9b0d7 361 #endif
fbb9b71b 362 }
363 float sinPhi, cosPhi;
364 if ( r.fNHits >= 10 && CAMath::Abs( tParam.SinPhi() ) < .99 ) {
365 sinPhi = tParam.SinPhi();
366 cosPhi = CAMath::Sqrt( 1 - sinPhi * sinPhi );
367 } else {
368 sinPhi = dy * ri;
369 cosPhi = dx * ri;
370 }
7be9b0d7 371 #ifdef DRAW
fbb9b71b 372 if ( drawFit ) std::cout << "sinPhi0 = " << sinPhi << ", cosPhi0 = " << cosPhi << std::endl;
7be9b0d7 373 #endif
91794c67 374 if ( !tParam.TransportToX( x, sinPhi, cosPhi, tracker.Param().ConstBz(), -1 ) ) {
7be9b0d7 375 #ifdef DRAW
fbb9b71b 376 if ( drawFit ) std::cout << " tracklet " << r.fItr << ", row " << iRow << ": can not transport!!" << std::endl;
7be9b0d7 377 #endif
b22af1bf 378#ifndef EXTERN_ROW_HITS
379 tracklet.SetRowHit( iRow, -1 );
380#else
381 tracker.TrackletRowHits()[iRow * s.fNTracklets + r.fItr] = -1;
382#endif
fbb9b71b 383 break;
384 }
385 //std::cout<<"mark1 "<<r.fItr<<std::endl;
386 //tParam.Print();
387 tracker.GetErrors2( iRow, tParam.GetZ(), sinPhi, cosPhi, tParam.GetDzDs(), err2Y, err2Z );
388 //std::cout<<"mark2"<<std::endl;
389
390 if ( drawFit ) {
7be9b0d7 391 #ifdef DRAW
fbb9b71b 392 std::cout << " params after transport=" << std::endl;
393 tParam.Print();
394 std::cout << "fit tracklet before filter: " << r.fItr << ", row " << iRow << " errs=" << err2Y << " " << err2Z << std::endl;
fbb9b71b 395 AliHLTTPCCADisplay::Instance().DrawTracklet( tParam, hitstore, kBlue, 2., 1 );
396 AliHLTTPCCADisplay::Instance().Ask();
7be9b0d7 397 #endif
fbb9b71b 398 }
399 if ( !tParam.Filter( y, z, err2Y, err2Z, .99 ) ) {
7be9b0d7 400 #ifdef DRAW
fbb9b71b 401 if ( drawFit ) std::cout << " tracklet " << r.fItr << ", row " << iRow << ": can not filter!!" << std::endl;
7be9b0d7 402 #endif
b22af1bf 403#ifndef EXTERN_ROW_HITS
404 tracklet.SetRowHit( iRow, -1 );
405#else
406 tracker.TrackletRowHits()[iRow * s.fNTracklets + r.fItr] = -1;
407#endif
fbb9b71b 408 break;
409 }
00d07bcd 410 }
b22af1bf 411#ifndef EXTERN_ROW_HITS
412 tracklet.SetRowHit( iRow, oldIH );
413#else
414 tracker.TrackletRowHits()[iRow * s.fNTracklets + r.fItr] = oldIH;
415#endif
fbb9b71b 416 if ( drawFit ) {
7be9b0d7 417 #ifdef DRAW
fbb9b71b 418 std::cout << "fit tracklet after filter " << r.fItr << ", row " << iRow << std::endl;
419 tParam.Print();
fbb9b71b 420 AliHLTTPCCADisplay::Instance().DrawTracklet( tParam, hitstore, kGreen, 2. );
421 AliHLTTPCCADisplay::Instance().Ask();
7be9b0d7 422 #endif
693d2443 423 }
00d07bcd 424 r.fNHits++;
425 r.fLastRow = iRow;
ce565086 426 r.fEndRow = iRow;
00d07bcd 427 break;
fbb9b71b 428 } while ( 0 );
429
430 if ( r.fCurrIH < 0 ) {
7be9b0d7 431 #ifdef DRAW
fbb9b71b 432 if ( drawFitted ) std::cout << "fitted tracklet " << r.fItr << ", nhits=" << r.fNHits << std::endl;
7be9b0d7 433 #endif
693d2443 434 r.fStage = 1;
435 //AliHLTTPCCAPerformance::Instance().HNHitsPerSeed()->Fill(r.fNHits);
fbb9b71b 436 if ( r.fNHits < 3 ) { r.fNHits = 0; r.fGo = 0;}//SG!!!
437 if ( CAMath::Abs( tParam.SinPhi() ) > .999 ) {
7be9b0d7 438 #ifdef DRAW
fbb9b71b 439 if ( drawFitted ) std::cout << " fitted tracklet error: sinPhi=" << tParam.SinPhi() << std::endl;
7be9b0d7 440 #endif
fbb9b71b 441 r.fNHits = 0; r.fGo = 0;
442 } else {
443 //tParam.SetCosPhi( CAMath::Sqrt(1-tParam.SinPhi()*tParam.SinPhi()) );
693d2443 444 }
fbb9b71b 445 if ( drawFitted ) {
7be9b0d7 446 #ifdef DRAW
fbb9b71b 447 std::cout << "fitted tracklet " << r.fItr << " miss=" << r.fNMissed << " go=" << r.fGo << std::endl;
448 tParam.Print();
fbb9b71b 449 AliHLTTPCCADisplay::Instance().DrawTracklet( tParam, hitstore, kBlue );
450 AliHLTTPCCADisplay::Instance().Ask();
7be9b0d7 451 #endif
693d2443 452 }
fbb9b71b 453 }
454 } else { // forward/backward searching part
455 do {
456 if ( drawSearch ) {
7be9b0d7 457 #ifdef DRAW
fbb9b71b 458 std::cout << "search tracklet " << r.fItr << " row " << iRow << " miss=" << r.fNMissed << " go=" << r.fGo << " stage=" << r.fStage << std::endl;
7be9b0d7 459 #endif
fbb9b71b 460 }
461
462 if ( r.fStage == 2 && ( ( iRow >= r.fEndRow ) ||
463 ( iRow >= r.fStartRow && ( iRow - r.fStartRow ) % 2 == 0 )
464 ) ) break;
465 if ( r.fNMissed > kMaxRowGap ) {
466 break;
467 }
468
469 r.fNMissed++;
470
471 float x = row.X();
472 float err2Y, err2Z;
473 if ( drawSearch ) {
7be9b0d7 474 #ifdef DRAW
fbb9b71b 475 std::cout << "tracklet " << r.fItr << " before transport to row " << iRow << " : " << std::endl;
476 tParam.Print();
7be9b0d7 477 #endif
fbb9b71b 478 }
91794c67 479 if ( !tParam.TransportToX( x, tParam.SinPhi(), tParam.GetCosPhi(), tracker.Param().ConstBz(), .99 ) ) {
7be9b0d7 480 #ifdef DRAW
fbb9b71b 481 if ( drawSearch ) std::cout << " tracklet " << r.fItr << ", row " << iRow << ": can not transport!!" << std::endl;
7be9b0d7 482 #endif
b22af1bf 483#ifndef EXTERN_ROW_HITS
484 tracklet.SetRowHit(iRow, -1);
485#else
486 tracker.TrackletRowHits()[iRow * s.fNTracklets + r.fItr] = -1;
487#endif
fbb9b71b 488 break;
489 }
4acc2401 490 if ( row.NHits() < 1 ) {
491 // skip empty row
b22af1bf 492#ifndef EXTERN_ROW_HITS
493 tracklet.SetRowHit(iRow, -1);
494#else
495 tracker.TrackletRowHits()[iRow * s.fNTracklets + r.fItr] = -1;
496#endif
4acc2401 497 break;
498 }
fbb9b71b 499 if ( drawSearch ) {
7be9b0d7 500 #ifdef DRAW
fbb9b71b 501 std::cout << "tracklet " << r.fItr << " after transport to row " << iRow << " : " << std::endl;
502 tParam.Print();
fbb9b71b 503 AliHLTTPCCADisplay::Instance().DrawTracklet( tParam, hitstore, kBlue, 2., 1 );
504 AliHLTTPCCADisplay::Instance().Ask();
7be9b0d7 505 #endif
fbb9b71b 506 }
b22af1bf 507#ifdef HLTCA_GPU_PREFETCHDATA
fbb9b71b 508 uint4 *tmpint4 = s.fData[r.fCurrentData];
b22af1bf 509#endif
fbb9b71b 510
b22af1bf 511//#ifdef HLTCA_GPU_REORDERHITDATA
512// const ushort2 *hits = reinterpret_cast<ushort2*>( tmpint4 );
513//#else
514//#ifdef HLTCA_GPU_PREFETCHDATA
515// const ushort_v *hitsx = reinterpret_cast<ushort_v*>( tmpint4 );
516// const ushort_v *hitsy = reinterpret_cast<ushort_v*>( tmpint4 ) + NextMultipleOf<sizeof(HLTCA_GPU_ROWALIGNMENT) / sizeof(ushort_v)>(row.NHits());
517//#else
518#ifndef HLTCA_GPU_TEXTURE_FETCH
519 const ushort2 *hits = tracker.HitData(row);
520#endif
521//#endif
522//#endif
fbb9b71b 523
524 float fY = tParam.GetY();
525 float fZ = tParam.GetZ();
526 int best = -1;
527
528 { // search for the closest hit
4acc2401 529 const int fIndYmin = row.Grid().GetBinBounded( fY - 1.f, fZ - 1.f );
530 assert( fIndYmin >= 0 );
fbb9b71b 531
532 int ds;
533 int fY0 = ( int ) ( ( fY - y0 ) * stepYi );
534 int fZ0 = ( int ) ( ( fZ - z0 ) * stepZi );
535 int ds0 = ( ( ( int )1 ) << 30 );
536 ds = ds0;
537
fbb9b71b 538 unsigned int fHitYfst = 1, fHitYlst = 0, fHitYfst1 = 1, fHitYlst1 = 0;
539
fbb9b71b 540 if ( drawSearch ) {
693d2443 541#ifdef DRAW
fbb9b71b 542 std::cout << " tracklet " << r.fItr << ", row " << iRow << ": grid N=" << row.Grid().N() << std::endl;
543 std::cout << " tracklet " << r.fItr << ", row " << iRow << ": minbin=" << fIndYmin << std::endl;
693d2443 544#endif
fbb9b71b 545 }
546 {
547 int nY = row.Grid().Ny();
548
b22af1bf 549//#ifdef HLTCA_GPU_PREFETCHDATA
550// const unsigned short *sGridP = ( reinterpret_cast<unsigned short*>( tmpint4 ) );
551//#else
552#ifndef HLTCA_GPU_TEXTURE_FETCH
553 const unsigned short *sGridP = tracker.FirstHitInBin(row);
554#endif
555//#endif
556
557#ifdef HLTCA_GPU_TEXTURE_FETCH
558 fHitYfst = tex1Dfetch(gAliTexRefu, ((char*) tracker.Data().FirstHitInBin(row) - tracker.pData()->GPUTextureBase()) / sizeof(unsigned short) + fIndYmin);
559 fHitYlst = tex1Dfetch(gAliTexRefu, ((char*) tracker.Data().FirstHitInBin(row) - tracker.pData()->GPUTextureBase()) / sizeof(unsigned short) + fIndYmin+2);
560 fHitYfst1 = tex1Dfetch(gAliTexRefu, ((char*) tracker.Data().FirstHitInBin(row) - tracker.pData()->GPUTextureBase()) / sizeof(unsigned short) + fIndYmin+nY);
561 fHitYlst1 = tex1Dfetch(gAliTexRefu, ((char*) tracker.Data().FirstHitInBin(row) - tracker.pData()->GPUTextureBase()) / sizeof(unsigned short) + fIndYmin+nY+2);
562#else
fbb9b71b 563 fHitYfst = sGridP[fIndYmin];
564 fHitYlst = sGridP[fIndYmin+2];
565 fHitYfst1 = sGridP[fIndYmin+nY];
566 fHitYlst1 = sGridP[fIndYmin+nY+2];
b22af1bf 567#endif
568 assert( (signed) fHitYfst <= row.NHits() );
569 assert( (signed) fHitYlst <= row.NHits() );
570 assert( (signed) fHitYfst1 <= row.NHits() );
571 assert( (signed) fHitYlst1 <= row.NHits() );
fbb9b71b 572 if ( drawSearch ) {
693d2443 573#ifdef DRAW
fbb9b71b 574 std::cout << " Grid, row " << iRow << ": nHits=" << row.NHits() << ", grid n=" << row.Grid().N() << ", c[n]=" << sGridP[row.Grid().N()] << std::endl;
575 std::cout << "hit steps = " << stepY << " " << stepZ << std::endl;
576 std::cout << " Grid bins:" << std::endl;
577 for ( unsigned int i = 0; i < row.Grid().N(); i++ ) {
578 std::cout << " bin " << i << ": ";
579 for ( int j = sGridP[i]; j < sGridP[i+1]; j++ ) {
580 ushort2 hh = hits[j];
581 float y = y0 + hh.x * stepY;
582 float z = z0 + hh.y * stepZ;
583 std::cout << "[" << j << "|" << y << "," << z << "] ";
584 }
585 std::cout << std::endl;
586 }
693d2443 587#endif
fbb9b71b 588 }
693d2443 589#ifdef DRAW
b22af1bf 590 if ( sGridP[row.Grid().N()] != row.NHits() ) {
fbb9b71b 591 std::cout << " grid, row " << iRow << ": nHits=" << row.NHits() << ", grid n=" << row.Grid().N() << ", c[n]=" << sGridP[row.Grid().N()] << std::endl;
592 //exit(0);
fbb9b71b 593 }
b22af1bf 594#endif
fbb9b71b 595 }
b22af1bf 596#ifdef DRAW
fbb9b71b 597 if ( drawSearch ) {
fbb9b71b 598 std::cout << " tracklet " << r.fItr << ", row " << iRow << ", yz= " << fY << "," << fZ << ": search hits=" << fHitYfst << " " << fHitYlst << " / " << fHitYfst1 << " " << fHitYlst1 << std::endl;
599 std::cout << " hit search :" << std::endl;
fbb9b71b 600 }
b22af1bf 601#endif
fbb9b71b 602 for ( unsigned int fIh = fHitYfst; fIh < fHitYlst; fIh++ ) {
b22af1bf 603 assert( (signed) fIh < row.NHits() );
604 ushort2 hh;
605#if defined(HLTCA_GPU_TEXTURE_FETCH)
606 hh = tex1Dfetch(gAliTexRefu2, ((char*) tracker.Data().HitData() - tracker.pData()->GPUTextureBase()) / sizeof(ushort2) + row.HitNumberOffset() + fIh);
607#else
608 hh = hits[fIh];
609#endif
fbb9b71b 610 int ddy = ( int )( hh.x ) - fY0;
611 int ddz = ( int )( hh.y ) - fZ0;
612 int dds = CAMath::Abs( ddy ) + CAMath::Abs( ddz );
613 if ( drawSearch ) {
7be9b0d7 614 #ifdef DRAW
fbb9b71b 615 std::cout << fIh << ": hityz= " << hh.x << " " << hh.y << "(" << hh.x*stepY << " " << hh.y*stepZ << "), trackyz=" << fY0 << " " << fZ0 << "(" << fY0*stepY << " " << fZ0*stepZ << "), dy,dz,ds= " << ddy << " " << ddz << " " << dds << "(" << ddy*stepY << " " << ddz*stepZ << std::endl;
7be9b0d7 616 #endif
fbb9b71b 617 }
618 if ( dds < ds ) {
619 ds = dds;
620 best = fIh;
621 }
622 }
623
b22af1bf 624 for ( unsigned int fIh = fHitYfst1; fIh < fHitYlst1; fIh++ ) {
625 ushort2 hh;
626#if defined(HLTCA_GPU_TEXTURE_FETCH)
627 hh = tex1Dfetch(gAliTexRefu2, ((char*) tracker.Data().HitData() - tracker.pData()->GPUTextureBase()) / sizeof(ushort2) + row.HitNumberOffset() + fIh);
628#else
629 hh = hits[fIh];
630#endif
fbb9b71b 631 int ddy = ( int )( hh.x ) - fY0;
632 int ddz = ( int )( hh.y ) - fZ0;
633 int dds = CAMath::Abs( ddy ) + CAMath::Abs( ddz );
634 if ( drawSearch ) {
7be9b0d7 635 #ifdef DRAW
fbb9b71b 636 std::cout << fIh << ": hityz= " << hh.x << " " << hh.y << "(" << hh.x*stepY << " " << hh.y*stepZ << "), trackyz=" << fY0 << " " << fZ0 << "(" << fY0*stepY << " " << fZ0*stepZ << "), dy,dz,ds= " << ddy << " " << ddz << " " << dds << "(" << ddy*stepY << " " << ddz*stepZ << std::endl;
7be9b0d7 637 #endif
fbb9b71b 638 }
639 if ( dds < ds ) {
640 ds = dds;
641 best = fIh;
642 }
643 }
644 }// end of search for the closest hit
645
b22af1bf 646 if ( best < 0 )
647 {
648#ifndef EXTERN_ROW_HITS
649 tracklet.SetRowHit(iRow, -1);
650#else
651 tracker.TrackletRowHits()[iRow * s.fNTracklets + r.fItr] = -1;
652#endif
653 break;
654 }
fbb9b71b 655 if ( drawSearch ) {
7be9b0d7 656 #ifdef DRAW
fbb9b71b 657 std::cout << "hit search " << r.fItr << ", row " << iRow << " hit " << best << " found" << std::endl;
fbb9b71b 658 AliHLTTPCCADisplay::Instance().DrawSliceHit( iRow, best, kRed, 1. );
659 AliHLTTPCCADisplay::Instance().Ask();
660 AliHLTTPCCADisplay::Instance().DrawSliceHit( iRow, best, kWhite, 1 );
661 AliHLTTPCCADisplay::Instance().DrawSliceHit( iRow, best );
7be9b0d7 662 #endif
fbb9b71b 663 }
664
b22af1bf 665 ushort2 hh;
666#if defined(HLTCA_GPU_TEXTURE_FETCH)
667 hh = tex1Dfetch(gAliTexRefu2, ((char*) tracker.Data().HitData() - tracker.pData()->GPUTextureBase()) / sizeof(ushort2) + row.HitNumberOffset() + best);
668#else
669 hh = hits[best];
670#endif
fbb9b71b 671
672 //std::cout<<"mark 3, "<<r.fItr<<std::endl;
673 //tParam.Print();
674 tracker.GetErrors2( iRow, *( ( AliHLTTPCCATrackParam* )&tParam ), err2Y, err2Z );
675 //std::cout<<"mark 4"<<std::endl;
676
677 float y = y0 + hh.x * stepY;
678 float z = z0 + hh.y * stepZ;
679 float dy = y - fY;
680 float dz = z - fZ;
681
682 const float kFactor = tracker.Param().HitPickUpFactor() * tracker.Param().HitPickUpFactor() * 3.5 * 3.5;
683 float sy2 = kFactor * ( tParam.GetErr2Y() + err2Y );
684 float sz2 = kFactor * ( tParam.GetErr2Z() + err2Z );
685 if ( sy2 > 2. ) sy2 = 2.;
686 if ( sz2 > 2. ) sz2 = 2.;
687
688 if ( drawSearch ) {
7be9b0d7 689 #ifdef DRAW
fbb9b71b 690 std::cout << "dy,sy= " << dy << " " << CAMath::Sqrt( sy2 ) << ", dz,sz= " << dz << " " << CAMath::Sqrt( sz2 ) << std::endl;
691 std::cout << "dy,dz= " << dy << " " << dz << ", sy,sz= " << CAMath::Sqrt( sy2 ) << " " << CAMath::Sqrt( sz2 ) << ", sy,sz= " << CAMath::Sqrt( kFactor*( tParam.GetErr2Y() + err2Y ) ) << " " << CAMath::Sqrt( kFactor*( tParam.GetErr2Z() + err2Z ) ) << std::endl;
7be9b0d7 692 #endif
fbb9b71b 693 }
694 if ( CAMath::FMulRZ( dy, dy ) > sy2 || CAMath::FMulRZ( dz, dz ) > sz2 ) {
695 if ( drawSearch ) {
7be9b0d7 696 #ifdef DRAW
fbb9b71b 697 std::cout << "found hit is out of the chi2 window\n " << std::endl;
7be9b0d7 698 #endif
fbb9b71b 699 }
b22af1bf 700#ifndef EXTERN_ROW_HITS
701 tracklet.SetRowHit(iRow, -1);
702#else
703 tracker.TrackletRowHits()[iRow * s.fNTracklets + r.fItr] = -1;
704#endif
fbb9b71b 705 break;
706 }
693d2443 707#ifdef DRAW
fbb9b71b 708 //if( SAVE() ) hitstore[ iRow ] = best;
709 //std::cout<<"hit search before filter: "<<r.fItr<<", row "<<iRow<<std::endl;
710 //AliHLTTPCCADisplay::Instance().DrawTracklet(tParam, hitstore, kBlue);
711 //AliHLTTPCCADisplay::Instance().Ask();
693d2443 712#endif
fbb9b71b 713 if ( !tParam.Filter( y, z, err2Y, err2Z, .99 ) ) {
714 if ( drawSearch ) {
7be9b0d7 715 #ifdef DRAW
fbb9b71b 716 std::cout << "tracklet " << r.fItr << " at row " << iRow << " : can not filter!!!! " << std::endl;
7be9b0d7 717 #endif
fbb9b71b 718 }
719 break;
720 }
b22af1bf 721#ifndef EXTERN_ROW_HITS
722 tracklet.SetRowHit( iRow, best );
723#else
724 tracker.TrackletRowHits()[iRow * s.fNTracklets + r.fItr] = best;
725#endif
fbb9b71b 726 if ( drawSearch ) {
7be9b0d7 727 #ifdef DRAW
fbb9b71b 728 std::cout << "tracklet " << r.fItr << " after filter at row " << iRow << " : " << std::endl;
729 tParam.Print();
fbb9b71b 730 AliHLTTPCCADisplay::Instance().DrawTracklet( tParam, hitstore, kRed );
731 AliHLTTPCCADisplay::Instance().Ask();
7be9b0d7 732 #endif
fbb9b71b 733 }
734 r.fNHits++;
735 r.fNMissed = 0;
736 if ( r.fStage == 1 ) r.fLastRow = iRow;
737 else r.fFirstRow = iRow;
738 } while ( 0 );
739 }
00d07bcd 740}
741
b22af1bf 742#ifdef HLTCA_GPUCODE
743GPUd() void AliHLTTPCCATrackletConstructor::CopyTrackletTempData( AliHLTTPCCAThreadMemory &rMemSrc, AliHLTTPCCAThreadMemory &rMemDst, AliHLTTPCCATrackParam &tParamSrc, AliHLTTPCCATrackParam &tParamDst)
744{
745 //Copy Temporary Tracklet data from registers to global mem and vice versa
746 rMemDst.fStartRow = rMemSrc.fStartRow;
747 rMemDst.fEndRow = rMemSrc.fEndRow;
748 rMemDst.fFirstRow = rMemSrc.fFirstRow;
749 rMemDst.fLastRow = rMemSrc.fLastRow;
750 rMemDst.fCurrIH = rMemSrc.fCurrIH;
751 rMemDst.fGo = rMemSrc.fGo;
752 rMemDst.fStage = rMemSrc.fStage;
753 rMemDst.fNHits = rMemSrc.fNHits;
754 rMemDst.fNMissed = rMemSrc.fNMissed;
755 rMemDst.fLastY = rMemSrc.fLastY;
756 rMemDst.fLastZ = rMemSrc.fLastZ;
757
758 tParamDst.SetSinPhi( tParamSrc.GetSinPhi() );
759 tParamDst.SetDzDs( tParamSrc.GetDzDs() );
760 tParamDst.SetQPt( tParamSrc.GetQPt() );
761 tParamDst.SetSignCosPhi( tParamSrc.GetSignCosPhi() );
762 tParamDst.SetChi2( tParamSrc.GetChi2() );
763 tParamDst.SetNDF( tParamSrc.GetNDF() );
764 tParamDst.SetCov( 0, tParamSrc.GetCov(0) );
765 tParamDst.SetCov( 1, tParamSrc.GetCov(1) );
766 tParamDst.SetCov( 2, tParamSrc.GetCov(2) );
767 tParamDst.SetCov( 3, tParamSrc.GetCov(3) );
768 tParamDst.SetCov( 4, tParamSrc.GetCov(4) );
769 tParamDst.SetCov( 5, tParamSrc.GetCov(5) );
770 tParamDst.SetCov( 6, tParamSrc.GetCov(6) );
771 tParamDst.SetCov( 7, tParamSrc.GetCov(7) );
772 tParamDst.SetCov( 8, tParamSrc.GetCov(8) );
773 tParamDst.SetCov( 9, tParamSrc.GetCov(9) );
774 tParamDst.SetCov( 10, tParamSrc.GetCov(10) );
775 tParamDst.SetCov( 11, tParamSrc.GetCov(11) );
776 tParamDst.SetCov( 12, tParamSrc.GetCov(12) );
777 tParamDst.SetCov( 13, tParamSrc.GetCov(13) );
778 tParamDst.SetCov( 14, tParamSrc.GetCov(14) );
779 tParamDst.SetX( tParamSrc.GetX() );
780 tParamDst.SetY( tParamSrc.GetY() );
781 tParamDst.SetZ( tParamSrc.GetZ() );
782}
00d07bcd 783
b22af1bf 784GPUd() int AliHLTTPCCATrackletConstructor::FetchTracklet(AliHLTTPCCATracker &tracker, AliHLTTPCCASharedMemory &sMem, int Reverse, int RowBlock, int &mustInit)
785{
786 //Fetch a new trackled to be processed by this thread
787 __syncthreads();
788 int nextTracketlFirstRun = sMem.fNextTrackletFirstRun;
789 if (threadIdx.x == TRACKLET_CONSTRUCTOR_NMEMTHREDS)
790 {
791 sMem.fNTracklets = *tracker.NTracklets();
792 if (sMem.fNextTrackletFirstRun)
793 {
794#ifdef HLTCA_GPU_SCHED_FIXED_START
795 const int iSlice = tracker.GPUParametersConst()->fGPUnSlices * (blockIdx.x + (HLTCA_GPU_BLOCK_COUNT % tracker.GPUParametersConst()->fGPUnSlices != 0 && tracker.GPUParametersConst()->fGPUnSlices * (blockIdx.x + 1) % HLTCA_GPU_BLOCK_COUNT != 0)) / HLTCA_GPU_BLOCK_COUNT;
796 const int nSliceBlockOffset = HLTCA_GPU_BLOCK_COUNT * iSlice / tracker.GPUParametersConst()->fGPUnSlices;
797 const uint2 &nTracklet = tracker.BlockStartingTracklet()[blockIdx.x - nSliceBlockOffset];
798
799 sMem.fNextTrackletCount = nTracklet.y;
800 if (sMem.fNextTrackletCount == 0)
801 {
802 sMem.fNextTrackletFirstRun = 0;
803 }
804 else
805 {
806 if (tracker.TrackletStartHits()[nTracklet.x].RowIndex() / HLTCA_GPU_SCHED_ROW_STEP != RowBlock)
807 {
808 sMem.fNextTrackletCount = 0;
809 }
810 else
811 {
812 sMem.fNextTrackletFirst = nTracklet.x;
813 sMem.fNextTrackletNoDummy = 1;
814 }
815 }
816#endif
817 }
818 else
819 {
820 const int nFetchTracks = CAMath::Max(CAMath::Min((*tracker.RowBlockPos(Reverse, RowBlock)).x - (*tracker.RowBlockPos(Reverse, RowBlock)).y, HLTCA_GPU_THREAD_COUNT - TRACKLET_CONSTRUCTOR_NMEMTHREDS), 0);
821 sMem.fNextTrackletCount = nFetchTracks;
822 const int nUseTrack = nFetchTracks ? CAMath::AtomicAdd(&(*tracker.RowBlockPos(Reverse, RowBlock)).y, nFetchTracks) : 0;
823 sMem.fNextTrackletFirst = nUseTrack;
824
825 const int nFillTracks = CAMath::Min(nFetchTracks, nUseTrack + nFetchTracks - (*((volatile int2*) (tracker.RowBlockPos(Reverse, RowBlock)))).x);
826 if (nFillTracks > 0)
827 {
828 const int nStartFillTrack = CAMath::AtomicAdd(&(*tracker.RowBlockPos(Reverse, RowBlock)).x, nFillTracks);\r
829 if (nFillTracks + nStartFillTrack >= HLTCA_GPU_MAX_TRACKLETS)\r
830 {\r
831 tracker.GPUParameters()->fGPUError = HLTCA_GPU_ERROR_ROWBLOCK_TRACKLET_OVERFLOW;\r
832 }\r
833 for (int i = 0;i < nFillTracks;i++)\r
834 {\r
835 tracker.RowBlockTracklets(Reverse, RowBlock)[(nStartFillTrack + i) % HLTCA_GPU_MAX_TRACKLETS] = -3; //Dummy filling track\r
836 }\r
837 }
838 sMem.fNextTrackletNoDummy = 0;\r
839 }
840 }
841 __syncthreads();
842 mustInit = 0;
843 if (sMem.fNextTrackletCount == 0)
844 {
845 return(-2); //No more track in this RowBlock
846 }
847#if HLTCA_GPU_TRACKLET_CONSTRUCTOR_NMEMTHREDS > 0
848 else if (threadIdx.x < TRACKLET_CONSTRUCTOR_NMEMTHREDS)
849 {
850 return(-1);
851 }
852#endif
853 else if (threadIdx.x - TRACKLET_CONSTRUCTOR_NMEMTHREDS >= sMem.fNextTrackletCount)
854 {
855 return(-1); //No track in this RowBlock for this thread
856 }
857 else if (nextTracketlFirstRun)
858 {
859 if (threadIdx.x == TRACKLET_CONSTRUCTOR_NMEMTHREDS) sMem.fNextTrackletFirstRun = 0;
860 mustInit = 1;
861 return(sMem.fNextTrackletFirst + threadIdx.x - TRACKLET_CONSTRUCTOR_NMEMTHREDS);
862 }
863 else
864 {
865 const int nTrackPos = sMem.fNextTrackletFirst + threadIdx.x - TRACKLET_CONSTRUCTOR_NMEMTHREDS;\r
866 mustInit = (nTrackPos < tracker.RowBlockPos(Reverse, RowBlock)->w);\r
867 volatile int* const ptrTracklet = &tracker.RowBlockTracklets(Reverse, RowBlock)[nTrackPos % HLTCA_GPU_MAX_TRACKLETS];\r
868 int nTracklet;\r
869 int nTryCount = 0;\r
870 while ((nTracklet = *ptrTracklet) == -1)\r
871 {\r
872 for (int i = 0;i < 10000;i++)\r
873 sMem.fNextTrackletStupidDummy++;\r
874 nTryCount++;\r
875 if (nTryCount > 20)\r
876 {\r
877 tracker.GPUParameters()->fGPUError = HLTCA_GPU_ERROR_SCHEDULE_COLLISION;\r
878 return(-1);\r
879 }\r
880 };\r
881 return(nTracklet);\r
882 }
883}
00d07bcd 884
b22af1bf 885GPUd() void AliHLTTPCCATrackletConstructor::AliHLTTPCCATrackletConstructorNewGPU(AliHLTTPCCATracker *pTracker)
fbb9b71b 886{
b22af1bf 887 //Main Tracklet construction function that calls the scheduled (FetchTracklet) and then Processes the tracklet (mainly UpdataTracklet) and at the end stores the tracklet.
888 //Can also dispatch a tracklet to be rescheduled
889#ifdef HLTCA_GPU_EMULATION_SINGLE_TRACKLET
890 pTracker[0].BlockStartingTracklet()[0].x = HLTCA_GPU_EMULATION_SINGLE_TRACKLET;
891 pTracker[0].BlockStartingTracklet()[0].y = 1;
892 for (int i = 1;i < HLTCA_GPU_BLOCK_COUNT;i++)
893 {
894 pTracker[0].BlockStartingTracklet()[i].x = pTracker[0].BlockStartingTracklet()[i].y = 0;
895 }
896#endif
00d07bcd 897
b22af1bf 898 GPUshared() AliHLTTPCCASharedMemory sMem;
fbb9b71b 899
b22af1bf 900#ifdef HLTCA_GPU_SCHED_FIXED_START
901 if (threadIdx.x == TRACKLET_CONSTRUCTOR_NMEMTHREDS)
902 {
903 sMem.fNextTrackletFirstRun = 1;
904 }
905 __syncthreads();
906#endif
fbb9b71b 907
b22af1bf 908#ifdef HLTCA_GPU_TRACKLET_CONSTRUCTOR_DO_PROFILE
909 if (threadIdx.x == TRACKLET_CONSTRUCTOR_NMEMTHREDS)
910 {
911 sMem.fMaxSync = 0;
912 }
913 int threadSync = 0;
914#endif
fbb9b71b 915
b22af1bf 916 for (int iReverse = 0;iReverse < 2;iReverse++)
917 {
918 for (int iRowBlock = 0;iRowBlock < HLTCA_ROW_COUNT / HLTCA_GPU_SCHED_ROW_STEP + 1;iRowBlock++)
919 {
920#ifdef HLTCA_GPU_SCHED_FIXED_SLICE
921 int iSlice = tracker.GPUParametersConst()->fGPUnSlices * (blockIdx.x + (HLTCA_GPU_BLOCK_COUNT % tracker.GPUParametersConst()->fGPUnSlices != 0 && tracker.GPUParametersConst()->fGPUnSlices * (blockIdx.x + 1) % HLTCA_GPU_BLOCK_COUNT != 0)) / HLTCA_GPU_BLOCK_COUNT;
922#else
923 for (int iSlice = 0;iSlice < pTracker[0].GPUParametersConst()->fGPUnSlices;iSlice++)
924#endif
925 {
926 AliHLTTPCCATracker &tracker = pTracker[iSlice];
927 if (sMem.fNextTrackletFirstRun && iSlice != tracker.GPUParametersConst()->fGPUnSlices * (blockIdx.x + (HLTCA_GPU_BLOCK_COUNT % tracker.GPUParametersConst()->fGPUnSlices != 0 && tracker.GPUParametersConst()->fGPUnSlices * (blockIdx.x + 1) % HLTCA_GPU_BLOCK_COUNT != 0)) / HLTCA_GPU_BLOCK_COUNT)
928 {
929 continue;
930 }
931 /*if (!sMem.fNextTrackletFirstRun && tracker.RowBlockPos(1, HLTCA_ROW_COUNT / HLTCA_GPU_SCHED_ROW_STEP)->x <= tracker.RowBlockPos(1, HLTCA_ROW_COUNT / HLTCA_GPU_SCHED_ROW_STEP)->y)
932 {
933 continue;
934 }*/
935 int sharedRowsInitialized = 0;
936
937 int iTracklet;
938 int mustInit;
939 while ((iTracklet = FetchTracklet(tracker, sMem, iReverse, iRowBlock, mustInit)) != -2)
940 {
941#ifdef HLTCA_GPU_TRACKLET_CONSTRUCTOR_DO_PROFILE
942 CAMath::AtomicMax(&sMem.fMaxSync, threadSync);
943 __syncthreads();
944 threadSync = CAMath::Min(sMem.fMaxSync, 100000000 / blockDim.x / gridDim.x);
945#endif
946#ifndef HLTCA_GPU_PREFETCHDATA
947 if (!sharedRowsInitialized)
948 {
949#ifdef HLTCA_GPU_PREFETCH_ROWBLOCK_ONLY
950 if (iReverse)
951 {
952 for (int i = CAMath::Max(0, HLTCA_ROW_COUNT - (iRowBlock + 1) * HLTCA_GPU_SCHED_ROW_STEP) * sizeof(AliHLTTPCCARow) / sizeof(int) + threadIdx.x;i < (HLTCA_ROW_COUNT - iRowBlock * HLTCA_GPU_SCHED_ROW_STEP) * sizeof(AliHLTTPCCARow) / sizeof(int);i += blockDim.x)
953 {
954 reinterpret_cast<int*>(&sMem.fRows)[i] = reinterpret_cast<int*>(tracker.SliceDataRows())[i];
955 }
956 }
957 else
958 {
959 for (int i = CAMath::Max(1, iRowBlock * HLTCA_GPU_SCHED_ROW_STEP) * sizeof(AliHLTTPCCARow) / sizeof(int) + threadIdx.x;i < CAMath::Min((iRowBlock + 1) * HLTCA_GPU_SCHED_ROW_STEP, HLTCA_ROW_COUNT) * sizeof(AliHLTTPCCARow) / sizeof(int);i += blockDim.x)
960 {
961 reinterpret_cast<int*>(&sMem.fRows)[i] = reinterpret_cast<int*>(tracker.SliceDataRows())[i];
962 }
963 }
964#else
965 for (int i = threadIdx.x;i < HLTCA_ROW_COUNT * sizeof(AliHLTTPCCARow) / sizeof(int);i += blockDim.x)
966 {
967 reinterpret_cast<int*>(&sMem.fRows)[i] = reinterpret_cast<int*>(tracker.SliceDataRows())[i];
968 }
969#endif
970 sharedRowsInitialized = 1;
971 }
972#endif
973#ifdef HLTCA_GPU_RESCHED
974 short2 storeToRowBlock;
975 int storePosition = 0;
976 if (threadIdx.x - TRACKLET_CONSTRUCTOR_NMEMTHREDS < 2 * (HLTCA_ROW_COUNT / HLTCA_GPU_SCHED_ROW_STEP + 1))
977 {
978 const int nReverse = (threadIdx.x - TRACKLET_CONSTRUCTOR_NMEMTHREDS) / (HLTCA_ROW_COUNT / HLTCA_GPU_SCHED_ROW_STEP + 1);
979 const int nRowBlock = (threadIdx.x - TRACKLET_CONSTRUCTOR_NMEMTHREDS) % (HLTCA_ROW_COUNT / HLTCA_GPU_SCHED_ROW_STEP + 1);
980 sMem.fTrackletStoreCount[nReverse][nRowBlock] = 0;
981 }
982#endif
983 __syncthreads();
984 AliHLTTPCCATrackParam tParam;
985 AliHLTTPCCAThreadMemory rMem;
fbb9b71b 986
b22af1bf 987 rMem.fCurrentData = 0;
988
989#ifdef HLTCA_GPU_EMULATION_DEBUG_TRACKLET
990 if (iTracklet == HLTCA_GPU_EMULATION_DEBUG_TRACKLET)
991 {
992 tracker.GPUParameters()->fGPUError = 1;
993 }
994#endif
995 AliHLTTPCCAThreadMemory &rMemGlobal = tracker.GPUTrackletTemp()[iTracklet].fThreadMem;
996 AliHLTTPCCATrackParam &tParamGlobal = tracker.GPUTrackletTemp()[iTracklet].fParam;
997 if (mustInit)
998 {
999 AliHLTTPCCAHitId id = tracker.TrackletStartHits()[iTracklet];
1000
1001 rMem.fStartRow = rMem.fEndRow = rMem.fFirstRow = rMem.fLastRow = id.RowIndex();
1002 rMem.fCurrIH = id.HitIndex();
1003 rMem.fStage = 0;
1004 rMem.fNHits = 0;
1005 rMem.fNMissed = 0;
1006
1007 AliHLTTPCCATrackletConstructor::InitTracklet(tParam);
1008 }
1009 else if (iTracklet >= 0)
1010 {
1011 CopyTrackletTempData( rMemGlobal, rMem, tParamGlobal, tParam );
1012 }
1013#ifdef HLTCA_GPU_PREFETCHDATA
1014 else if (threadIdx.x < TRACKLET_CONSTRUCTOR_NMEMTHREDS)
1015 {
1016 ReadData(threadIdx.x, sMem, rMem, tracker, iReverse ? (HLTCA_ROW_COUNT - 1 - iRowBlock * HLTCA_GPU_SCHED_ROW_STEP) : (CAMath::Max(1, iRowBlock * HLTCA_GPU_SCHED_ROW_STEP)));
1017 }
1018#endif
1019 rMem.fItr = iTracklet;
1020 rMem.fGo = (iTracklet >= 0);
1021
1022#ifdef HLTCA_GPU_RESCHED
1023 storeToRowBlock.x = iRowBlock + 1;
1024 storeToRowBlock.y = iReverse;
1025#ifdef HLTCA_GPU_PREFETCHDATA
1026 rMem.fCurrentData ^= 1;
1027 __syncthreads();
1028#endif
1029 if (iReverse)
1030 {
1031 for (int j = HLTCA_ROW_COUNT - 1 - iRowBlock * HLTCA_GPU_SCHED_ROW_STEP;j >= CAMath::Max(0, HLTCA_ROW_COUNT - (iRowBlock + 1) * HLTCA_GPU_SCHED_ROW_STEP);j--)
1032 {
1033#ifdef HLTCA_GPU_TRACKLET_CONSTRUCTOR_DO_PROFILE
1034 if (rMem.fNMissed <= kMaxRowGap && rMem.fGo && !(j >= rMem.fEndRow || ( j >= rMem.fStartRow && j - rMem.fStartRow % 2 == 0)))
1035 pTracker[0].fStageAtSync[threadSync++ * blockDim.x * gridDim.x + blockIdx.x * blockDim.x + threadIdx.x] = rMem.fStage + 1;
1036#endif
1037#ifdef HLTCA_GPU_PREFETCHDATA
1038 if (threadIdx.x < TRACKLET_CONSTRUCTOR_NMEMTHREDS && j > CAMath::Max(0, HLTCA_ROW_COUNT - (iRowBlock + 1) * HLTCA_GPU_SCHED_ROW_STEP))
1039 {
1040 ReadData(threadIdx.x, sMem, rMem, tracker, j - 1);
1041 }
1042 else
1043#endif
1044 if (iTracklet >= 0)
1045 {
1046 UpdateTracklet(gridDim.x, blockDim.x, blockIdx.x, threadIdx.x, sMem, rMem, tracker, tParam, j);
1047 if (rMem.fNMissed > kMaxRowGap)
1048 {
1049 rMem.fGo = 0;
1050#ifndef HLTCA_GPU_PREFETCHDATA
1051 break;
1052#endif
1053 }
1054 }
1055#ifdef HLTCA_GPU_PREFETCHDATA
1056 __syncthreads();
1057 rMem.fCurrentData ^= 1;
1058#endif
1059 }
1060
1061 if (iTracklet >= 0 && (!rMem.fGo || iRowBlock == HLTCA_ROW_COUNT / HLTCA_GPU_SCHED_ROW_STEP))
1062 {
1063 StoreTracklet( gridDim.x, blockDim.x, blockIdx.x, threadIdx.x, sMem, rMem, tracker, tParam );
1064 }
1065 }
1066 else
1067 {
1068 for (int j = CAMath::Max(1, iRowBlock * HLTCA_GPU_SCHED_ROW_STEP);j < CAMath::Min((iRowBlock + 1) * HLTCA_GPU_SCHED_ROW_STEP, HLTCA_ROW_COUNT);j++)
1069 {
1070#ifdef HLTCA_GPU_TRACKLET_CONSTRUCTOR_DO_PROFILE
1071 if (rMem.fNMissed <= kMaxRowGap && rMem.fGo && j >= rMem.fStartRow && (rMem.fStage > 0 || rMem.fCurrIH >= 0 || (j - rMem.fStartRow) % 2 == 0 ))
1072 pTracker[0].fStageAtSync[threadSync++ * blockDim.x * gridDim.x + blockIdx.x * blockDim.x + threadIdx.x] = rMem.fStage + 1;
1073#endif
1074#ifdef HLTCA_GPU_PREFETCHDATA
1075 if (threadIdx.x < TRACKLET_CONSTRUCTOR_NMEMTHREDS && j < CAMath::Min((iRowBlock + 1) * HLTCA_GPU_SCHED_ROW_STEP, HLTCA_ROW_COUNT) - 1)
1076 {
1077 ReadData(threadIdx.x, sMem, rMem, tracker, j + 1);
1078 }
1079 else
1080#endif
1081 if (iTracklet >= 0)
1082 {
1083 UpdateTracklet( gridDim.x, blockDim.x, blockIdx.x, threadIdx.x, sMem, rMem, tracker, tParam, j);
1084#ifndef HLTCA_GPU_PREFETCHDATA
1085 //if (rMem.fNMissed > kMaxRowGap || rMem.fGo == 0) break; //DR!!! CUDA Crashes with this enabled
1086#endif
1087 }
1088#ifdef HLTCA_GPU_PREFETCHDATA
1089 __syncthreads();
1090 rMem.fCurrentData ^= 1;
1091#endif
1092 }
1093 if (rMem.fGo && (rMem.fNMissed > kMaxRowGap || iRowBlock == HLTCA_ROW_COUNT / HLTCA_GPU_SCHED_ROW_STEP))
1094 {
1095#if defined(HLTCA_GPU_PREFETCHDATA) | !defined(HLTCA_GPU_PREFETCH_ROWBLOCK_ONLY)
1096 if ( !tParam.TransportToX( tracker.Row( rMem.fEndRow ).X(), tracker.Param().ConstBz(), .999 ) )
1097#else
1098 if ( !tParam.TransportToX( sMem.fRows[ rMem.fEndRow ].X(), tracker.Param().ConstBz(), .999 ) )
1099#endif
1100 {
1101 rMem.fGo = 0;
1102 }
1103 else
1104 {
1105 storeToRowBlock.x = (HLTCA_ROW_COUNT - rMem.fEndRow) / HLTCA_GPU_SCHED_ROW_STEP;
1106 storeToRowBlock.y = 1;
1107 rMem.fNMissed = 0;
1108 rMem.fStage = 2;
1109 }
1110 }
1111
1112 if (iTracklet >= 0 && !rMem.fGo)
1113 {
1114 StoreTracklet( gridDim.x, blockDim.x, blockIdx.x, threadIdx.x, sMem, rMem, tracker, tParam );
1115 }
1116 }
1117
1118 if (rMem.fGo && (iRowBlock != HLTCA_ROW_COUNT / HLTCA_GPU_SCHED_ROW_STEP || iReverse == 0))
1119 {
1120 CopyTrackletTempData( rMem, rMemGlobal, tParam, tParamGlobal );
1121 storePosition = CAMath::AtomicAdd(&sMem.fTrackletStoreCount[storeToRowBlock.y][storeToRowBlock.x], 1);
1122 }
1123#else
1124 if (iTracklet >= 0)
1125 {
1126 for (int j = rMem.fStartRow;j < HLTCA_ROW_COUNT;j++)\r
1127 {\r
1128 UpdateTracklet(gridDim.x, blockDim.x, blockIdx.x, threadIdx.x, sMem, rMem, tracker, tParam, j);\r
1129 if (!rMem.fGo) break;\r
1130 }\r
1131\r
1132 rMem.fNMissed = 0;\r
1133 rMem.fStage = 2;\r
1134 if ( rMem.fGo )\r
1135 {\r
1136 if ( !tParam.TransportToX( tracker.Row( rMem.fEndRow ).X(), tracker.Param().ConstBz(), .999 ) ) rMem.fGo = 0;\r
1137 }\r
1138\r
1139 for (int j = rMem.fEndRow;j >= 0;j--)\r
1140 {\r
1141 if (!rMem.fGo) break;\r
1142 UpdateTracklet( gridDim.x, blockDim.x, blockIdx.x, threadIdx.x, sMem, rMem, tracker, tParam, j);\r
1143 }\r
1144\r
1145 StoreTracklet( gridDim.x, blockDim.x, blockIdx.x, threadIdx.x, sMem, rMem, tracker, tParam );
1146 }
1147#endif
1148#ifdef HLTCA_GPU_RESCHED
1149 __syncthreads();
1150 if (threadIdx.x - TRACKLET_CONSTRUCTOR_NMEMTHREDS < 2 * (HLTCA_ROW_COUNT / HLTCA_GPU_SCHED_ROW_STEP + 1))
1151 {
1152 const int nReverse = (threadIdx.x - TRACKLET_CONSTRUCTOR_NMEMTHREDS) / (HLTCA_ROW_COUNT / HLTCA_GPU_SCHED_ROW_STEP + 1);
1153 const int nRowBlock = (threadIdx.x - TRACKLET_CONSTRUCTOR_NMEMTHREDS) % (HLTCA_ROW_COUNT / HLTCA_GPU_SCHED_ROW_STEP + 1);
1154 if (sMem.fTrackletStoreCount[nReverse][nRowBlock])
1155 {
1156 sMem.fTrackletStoreCount[nReverse][nRowBlock] = CAMath::AtomicAdd(&tracker.RowBlockPos(nReverse, nRowBlock)->x, sMem.fTrackletStoreCount[nReverse][nRowBlock]);
1157 }
1158 }
1159 __syncthreads();
1160 if (iTracklet >= 0 && rMem.fGo && (iRowBlock != HLTCA_ROW_COUNT / HLTCA_GPU_SCHED_ROW_STEP || iReverse == 0))
1161 {
1162 tracker.RowBlockTracklets(storeToRowBlock.y, storeToRowBlock.x)[sMem.fTrackletStoreCount[storeToRowBlock.y][storeToRowBlock.x] + storePosition] = iTracklet;
1163 }
1164 __syncthreads();
1165#endif
1166 }
1167 }
1168 }
1169 }
1170#ifdef HLTCA_GPU_TRACKLET_CONSTRUCTOR_DO_PROFILE
1171
1172#endif
00d07bcd 1173}
1174
b22af1bf 1175GPUd() void AliHLTTPCCATrackletConstructor::AliHLTTPCCATrackletConstructorInit(int iTracklet, AliHLTTPCCATracker &tracker)
1176{
1177 //Initialize Row Blocks
1178
1179#ifndef HLTCA_GPU_EMULATION_SINGLE_TRACKLET
1180AliHLTTPCCAHitId id = tracker.TrackletStartHits()[iTracklet];
1181#ifdef HLTCA_GPU_SCHED_FIXED_START
1182 const int firstDynamicTracklet = tracker.GPUParameters()->fScheduleFirstDynamicTracklet;
1183 if (iTracklet >= firstDynamicTracklet)
1184#endif
1185 {
1186 const int firstTrackletInRowBlock = CAMath::Max(firstDynamicTracklet, tracker.RowStartHitCountOffset()[CAMath::Max(id.RowIndex() / HLTCA_GPU_SCHED_ROW_STEP * HLTCA_GPU_SCHED_ROW_STEP, 1)].z);
1187 if (iTracklet == firstTrackletInRowBlock)
1188 {
1189 const int firstRowInNextBlock = (id.RowIndex() / HLTCA_GPU_SCHED_ROW_STEP + 1) * HLTCA_GPU_SCHED_ROW_STEP;
1190 int trackletsInRowBlock;
1191 if (firstRowInNextBlock >= HLTCA_ROW_COUNT - 3)
1192 trackletsInRowBlock = *tracker.NTracklets() - firstTrackletInRowBlock;
1193 else
1194 trackletsInRowBlock = CAMath::Max(firstDynamicTracklet, tracker.RowStartHitCountOffset()[firstRowInNextBlock].z) - firstTrackletInRowBlock;
1195
1196 tracker.RowBlockPos(0, id.RowIndex() / HLTCA_GPU_SCHED_ROW_STEP)->x = trackletsInRowBlock;
1197 tracker.RowBlockPos(0, id.RowIndex() / HLTCA_GPU_SCHED_ROW_STEP)->w = trackletsInRowBlock;
1198 }
1199 tracker.RowBlockTracklets(0, id.RowIndex() / HLTCA_GPU_SCHED_ROW_STEP)[iTracklet - firstTrackletInRowBlock] = iTracklet;
1200 }
1201#endif
1202}
1203
1204GPUg() void AliHLTTPCCATrackletConstructorInit(int iSlice)
1205{
1206 //GPU Wrapper for AliHLTTPCCATrackletConstructor::AliHLTTPCCATrackletConstructorInit
1207 AliHLTTPCCATracker &tracker = ( ( AliHLTTPCCATracker* ) gAliHLTTPCCATracker )[iSlice];
1208 int i = blockIdx.x * blockDim.x + threadIdx.x;
1209 if (i >= *tracker.NTracklets()) return;
1210 AliHLTTPCCATrackletConstructor::AliHLTTPCCATrackletConstructorInit(i, tracker);
1211}
1212
1213GPUg() void AliHLTTPCCATrackletConstructorNewGPU()
1214{
1215 //GPU Wrapper for AliHLTTPCCATrackletConstructor::AliHLTTPCCATrackletConstructorNewGPU
1216 AliHLTTPCCATracker *pTracker = ( ( AliHLTTPCCATracker* ) gAliHLTTPCCATracker );
1217 AliHLTTPCCATrackletConstructor::AliHLTTPCCATrackletConstructorNewGPU(pTracker);
1218}
1219
1220#else
1221GPUd() void AliHLTTPCCATrackletConstructor::AliHLTTPCCATrackletConstructorNewCPU(AliHLTTPCCATracker &tracker)
1222{
1223 //Tracklet constructor simple CPU Function that does not neew a scheduler
1224 GPUshared() AliHLTTPCCASharedMemory sMem;
1225 sMem.fNTracklets = *tracker.NTracklets();
1226 for (int iTracklet = 0;iTracklet < *tracker.NTracklets();iTracklet++)
1227 {
1228 AliHLTTPCCATrackParam tParam;
1229 AliHLTTPCCAThreadMemory rMem;
1230
1231 AliHLTTPCCAHitId id = tracker.TrackletStartHits()[iTracklet];
1232
1233 rMem.fStartRow = rMem.fEndRow = rMem.fFirstRow = rMem.fLastRow = id.RowIndex();
1234 rMem.fCurrIH = id.HitIndex();
1235 rMem.fStage = 0;
1236 rMem.fNHits = 0;
1237 rMem.fNMissed = 0;
1238
1239 AliHLTTPCCATrackletConstructor::InitTracklet(tParam);
1240
1241 rMem.fItr = iTracklet;
1242 rMem.fGo = 1;
1243
1244 for (int j = rMem.fStartRow;j < tracker.Param().NRows();j++)\r
1245 {\r
1246 UpdateTracklet(1, 1, 0, iTracklet, sMem, rMem, tracker, tParam, j);\r
1247 if (!rMem.fGo) break;\r
1248 }\r
1249\r
1250 rMem.fNMissed = 0;\r
1251 rMem.fStage = 2;\r
1252 if ( rMem.fGo )\r
1253 {\r
1254 if ( !tParam.TransportToX( tracker.Row( rMem.fEndRow ).X(), tracker.Param().ConstBz(), .999 ) ) rMem.fGo = 0;\r
1255 }\r
1256\r
1257 for (int j = rMem.fEndRow;j >= 0;j--)\r
1258 {\r
1259 if (!rMem.fGo) break;\r
1260 UpdateTracklet( 1, 1, 0, iTracklet, sMem, rMem, tracker, tParam, j);\r
1261 }\r
1262\r
1263 StoreTracklet( 1, 1, 0, iTracklet, sMem, rMem, tracker, tParam );
1264 }
1265}
1266#endif