HLT/TPCLib/tracking-ca/AliHLTTPCCATracker.cxx

   1 // @(#) $Id$
   2 //***************************************************************************
   3 // This file is property of and copyright by the ALICE HLT Project          *
   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.                    *
  17 //***************************************************************************
  18
  19 #include "AliHLTTPCCATracker.h"
  20 #include "AliHLTTPCCAOutTrack.h"
  21 #include "AliHLTTPCCAGrid.h"
  22 #include "AliHLTTPCCARow.h"
  23 #include "AliHLTTPCCATrack.h"
  24 #include "AliHLTTPCCAMath.h"
  25 #include "AliHLTTPCCAHit.h"
  26
  27 #include "TStopwatch.h"
  28 #include "AliHLTTPCCAHitArea.h"
  29 #include "AliHLTTPCCANeighboursFinder.h"
  30 #include "AliHLTTPCCANeighboursCleaner.h"
  31 #include "AliHLTTPCCAStartHitsFinder.h"
  32 #include "AliHLTTPCCATrackletConstructor.h"
  33 #include "AliHLTTPCCATrackletSelector.h"
  34 #include "AliHLTTPCCAProcess.h"
  35 #include "AliHLTTPCCALinksWriter.h"
  36 #include "AliHLTTPCCAUsedHitsInitialiser.h"
  37
  38 #include "AliHLTTPCCATrackParam.h"
  39 #include "AliHLTTPCCATrackParam1.h"
  40
  41 #if !defined(HLTCA_GPUCODE)
  42 #if defined( HLTCA_STANDALONE )
  43 #include <iostream.h>
  44 #else
  45 #include "Riostream.h"
  46 #endif
  47 #endif
  48
  49 //#define DRAW
  50
  51 #ifdef DRAW
  52   #include "AliHLTTPCCADisplay.h"
  53   #include "TApplication.h"
  54 #endif //DRAW
  55
  56 ClassImp(AliHLTTPCCATracker)
  57
  58 #if !defined(HLTCA_GPUCODE)
  59
  60 AliHLTTPCCATracker::AliHLTTPCCATracker()
  61   :
  62   fParam(),
  63   fNHitsTotal(0),
  64   fGridSizeTotal(0),
  65   fGrid1SizeTotal(0),
  66   fHits(0),
  67   fHits1(0),
  68   fGridContents(0),
  69   fGrid1Contents(0),
  70   fHitsID(0),
  71   fHitLinkUp(0),
  72   fHitLinkDown(0),
  73   fHitIsUsed(0),
  74   fStartHits(0),
  75   fTracklets(0),
  76   fNTracks(0),
  77   fTracks(0),
  78   fTrackHits(0),
  79   fNOutTracks(0),
  80   fNOutTrackHits(0),
  81   fOutTracks(0),
  82   fOutTrackHits(0),
  83   fEventMemory(0),
  84   fEventMemSize(0),
  85   fTexHitsFullData(0),
  86   fTexHitsFullSize(0)
  87 {
  88   // constructor
  89   //fRows = new AliHLTTPCCARow[fParam.NRows()];
  90   //Initialize( fParam );
  91 }
  92
  93 AliHLTTPCCATracker::AliHLTTPCCATracker( const AliHLTTPCCATracker& )
  94   :
  95   fParam(),
  96   fNHitsTotal(0),
  97   fGridSizeTotal(0),
  98   fGrid1SizeTotal(0),
  99   fHits(0),
 100   fHits1(0),
 101   fGridContents(0),
 102   fGrid1Contents(0),
 103   fHitsID(0),
 104   fHitLinkUp(0),
 105   fHitLinkDown(0),
 106   fHitIsUsed(0),
 107   fStartHits(0),
 108   fTracklets(0),
 109   fNTracks(0),
 110   fTracks(0),
 111   fTrackHits(0),
 112   fNOutTracks(0),
 113   fNOutTrackHits(0),
 114   fOutTracks(0),
 115   fOutTrackHits(0),
 116   fEventMemory(0),
 117   fEventMemSize(0),
 118   fTexHitsFullData(0),
 119   fTexHitsFullSize(0)
 120 {
 121   // dummy
 122 }
 123
 124 AliHLTTPCCATracker &AliHLTTPCCATracker::operator=( const AliHLTTPCCATracker& )
 125 {
 126   // dummy
 127   fOutTrackHits=0;
 128   fOutTracks=0;
 129   fNOutTracks=0;
 130   fTrackHits = 0;
 131   fEventMemory = 0;
 132   return *this;
 133 }
 134
 135 GPUd() AliHLTTPCCATracker::~AliHLTTPCCATracker()
 136 {
 137   // destructor
 138   StartEvent();
 139 }
 140 #endif
 141
 142
 143 GPUd() UChar_t AliHLTTPCCATracker::GetGridContent(  UInt_t i ) const
 144 {
 145   //* get grid content
 146 #if defined(HLTCA_GPUSTEP)
 147   return (UChar_t) tex1Dfetch(texGrid,i).x;
 148 #else
 149   return fGridContents[i];
 150 #endif
 151 }
 152
 153
 154 GPUd() AliHLTTPCCAHit AliHLTTPCCATracker::GetHit(  UInt_t i ) const
 155 {
 156   //* get hit
 157 #if defined(HLTCA_USE_GPU)
 158   AliHLTTPCCAHit h;
 159   float2 f = tex1Dfetch(texHits,i);
 160   h.Y() = f.x;
 161   h.Z() = f.y;
 162   return h;
 163 #else
 164   return fHits[i];
 165 #endif
 166 }
 167
 168
 169
 170
 171 // ----------------------------------------------------------------------------------
 172 GPUd() void AliHLTTPCCATracker::Initialize( AliHLTTPCCAParam &param )
 173 {
 174   // initialisation
 175   StartEvent();
 176   fParam = param;
 177   fParam.Update();
 178   for( Int_t irow=0; irow<fParam.NRows(); irow++ ){
 179     fRows[irow].X() = fParam.RowX(irow);
 180     fRows[irow].MaxY() = CAMath::Tan( fParam.DAlpha()/2.)*fRows[irow].X();
 181   }
 182   StartEvent();
 183 }
 184
 185 GPUd() void AliHLTTPCCATracker::StartEvent()
 186 {
 187   // start new event and fresh the memory
 188
 189   if( fEventMemory ) delete[] fEventMemory;
 190   if( fOutTracks ) delete[] fOutTracks;
 191   if( fOutTrackHits ) delete[] fOutTrackHits;
 192   fEventMemory = 0;
 193   fHits = 0;
 194   fHits1 = 0;
 195   fHitsID = 0;
 196   fTrackHits = 0;
 197   fTracks = 0;
 198   fOutTrackHits = 0;
 199   fOutTracks = 0;
 200   fNTracks = 0;
 201   fNOutTrackHits = 0;
 202   fNOutTracks = 0;
 203   fNHitsTotal = 0;
 204 }
 205
 206 GPUhd() void  AliHLTTPCCATracker::SetPointers()
 207 {
 208   // set all pointers to the event memory
 209
 210   fEventMemSize = 0;
 211   UInt_t &size = fEventMemSize;
 212   fHits = (AliHLTTPCCAHit*) (fEventMemory+ size);
 213   size+= sizeof(AliHLTTPCCAHit)*fNHitsTotal;
 214   fHits1 = (ushort2*) (fEventMemory+ size);
 215   size+= sizeof(ushort2)*fNHitsTotal;
 216   fGridContents = (UChar_t *) (fEventMemory + size);
 217   size+= sizeof(Int_t)*fGridSizeTotal;
 218   fGrid1Contents = (UInt_t *) (fEventMemory + size);
 219   size+= sizeof(UInt_t)*fGrid1SizeTotal;
 220   fHitsID = (Int_t *) (fEventMemory + size);
 221   size+= sizeof(Int_t) * fNHitsTotal;
 222   fHitLinkUp = (Short_t *) (fEventMemory + size);
 223   size+= sizeof(Int_t) * fNHitsTotal;
 224   fHitLinkDown = (Short_t *) (fEventMemory + size);
 225   size+= sizeof(Int_t) * fNHitsTotal;
 226   fHitIsUsed = (Int_t *) (fEventMemory + size);
 227   size+= sizeof(Int_t) * fNHitsTotal;
 228   fStartHits = (Int_t *) (fEventMemory + size);
 229   size+= sizeof(Int_t) * (fNHitsTotal+1);
 230   size = (size/16+1)*16;
 231   fTexHitsFullData = (uint4*)(fEventMemory+ size);
 232   size+= ((sizeof(UShort_t)*6*fNHitsTotal + sizeof(UShort_t)*2*fGrid1SizeTotal )/16+1)*16;
 233
 234   fTracklets = (Int_t *) (fEventMemory + size);
 235   size+= sizeof(Int_t) * (1 + fNHitsTotal*(5+ sizeof(AliHLTTPCCATrackParam)/4 + 160 ));
 236   fNTracks = (Int_t *) (fEventMemory + size);
 237   size+= sizeof(Int_t);
 238   fTracks = (AliHLTTPCCATrack* )(fEventMemory + size);
 239   size+= sizeof(AliHLTTPCCATrack) * (fNHitsTotal+1);
 240   fTrackHits = ( Int_t *)(fEventMemory + size);
 241   size+= sizeof(Int_t) * (10*fNHitsTotal+1);
 242
 243   fOutTrackHits = 0;
 244   fOutTracks = 0;
 245 }
 246
 247 GPUd() void AliHLTTPCCATracker::ReadEvent( Int_t *RowFirstHit, Int_t *RowNHits, Float_t *Y, Float_t *Z, Int_t NHits )
 248 {
 249   //* Read event
 250
 251   fNHitsTotal = NHits;
 252   fGridSizeTotal = 0;
 253   fGrid1SizeTotal = 0;
 254   fTexHitsFullSize = 0;
 255
 256   //cout<<"event mem = "<<fEventMemory<<endl;
 257   for( Int_t iRow=0; iRow<fParam.NRows(); iRow++ ){
 258     //cout<<"row, nhits="<<iRow<<" "<<RowNHits[iRow]<<endl;
 259     //cout<<"row, firsthit="<<iRow<<" "<<RowFirstHit[iRow]<<endl;
 260     AliHLTTPCCARow &row = fRows[iRow];
 261     row.FirstHit() = RowFirstHit[iRow];
 262     row.NHits() = RowNHits[iRow];
 263     Float_t yMin=1.e3, yMax=-1.e3, zMin=1.e3, zMax=-1.e3;
 264     Int_t nGrid =  row.NHits();
 265     for( Int_t i=0; i<row.NHits(); i++ ){
 266       Int_t j = RowFirstHit[iRow]+i;
 267       if( yMax < Y[j] ) yMax = Y[j];
 268       if( yMin > Y[j] ) yMin = Y[j];
 269       if( zMax < Z[j] ) zMax = Z[j];
 270       if( zMin > Z[j] ) zMin = Z[j];
 271     }
 272     if( nGrid == 0 ){
 273       yMin = yMax = zMin = zMax = 0;
 274       nGrid = 1;
 275     }
 276
 277     row.Grid().Create( yMin, yMax, zMin, zMax, nGrid );
 278
 279     float sy = ( CAMath::Abs( row.Grid().StepYInv() ) >1.e-4 ) ?1./row.Grid().StepYInv() :1;
 280     float sz = ( CAMath::Abs( row.Grid().StepZInv() ) >1.e-4 ) ?1./row.Grid().StepZInv() :1;
 281
 282     //cout<<"grid n = "<<row.Grid().N()<<" "<<sy<<" "<<sz<<" "<<yMin<<" "<<yMax<<" "<<zMin<<" "<<zMax<<endl;
 283
 284     bool recreate=0;
 285     if( sy < 2. ) { recreate = 1; sy = 2; }
 286     if( sz < 2. ) { recreate = 1; sz = 2; }
 287     if( recreate ) row.Grid().Create( yMin, yMax, zMin, zMax, sy, sz );
 288
 289     fGridSizeTotal+=row.Grid().N()+3+10;
 290     //cout<<"grid n = "<<row.Grid().N()<<endl;
 291   }
 292
 293   fGrid1SizeTotal = fGridSizeTotal+10;
 294
 295   SetPointers();
 296
 297   fEventMemory = (char*) ( new uint4 [ fEventMemSize/sizeof(uint4) + 100]);
 298   SetPointers();
 299
 300   fGridSizeTotal = 0;
 301   fGrid1SizeTotal = 0;
 302
 303   for( Int_t iRow=0; iRow<fParam.NRows(); iRow++ ){
 304     AliHLTTPCCARow &row = fRows[iRow];
 305     AliHLTTPCCAGrid &grid = row.Grid();
 306
 307     Int_t c[grid.N()+3+10];
 308     Int_t bins[row.NHits()];
 309     Int_t filled[ row.Grid().N() +3+10 ];
 310
 311     for( UInt_t bin=0; bin<row.Grid().N()+3; bin++ ) filled[bin] = 0;
 312
 313     for( Int_t i=0; i<row.NHits(); i++ ){
 314       Int_t j = RowFirstHit[iRow]+i;
 315       Int_t bin = row.Grid().GetBin( Y[j], Z[j] );
 316       bins[i] = bin;
 317       filled[bin]++;
 318     }
 319
 320     {
 321       Int_t n=0;
 322       for( UInt_t bin=0; bin<row.Grid().N()+3; bin++ ){
 323         c[bin] = n;
 324         n+=filled[bin];
 325       }
 326     }
 327     for( Int_t i=0; i<row.NHits(); i++ ){
 328       Int_t bin = bins[i];
 329       Int_t ind = c[bin] + filled[bin]-1;
 330       AliHLTTPCCAHit &h = fHits[RowFirstHit[iRow]+ind];
 331       fHitsID[RowFirstHit[iRow]+ind] = RowFirstHit[iRow]+i;
 332       h.Y() = Y[row.FirstHit()+i];
 333       h.Z() = Z[row.FirstHit()+i];
 334       filled[bin]--;
 335     }
 336
 337     grid.Offset() = fGridSizeTotal;
 338     Int_t off= grid.N()+3+10;
 339     fGridSizeTotal+=off;
 340     Int_t n2 = grid.N()/2;
 341     grid.Content2() = c[n2];
 342     UChar_t *cnew = fGridContents + grid.Offset();
 343
 344     for( Int_t i=0; i<n2; i++ ){
 345       Int_t v = c[i];
 346       if( v>=256 ){
 347         //cout<<" ERROR!!! "<<v<<endl;
 348         v = 255;
 349       }else if( v<0 ){
 350         //cout<<" ERROR!!! "<<v<<endl;
 351         v = 0;
 352       }
 353       cnew[i] = (UChar_t ) v;
 354     }
 355     for( UInt_t i=n2; i<grid.N()+3; i++ ){
 356       Int_t v = c[i] - grid.Content2();
 357       if( v>=256 ){
 358         //cout<<" ERROR 1 !!! "<<v<<endl;
 359         v = 255;
 360       }else if( v<0 ){
 361         //cout<<" ERROR 1 !!! "<<v<<endl;
 362         v = 0;
 363       }
 364       cnew[i] = (UChar_t) v;
 365     }
 366
 367
 368     UInt_t *cnew1 = fGrid1Contents + grid.Offset();
 369
 370     for( UInt_t i=0; i<grid.N()+1; i++ ){
 371       UInt_t g0n = 0;
 372       UInt_t g1n = 0;
 373       UInt_t g1 = 0;
 374       UInt_t g0 = c[i];// max [gN]
 375       UInt_t g0e = c[i+2]; //max[gN+2]
 376       g0n = g0e - g0;
 377       if( i+grid.Ny()< grid.N()+1 ){// max [gN-gNy]
 378         g1 = c[i+grid.Ny()]; // max [gN]
 379         UInt_t g1e = c[i+grid.Ny()+2];//max [gN+2]
 380         g1n = g1e - g1;
 381       }
 382
 383       if( g0n > 63 ) g0n = 63;
 384       if( g1n > 63 ) g1n = 63;
 385       cnew1[i] = (g1n<<26) + (g1<<16) + (g0n<<10) + g0;
 386     }
 387     {
 388       float y0 = row.Grid().YMin();
 389       float stepY = (row.Grid().YMax() - y0)*(1./65535.);
 390       float z0 = row.Grid().ZMin();
 391       float stepZ = (row.Grid().ZMax() - z0)*(1./65535.);
 392       float stepYi = 1./stepY;
 393       float stepZi = 1./stepZ;
 394
 395       row.Hy0() = y0;
 396       row.Hz0() = z0;
 397       row.HstepY() = stepY;
 398       row.HstepZ() = stepZ;
 399       row.HstepYi() = stepYi;
 400       row.HstepZi() = stepZi;
 401
 402       for( Int_t ih=0; ih<row.NHits(); ih++ ){
 403         Int_t ihTot = RowFirstHit[iRow]+ih;
 404         AliHLTTPCCAHit &hh = fHits[ihTot];
 405         ushort2  &h = fHits1[ihTot];
 406         float xx = ((hh.Y() - y0)*stepYi); //SG!!!
 407         float yy = ((hh.Z() - z0)*stepZi);
 408         if( xx<0 || yy<0 || xx>=65536 || yy>= 65536 ){
 409           cout<<"!!!! hit packing error!!! "<<xx<<" "<<yy<<" "<<endl;
 410         }
 411         h.x = (UShort_t) xx;//((hh.Y() - y0)*stepYi);
 412         h.y = (UShort_t) yy;//((hh.Z() - z0)*stepZi);
 413       }
 414     }
 415
 416     if(1){
 417       row.FullOffset() = fTexHitsFullSize;
 418       ushort2 *p= (ushort2*)(fTexHitsFullData+row.FullOffset());
 419       for( Int_t ih=0; ih<row.NHits(); ih++ ){
 420         Int_t ihTot = RowFirstHit[iRow]+ih;
 421         p[ih] = fHits1[ihTot];
 422       }
 423       Int_t size = row.NHits()*sizeof(ushort2);
 424
 425       row.FullGridOffset() = row.NHits()*2;
 426       UShort_t *p1 = ((UShort_t *)p) + row.FullGridOffset();
 427
 428       Int_t n = grid.N();
 429       for( Int_t i=0; i<n; i++ ){
 430         p1[i] = c[i];
 431       }
 432       UShort_t a = c[n];
 433       Int_t nn = n+grid.Ny()+2;
 434       for( Int_t i=n; i<nn; i++ ) p1[i] = a;
 435
 436       size+= (nn)*sizeof(UShort_t);
 437       row.FullLinkOffset() = row.NHits()*2 + nn;
 438       size+= row.NHits()*2*sizeof(Short_t);
 439       if( size%16 ) size = size/16+1;
 440       else size = size/16;
 441       row.FullSize()=size;
 442       //cout<<iRow<<", "<<row.fNHits<<"= "<<size*16<<"b: "<<row.fFullOffset<<" "<<row.fFullSize<<" "<<row.fFullGridOffset<<" "<<row.fFullLinkOffset<<endl;
 443
 444       fTexHitsFullSize+=size;
 445     }
 446   }
 447   fGrid1SizeTotal = fGridSizeTotal+10;
 448 }
 449
 450
 451 GPUh() void AliHLTTPCCATracker::Reconstruct()
 452 {
 453   //* reconstruction of event
 454
 455 #ifdef DRAW
 456   if( !gApplication ){
 457     TApplication *myapp = new TApplication("myapp",0,0);
 458   }
 459   //AliHLTTPCCADisplay::Instance().Init();
 460
 461   AliHLTTPCCADisplay::Instance().SetCurrentSlice( this );
 462   AliHLTTPCCADisplay::Instance().SetSliceView();
 463   AliHLTTPCCADisplay::Instance().DrawSlice( this );
 464   //for( Int_t iRow=0; iRow<fParam.NRows(); iRow++ )
 465   //for (Int_t i = 0; i<fRows[iRow].NHits(); i++)
 466   //AliHLTTPCCADisplay::Instance().DrawHit( iRow, i );
 467   //AliHLTTPCCADisplay::Instance().Ask();
 468 #endif
 469
 470   fTimers[0] = 0; // find neighbours
 471   fTimers[1] = 0; // construct tracklets
 472   fTimers[2] = 0; // fit tracklets
 473   fTimers[3] = 0; // prolongation of tracklets
 474   fTimers[4] = 0; // selection
 475   fTimers[5] = 0; // write output
 476   fTimers[6] = 0;
 477   fTimers[7] = 0;
 478
 479   if( fNHitsTotal < 1 ) return;
 480   //if( fParam.ISlice()!=3 ) return;
 481   TStopwatch timer0;
 482   *fNTracks = 0;
 483 #if !defined(HLTCA_GPUCODE)
 484
 485   AliHLTTPCCAProcess<AliHLTTPCCANeighboursFinder>( Param().NRows(), 1, *this );
 486   AliHLTTPCCAProcess<AliHLTTPCCANeighboursCleaner>( Param().NRows()-2, 1, *this );
 487   AliHLTTPCCAProcess<AliHLTTPCCAStartHitsFinder>( Param().NRows()-4, 1, *this );
 488
 489   Int_t nStartHits = *fStartHits;
 490
 491   Int_t nThreads = 128;
 492   Int_t nBlocks = fNHitsTotal/nThreads + 1;
 493   if( nBlocks<12 ){
 494     nBlocks = 12;
 495     nThreads = fNHitsTotal/12+1;
 496     if( nThreads%32 ) nThreads = (nThreads/32+1)*32;
 497   }
 498
 499   nThreads = fNHitsTotal;
 500   nBlocks = 1;
 501
 502   AliHLTTPCCAProcess<AliHLTTPCCAUsedHitsInitialiser>(nBlocks, nThreads,*this);
 503
 504   nThreads = 256;
 505   nBlocks = 30;
 506
 507   nThreads = 1;
 508   nBlocks = 1;
 509
 510   AliHLTTPCCAProcess<AliHLTTPCCALinksWriter>(nBlocks, nThreads,*this);
 511
 512   Int_t nMemThreads = 128;
 513   nThreads = 256;//96;
 514   nBlocks = nStartHits/nThreads + 1;
 515   if( nBlocks<30 ){
 516     nBlocks = 30;
 517     nThreads = (nStartHits)/30+1;
 518     if( nThreads%32 ) nThreads = (nThreads/32+1)*32;
 519   }
 520
 521   nThreads = nStartHits;
 522   nBlocks = 1;
 523
 524   AliHLTTPCCAProcess1<AliHLTTPCCATrackletConstructor>(nBlocks, nMemThreads+nThreads,*this);
 525
 526   {
 527     nThreads = 128;
 528     nBlocks = nStartHits/nThreads + 1;
 529     if( nBlocks<12 ){
 530       nBlocks = 12;
 531       nThreads = nStartHits/12+1;
 532       nThreads = (nThreads/32+1)*32;
 533     }
 534     *fStartHits = 0;
 535     *fTrackHits = 0;
 536
 537     nThreads = nStartHits;
 538     nBlocks = 1;
 539
 540
 541     AliHLTTPCCAProcess<AliHLTTPCCATrackletSelector>(nBlocks, nThreads,*this);
 542
 543     //cudaMemcpy(cpuTrackerCopy.fNTracks, gpuTrackerCopy.fNTracks, sizeof(int), cudaMemcpyDeviceToHost);
 544     //cudaMemcpy(cpuTrackerCopy.fTrackHits, gpuTrackerCopy.fTrackHits, sizeof(int), cudaMemcpyDeviceToHost);
 545
 546     //Int_t size = sizeof(AliHLTTPCCATrack)*( *cpuTrackerCopy.fNTracks );
 547     //cudaMemcpy(cpuTrackerCopy.fTracks, gpuTrackerCopy.fTracks, size, cudaMemcpyDeviceToHost);
 548     //cout<<"Tracks size = "<<size<<endl;
 549
 550     //size = sizeof(Int_t)*( *cpuTrackerCopy.fTrackHits );
 551     //cudaMemcpy(cpuTrackerCopy.fTrackHits+1, gpuTrackerCopy.fTrackHits+1, size, cudaMemcpyDeviceToHost);
 552     //cout<<"Track hits size = "<<size<<endl;
 553     //cpuTrackerCopy.WriteOutput();
 554
 555     Int_t nTracklets = *fStartHits;
 556
 557     //cout<<"Slice "<<Param().ISlice()<<": N start hits/tracklets/tracks = "<<nStartHits<<" "<<nTracklets<<" "<<*fNTracks<<endl;
 558    WriteOutput();
 559   }
 560
 561 #endif
 562
 563   timer0.Stop();
 564   fTimers[0] = timer0.CpuTime();
 565
 566  }
 567
 568
 569
 570
 571 GPUh() void AliHLTTPCCATracker::WriteOutput()
 572 {
 573   // write output
 574
 575   TStopwatch timer;
 576   fOutTrackHits = new Int_t[fNHitsTotal*10];
 577   fOutTracks = new AliHLTTPCCAOutTrack[*fNTracks];
 578   fNOutTrackHits = 0;
 579   fNOutTracks = 0;
 580   //cout<<"NTracks = "<<*fNTracks<<endl;
 581   //cout<<"NHits = "<<fNHitsTotal<<endl;
 582   for( Int_t iTr=0; iTr<*fNTracks; iTr++){
 583     //cout<<"iTr = "<<iTr<<endl;
 584     AliHLTTPCCATrack &iTrack = fTracks[iTr];
 585     if( !iTrack.Alive() ) continue;
 586     if( iTrack.NHits()<3 ) continue;
 587     //cout<<10<<endl;
 588     AliHLTTPCCAOutTrack &out = fOutTracks[fNOutTracks];
 589     out.FirstHitRef() = fNOutTrackHits;
 590     out.NHits() = 0;
 591     out.OrigTrackID() = iTr;
 592     {
 593       out.StartPoint() = iTrack.Param();
 594       out.EndPoint() = iTrack.Param();
 595     }
 596     //cout<<11<<endl;
 597
 598     Int_t iID = iTrack.FirstHitID();
 599     Int_t fNOutTrackHitsOld = fNOutTrackHits;
 600     //cout<<12<<" "<<iID<<" "<<iTrack.NHits()<<endl;
 601     for( Int_t ith=0; ith<iTrack.NHits(); ith++ ){
 602       //cout<<ith<<":"<<endl;
 603       Int_t ic = (fTrackHits[iID+ith]);
 604       //cout<<ic<<endl;
 605       AliHLTTPCCARow &row = ID2Row(ic);
 606       Int_t ih = ID2IHit(ic);
 607       //cout<<"write row,hit="<<ID2IRow(ic)<<" "<<ih<<endl;
 608       fOutTrackHits[fNOutTrackHits] = fHitsID[row.FirstHit()+ih];
 609       fNOutTrackHits++;
 610       //cout<<"ok"<<endl;
 611       if( fNOutTrackHits>fNHitsTotal*10 ){
 612         cout<<"fNOutTrackHits>fNHitsTotal"<<endl;
 613         exit(0);//SG!!!
 614         return;
 615       }
 616       out.NHits()++;
 617     }
 618     //cout<<13<<endl;
 619     //cout<<fNOutTracks<<": itr = "<<iTr<<", n outhits = "<<out.NHits()<<endl;
 620     if( out.NHits() >= 2 ){
 621       fNOutTracks++;
 622     }else {
 623       fNOutTrackHits = fNOutTrackHitsOld;
 624     }
 625   }
 626   timer.Stop();
 627   fTimers[5]+=timer.CpuTime();
 628 }
 629
 630 GPUh() void AliHLTTPCCATracker::FitTrackFull( AliHLTTPCCATrack &/**/, Float_t * /**/ ) const
 631 {
 632   // fit track with material
 633 #ifdef XXX
 634   //* Fit the track
 635   FitTrack( iTrack, tt0 );
 636   if( iTrack.NHits()<=3 ) return;
 637
 638   AliHLTTPCCATrackParam &t = iTrack.Param();
 639   AliHLTTPCCATrackParam t0 = t;
 640
 641   t.Chi2() = 0;
 642   t.NDF() = -5;
 643   Bool_t first = 1;
 644
 645   Int_t iID = iTrack.FirstHitID();
 646   for( Int_t ih=0; ih<iTrack.NHits(); ih++, iID++ ){
 647     Int_t *ic = &(fTrackHits[iID]);
 648     Int_t iRow = ID2IRow(*ic);
 649     AliHLTTPCCARow &row = fRows[iRow];
 650     if( !t0.TransportToX( row.X() ) ) continue;
 651     Float_t dy, dz;
 652     AliHLTTPCCAHit &h = ID2Hit(*ic);
 653
 654     // check for wrong hits
 655     if(0){
 656       dy = t0.GetY() - h.Y();
 657       dz = t0.GetZ() - h.Z();
 658
 659       //if( dy*dy > 3.5*3.5*(/*t0.GetErr2Y() + */h.ErrY()*h.ErrY() ) ) continue;//SG!!!
 660       //if( dz*dz > 3.5*3.5*(/*t0.GetErr2Z() + */h.ErrZ()*h.ErrZ() ) ) continue;
 661     }
 662
 663     if( !t.TransportToX( row.X() ) ) continue;
 664
 665     //* Update the track
 666
 667     if( first ){
 668       t.Cov()[ 0] = .5*.5;
 669       t.Cov()[ 1] = 0;
 670       t.Cov()[ 2] = .5*.5;
 671       t.Cov()[ 3] = 0;
 672       t.Cov()[ 4] = 0;
 673       t.Cov()[ 5] = .2*.2;
 674       t.Cov()[ 6] = 0;
 675       t.Cov()[ 7] = 0;
 676       t.Cov()[ 8] = 0;
 677       t.Cov()[ 9] = .2*.2;
 678       t.Cov()[10] = 0;
 679       t.Cov()[11] = 0;
 680       t.Cov()[12] = 0;
 681       t.Cov()[13] = 0;
 682       t.Cov()[14] = .2*.2;
 683       t.Chi2() = 0;
 684       t.NDF() = -5;
 685     }
 686     Float_t err2Y, err2Z;
 687     GetErrors2( iRow, t, err2Y, err2Z );
 688
 689     if( !t.Filter2( h.Y(), h.Z(), err2Y, err2Z ) ) continue;
 690
 691     first = 0;
 692   }
 693   /*
 694   Float_t cosPhi = iTrack.Param().GetCosPhi();
 695   p0.Param().TransportToX(ID2Row( iTrack.PointID()[0] ).X());
 696   p2.Param().TransportToX(ID2Row( iTrack.PointID()[1] ).X());
 697   if( p0.Param().GetCosPhi()*cosPhi<0 ){ // change direction
 698   Float_t *par = p0.Param().Par();
 699   Float_t *cov = p0.Param().Cov();
 700   par[2] = -par[2]; // sin phi
 701   par[3] = -par[3]; // DzDs
 702     par[4] = -par[4]; // kappa
 703     cov[3] = -cov[3];
 704     cov[4] = -cov[4];
 705     cov[6] = -cov[6];
 706     cov[7] = -cov[7];
 707     cov[10] = -cov[10];
 708     cov[11] = -cov[11];
 709     p0.Param().CosPhi() = -p0.Param().GetCosPhi();
 710   }
 711   */
 712 #endif
 713 }
 714
 715 GPUh() void AliHLTTPCCATracker::FitTrack( AliHLTTPCCATrack &/*track*/, Float_t */*t0[]*/ ) const
 716 {
 717   //* Fit the track
 718 #ifdef XXX
 719   AliHLTTPCCAEndPoint &p2 = ID2Point(track.PointID()[1]);
 720   AliHLTTPCCAHit &c0 = ID2Hit(fTrackHits[p0.TrackHitID()].HitID());
 721   AliHLTTPCCAHit &c1 = ID2Hit(fTrackHits[track.HitID()[1]].HitID());
 722   AliHLTTPCCAHit &c2 = ID2Hit(fTrackHits[p2.TrackHitID()].HitID());
 723   AliHLTTPCCARow &row0 = ID2Row(fTrackHits[p0.TrackHitID()].HitID());
 724   AliHLTTPCCARow &row1 = ID2Row(fTrackHits[track.HitID()[1]].HitID());
 725   AliHLTTPCCARow &row2 = ID2Row(fTrackHits[p2.TrackHitID()].HitID());
 726   Float_t sp0[5] = {row0.X(), c0.Y(), c0.Z(), c0.ErrY(), c0.ErrZ() };
 727   Float_t sp1[5] = {row1.X(), c1.Y(), c1.Z(), c1.ErrY(), c1.ErrZ() };
 728   Float_t sp2[5] = {row2.X(), c2.Y(), c2.Z(), c2.ErrY(), c2.ErrZ() };
 729   //cout<<"Fit track, points ="<<sp0[0]<<" "<<sp0[1]<<" / "<<sp1[0]<<" "<<sp1[1]<<" / "<<sp2[0]<<" "<<sp2[1]<<endl;
 730   if( track.NHits()>=3 ){
 731     p0.Param().ConstructXYZ3(sp0,sp1,sp2,p0.Param().CosPhi(), t0);
 732     p2.Param().ConstructXYZ3(sp2,sp1,sp0,p2.Param().CosPhi(), t0);
 733     //p2.Param() = p0.Param();
 734     //p2.Param().TransportToX(row2.X());
 735     //p2.Param().Par()[1] = -p2.Param().Par()[1];
 736     //p2.Param().Par()[4] = -p2.Param().Par()[4];
 737   } else {
 738     p0.Param().X() = row0.X();
 739     p0.Param().Y() = c0.Y();
 740     p0.Param().Z() = c0.Z();
 741     p0.Param().Err2Y() = c0.ErrY()*c0.ErrY();
 742     p0.Param().Err2Z() = c0.ErrZ()*c0.ErrZ();
 743     p2.Param().X() = row2.X();
 744     p2.Param().Y() = c2.Y();
 745     p2.Param().Z() = c2.Z();
 746     p2.Param().Err2Y() = c2.ErrY()*c2.ErrY();
 747     p2.Param().Err2Z() = c2.ErrZ()*c2.ErrZ();
 748   }
 749 #endif
 750 }
 751
 752
 753
 754 GPUd() void AliHLTTPCCATracker::GetErrors2( Int_t iRow, const AliHLTTPCCATrackParam &t, Float_t &Err2Y, Float_t &Err2Z ) const
 755 {
 756   //
 757   // Use calibrated cluster error from OCDB
 758   //
 759
 760   Float_t z = CAMath::Abs((250.-0.275)-CAMath::Abs(t.GetZ()));
 761   Int_t    type = (iRow<63) ? 0: (iRow>126) ? 1:2;
 762   Float_t cosPhiInv = CAMath::Abs(t.GetCosPhi())>1.e-2 ?1./t.GetCosPhi() :0;
 763   Float_t angleY = t.GetSinPhi()*cosPhiInv ;
 764   Float_t angleZ = t.GetDzDs()*cosPhiInv ;
 765
 766   Err2Y = fParam.GetClusterError2(0,type, z,angleY);
 767   Err2Z = fParam.GetClusterError2(1,type, z,angleZ);
 768 }
 769
 770 GPUd() void AliHLTTPCCATracker::GetErrors2( Int_t iRow, const AliHLTTPCCATrackParam1 &t, Float_t &Err2Y, Float_t &Err2Z ) const
 771 {
 772   //
 773   // Use calibrated cluster error from OCDB
 774   //
 775
 776   Float_t z = CAMath::Abs((250.-0.275)-CAMath::Abs(t.GetZ()));
 777   Int_t    type = (iRow<63) ? 0: (iRow>126) ? 1:2;
 778   Float_t cosPhiInv = CAMath::Abs(t.GetCosPhi())>1.e-2 ?1./t.GetCosPhi() :0;
 779   Float_t angleY = t.GetSinPhi()*cosPhiInv ;
 780   Float_t angleZ = t.GetDzDs()*cosPhiInv ;
 781
 782   Err2Y = fParam.GetClusterError2(0,type, z,angleY);
 783   Err2Z = fParam.GetClusterError2(1,type, z,angleZ);
 784 }