From 00d07bcdb04ce4d1edd80de14119e146e662205b Mon Sep 17 00:00:00 2001 From: richterm Date: Tue, 13 Jan 2009 14:37:57 +0000 Subject: [PATCH] Completely reworked version of TPC CA tracker (Sergey) !!! Note: the ca tracking has been disabled for the moment due to various compilation problems. Will be corrected soon. --- HLT/TPCLib/AliHLTTPCAgent.cxx | 4 +- HLT/TPCLib/tracking-ca/AliHLTTPCCACell.h | 65 - HLT/TPCLib/tracking-ca/AliHLTTPCCADef.h | 127 ++ HLT/TPCLib/tracking-ca/AliHLTTPCCADisplay.cxx | 289 ++- HLT/TPCLib/tracking-ca/AliHLTTPCCADisplay.h | 34 +- HLT/TPCLib/tracking-ca/AliHLTTPCCAEndPoint.h | 48 - HLT/TPCLib/tracking-ca/AliHLTTPCCAGBHit.h | 9 +- HLT/TPCLib/tracking-ca/AliHLTTPCCAGBTrack.h | 4 +- .../tracking-ca/AliHLTTPCCAGBTracker.cxx | 225 ++- HLT/TPCLib/tracking-ca/AliHLTTPCCAGBTracker.h | 17 +- HLT/TPCLib/tracking-ca/AliHLTTPCCAGrid.cxx | 84 +- HLT/TPCLib/tracking-ca/AliHLTTPCCAGrid.h | 70 +- HLT/TPCLib/tracking-ca/AliHLTTPCCAHit.h | 42 +- HLT/TPCLib/tracking-ca/AliHLTTPCCAHitArea.cxx | 111 ++ HLT/TPCLib/tracking-ca/AliHLTTPCCAHitArea.h | 55 + ...ndPoint.cxx => AliHLTTPCCALinksWriter.cxx} | 27 +- .../tracking-ca/AliHLTTPCCALinksWriter.h | 31 + HLT/TPCLib/tracking-ca/AliHLTTPCCAMCPoint.h | 2 +- HLT/TPCLib/tracking-ca/AliHLTTPCCAMCTrack.cxx | 8 +- HLT/TPCLib/tracking-ca/AliHLTTPCCAMCTrack.h | 2 +- HLT/TPCLib/tracking-ca/AliHLTTPCCAMath.h | 223 +++ .../AliHLTTPCCANeighboursCleaner.cxx | 70 + .../AliHLTTPCCANeighboursCleaner.h | 51 + .../AliHLTTPCCANeighboursFinder.cxx | 150 ++ .../tracking-ca/AliHLTTPCCANeighboursFinder.h | 69 + ...l.cxx => AliHLTTPCCANeighboursFinder1.cxx} | 23 +- .../AliHLTTPCCANeighboursFinder1.h | 34 + .../tracking-ca/AliHLTTPCCAOutTrack.cxx | 2 +- HLT/TPCLib/tracking-ca/AliHLTTPCCAOutTrack.h | 13 +- HLT/TPCLib/tracking-ca/AliHLTTPCCAParam.cxx | 41 +- HLT/TPCLib/tracking-ca/AliHLTTPCCAParam.h | 87 +- .../tracking-ca/AliHLTTPCCAPerformance.cxx | 591 +++++-- .../tracking-ca/AliHLTTPCCAPerformance.h | 34 +- HLT/TPCLib/tracking-ca/AliHLTTPCCAProcess.h | 104 ++ HLT/TPCLib/tracking-ca/AliHLTTPCCARow.cxx | 39 +- HLT/TPCLib/tracking-ca/AliHLTTPCCARow.h | 75 +- .../AliHLTTPCCAStartHitsFinder.cxx | 72 + .../tracking-ca/AliHLTTPCCAStartHitsFinder.h | 53 + HLT/TPCLib/tracking-ca/AliHLTTPCCATrack.cxx | 7 +- HLT/TPCLib/tracking-ca/AliHLTTPCCATrack.h | 37 +- .../tracking-ca/AliHLTTPCCATrackConvertor.cxx | 8 +- .../tracking-ca/AliHLTTPCCATrackConvertor.h | 3 +- .../tracking-ca/AliHLTTPCCATrackParam.cxx | 274 +-- .../tracking-ca/AliHLTTPCCATrackParam.h | 113 +- .../tracking-ca/AliHLTTPCCATrackParam1.cxx | 949 ++++++++++ .../tracking-ca/AliHLTTPCCATrackParam1.h | 117 ++ HLT/TPCLib/tracking-ca/AliHLTTPCCATracker.cxx | 1545 +++++++---------- HLT/TPCLib/tracking-ca/AliHLTTPCCATracker.h | 132 +- .../AliHLTTPCCATrackerComponent.cxx | 98 +- .../tracking-ca/AliHLTTPCCATrackerComponent.h | 14 +- .../AliHLTTPCCATrackletConstructor.cxx | 472 +++++ .../AliHLTTPCCATrackletConstructor.h | 113 ++ .../AliHLTTPCCATrackletSelector.cxx | 80 + .../tracking-ca/AliHLTTPCCATrackletSelector.h | 38 + ...cxx => AliHLTTPCCAUsedHitsInitialiser.cxx} | 26 +- .../AliHLTTPCCAUsedHitsInitialiser.h | 51 + HLT/TPCLib/tracking-ca/AliTPCtrackerCA.cxx | 94 +- HLT/TPCLib/tracking-ca/AliTPCtrackerCA.h | 4 +- HLT/libAliHLTTPC.pkg | 50 +- 59 files changed, 5032 insertions(+), 2178 deletions(-) delete mode 100644 HLT/TPCLib/tracking-ca/AliHLTTPCCACell.h create mode 100644 HLT/TPCLib/tracking-ca/AliHLTTPCCADef.h delete mode 100644 HLT/TPCLib/tracking-ca/AliHLTTPCCAEndPoint.h create mode 100644 HLT/TPCLib/tracking-ca/AliHLTTPCCAHitArea.cxx create mode 100644 HLT/TPCLib/tracking-ca/AliHLTTPCCAHitArea.h rename HLT/TPCLib/tracking-ca/{AliHLTTPCCAEndPoint.cxx => AliHLTTPCCALinksWriter.cxx} (57%) create mode 100644 HLT/TPCLib/tracking-ca/AliHLTTPCCALinksWriter.h create mode 100644 HLT/TPCLib/tracking-ca/AliHLTTPCCAMath.h create mode 100644 HLT/TPCLib/tracking-ca/AliHLTTPCCANeighboursCleaner.cxx create mode 100644 HLT/TPCLib/tracking-ca/AliHLTTPCCANeighboursCleaner.h create mode 100644 HLT/TPCLib/tracking-ca/AliHLTTPCCANeighboursFinder.cxx create mode 100644 HLT/TPCLib/tracking-ca/AliHLTTPCCANeighboursFinder.h rename HLT/TPCLib/tracking-ca/{AliHLTTPCCACell.cxx => AliHLTTPCCANeighboursFinder1.cxx} (64%) create mode 100644 HLT/TPCLib/tracking-ca/AliHLTTPCCANeighboursFinder1.h create mode 100644 HLT/TPCLib/tracking-ca/AliHLTTPCCAProcess.h create mode 100644 HLT/TPCLib/tracking-ca/AliHLTTPCCAStartHitsFinder.cxx create mode 100644 HLT/TPCLib/tracking-ca/AliHLTTPCCAStartHitsFinder.h create mode 100644 HLT/TPCLib/tracking-ca/AliHLTTPCCATrackParam1.cxx create mode 100644 HLT/TPCLib/tracking-ca/AliHLTTPCCATrackParam1.h create mode 100644 HLT/TPCLib/tracking-ca/AliHLTTPCCATrackletConstructor.cxx create mode 100644 HLT/TPCLib/tracking-ca/AliHLTTPCCATrackletConstructor.h create mode 100644 HLT/TPCLib/tracking-ca/AliHLTTPCCATrackletSelector.cxx create mode 100644 HLT/TPCLib/tracking-ca/AliHLTTPCCATrackletSelector.h rename HLT/TPCLib/tracking-ca/{AliHLTTPCCAHit.cxx => AliHLTTPCCAUsedHitsInitialiser.cxx} (63%) create mode 100644 HLT/TPCLib/tracking-ca/AliHLTTPCCAUsedHitsInitialiser.h diff --git a/HLT/TPCLib/AliHLTTPCAgent.cxx b/HLT/TPCLib/AliHLTTPCAgent.cxx index a84d29f4aa1..f977f627649 100644 --- a/HLT/TPCLib/AliHLTTPCAgent.cxx +++ b/HLT/TPCLib/AliHLTTPCAgent.cxx @@ -46,7 +46,7 @@ AliHLTTPCAgent gAliHLTTPCAgent; #include "AliHLTTPCCalibCEComponent.h" #include "AliHLTTPCCalibPulserComponent.h" #include "AliHLTTPCCalibPedestalComponent.h" -#include "AliHLTTPCCATrackerComponent.h" +//#include "AliHLTTPCCATrackerComponent.h" #include "AliHLTTPCGlobalMergerComponent.h" #include "AliHLTTPCSliceTrackerComponent.h" #include "AliHLTTPCVertexFinderComponent.h" @@ -242,7 +242,7 @@ int AliHLTTPCAgent::RegisterComponents(AliHLTComponentHandler* pHandler) const pHandler->AddComponent(new AliHLTTPCCompModelDeconverterComponent); pHandler->AddComponent(new AliHLTTPCCompModelConverterComponent); pHandler->AddComponent(new AliHLTTPCCompDumpComponent); - pHandler->AddComponent(new AliHLTTPCCATrackerComponent); + // pHandler->AddComponent(new AliHLTTPCCATrackerComponent); pHandler->AddComponent(new AliHLTTPCGlobalMergerComponent); pHandler->AddComponent(new AliHLTTPCSliceTrackerComponent); pHandler->AddComponent(new AliHLTTPCVertexFinderComponent); diff --git a/HLT/TPCLib/tracking-ca/AliHLTTPCCACell.h b/HLT/TPCLib/tracking-ca/AliHLTTPCCACell.h deleted file mode 100644 index 0f5bfe6f690..00000000000 --- a/HLT/TPCLib/tracking-ca/AliHLTTPCCACell.h +++ /dev/null @@ -1,65 +0,0 @@ -//-*- Mode: C++ -*- -// @(#) $Id$ - -//* This file is property of and copyright by the ALICE HLT Project * -//* ALICE Experiment at CERN, All rights reserved. * -//* See cxx source for full Copyright notice * - -#ifndef ALIHLTTPCCACELL_H -#define ALIHLTTPCCACELL_H - - -#include "Rtypes.h" - -/** - * @class AliHLTTPCCACell - * - * The AliHLTTPCCACell class describes the "Cell" object --- - * the set of neghbouring clusters in the same TPC row. - * Cells are used as the minimal data units - * by the Cellular Automaton tracking algorithm. - * - */ -class AliHLTTPCCACell -{ - public: - - //AliHLTTPCCACell(): fY(0),fZ(0),fErrY(0),fErrZ(0),fFirstHitRef(0),fNHits(0),fLink(0),fTrackID(0){} - - //virtual ~AliHLTTPCCACell(){} - - Float_t &Y(){ return fY; } - Float_t &Z(){ return fZ; } - Float_t &ErrY(){ return fErrY; } - Float_t &ErrZ(){ return fErrZ; } - Float_t &ZMin(){ return fZMin; } - Float_t &ZMax(){ return fZMax; } - - Int_t &FirstHitRef(){ return fFirstHitRef; } - Int_t &NHits() { return fNHits; } - - Int_t &Link() { return fLink; } - Int_t &Status() { return fStatus; } - Int_t &TrackID() { return fTrackID; } - - protected: - - Float_t fY, fZ, fZMin,fZMax; //* Y and Z coordinates - Float_t fErrY, fErrZ; //* cell errors in Y and Z - - Int_t fFirstHitRef; //* index of the first cell hit in the cell->hit reference array - Int_t fNHits; //* number of hits in the cell - - Int_t fLink; //* link to the next cell on the track - Int_t fStatus; //* status flag - Int_t fTrackID; //* index of track - - private: - - void Dummy(); // to make rulechecker happy by having something in .cxx file - - //ClassDef(AliHLTTPCCACell,1); -}; - - -#endif diff --git a/HLT/TPCLib/tracking-ca/AliHLTTPCCADef.h b/HLT/TPCLib/tracking-ca/AliHLTTPCCADef.h new file mode 100644 index 00000000000..23010b9ab81 --- /dev/null +++ b/HLT/TPCLib/tracking-ca/AliHLTTPCCADef.h @@ -0,0 +1,127 @@ +//-*- Mode: C++ -*- + +//* This file is property of and copyright by the ALICE HLT Project * +//* ALICE Experiment at CERN, All rights reserved. * +//* See cxx source for full Copyright notice * + +#ifndef ALIHLTTPCCADEF_H +#define ALIHLTTPCCADEF_H + + +/** + * Definitions needed for AliHLTTPCCATracker + * + */ + +//#define HLTCA_STANDALONE // compilation w/o root + +#ifdef __CUDACC__ + +#define HLTCA_GPUCODE + +#endif + +#if defined(HLTCA_STANDALONE) || defined(HLTCA_GPUCODE) + +class TObject{}; + +#define ClassDef(name,id) +#define ClassImp(name) + +typedef char Char_t; //Signed Character 1 byte (char) +typedef unsigned char UChar_t; //Unsigned Character 1 byte (unsigned char) +typedef short Short_t; //Signed Short integer 2 bytes (short) +typedef unsigned short UShort_t; //Unsigned Short integer 2 bytes (unsigned short) +#ifdef R__INT16 +typedef long Int_t; //Signed integer 4 bytes +typedef unsigned long UInt_t; //Unsigned integer 4 bytes +#else +typedef int Int_t; //Signed integer 4 bytes (int) +typedef unsigned int UInt_t; //Unsigned integer 4 bytes (unsigned int) +#endif +#ifdef R__B64 // Note: Long_t and ULong_t are currently not portable types +typedef int Seek_t; //File pointer (int) +typedef long Long_t; //Signed long integer 8 bytes (long) +typedef unsigned long ULong_t; //Unsigned long integer 8 bytes (unsigned long) +#else +typedef int Seek_t; //File pointer (int) +typedef long Long_t; //Signed long integer 4 bytes (long) +typedef unsigned long ULong_t; //Unsigned long integer 4 bytes (unsigned long) +#endif +typedef float Float_t; //Float 4 bytes (float) +typedef float Float16_t; //Float 4 bytes written with a truncated mantissa +typedef double Double_t; //Double 8 bytes +typedef double Double32_t; //Double 8 bytes in memory, written as a 4 bytes float +typedef char Text_t; //General string (char) +typedef bool Bool_t; //Boolean (0=false, 1=true) (bool) +typedef unsigned char Byte_t; //Byte (8 bits) (unsigned char) +typedef short Version_t; //Class version identifier (short) +typedef const char Option_t; //Option string (const char) +typedef int Ssiz_t; //String size (int) +typedef float Real_t; //TVector and TMatrix element type (float) +#if defined(R__WIN32) && !defined(__CINT__) +typedef __int64 Long64_t; //Portable signed long integer 8 bytes +typedef unsigned __int64 ULong64_t; //Portable unsigned long integer 8 bytes +#else +typedef long long Long64_t; //Portable signed long integer 8 bytes +typedef unsigned long long ULong64_t;//Portable unsigned long integer 8 bytes +#endif +typedef double Axis_t; //Axis values type (double) +typedef double Stat_t; //Statistics type (double) +typedef short Font_t; //Font number (short) +typedef short Style_t; //Style number (short) +typedef short Marker_t; //Marker number (short) +typedef short Width_t; //Line width (short) +typedef short Color_t; //Color number (short) +typedef short SCoord_t; //Screen coordinates (short) +typedef double Coord_t; //Pad world coordinates (double) +typedef float Angle_t; //Graphics angle (float) +typedef float Size_t; //Attribute size (float) + +#else + +#include "Rtypes.h" +#include "TObject.h" + +#endif + + + +#ifdef HLTCA_GPUCODE + +#define GPUd() __device__ +#define GPUhd() __host__ __device__ +#define GPUh() __host__ inline +#define GPUg() __global__ + +#define GPUshared() __shared__ +#define GPUsync() __syncthreads() + +__constant__ float4 cTracker[30000/sizeof(float4)]; +texture texGrid; +texture texHits; + +#else + +#define GPUd() +#define GPUhd() +#define GPUg() +#define GPUh() +#define GPUshared() +#define GPUsync() + +struct float2{ float x; float y; }; +struct uchar2{ unsigned char x; unsigned char y; }; +struct ushort2{ unsigned short x; unsigned short y; }; +struct uint1{ unsigned int x; }; +struct uint4{ unsigned int x,y,z,w; }; + +float2 inline make_float2(float x, float y ){ + float2 ret = {x,y}; + return ret; +} + +#endif + + +#endif diff --git a/HLT/TPCLib/tracking-ca/AliHLTTPCCADisplay.cxx b/HLT/TPCLib/tracking-ca/AliHLTTPCCADisplay.cxx index 76b1f2110e1..f12d509cf3e 100644 --- a/HLT/TPCLib/tracking-ca/AliHLTTPCCADisplay.cxx +++ b/HLT/TPCLib/tracking-ca/AliHLTTPCCADisplay.cxx @@ -16,20 +16,21 @@ // provided "as is" without express or implied warranty. * //*************************************************************************** -#include "AliHLTTPCCADisplay.h" -#include "AliHLTTPCCATracker.h" -#include "AliHLTTPCCAEndPoint.h" -#include "AliHLTTPCCARow.h" -#include "AliHLTTPCCATrack.h" +//#include "AliHLTTPCCADisplay.h" + +#ifdef XXXX + +//#include "AliHLTTPCCATracker.h" +//#include "AliHLTTPCCARow.h" +//#include "AliHLTTPCCATrack.h" //#include "TString.h" -#include "Riostream.h" -#include "TMath.h" -#include "TStyle.h" -#include "TCanvas.h" +//#include "Riostream.h" +//#include "TMath.h" +//#include "TStyle.h" +//#include "TCanvas.h" -ClassImp(AliHLTTPCCADisplay) AliHLTTPCCADisplay &AliHLTTPCCADisplay::Instance() { @@ -147,7 +148,9 @@ void AliHLTTPCCADisplay::SetCurrentSlice( AliHLTTPCCATracker *slice ) fYX->Range(cx-dr, cy-dr*1.05, cx+dr, cy+dr); fZX->Range(cz-dz, cy-dr*1.05, cz+dz, cy+dr); - //fYX->Range(cx-dr/6, cy-dr, cx+dr/8, cy-dr + dr/8); + fYX->Range(cx, cy-dr*0.8, cx+dr/2, cy+dr*0.1); + fZX->Range(cz, cy-dr*.8, cz+dz, cy+dr*0.1); + //fYX->Range(cx-dr/6, cy-dr, cx+dr/8, cy-dr + dr/8); //fZX->Range(cz+dz/2+dz/6, cy-dr , cz+dz-dz/8, cy-dr + dr/8); //fYX->Range(cx-dr/3, cy-dr/3, cx+dr, cy+dr); @@ -295,9 +298,10 @@ void AliHLTTPCCADisplay::DrawHit( Int_t iRow, Int_t iHit, Int_t color ) //Double_t dgy = 3.5*TMath::Abs(h->ErrY()*fSlice->Param().CosAlpha() - fSlice->Param().ErrX()*fSlice->Param().SinAlpha() ); Double_t dx = 0.1;//fSlice->Param().ErrX()*TMath::Sqrt(12.)/2.; - Double_t dy = h->ErrY()*3.5; + Double_t dy = 0.35;//h->ErrY()*3.5; //Double_t dz = h->ErrZ()*3.5; - fMarker.SetMarkerSize(.3); + fMarker.SetMarkerSize(.5); + //fMarker.SetMarkerSize(.3); fMarker.SetMarkerColor(color); fArc.SetLineColor(color); fArc.SetFillStyle(0); @@ -315,173 +319,108 @@ void AliHLTTPCCADisplay::DrawHit( Int_t iRow, Int_t iHit, Int_t color ) fMarker.DrawMarker(h->Z(), vy); } -void AliHLTTPCCADisplay::DrawCell( Int_t iRow, AliHLTTPCCACell &cell, Int_t width, Int_t color ) -{ - // draw cell - AliHLTTPCCARow &row = fSlice->Rows()[iRow]; - Double_t vx, vy, vdx, vdy, vz = cell.Z(), vdz = cell.ErrZ()*3.5; - Slice2View(row.X(), cell.Y(), &vx, &vy); - Slice2View(0.2, cell.ErrY()*3.5, &vdx, &vdy); - if( color<0 ) color = GetColor(cell.Z()); - fLine.SetLineColor(color); - fLine.SetLineWidth(width); - fArc.SetLineColor(color); - fArc.SetFillStyle(0); - fYX->cd(); - //fLine.DrawLine(vx-vdx,vy-vdy, vx+vdx, vy+vdy ); - fArc.DrawEllipse(vx, vy, vdx, vdy, 0,360, 0); - fZX->cd(); - //fLine.DrawLine(cell.Z()-3*cell.ErrZ(),vy-vdy, cell.Z()+3*cell.ErrZ(), vy+vdy ); - fArc.DrawEllipse(vz, vy, vdz, vdy, 0,360, 0); - fLine.SetLineWidth(1); -} -void AliHLTTPCCADisplay::DrawCell( Int_t iRow, Int_t iCell, Int_t width, Int_t color ) -{ - // draw cell - AliHLTTPCCARow &row = fSlice->Rows()[iRow]; - DrawCell( iRow, row.Cells()[iCell], width, color ); -} -void AliHLTTPCCADisplay::DrawEndPoint( Int_t ID, Float_t R, Int_t width, Int_t color) -{ - // draw endpoint - if( !fSlice ) return; - AliHLTTPCCARow &row = fSlice->ID2Row(ID); - AliHLTTPCCAEndPoint &p = fSlice->ID2Point(ID); - AliHLTTPCCACell &c = fSlice->ID2Cell(p.CellID()); - if( color<0 ) color = GetColor( c.Z() ); - fArc.SetLineColor(color); - fArc.SetFillStyle(0); - fArc.SetLineWidth(width); + - Double_t vx, vy; - Slice2View( row.X(), c.Y(), &vx, &vy ); - fYX->cd(); - fArc.DrawEllipse( vx, vy, R, R, 0,360, 90); - fZX->cd(); - fArc.DrawEllipse( c.Z(), vy, R, R, 0,360, 90); - fArc.SetLineWidth(1); -} - -void AliHLTTPCCADisplay::ConnectEndPoints( Int_t iID, Int_t jID, Float_t R, Int_t width, Int_t color ) + +void AliHLTTPCCADisplay::DrawMergedHit( Int_t iRow, Int_t iHit, Int_t color ) { - // connect endpoints - if( !fSlice ) return; - AliHLTTPCCARow &irow = fSlice->ID2Row(iID); - AliHLTTPCCAEndPoint &ip = fSlice->ID2Point(iID); - AliHLTTPCCACell &ic = fSlice->ID2Cell(ip.CellID()); - AliHLTTPCCARow &jrow = fSlice->ID2Row(jID); - AliHLTTPCCAEndPoint &jp = fSlice->ID2Point(jID); - AliHLTTPCCACell &jc = fSlice->ID2Cell(jp.CellID()); - if( color<0 ) color = GetColor( ic.Z() ); +#ifdef XXX + // connect two cells on display, kind of row is drawing - fArc.SetLineColor(color); - fArc.SetFillStyle(0); - fArc.SetLineWidth(width); - fLine.SetLineWidth(width); - fLine.SetLineColor(color); - Double_t ivx, ivy; - Slice2View( irow.X(), ic.Y(), &ivx, &ivy ); - Double_t jvx, jvy; - Slice2View( jrow.X(), jc.Y(), &jvx, &jvy ); + AliHLTTPCCARow &row = fSlice->Rows()[iRow]; + AliHLTTPCCAHit &h = row.Hits()[iHit]; + AliHLTTPCCAHit &hyz = row.HitsYZ()[iHit]; + + Double_t x = row.X(); + Double_t y = hyz.Y(); + Double_t z = hyz.Z(); + Double_t x1 = x, x2 = x; + Double_t y1 = y, y2 = y; + Double_t z1 = z, z2 = z; + Int_t iRow1 = iRow, iHit1 = iHit; + Int_t iRow2 = iRow, iHit2 = iHit; + + if( fSlice->HitLinksDown()[]>=0 ){ + iRow1 = iRow - 1; + iHit1 = h.LinkDown(); + AliHLTTPCCARow &row1 = fSlice->Rows()[iRow1]; + AliHLTTPCCAHitYZ &h1 = row1.HitsYZ()[iHit1]; + x1 = row1.X(); + y1 = h1.Y(); + z1 = h1.Z(); + } + if( h.LinkUp()>=0 ){ + iRow2 = iRow+1; + iHit2 = h.LinkUp(); + AliHLTTPCCARow &row2 = fSlice->Rows()[iRow2]; + AliHLTTPCCAHitYZ &h2 = row2.HitsYZ()[iHit2]; + x2 = row2.X(); + y2 = h2.Y(); + z2 = h2.Z(); + } + if( color<0 ) color = GetColor( (z+z1+z2)/3. ); - fYX->cd(); - fArc.DrawEllipse( ivx, ivy, R, R, 0,360, 90); - fArc.DrawEllipse( jvx, jvy, R, R, 0,360, 90); - fLine.DrawLine(ivx, ivy,jvx, jvy); - fZX->cd(); - fArc.DrawEllipse( ic.Z(), ivy, R, R, 0,360, 90); - fArc.DrawEllipse( jc.Z(), jvy, R, R, 0,360, 90); - fLine.DrawLine(ic.Z(), ivy, jc.Z(), jvy); - fArc.SetLineWidth(1); - fLine.SetLineWidth(1); -} -void AliHLTTPCCADisplay::ConnectCells( Int_t iRow1, AliHLTTPCCACell &cell1, - Int_t iRow2, AliHLTTPCCACell &cell2, Int_t color ) -{ - // connect two cells on display, kind of row is drawing - AliHLTTPCCARow &row1 = fSlice->Rows()[iRow1]; - AliHLTTPCCARow &row2 = fSlice->Rows()[iRow2]; - - AliHLTTPCCAHit &h11 = row1.GetCellHit(cell1,0); - AliHLTTPCCAHit &h12 = row1.GetCellHit(cell1,cell1.NHits()-1); - AliHLTTPCCAHit &h21 = row2.GetCellHit(cell2,0); - AliHLTTPCCAHit &h22= row2.GetCellHit(cell2,cell2.NHits()-1); - - Double_t x11 = row1.X(); - Double_t x12 = row1.X(); - Double_t y11 = h11.Y() - h11.ErrY()*3; - Double_t y12 = h12.Y() + h12.ErrY()*3; - Double_t z11 = h11.Z() - h11.ErrZ()*3; - Double_t z12 = h12.Z() + h12.ErrZ()*3; - Double_t x21 = row2.X(); - Double_t x22 = row2.X(); - Double_t y21 = h21.Y() - h21.ErrY()*3; - Double_t y22 = h22.Y() + h22.ErrY()*3; - Double_t z21 = h21.Z() - h21.ErrZ()*3; - Double_t z22 = h22.Z() + h22.ErrZ()*3; - - Double_t vx11, vx12, vy11, vy12, vx21, vx22, vy21, vy22; - - Slice2View(x11,y11, &vx11, &vy11 ); - Slice2View(x12,y12, &vx12, &vy12 ); - Slice2View(x21,y21, &vx21, &vy21 ); - Slice2View(x22,y22, &vx22, &vy22 ); - - Double_t lx[] = { vx11, vx12, vx22, vx21, vx11 }; - Double_t ly[] = { vy11, vy12, vy22, vy21, vy11 }; - Double_t lz[] = { z11, z12, z22, z21, z11 }; - - if( color<0 ) color = GetColor( (z11+z12+z22+z21)/4. ); + Slice2View(x,y, &x, &y ); + Slice2View(x1,y1, &x1, &y1 ); + Slice2View(x2,y2, &x2, &y2 ); + + Double_t lx[] = { x1, x, x2 }; + Double_t ly[] = { y1, y, y2 }; + Double_t lz[] = { z1, z, z2 }; + fPLine.SetLineColor(color); fPLine.SetLineWidth(1); //fPLine.SetFillColor(color); fPLine.SetFillStyle(-1); fYX->cd(); - fPLine.DrawPolyLine(5, lx, ly ); + fPLine.DrawPolyLine(3, lx, ly ); fZX->cd(); - fPLine.DrawPolyLine(5, lz, ly ); - DrawCell( iRow1, cell1, 1, color ); - DrawCell( iRow2, cell2, 1, color ); + fPLine.DrawPolyLine(3, lz, ly ); + DrawHit( iRow, iHit, color ); + DrawHit( iRow1, iHit1, color ); + DrawHit( iRow2, iHit2, color ); +#endif } - - -void AliHLTTPCCADisplay::DrawTrack( AliHLTTPCCATrack &track, Int_t color, Bool_t DrawCells ) +void AliHLTTPCCADisplay::DrawTrack( AliHLTTPCCATrack &track, Int_t color, Bool_t DrawHits ) { // draw track - - if( track.NCells()<2 ) return; - int width = 1; - - AliHLTTPCCADisplayTmpCell *vCells = new AliHLTTPCCADisplayTmpCell[track.NCells()]; - AliHLTTPCCATrackParam &t = fSlice->ID2Point(track.PointID()[0]).Param(); - - Int_t iID = track.FirstCellID(); - { - Int_t iCell=0; - while( iID>=0 ){ - AliHLTTPCCACell *c = &(fSlice->ID2Cell( iID )); - AliHLTTPCCARow &row = fSlice->ID2Row(iID); - vCells[iCell].ID() = iID; - vCells[iCell].S() = t.GetS( row.X(), c->Y() ); - vCells[iCell].Z() = c->Z(); - iCell++; - iID = c->Link(); + + if( track.NHits()<2 ) return; + int width = 2; + + AliHLTTPCCADisplayTmpHit *vHits = new AliHLTTPCCADisplayTmpHit[track.NHits()]; + AliHLTTPCCATrackParam &t = track.Param(); + + Int_t iID = track.FirstHitID(); + int nhits = 0; + { + Int_t iHit = 0; + for( int ih=0; ihTrackHits()[iID]; + AliHLTTPCCAHit *h = &(fSlice->ID2Hit( i )); + AliHLTTPCCARow &row = fSlice->ID2Row(i); + vHits[iHit].ID() = i; + vHits[iHit].S() = t.GetS( row.X(), h->Y() ); + vHits[iHit].Z() = h->Z(); + iHit++; + nhits++; + iID++; } } - sort(vCells, vCells + track.NCells(), AliHLTTPCCADisplayTmpCell::CompareCellZ ); - + sort(vHits, vHits + track.NHits(), AliHLTTPCCADisplayTmpHit::CompareHitZ ); + cout<<"Draw track, nhits = "<ID2Cell(vCells[0].ID()); - AliHLTTPCCACell &c2 = fSlice->ID2Cell(vCells[track.NCells()-1].ID()); + AliHLTTPCCAHit &c1 = fSlice->ID2Hit(vHits[0].ID()); + AliHLTTPCCAHit &c2 = fSlice->ID2Hit(vHits[track.NHits()-1].ID()); if( color<0 ) color = GetColor( (c1.Z()+c2.Z())/2. ); } @@ -489,8 +428,8 @@ void AliHLTTPCCADisplay::DrawTrack( AliHLTTPCCATrack &track, Int_t color, Bool_t fMarker.SetMarkerSize(1.); /* for( Int_t i=0; i<3; i++){ - AliHLTTPCCACell &c1 = fSlice->ID2Cell(track.CellID()[i]); - AliHLTTPCCARow &row1 = fSlice->ID2Row(track.CellID()[i]); + AliHLTTPCCAHit &c1 = fSlice->ID2Hit(track.HitID()[i]); + AliHLTTPCCARow &row1 = fSlice->ID2Row(track.HitID()[i]); Double_t vx1, vy1; Slice2View(row1.X(), c1.Y(), &vx1, &vy1 ); fYX->cd(); @@ -499,15 +438,16 @@ void AliHLTTPCCADisplay::DrawTrack( AliHLTTPCCATrack &track, Int_t color, Bool_t fMarker.DrawMarker(c1.Z(),vy1); } */ - DrawTrackletPoint( fSlice->ID2Point(track.PointID()[0]).Param(), kBlack);//color ); - DrawTrackletPoint( fSlice->ID2Point(track.PointID()[1]).Param(), kBlack);//color ); - - for( Int_t iCell=0; iCellID2Point(track.PointID()[0]).Param(), kBlack);//color ); + //DrawTrackletPoint( fSlice->ID2Point(track.PointID()[1]).Param(), kBlack);//color ); + //cout<<"DrawTrack end points x = "<ID2Point(track.PointID()[0]).Param().GetX()<<" "<ID2Point(track.PointID()[1]).Param().GetX()<ID2Cell(vCells[iCell].ID()); - AliHLTTPCCACell &c2 = fSlice->ID2Cell(vCells[iCell+1].ID()); - AliHLTTPCCARow &row1 = fSlice->ID2Row(vCells[iCell].ID()); - AliHLTTPCCARow &row2 = fSlice->ID2Row(vCells[iCell+1].ID()); + AliHLTTPCCAHit &c1 = fSlice->ID2Hit(vHits[iHit].ID()); + AliHLTTPCCAHit &c2 = fSlice->ID2Hit(vHits[iHit+1].ID()); + AliHLTTPCCARow &row1 = fSlice->ID2Row(vHits[iHit].ID()); + AliHLTTPCCARow &row2 = fSlice->ID2Row(vHits[iHit+1].ID()); Float_t x1, y1, z1, x2, y2, z2; t.GetDCAPoint( row1.X(), c1.Y(), c1.Z(), x1, y1, z1 ); t.GetDCAPoint( row2.X(), c2.Y(), c2.Z(), x2, y2, z2 ); @@ -563,14 +503,14 @@ void AliHLTTPCCADisplay::DrawTrack( AliHLTTPCCATrack &track, Int_t color, Bool_t } } - for( Int_t iCell=0; iCellID2Cell(vCells[iCell].ID()); - AliHLTTPCCACell &c2 = fSlice->ID2Cell(vCells[iCell+1].ID()); - AliHLTTPCCARow &row1 = fSlice->ID2Row(vCells[iCell].ID()); - AliHLTTPCCARow &row2 = fSlice->ID2Row(vCells[iCell+1].ID()); + for( Int_t iHit=0; iHitID2Hit(vHits[iHit].ID()); + AliHLTTPCCAHit &c2 = fSlice->ID2Hit(vHits[iHit+1].ID()); + AliHLTTPCCARow &row1 = fSlice->ID2Row(vHits[iHit].ID()); + AliHLTTPCCARow &row2 = fSlice->ID2Row(vHits[iHit+1].ID()); - if( DrawCells ) ConnectCells( fSlice->ID2IRow(vCells[iCell].ID()),c1, - fSlice->ID2IRow(vCells[iCell+1].ID()),c2, color ); + //if( DrawHits ) ConnectHits( fSlice->ID2IRow(vHits[iHit].ID()),c1, + //fSlice->ID2IRow(vHits[iHit+1].ID()),c2, color ); Float_t x1, y1, z1, x2, y2, z2; t.GetDCAPoint( row1.X(), c1.Y(), c1.Z(), x1, y1, z1 ); t.GetDCAPoint( row2.X(), c2.Y(), c2.Z(), x2, y2, z2 ); @@ -586,7 +526,7 @@ void AliHLTTPCCADisplay::DrawTrack( AliHLTTPCCATrack &track, Int_t color, Bool_t fLine.DrawLine(z1,vy1, z2, vy2 ); } fLine.SetLineWidth(1); - delete[] vCells; + delete[] vHits; } @@ -602,7 +542,7 @@ void AliHLTTPCCADisplay::DrawTrackletPoint( AliHLTTPCCATrackParam &t, Int_t colo Double_t ex = t.GetCosPhi(); Double_t ey = sinPhi; - int width = 3; + int width = 1; if( color<0 ) color = GetColor( t.GetZ() ); @@ -629,3 +569,4 @@ void AliHLTTPCCADisplay::DrawTrackletPoint( AliHLTTPCCATrackParam &t, Int_t colo fLine.DrawLine(z-3.5*dz,vy-vdy, z+3.5*dz, vy+vdy ); fLine.SetLineWidth(1); } +#endif //XXXX diff --git a/HLT/TPCLib/tracking-ca/AliHLTTPCCADisplay.h b/HLT/TPCLib/tracking-ca/AliHLTTPCCADisplay.h index 79eb0d41344..ac1df77d5b8 100644 --- a/HLT/TPCLib/tracking-ca/AliHLTTPCCADisplay.h +++ b/HLT/TPCLib/tracking-ca/AliHLTTPCCADisplay.h @@ -13,12 +13,12 @@ #ifndef ALIHLTTPCCADISPLAY_H #define ALIHLTTPCCADISPLAY_H + +#ifdef XXXX class AliHLTTPCCATracker; -class AliHLTTPCCACell; class AliHLTTPCCATrack; class AliHLTTPCCATrackParam; class TCanvas; -#include "TObject.h" #include "TArc.h" #include "TLine.h" #include "TPolyLine.h" @@ -26,15 +26,15 @@ class TCanvas; #include "TCrown.h" #include "TMarker.h" #include "TLatex.h" - +#endif /** * @class AliHLTTPCCADisplay */ -class AliHLTTPCCADisplay:public TObject +class AliHLTTPCCADisplay { public: - +#ifdef XXXX static AliHLTTPCCADisplay &Instance(); AliHLTTPCCADisplay(); @@ -60,13 +60,8 @@ class AliHLTTPCCADisplay:public TObject void DrawSlice( AliHLTTPCCATracker *slice ); void DrawHit( Int_t iRow,Int_t iHit, Int_t color=-1 ); - void DrawCell( Int_t iRow, AliHLTTPCCACell &cell, Int_t width=1, Int_t color=-1 ); - void DrawCell( Int_t iRow, Int_t iCell, Int_t width=1, Int_t color=-1 ); - - void DrawEndPoint( Int_t ID, Float_t R, Int_t width=1, Int_t color=-1 ); - void ConnectEndPoints( Int_t iID, Int_t jID, Float_t R, Int_t width=1, Int_t color=-1 ); - void ConnectCells( Int_t iRow1, AliHLTTPCCACell &cell1, Int_t iRow2, AliHLTTPCCACell &cell2, Int_t color=-1 ); + void DrawMergedHit( Int_t iRow, Int_t iHit, Int_t color=-1 ); void DrawTrack( AliHLTTPCCATrack &track, Int_t color=-1, Bool_t DrawCells=1 ); void DrawTrackletPoint( AliHLTTPCCATrackParam &t, Int_t color=-1 ); @@ -94,31 +89,30 @@ class AliHLTTPCCADisplay:public TObject TCrown fCrown; //! TLatex fLatex; //! - class AliHLTTPCCADisplayTmpCell{ + class AliHLTTPCCADisplayTmpHit{ public: - Int_t &ID(){ return fCellID; } + Int_t &ID(){ return fHitID; } Double_t &S(){ return fS; } Double_t &Z(){ return fZ; } - static Bool_t CompareCellDS( const AliHLTTPCCADisplayTmpCell &a, - const AliHLTTPCCADisplayTmpCell &b ) + static Bool_t CompareHitDS( const AliHLTTPCCADisplayTmpHit &a, + const AliHLTTPCCADisplayTmpHit &b ) { return (a.fS < b.fS); } - static Bool_t CompareCellZ( const AliHLTTPCCADisplayTmpCell &a, - const AliHLTTPCCADisplayTmpCell &b ) + static Bool_t CompareHitZ( const AliHLTTPCCADisplayTmpHit &a, + const AliHLTTPCCADisplayTmpHit &b ) { return (a.fZ < b.fZ); } protected: - Int_t fCellID; // cell ID + Int_t fHitID; // cell ID Double_t fS; // cell position on the XY track curve Double_t fZ; // cell Z position }; - - ClassDef(AliHLTTPCCADisplay,1); +#endif // XXXX }; diff --git a/HLT/TPCLib/tracking-ca/AliHLTTPCCAEndPoint.h b/HLT/TPCLib/tracking-ca/AliHLTTPCCAEndPoint.h deleted file mode 100644 index 7ed1ac665d4..00000000000 --- a/HLT/TPCLib/tracking-ca/AliHLTTPCCAEndPoint.h +++ /dev/null @@ -1,48 +0,0 @@ -//-*- Mode: C++ -*- -// $Id$ - -//* This file is property of and copyright by the ALICE HLT Project * -//* ALICE Experiment at CERN, All rights reserved. * -//* See cxx source for full Copyright notice * - -//* * -//* The AliHLTTPCCAEndPoint class describes the end point of the track * -//* it contains the fitted track parameters at this point * -//* The class is used for the matching of tracks withing one TPC slice * -//* * - -#ifndef ALIHLTTPCCAENDPOINT_H -#define ALIHLTTPCCAENDPOINT_H - - -#include "Rtypes.h" -#include "AliHLTTPCCATrackParam.h" - -/** - * @class AliHLTTPCCAEndPoint - */ -class AliHLTTPCCAEndPoint -{ - public: - - AliHLTTPCCAEndPoint() :fCellID(0),fTrackID(0),fLink(0),fParam(){} - Int_t &CellID() { return fCellID; } - Int_t &TrackID() { return fTrackID; } - Int_t &Link() { return fLink; } - - AliHLTTPCCATrackParam &Param(){ return fParam; }; - - protected: - - Int_t fCellID; //* index of the cell - Int_t fTrackID; //* index of the track - Int_t fLink; //* link to the neighbour (if found) - AliHLTTPCCATrackParam fParam; //* track parameters at the end point - -private: - void Dummy(); //* to make rulechecker happy by having something in .cxx file - -}; - - -#endif diff --git a/HLT/TPCLib/tracking-ca/AliHLTTPCCAGBHit.h b/HLT/TPCLib/tracking-ca/AliHLTTPCCAGBHit.h index aee3f0e0b92..864b03fa668 100644 --- a/HLT/TPCLib/tracking-ca/AliHLTTPCCAGBHit.h +++ b/HLT/TPCLib/tracking-ca/AliHLTTPCCAGBHit.h @@ -8,8 +8,7 @@ #ifndef ALIHLTTPCCAGBHIT_H #define ALIHLTTPCCAGBHIT_H -#include "Rtypes.h" -#include "AliHLTTPCCAHit.h" +#include "AliHLTTPCCADef.h" /** * @class AliHLTTPCCAGBHit @@ -23,7 +22,7 @@ class AliHLTTPCCAGBHit public: AliHLTTPCCAGBHit() :fX(0),fY(0),fZ(0),fErrX(0),fErrY(0),fErrZ(0),fAmp(0), - fISlice(0), fIRow(0), fID(0), fIsUsed(0),fSliceHit(){} + fISlice(0), fIRow(0), fID(0), fIsUsed(0){} virtual ~AliHLTTPCCAGBHit(){} @@ -41,8 +40,6 @@ class AliHLTTPCCAGBHit Int_t &ID(){ return fID; } Bool_t &IsUsed(){ return fIsUsed; }; - AliHLTTPCCAHit &SliceHit(){ return fSliceHit; } - static bool Compare(const AliHLTTPCCAGBHit &a, const AliHLTTPCCAGBHit &b); @@ -69,8 +66,6 @@ class AliHLTTPCCAGBHit Int_t fID; //* external ID (id of AliTPCcluster) Bool_t fIsUsed; //* is used by GBTracks - AliHLTTPCCAHit fSliceHit; //* corresponding slice hit - //ClassDef(AliHLTTPCCAGBHit,1); }; diff --git a/HLT/TPCLib/tracking-ca/AliHLTTPCCAGBTrack.h b/HLT/TPCLib/tracking-ca/AliHLTTPCCAGBTrack.h index 42eddac2561..d457f7ad682 100644 --- a/HLT/TPCLib/tracking-ca/AliHLTTPCCAGBTrack.h +++ b/HLT/TPCLib/tracking-ca/AliHLTTPCCAGBTrack.h @@ -8,8 +8,8 @@ #ifndef ALIHLTTPCCAGBTRACK_H #define ALIHLTTPCCAGBTRACK_H -#include "Rtypes.h" +#include "AliHLTTPCCADef.h" #include "AliHLTTPCCATrackParam.h" /** @@ -45,7 +45,7 @@ class AliHLTTPCCAGBTrack void Dummy(); // to make rulechecker happy by having something in .cxx file - ClassDef(AliHLTTPCCAGBTrack,1); + ClassDef(AliHLTTPCCAGBTrack,1) }; diff --git a/HLT/TPCLib/tracking-ca/AliHLTTPCCAGBTracker.cxx b/HLT/TPCLib/tracking-ca/AliHLTTPCCAGBTracker.cxx index 84cd47c36ee..9e29c09a253 100644 --- a/HLT/TPCLib/tracking-ca/AliHLTTPCCAGBTracker.cxx +++ b/HLT/TPCLib/tracking-ca/AliHLTTPCCAGBTracker.cxx @@ -21,10 +21,11 @@ #include "AliHLTTPCCAOutTrack.h" #include "AliHLTTPCCATracker.h" #include "AliHLTTPCCAGBTrack.h" +#include "AliHLTTPCCATrackParam.h" -#include "TMath.h" +#include "AliHLTTPCCAMath.h" #include "TStopwatch.h" -#include "Riostream.h" +#include //#define DRAW @@ -33,10 +34,9 @@ #include "TApplication.h" #endif //DRAW -ClassImp(AliHLTTPCCAGBTracker) AliHLTTPCCAGBTracker::AliHLTTPCCAGBTracker() - : TObject(), + : fSlices(0), fNSlices(0), fHits(0), @@ -53,7 +53,7 @@ AliHLTTPCCAGBTracker::AliHLTTPCCAGBTracker() } AliHLTTPCCAGBTracker::AliHLTTPCCAGBTracker(const AliHLTTPCCAGBTracker&) - : TObject(), + : fSlices(0), fNSlices(0), fHits(0), @@ -148,12 +148,6 @@ void AliHLTTPCCAGBTracker::ReadHit( Float_t x, Float_t y, Float_t z, hit.ISlice()=iSlice; hit.IRow() = iRow; hit.IsUsed() = 0; - - hit.SliceHit().Y() = y; - hit.SliceHit().Z() = z; - hit.SliceHit().ErrY() = errY; - hit.SliceHit().ErrZ() = errZ; - hit.SliceHit().ID() = 0; fNHits++; } @@ -161,8 +155,7 @@ void AliHLTTPCCAGBTracker::FindTracks() { //* main tracking routine fTime = 0; - fStatNEvents++; - + fStatNEvents++; #ifdef DRAW if( fStatNEvents<=1 ){ if( !gApplication ){ @@ -173,49 +166,66 @@ void AliHLTTPCCAGBTracker::FindTracks() AliHLTTPCCADisplay::Instance().SetSliceView(); //AliHLTTPCCADisplay::Instance().SetTPCView(); //AliHLTTPCCADisplay::Instance().DrawTPC(); -#endif //DRAW +#endif //DRAW if( fNHits<=0 ) return; - sort(fHits,fHits+fNHits, AliHLTTPCCAGBHit::Compare ); + std::sort(fHits,fHits+fNHits, AliHLTTPCCAGBHit::Compare ); // Read hits, row by row - Int_t oldRow = -1; - Int_t oldSlice = -1; - Int_t nRowHits = 0; - Int_t firstRowHit = 0; int nHitsTotal = fNHits; - AliHLTTPCCAHit *vHits = new AliHLTTPCCAHit[nHitsTotal]; // CA hit array + Float_t *hitY = new Float_t [nHitsTotal]; + Float_t *hitZ = new Float_t [nHitsTotal]; + + Int_t sliceNHits[fNSlices]; + Int_t rowNHits[fNSlices][200]; + for( Int_t is=0; is=0 && oldSlice>=0 ){ - fSlices[oldSlice].ReadHitRow( oldRow, vHits+firstRowHit, nRowHits ); + AliHLTTPCCAGBHit &h = fHits[ih]; + sliceNHits[h.ISlice()]++; + rowNHits[h.ISlice()][h.IRow()]++; + } + + Int_t firstSliceHit = 0; + for( Int_t is=0; is=0 && oldSlice>=0 ){ - fSlices[oldSlice].ReadHitRow( oldRow, vHits+firstRowHit, nRowHits ); + //if( is==24 ){//SG!!! + fSlices[is].ReadEvent( rowFirstHits, rowNHits[is], hitY, hitZ, sliceNHits[is] ); + //} + firstSliceHit+=sliceNHits[is]; } - delete[] vHits; - - //cout<<"Start CA reconstruction"< iT1.Y() || iS.Param().RMax() < iT1.Y() ) iOK[0][itr]=0; } if( iOK[1][itr] ){ - iOK[1][itr] = iT2.TransportToX( 0 ); + iOK[1][itr] = iT2.TransportToX( 0, .99 ); if( iS.Param().RMin() > iT2.Y() || iS.Param().RMax() < iT2.Y() ) iOK[1][itr]=0; } } @@ -328,14 +342,14 @@ void AliHLTTPCCAGBTracker::Merging() AliHLTTPCCATrackParam &jT2 = jTrParams[1][jtr]; jT1 = jS.OutTracks()[jtr].StartPoint(); jT2 = jS.OutTracks()[jtr].EndPoint(); - jOK[0][jtr] = jT1.Rotate( -dalpha/2 - TMath::Pi()/2 ); - jOK[1][jtr] = jT2.Rotate( -dalpha/2 - TMath::Pi()/2 ); + jOK[0][jtr] = jT1.Rotate( -dalpha/2 - CAMath::Pi()/2 ); + jOK[1][jtr] = jT2.Rotate( -dalpha/2 - CAMath::Pi()/2 ); if( jOK[0][jtr] ){ - jOK[0][jtr] = jT1.TransportToX( 0 ); + jOK[0][jtr] = jT1.TransportToX( 0, .99 ); if( jS.Param().RMin() > jT1.Y() || jS.Param().RMax() < jT1.Y() ) jOK[0][jtr]=0; } if( jOK[1][jtr] ){ - jOK[1][jtr] = jT2.TransportToX( 0 ); + jOK[1][jtr] = jT2.TransportToX( 0, .99 ); if( jS.Param().RMin() > jT2.Y() || jS.Param().RMax() < jT2.Y() ) jOK[1][jtr]=0; } } @@ -420,6 +434,7 @@ void AliHLTTPCCAGBTracker::Merging() if( iOK[i] ) delete[] iOK[i]; if( jOK[i] ) delete[] jOK[i]; } + timerMerge1.Stop(); fStatTime[10]+=timerMerge1.CpuTime(); @@ -476,7 +491,7 @@ void AliHLTTPCCAGBTracker::Merging() AliHLTTPCCATracker &jslice = fSlices[jSlice]; AliHLTTPCCAOutTrack &jTr = jslice.OutTracks()[jtr]; for( int jhit=0; jhit=0 ) continue; @@ -485,8 +500,8 @@ void AliHLTTPCCAGBTracker::Merging() p.fX = jslice.Rows()[h.IRow()].X(); p.fY = h.Y(); p.fZ = h.Z(); - p.fErr2Y = h.ErrY()*h.ErrY(); - p.fErr2Z = h.ErrZ()*h.ErrZ(); + //p.fErr2Y = h.ErrY()*h.ErrY(); + //p.fErr2Z = h.ErrZ()*h.ErrZ(); p.fAmp = h.Amp(); nHits++; } @@ -494,8 +509,7 @@ void AliHLTTPCCAGBTracker::Merging() jSlice = nextSlice[jSlice]; } while( jtr >=0 ); - if( nHits < 30 ) continue; //SG!!! - + if( nHits < 10 ) continue; //SG!!! Int_t firstRow = 0, lastRow = maxNRows-1; for( firstRow=0; firstRow=TMath::Abs(midRow-mmidRow) ) continue; + if( CAMath::Abs(i-mmidRow)>=CAMath::Abs(midRow-mmidRow) ) continue; midRow = i; } if( midRow==firstRow || midRow==lastRow ) continue; @@ -526,7 +540,7 @@ void AliHLTTPCCAGBTracker::Merging() { - { + { FitPoint &p0 = fitPoints[firstRow]; FitPoint &p1 = fitPoints[midRow]; FitPoint &p2 = fitPoints[lastRow]; @@ -535,15 +549,15 @@ void AliHLTTPCCAGBTracker::Merging() Float_t x2=p2.fX, y2=p2.fY, z2=p2.fZ; if( p1.fISlice!=p0.fISlice ){ Float_t dAlpha = fSlices[p0.fISlice].Param().Alpha() - fSlices[p1.fISlice].Param().Alpha(); - Float_t c = TMath::Cos(dAlpha); - Float_t s = TMath::Sin(dAlpha); + Float_t c = CAMath::Cos(dAlpha); + Float_t s = CAMath::Sin(dAlpha); x1 = p1.fX*c + p1.fY*s; y1 = p1.fY*c - p1.fX*s; } if( p2.fISlice!=p0.fISlice ){ Float_t dAlpha = fSlices[p0.fISlice].Param().Alpha() - fSlices[p2.fISlice].Param().Alpha(); - Float_t c = TMath::Cos(dAlpha); - Float_t s = TMath::Sin(dAlpha); + Float_t c = CAMath::Cos(dAlpha); + Float_t s = CAMath::Sin(dAlpha); x2 = p2.fX*c + p2.fY*s; y2 = p2.fY*c - p2.fX*s; } @@ -572,7 +586,7 @@ void AliHLTTPCCAGBTracker::Merging() } //* Transport to the new row - if( !t0.TransportToX( p.fX ) ) continue; + if( !t0.TransportToX( p.fX, .99 ) ) continue; //* Calculate hit errors @@ -586,7 +600,7 @@ void AliHLTTPCCAGBTracker::Merging() AliHLTTPCCATracker *cslice = &(fSlices[currslice]); AliHLTTPCCARow *row = &(cslice->Rows()[iRow]); - if( !t0.TransportToX( row->X() ) ) continue; + if( !t0.TransportToX( row->X(), .99 ) ) continue; if( t0.GetY() > row->MaxY() ){ //next slice @@ -598,8 +612,8 @@ void AliHLTTPCCAGBTracker::Merging() currslice = j; cslice = &(fSlices[currslice]); row = &(cslice->Rows()[iRow]); - if( !t0.TransportToX( row->X() ) ) continue; - if( TMath::Abs(t0.GetY()) > row->MaxY() ) continue; + if( !t0.TransportToX( row->X(), .99 ) ) continue; + if( CAMath::Abs(t0.GetY()) > row->MaxY() ) continue; }else if( t0.GetY() < -row->MaxY() ){ //prev slice Int_t j = prevSlice[currslice]; @@ -608,14 +622,14 @@ void AliHLTTPCCAGBTracker::Merging() currslice = j; cslice = &(fSlices[currslice]); row = &(cslice->Rows()[iRow]); - if( !t0.TransportToX( row->X() ) ) continue; - if( TMath::Abs(t0.GetY()) > row->MaxY() ) continue; + if( !t0.TransportToX( row->X(), .99 ) ) continue; + if( CAMath::Abs(t0.GetY()) > row->MaxY() ) continue; } Int_t bestsh = -1; Float_t ds = 1.e10; for( Int_t ish=0; ishNHits(); ish++ ){ - AliHLTTPCCAHit &sh = row->Hits()[ish]; + AliHLTTPCCAHit &sh = cslice->Hits()[row->FirstHit()+ish]; Float_t dy = sh.Y() - t0.GetY(); Float_t dz = sh.Z() - t0.GetZ(); Float_t dds = dy*dy+dz*dz; @@ -630,7 +644,7 @@ void AliHLTTPCCAGBTracker::Merging() GetErrors2( currslice, iRow, t0, p.fErr2Y, p.fErr2Z ); - AliHLTTPCCAHit &sh = row->Hits()[bestsh]; + AliHLTTPCCAHit &sh = cslice->Hits()[row->FirstHit()+bestsh]; Float_t dy = sh.Y() - t0.GetY(); Float_t dz = sh.Z() - t0.GetZ(); Float_t s2z = /*t0.GetErr2Z() + */ p.fErr2Z; @@ -639,7 +653,7 @@ void AliHLTTPCCAGBTracker::Merging() if( dy*dy>factor2*s2y ) continue; p.fISlice = currslice; - p.fHitID = sh.ID(); + p.fHitID = fFirstSliceHit[p.fISlice] + cslice->HitsID()[row->FirstHit() + bestsh]; p.fX = row->X(); p.fY = sh.Y(); p.fZ = sh.Z(); @@ -648,17 +662,18 @@ void AliHLTTPCCAGBTracker::Merging() //* Update the track - t0.Filter2( p.fY, p.fZ, p.fErr2Y, p.fErr2Z ); + t0.Filter2( p.fY, p.fZ, p.fErr2Y, p.fErr2Z, .99 ); } //* final refit, dE/dx calculation //cout<<"\n\nstart refit..\n"<1.e-4 ){ - Float_t dLdX = TMath::Sqrt(1.+t0.DzDs()*t0.DzDs())/TMath::Abs(t0.CosPhi()); + if( CAMath::Abs( t0.CosPhi() )>1.e-4 ){ + Float_t dLdX = CAMath::Sqrt(1.+t0.DzDs()*t0.DzDs())/CAMath::Abs(t0.CosPhi()); sumDeDx+=p.fAmp/dLdX; nDeDx++; } @@ -753,6 +768,21 @@ void AliHLTTPCCAGBTracker::Merging() } if( t.NHits()<30 ) continue;//SG!! + Double_t dAlpha = 0; + { + Double_t xTPC=83.65; + Double_t ddAlpha = 0.00609235; + + if( t0.TransportToXWithMaterial( xTPC, fitPar ) ){ + Double_t y=t0.GetY(); + Double_t ymax=xTPC*CAMath::Tan(dAlpha/2.); + if (y > ymax) { + if( t0.Rotate( ddAlpha ) ){ dAlpha=ddAlpha; t0.TransportToXWithMaterial( xTPC, fitPar ); } + } else if (y <-ymax) { + if( t0.Rotate( -ddAlpha ) ){ dAlpha=-ddAlpha; t0.TransportToXWithMaterial( xTPC, fitPar );} + } + } + } { Bool_t ok=1; @@ -771,15 +801,15 @@ void AliHLTTPCCAGBTracker::Merging() } } - if( TMath::Abs(t0.Kappa())<1.e-8 ) t0.Kappa() = 1.e-8; + if( CAMath::Abs(t0.Kappa())<1.e-8 ) t0.Kappa() = 1.e-8; t.Param() = t0; - t.Alpha() = fSlices[currslice].Param().Alpha(); + t.Alpha() = fSlices[currslice].Param().Alpha() + dAlpha; nTrackHits+= t.NHits(); fNTracks++; } } } - cout<<"\n\nRejected: "<126) ? 1:2; - Float_t cosPhiInv = TMath::Abs(t.GetCosPhi())>1.e-2 ?1./t.GetCosPhi() :0; + Float_t cosPhiInv = CAMath::Abs(t.GetCosPhi())>1.e-2 ?1./t.GetCosPhi() :0; Float_t angleY = t.GetSinPhi()*cosPhiInv ; Float_t angleZ = t.GetDzDs()*cosPhiInv ; @@ -856,8 +886,9 @@ void AliHLTTPCCAGBTracker::GetErrors2( AliHLTTPCCAGBHit &h, AliHLTTPCCATrackPara GetErrors2( h.ISlice(), h.IRow(), t, Err2Y, Err2Z ); } -void AliHLTTPCCAGBTracker::WriteSettings( ostream &out ) +void AliHLTTPCCAGBTracker::WriteSettings( ostream &out ) const { + //* write settings to the file out<< NSlices()<> nSlices; SetNSlices( nSlices ); @@ -876,8 +908,10 @@ void AliHLTTPCCAGBTracker::ReadSettings( istream &in ) } } -void AliHLTTPCCAGBTracker::WriteEvent( ostream &out ) +void AliHLTTPCCAGBTracker::WriteEvent( ostream &out ) const { + // write event to the file + out<> nHits; @@ -908,23 +944,24 @@ void AliHLTTPCCAGBTracker::ReadEvent( istream &in ) } } -void AliHLTTPCCAGBTracker::WriteTracks( ostream &out ) +void AliHLTTPCCAGBTracker::WriteTracks( ostream &out ) const { + //* Write tracks to file + out<>fTime; fStatTime[0]+=fTime; fStatNEvents++; diff --git a/HLT/TPCLib/tracking-ca/AliHLTTPCCAGBTracker.h b/HLT/TPCLib/tracking-ca/AliHLTTPCCAGBTracker.h index 989f4ce03e6..bb4e6d80a1c 100644 --- a/HLT/TPCLib/tracking-ca/AliHLTTPCCAGBTracker.h +++ b/HLT/TPCLib/tracking-ca/AliHLTTPCCAGBTracker.h @@ -8,7 +8,8 @@ #ifndef ALIHLTTPCCAGBTRACKER_H #define ALIHLTTPCCAGBTRACKER_H -#include "TObject.h" +#include "AliHLTTPCCADef.h" +#include class AliHLTTPCCATracker; class AliHLTTPCCAGBTrack; @@ -31,7 +32,7 @@ class AliHLTTPCCATrackParam; * The class is under construction. * */ -class AliHLTTPCCAGBTracker:public TObject +class AliHLTTPCCAGBTracker { public: @@ -40,7 +41,7 @@ public: AliHLTTPCCAGBTracker(const AliHLTTPCCAGBTracker&); AliHLTTPCCAGBTracker &operator=(const AliHLTTPCCAGBTracker&); - virtual ~AliHLTTPCCAGBTracker(); + ~AliHLTTPCCAGBTracker(); void StartEvent(); void SetNSlices( Int_t N ); @@ -65,11 +66,11 @@ public: void GetErrors2( AliHLTTPCCAGBHit &h, AliHLTTPCCATrackParam &t, Float_t &Err2Y, Float_t &Err2Z ); void GetErrors2( Int_t iSlice, Int_t iRow, AliHLTTPCCATrackParam &t, Float_t &Err2Y, Float_t &Err2Z ); - void WriteSettings( ostream &out ); + void WriteSettings( ostream &out ) const; void ReadSettings( istream &in ); - void WriteEvent( ostream &out ); + void WriteEvent( ostream &out ) const; void ReadEvent( istream &in ); - void WriteTracks( ostream &out ); + void WriteTracks( ostream &out ) const; void ReadTracks( istream &in ); protected: @@ -89,11 +90,11 @@ protected: }; AliHLTTPCCAGBSliceTrackInfo **fSliceTrackInfos; //* additional information for slice tracks - Double_t fTime; + Double_t fTime; //* total time Double_t fStatTime[20]; //* timers Int_t fStatNEvents; //* n events proceed + Int_t fFirstSliceHit[100]; // hit array - ClassDef(AliHLTTPCCAGBTracker,1) //Base class for conformal mapping tracking }; #endif diff --git a/HLT/TPCLib/tracking-ca/AliHLTTPCCAGrid.cxx b/HLT/TPCLib/tracking-ca/AliHLTTPCCAGrid.cxx index 571c85dbc5c..aeb886ef33f 100644 --- a/HLT/TPCLib/tracking-ca/AliHLTTPCCAGrid.cxx +++ b/HLT/TPCLib/tracking-ca/AliHLTTPCCAGrid.cxx @@ -17,51 +17,69 @@ //*************************************************************************** #include "AliHLTTPCCAGrid.h" -#include "TMath.h" +#include "AliHLTTPCCAMath.h" -AliHLTTPCCAGrid::AliHLTTPCCAGrid(const AliHLTTPCCAGrid&) - :fGrid(0),fNy(0),fNz(0),fN(0), - fYMin(0),fYMax(0),fZMin(0),fZMax(0),fStepYInv(0),fStepZInv(0) +GPUd() void AliHLTTPCCAGrid::Create( Float_t yMin, Float_t yMax, Float_t zMin, Float_t zMax, UInt_t n ) { - //* dummy + //* Create the grid + + fYMin = CAMath::Min(yMin,yMax); + fYMax = CAMath::Max(yMin,yMax); + fZMin = CAMath::Min(zMin,zMax); + fZMax = CAMath::Max(zMin,zMax); + fNy = (UInt_t) CAMath::Sqrt( CAMath::Abs( (Float_t) n ) ); + fNy = CAMath::Max(fNy,1); + fNz = fNy; + fN = fNy*fNz; + fStepYInv = (fYMax - fYMin); + fStepZInv = (fZMax - fZMin); + Int_t ky = (fNy>1) ?fNy-1 :1; + Int_t kz = (fNz>1) ?fNz-1 :1; + fStepYInv = ( fStepYInv>1.e-4 ) ?ky/fStepYInv :1; + fStepZInv = ( fStepZInv>1.e-4 ) ?kz/fStepZInv :1; } -AliHLTTPCCAGrid& AliHLTTPCCAGrid::operator=(const AliHLTTPCCAGrid&) +GPUd() void AliHLTTPCCAGrid::Create( Float_t yMin, Float_t yMax, Float_t zMin, Float_t zMax, float_t sy, Float_t sz ) { - //* dummy - return *this; + //* Create the grid + + fYMin = CAMath::Min(yMin,yMax); + fYMax = CAMath::Max(yMin,yMax); + fZMin = CAMath::Min(zMin,zMax); + fZMax = CAMath::Max(zMin,zMax); + fStepYInv = 1./sy; + fStepZInv = 1./sz; + + fNy = (UInt_t) ( (fYMax - fYMin)*fStepYInv + 1 ); + fNz = (UInt_t) ( (fZMax - fZMin)*fStepZInv + 1 ); + fYMax = fYMin + fNy*sy; + fZMax = fZMin + fNz*sz; + fN = fNy*fNz; } -void AliHLTTPCCAGrid::Create( Float_t yMin, Float_t yMax, Float_t zMin, Float_t zMax, Int_t n ) +GPUd() UInt_t AliHLTTPCCAGrid::GetBin( Float_t Y, Float_t Z ) const { - //* Create the grid + //* get the bin pointer - fYMin = TMath::Min(yMin,yMax); - fYMax = TMath::Max(yMin,yMax); - fZMin = TMath::Min(zMin,zMax); - fZMax = TMath::Max(zMin,zMax); - fNy = fNz = (Int_t) TMath::Sqrt( (Float_t) n ); - fNy = TMath::Max(fNy,1); - fNz = TMath::Max(fNz,1); - fN = fNy*fNz; - if( fGrid ) delete[] fGrid; - fGrid = new void*[fN]; - for( Int_t i=0; i1.e-4 ) ?(fNy-1)/fStepYInv :0; - fStepZInv = ( fStepZInv>1.e-4 ) ?(fNz-1)/fStepZInv :0; + Int_t yBin = (Int_t) CAMath::fmul_rz( Y-fYMin, fStepYInv ); + Int_t zBin = (Int_t) CAMath::fmul_rz( Z-fZMin, fStepZInv ); + Int_t bin = CAMath::mul24(zBin,fNy) + yBin; + if( bin<0 ) return 0; + if( bin>=(Int_t) fN ) return fN - 1; + return (UInt_t) bin; } -void **AliHLTTPCCAGrid::Get( Float_t Y, Float_t Z ) const +GPUd() void AliHLTTPCCAGrid::GetBin( Float_t Y, Float_t Z, UInt_t &bY, UInt_t &bZ ) const { //* get the bin pointer - Int_t yBin = (Int_t) ( (Y-fYMin)*fStepYInv ); - Int_t zBin = (Int_t) ( (Z-fZMin)*fStepZInv ); - if( yBin<0 ) yBin = 0; - else if( yBin>=fNy ) yBin = fNy - 1; - if( zBin<0 ) zBin = 0; - else if( zBin>=fNz ) zBin = fNz - 1; - return fGrid + zBin*fNy + yBin; + Int_t bbY = (Int_t) ( (Y-fYMin)*fStepYInv ); + Int_t bbZ = (Int_t) ( (Z-fZMin)*fStepZInv ); + + if( bbY<0 ) bbY = 0; + else if( bbY>=(Int_t)fNy ) bbY = fNy - 1; + if( bbZ<0 ) bbZ = 0; + else if( bbZ>=(Int_t)fNz ) bbZ = fNz - 1; + bY = (UInt_t) bbY; + bZ = (UInt_t) bbZ; } diff --git a/HLT/TPCLib/tracking-ca/AliHLTTPCCAGrid.h b/HLT/TPCLib/tracking-ca/AliHLTTPCCAGrid.h index 7e3ec8040fb..86e2eb10e02 100644 --- a/HLT/TPCLib/tracking-ca/AliHLTTPCCAGrid.h +++ b/HLT/TPCLib/tracking-ca/AliHLTTPCCAGrid.h @@ -8,7 +8,7 @@ #ifndef ALIHLTTPCCAGRID_H #define ALIHLTTPCCAGRID_H -#include "Rtypes.h" +#include "AliHLTTPCCADef.h" /** * @class AliHLTTPCCAGrid @@ -21,51 +21,51 @@ class AliHLTTPCCAGrid { public: - AliHLTTPCCAGrid():fGrid(0),fNy(0),fNz(0),fN(0), - fYMin(0),fYMax(0),fZMin(0),fZMax(0),fStepYInv(0),fStepZInv(0){} - - AliHLTTPCCAGrid(const AliHLTTPCCAGrid&); - AliHLTTPCCAGrid &operator=(const AliHLTTPCCAGrid&); - - virtual ~AliHLTTPCCAGrid(){ - if( fGrid ) delete[] fGrid; + + GPUd() void Create( Float_t yMin, Float_t yMax, Float_t zMin, Float_t zMax, UInt_t n ); + GPUd() void Create( Float_t yMin, Float_t yMax, Float_t zMin, Float_t zMax, Float_t sy, Float_t sz ); + + GPUd() UInt_t GetBin( Float_t Y, Float_t Z ) const; + GPUd() void GetBin( Float_t Y, Float_t Z, UInt_t &bY, UInt_t &bZ ) const ; + + GPUd() UInt_t GetBinNoCheck( Float_t Y, Float_t Z ) const { + UInt_t bY = (UInt_t) ( (Y-fYMin)*fStepYInv ); + UInt_t bZ = (UInt_t) ( (Z-fZMin)*fStepZInv ); + return bZ*fNy + bY; } - - void Create( Float_t yMin, Float_t yMax, Float_t zMin, Float_t zMax, Int_t n ); - - void **Get( Float_t Y, Float_t Z ) const; - - void **GetNoCheck( Float_t Y, Float_t Z ) const { - Int_t yBin = (Int_t) ( (Y-fYMin)*fStepYInv ); - Int_t zBin = (Int_t) ( (Z-fZMin)*fStepZInv ); - return fGrid + zBin*fNy + yBin; + + GPUd() void GetBinNoCheck( Float_t Y, Float_t Z, UInt_t &bY, UInt_t &bZ ) const { + bY = (UInt_t) ( (Y-fYMin)*fStepYInv ); + bZ = (UInt_t) ( (Z-fZMin)*fStepZInv ); } - Int_t N() const { return fN; } - Int_t Ny() const { return fNy; } - Int_t Nz() const { return fNz; } - Float_t YMin() const { return fYMin; } - Float_t YMax() const { return fYMax; } - Float_t ZMin() const { return fZMin; } - Float_t ZMax() const { return fZMax; } - Float_t StepYInv() const { return fStepYInv; } - Float_t StepZInv() const { return fStepZInv; } - void **Grid(){ return fGrid; } - protected: + GPUd() UInt_t N() const { return fN; } + GPUd() UInt_t Ny() const { return fNy; } + GPUd() UInt_t Nz() const { return fNz; } + GPUd() Float_t YMin() const { return fYMin; } + GPUd() Float_t YMax() const { return fYMax; } + GPUd() Float_t ZMin() const { return fZMin; } + GPUd() Float_t ZMax() const { return fZMax; } + GPUd() Float_t StepYInv() const { return fStepYInv; } + GPUd() Float_t StepZInv() const { return fStepZInv; } + GPUhd() UInt_t &Content2() { return fContent2;} + GPUhd() UInt_t &Offset() { return fOffset;} + + private: - void **fGrid; //* the grid as 1-d array - Int_t fNy; //* N bins in Y - Int_t fNz; //* N bins in Z - Int_t fN; //* total N bins + UInt_t fNy; //* N bins in Y + UInt_t fNz; //* N bins in Z + UInt_t fN; //* total N bins Float_t fYMin; //* minimal Y value Float_t fYMax; //* maximal Y value Float_t fZMin; //* minimal Z value Float_t fZMax; //* maximal Z value Float_t fStepYInv; //* inverse bin size in Y Float_t fStepZInv; //* inverse bin size in Z - -}; + UInt_t fContent2; //* content of the fN/2 bin [4 bytes ] + UInt_t fOffset; //* offset of this Grid content in the common content array +}; #endif diff --git a/HLT/TPCLib/tracking-ca/AliHLTTPCCAHit.h b/HLT/TPCLib/tracking-ca/AliHLTTPCCAHit.h index f414bea564b..5796d162b37 100644 --- a/HLT/TPCLib/tracking-ca/AliHLTTPCCAHit.h +++ b/HLT/TPCLib/tracking-ca/AliHLTTPCCAHit.h @@ -8,7 +8,7 @@ #ifndef ALIHLTTPCCAHIT_H #define ALIHLTTPCCAHIT_H -#include "Rtypes.h" +#include "AliHLTTPCCADef.h" /** * @class AliHLTTPCCAHit @@ -19,44 +19,16 @@ */ class AliHLTTPCCAHit { - public: - AliHLTTPCCAHit(): fY(0),fZ(0),fErrY(0),fErrZ(0),fID(0){;} - virtual ~AliHLTTPCCAHit(){} - - Float_t &Y() { return fY; } - Float_t &Z() { return fZ; } - Float_t &ErrY(){ return fErrY; } - Float_t &ErrZ(){ return fErrZ; } +public: + + GPUhd() Float_t &Y() { return fY; } + GPUhd() Float_t &Z() { return fZ; } + + //protected: - Int_t &ID(){ return fID; } - - void Set( Int_t id, Float_t y, Float_t z, - Float_t errY, Float_t errZ ); - - protected: - Float_t fY, fZ; // Y and Z position of the TPC cluster - Float_t fErrY, fErrZ; // position errors - Int_t fID; // external unique ID of this hit, - // used as cluster index in track->hit reference array - private: - - void Dummy(); // to make rulechecker happy by having something in .cxx file }; - -inline void AliHLTTPCCAHit::Set( Int_t id, Float_t y, Float_t z, - Float_t errY, Float_t errZ ) -{ - //* set parameters - fID = id; - fY = y; - fZ = z; - fErrY = errY; - fErrZ = errZ; -} - - #endif diff --git a/HLT/TPCLib/tracking-ca/AliHLTTPCCAHitArea.cxx b/HLT/TPCLib/tracking-ca/AliHLTTPCCAHitArea.cxx new file mode 100644 index 00000000000..21a499f739c --- /dev/null +++ b/HLT/TPCLib/tracking-ca/AliHLTTPCCAHitArea.cxx @@ -0,0 +1,111 @@ +//*************************************************************************** +// This file is property of and copyright by the ALICE HLT Project * +// ALICE Experiment at CERN, All rights reserved. * +// * +// Primary Authors: Sergey Gorbunov * +// Ivan Kisel * +// for The ALICE HLT Project. * +// * +// Permission to use, copy, modify and distribute this software and its * +// documentation strictly for non-commercial purposes is hereby granted * +// without fee, provided that the above copyright notice appears in all * +// copies and that both the copyright notice and this permission notice * +// appear in the supporting documentation. The authors make no claims * +// about the suitability of this software for any purpose. It is * +// provided "as is" without express or implied warranty. * +//*************************************************************************** + +#include "AliHLTTPCCAHitArea.h" +#include "AliHLTTPCCATracker.h" +#include "AliHLTTPCCAGrid.h" +#include "AliHLTTPCCAHit.h" + + +GPUd() void AliHLTTPCCAHitAreaInit( AliHLTTPCCAHitArea &a, AliHLTTPCCATracker &tracker, AliHLTTPCCAGrid &grid, UInt_t hitoffset, Float_t y, Float_t z, Float_t dy, Float_t dz ) +{ + // initialisation + + UInt_t gridOffset = grid.Offset(); + a.HitOffset() = hitoffset; + a.Y() = y; + a.Z() = z; + a.MinZ() = z-dz; + a.MaxZ() = z+dz; + a.MinY() = y-dy; + a.MaxY() = y+dy; + UInt_t bYmin, bZmin, bYmax; + grid.GetBin(a.MinY(), a.MinZ(), bYmin, bZmin); + grid.GetBin(a.MaxY(), a.MaxZ(), bYmax, a.BZmax()); + a.BDY() = bYmax - bYmin + 1; + a.Ny() = grid.Ny(); + a.N2() = gridOffset + (grid.N()>>1); + a.C2() = grid.Content2(); + a.IndYmin() = gridOffset + bZmin*a.Ny() + bYmin; + a.Iz() = bZmin; + a.HitYfst() = tracker.GetGridContent((UInt_t)( a.IndYmin())); + a.HitYlst() = tracker.GetGridContent((UInt_t)( a.IndYmin() + a.BDY())); + if( a.IndYmin()>=a.N2() ) a.HitYfst()+=a.C2(); + if( a.IndYmin()+ a.BDY() >=a.N2() ) a.HitYlst()+=a.C2(); + a.Ih() = a.HitYfst(); +} + + +GPUd() void AliHLTTPCCAHitArea::Init( AliHLTTPCCATracker &tracker,AliHLTTPCCAGrid &grid, UInt_t hitoffset, Float_t y, Float_t z, Float_t dy, Float_t dz ) +{ + //initialisation + AliHLTTPCCAHitAreaInit(*this, tracker, grid, hitoffset, y, z, dy, dz); +} + +GPUd() Int_t AliHLTTPCCAHitArea::GetNext(AliHLTTPCCATracker &tracker, AliHLTTPCCAHit &h) +{ + // get next hit index + Int_t ret = -1; + do{ + while( fIh>=fHitYlst ){ + if( fIz>=fBZmax ) return -1; + fIz++; + fIndYmin += fNy; + fHitYfst = tracker.GetGridContent((UInt_t)( fIndYmin)); + fHitYlst = tracker.GetGridContent((UInt_t)( fIndYmin + fBDY)); + if( fIndYmin>=fn2 ) fHitYfst+=fc2; + if( fIndYmin+ fBDY>=fn2 ) fHitYlst+=fc2; + fIh = fHitYfst; + } + + h = tracker.GetHit( fHitOffset + fIh ); + + if( h.fZ>fMaxZ || h.fZfMaxY ){ + fIh++; + continue; + } + ret = fIh; + fIh++; + break; + } while(1); + return ret; +} + + + +GPUd() Int_t AliHLTTPCCAHitArea::GetBest(AliHLTTPCCATracker &tracker, AliHLTTPCCAHit &h) +{ + // get closest hit in the area + Int_t best = -1; + Float_t ds = 1.e10; + do{ + AliHLTTPCCAHit hh; + Int_t ih=GetNext( tracker, hh ); + if( ih<0 ) break; + Float_t dy = hh.fY - fY; + Float_t dz = hh.fZ - fZ; + Float_t dds = dy*dy+dz*dz; + if( ddsPz()*pi; fPar[6] = 0; fPDG = part->GetPdgCode(); - if ( TMath::Abs(fPDG) < 100000 ){ + if ( CAMath::Abs(fPDG) < 100000 ){ TParticlePDG *pPDG = TDatabasePDG::Instance()->GetParticle(fPDG); if( pPDG ) fPar[6] = pPDG->Charge()/3.0*pi; } @@ -70,13 +70,13 @@ void AliHLTTPCCAMCTrack::SetTPCPar( Float_t X, Float_t Y, Float_t Z, fTPCPar[0] = X; fTPCPar[1] = Y; fTPCPar[2] = Z; - Double_t p = TMath::Sqrt(Px*Px + Py*Py + Pz*Pz); + Double_t p = CAMath::Sqrt(Px*Px + Py*Py + Pz*Pz); Double_t pi = ( p >1.e-4 ) ?1./p :0; fTPCPar[3] = Px*pi; fTPCPar[4] = Py*pi; fTPCPar[5] = Pz*pi; fTPCPar[6] = 0; - if ( TMath::Abs(fPDG) < 100000 ){ + if ( CAMath::Abs(fPDG) < 100000 ){ TParticlePDG *pPDG = TDatabasePDG::Instance()->GetParticle(fPDG); if( pPDG ) fTPCPar[6] = pPDG->Charge()/3.0*pi; } diff --git a/HLT/TPCLib/tracking-ca/AliHLTTPCCAMCTrack.h b/HLT/TPCLib/tracking-ca/AliHLTTPCCAMCTrack.h index cd3d793bd4e..a2d97d50450 100644 --- a/HLT/TPCLib/tracking-ca/AliHLTTPCCAMCTrack.h +++ b/HLT/TPCLib/tracking-ca/AliHLTTPCCAMCTrack.h @@ -8,7 +8,7 @@ #ifndef ALIHLTTPCCAMCTRACK_H #define ALIHLTTPCCAMCTRACK_H -#include "Rtypes.h" +#include "AliHLTTPCCADef.h" class TParticle; diff --git a/HLT/TPCLib/tracking-ca/AliHLTTPCCAMath.h b/HLT/TPCLib/tracking-ca/AliHLTTPCCAMath.h new file mode 100644 index 00000000000..ca6f9f77a35 --- /dev/null +++ b/HLT/TPCLib/tracking-ca/AliHLTTPCCAMath.h @@ -0,0 +1,223 @@ +//-*- Mode: C++ -*- +// @(#) $Id: AliHLTTPCCARow.h 27042 2008-07-02 12:06:02Z richterm $ + +//* This file is property of and copyright by the ALICE HLT Project * +//* ALICE Experiment at CERN, All rights reserved. * +//* See cxx source for full Copyright notice * + +#ifndef ALIHLTTPCCAMATH_H +#define ALIHLTTPCCAMATH_H + +#include "AliHLTTPCCADef.h" + +#if defined(HLTCA_STANDALONE) || defined(HLTCA_GPUCODE) +#include +#else +#include "TMath.h" +#endif + +/** + * @class ALIHLTTPCCAMath + * + * + */ +class AliHLTTPCCAMath +{ +public: + GPUd() static Float_t Min(Float_t x, Float_t y); + GPUd() static Float_t Max(Float_t x, Float_t y); + GPUd() static Int_t Min(Int_t x, Int_t y); + GPUd() static Int_t Max(Int_t x, Int_t y); + GPUd() static Float_t Sqrt(Float_t x ); + GPUd() static Float_t Abs(Float_t x ); + GPUd() static Double_t Abs(Double_t x ); + GPUd() static Int_t Abs(Int_t x ); + GPUd() static Float_t ASin(Float_t x ); + GPUd() static Float_t ATan2( Float_t y, Float_t x ); + GPUd() static Float_t Sin( Float_t x ); + GPUd() static Float_t Cos( Float_t x ); + GPUd() static Float_t Tan( Float_t x ); + GPUd() static Float_t Copysign( Float_t x, Float_t y ); + GPUd() static Float_t TwoPi(){ return 6.28319; } + GPUd() static Float_t Pi(){ return 3.1415926535897; } + GPUd() static Int_t Nint(Float_t x ); + GPUd() static Bool_t Finite(Float_t x ); + + GPUd() static Int_t atomicExch( Int_t *addr, Int_t val); + GPUd() static Int_t atomicAdd ( Int_t *addr, Int_t val); + GPUd() static Int_t atomicMax ( Int_t *addr, Int_t val); + GPUd() static Int_t atomicMin ( Int_t *addr, Int_t val); + GPUd() static Int_t mul24( Int_t a, Int_t b ); + GPUd() static Float_t fmul_rz( Float_t a, Float_t b ); +}; + +typedef AliHLTTPCCAMath CAMath; + + +#if defined( HLTCA_GPUCODE ) +#define choice(c1,c2,c3) c1 +#elif defined( HLTCA_STANDALONE ) +#define choice(c1,c2,c3) c2 +#else +#define choice(c1,c2,c3) c3 +#endif + + +GPUd() inline Int_t AliHLTTPCCAMath::Nint(Float_t x) +{ +#if defined(HLTCA_STANDALONE) || defined( HLTCA_GPUCODE ) + int i; + if (x >= 0) { + i = int(x + 0.5); + if (x + 0.5 == Float_t(i) && i & 1) i--; + } else { + i = int(x - 0.5); + if (x - 0.5 == Float_t(i) && i & 1) i++; + } + return i; +#else + return TMath::Nint(x); +#endif +} + +GPUd() inline Bool_t AliHLTTPCCAMath::Finite(Float_t x) +{ + return choice( 1, finite(x), finite(x) ); +} + +GPUd() inline Float_t AliHLTTPCCAMath::ATan2(Float_t y, Float_t x) +{ + return choice(atan2f(y,x), atan2(y,x), TMath::ATan2(y,x) ); +} + + +GPUd() inline Float_t AliHLTTPCCAMath::Copysign(Float_t x, Float_t y) +{ +#if defined( HLTCA_GPUCODE ) + return copysignf(x,y); +#else + x = CAMath::Abs(x); + return (y>=0) ?x : -x; +#endif +} + + +GPUd() inline Float_t AliHLTTPCCAMath::Sin(Float_t x) +{ + return choice( sinf(x), sin(x), TMath::Sin(x) ); +} + +GPUd() inline Float_t AliHLTTPCCAMath::Cos(Float_t x) +{ + return choice( cosf(x), cos(x), TMath::Cos(x) ); +} + +GPUd() inline Float_t AliHLTTPCCAMath::Tan(Float_t x) +{ + return choice( tanf(x), tan(x), TMath::Tan(x) ); +} + +GPUd() inline Float_t AliHLTTPCCAMath::Min(Float_t x, Float_t y) +{ + return choice( fminf(x,y), (xy ?x :y), TMath::Max(x,y) ); +} + +GPUd() inline Int_t AliHLTTPCCAMath::Min(Int_t x, Int_t y) +{ + return choice( min(x,y), (xy ?x :y), TMath::Max(x,y) ); +} + +GPUd() inline Float_t AliHLTTPCCAMath::Sqrt(Float_t x ) +{ + return choice( sqrtf(x), sqrt(x), TMath::Sqrt(x) ); +} + +GPUd() inline Float_t AliHLTTPCCAMath::Abs(Float_t x ) +{ + return choice( fabsf(x), fabs(x), TMath::Abs(x) ); +} + +GPUd() inline Double_t AliHLTTPCCAMath::Abs(Double_t x ) +{ + return choice( fabs(x), fabs(x), TMath::Abs(x) ); +} + +GPUd() inline Int_t AliHLTTPCCAMath::Abs(Int_t x ) +{ + return choice( abs(x), (x>=0 ?x :-x), TMath::Abs(x) ); +} + +GPUd() inline Float_t AliHLTTPCCAMath::ASin(Float_t x ) +{ + return choice( asinf(x), asin(x), TMath::ASin(x) ); +} + + +GPUd() inline Int_t AliHLTTPCCAMath::mul24( Int_t a, Int_t b ) +{ + return choice( __mul24(a,b), a*b, a*b ); +} + +GPUd() inline Float_t AliHLTTPCCAMath::fmul_rz( Float_t a, Float_t b ) +{ + return choice( __fmul_rz(a,b), a*b, a*b ); +} + + +GPUd() inline Int_t AliHLTTPCCAMath::atomicExch( Int_t *addr, Int_t val) +{ +#if defined( HLTCA_GPUCODE ) + return ::atomicExch(addr, val ); +#else + Int_t old = *addr; + *addr = val; + return old; +#endif +} + +GPUd() inline Int_t AliHLTTPCCAMath::atomicAdd ( Int_t *addr, Int_t val) +{ +#if defined( HLTCA_GPUCODE ) + return ::atomicAdd(addr, val ); +#else + Int_t old = *addr; + *addr += val; + return old; +#endif +} + +GPUd() inline Int_t AliHLTTPCCAMath::atomicMax ( Int_t *addr, Int_t val) +{ +#if defined( HLTCA_GPUCODE ) + return ::atomicMax(addr, val ); +#else + Int_t old = *addr; + if( *addr< val ) *addr = val; + return old; +#endif +} + +GPUd() inline Int_t AliHLTTPCCAMath::atomicMin ( Int_t *addr, Int_t val) +{ +#if defined( HLTCA_GPUCODE ) + return ::atomicMin(addr, val ); +#else + Int_t old = *addr; + if( *addr> val ) *addr = val; + return old; +#endif +} + +#undef CHOICE + +#endif diff --git a/HLT/TPCLib/tracking-ca/AliHLTTPCCANeighboursCleaner.cxx b/HLT/TPCLib/tracking-ca/AliHLTTPCCANeighboursCleaner.cxx new file mode 100644 index 00000000000..ba363880b00 --- /dev/null +++ b/HLT/TPCLib/tracking-ca/AliHLTTPCCANeighboursCleaner.cxx @@ -0,0 +1,70 @@ +// @(#) $Id: AliHLTTPCCANeighboursCleaner.cxx 27042 2008-07-02 12:06:02Z richterm $ +//*************************************************************************** +// This file is property of and copyright by the ALICE HLT Project * +// ALICE Experiment at CERN, All rights reserved. * +// * +// Primary Authors: Sergey Gorbunov * +// Ivan Kisel * +// for The ALICE HLT Project. * +// * +// Permission to use, copy, modify and distribute this software and its * +// documentation strictly for non-commercial purposes is hereby granted * +// without fee, provided that the above copyright notice appears in all * +// copies and that both the copyright notice and this permission notice * +// appear in the supporting documentation. The authors make no claims * +// about the suitability of this software for any purpose. It is * +// provided "as is" without express or implied warranty. * +//*************************************************************************** + +#include "AliHLTTPCCANeighboursCleaner.h" +#include "AliHLTTPCCAMath.h" +#include "AliHLTTPCCATracker.h" + +GPUd() void AliHLTTPCCANeighboursCleaner::Thread +( Int_t nBlocks, Int_t nThreads, Int_t iBlock, Int_t iThread, Int_t iSync, + AliHLTTPCCASharedMemory &s, AliHLTTPCCATracker &tracker ) +{ + // * + // * kill link to the neighbour if the neighbour is not pointed to the hit + // * + + if( iSync==0 ) + { + if( iThread==0 ){ + s.fNRows = tracker.Param().NRows(); + s.fIRow = iBlock+1; + if( s.fIRow <= s.fNRows-2 ){ + int iRowUp = s.fIRow+1; + int iRowDn = s.fIRow-1; + int firstDn = tracker.Rows()[iRowDn].FirstHit(); + s.fFirstHit = tracker.Rows()[s.fIRow].FirstHit(); + int firstUp = tracker.Rows()[iRowUp].FirstHit(); + Short_t *hhitLinkUp = tracker.HitLinkUp(); + Short_t *hhitLinkDown = tracker.HitLinkDown(); + s.fHitLinkUp = hhitLinkUp + s.fFirstHit; + s.fHitLinkDown = hhitLinkDown + s.fFirstHit; + s.fNHits = firstUp - s.fFirstHit; + s.fUpHitLinkDown = hhitLinkDown + firstUp; + s.fDnHitLinkUp = hhitLinkUp + firstDn; + } + } + } + else if( iSync==1 ) + { + if( s.fIRow <= s.fNRows-2 ){ + for( int ih=iThread; ih=0) && ( s.fUpHitLinkDown[up] != ih) ){ + s.fHitLinkUp[ih]= -1; + //HLTCA_GPU_SUFFIX(CAMath::atomicExch)( tracker.HitLinkUp() + s.fFirstHit+ih, -1 ); + } + if( (dn>=0) && ( s.fDnHitLinkUp [dn] != ih) ){ + s.fHitLinkDown[ih]=-1; + //HLTCA_GPU_SUFFIX(CAMath::atomicExch)( tracker.HitLinkDown() + s.fFirstHit+ih,-1); + } + } + } + } +} + diff --git a/HLT/TPCLib/tracking-ca/AliHLTTPCCANeighboursCleaner.h b/HLT/TPCLib/tracking-ca/AliHLTTPCCANeighboursCleaner.h new file mode 100644 index 00000000000..5167a1a2038 --- /dev/null +++ b/HLT/TPCLib/tracking-ca/AliHLTTPCCANeighboursCleaner.h @@ -0,0 +1,51 @@ +//-*- Mode: C++ -*- + +//* This file is property of and copyright by the ALICE HLT Project * +//* ALICE Experiment at CERN, All rights reserved. * +//* See cxx source for full Copyright notice * + +#ifndef ALIHLTTPCCANEIGHBOURSCLEANER_H +#define ALIHLTTPCCANEIGHBOURSCLEANER_H + + +#include "AliHLTTPCCADef.h" + +class AliHLTTPCCATracker; + +/** + * @class AliHLTTPCCANeighboursCleaner + * + */ +class AliHLTTPCCANeighboursCleaner +{ + public: + class AliHLTTPCCASharedMemory + { + friend class AliHLTTPCCANeighboursCleaner; + public: +#if !defined(HLTCA_GPUCODE) + AliHLTTPCCASharedMemory() + :fIRow(0),fNRows(0),fNHits(0),fHitLinkDown(0),fHitLinkUp(0),fUpHitLinkDown(0),fDnHitLinkUp(0),fFirstHit(0){} + AliHLTTPCCASharedMemory( const AliHLTTPCCASharedMemory& /*dummy*/) + :fIRow(0),fNRows(0),fNHits(0),fHitLinkDown(0),fHitLinkUp(0),fUpHitLinkDown(0),fDnHitLinkUp(0),fFirstHit(0) {} + AliHLTTPCCASharedMemory& operator=(const AliHLTTPCCASharedMemory& /*dummy*/){ return *this; } +#endif + protected: + Int_t fIRow; // current row index + Int_t fNRows; // number of rows + Int_t fNHits; // number of hits + Short_t *fHitLinkDown; // links to the previous row + Short_t *fHitLinkUp; // links to the next row + Short_t *fUpHitLinkDown; // links from next row + Short_t *fDnHitLinkUp; // links from previous row + Int_t fFirstHit; // index of the first row hit in global arrays + }; + + GPUd() static Int_t NThreadSyncPoints(){ return 1; } + + GPUd() static void Thread( Int_t nBlocks, Int_t nThreads, Int_t iBlock, Int_t iThread, Int_t iSync, + AliHLTTPCCASharedMemory &smem, AliHLTTPCCATracker &tracker ); +}; + + +#endif diff --git a/HLT/TPCLib/tracking-ca/AliHLTTPCCANeighboursFinder.cxx b/HLT/TPCLib/tracking-ca/AliHLTTPCCANeighboursFinder.cxx new file mode 100644 index 00000000000..3991fb18bc2 --- /dev/null +++ b/HLT/TPCLib/tracking-ca/AliHLTTPCCANeighboursFinder.cxx @@ -0,0 +1,150 @@ +// @(#) $Id: AliHLTTPCCANeighboursFinder.cxx 27042 2008-07-02 12:06:02Z richterm $ +//*************************************************************************** +// This file is property of and copyright by the ALICE HLT Project * +// ALICE Experiment at CERN, All rights reserved. * +// * +// Primary Authors: Sergey Gorbunov * +// Ivan Kisel * +// for The ALICE HLT Project. * +// * +// Permission to use, copy, modify and distribute this software and its * +// documentation strictly for non-commercial purposes is hereby granted * +// without fee, provided that the above copyright notice appears in all * +// copies and that both the copyright notice and this permission notice * +// appear in the supporting documentation. The authors make no claims * +// about the suitability of this software for any purpose. It is * +// provided "as is" without express or implied warranty. * +//*************************************************************************** + +#include "AliHLTTPCCANeighboursFinder.h" +#include "AliHLTTPCCAMath.h" +#include "AliHLTTPCCAHitArea.h" +#include "AliHLTTPCCATracker.h" + +GPUd() void AliHLTTPCCANeighboursFinder::Thread +( Int_t /*nBlocks*/, Int_t nThreads, Int_t iBlock, Int_t iThread, Int_t iSync, + AliHLTTPCCASharedMemory &s, AliHLTTPCCATracker &tracker ) +{ + //* find neighbours + + if( iSync==0 ) + { + if( iThread==0 ){ + s.fNRows = tracker.Param().NRows(); + s.fIRow = iBlock; + if( s.fIRow < s.fNRows ){ + s.fFirst = tracker.Rows()[s.fIRow].FirstHit(); + s.fHitLinkUp = tracker.HitLinkUp() + s.fFirst; + s.fHitLinkDn = tracker.HitLinkDown() + s.fFirst; + s.fNHits = tracker.Rows()[s.fIRow].NHits(); + if( (s.fIRow>0) && (s.fIRow= s.fNRows-1) ) return; + + const float kAreaSize = 3; + float chi2Cut = 3.*3.*4*(s.fUpDx*s.fUpDx + s.fDnDx*s.fDnDx ); + const Int_t kMaxN = 5; + + for( int ih=iThread; ih=1 && s.fUpNHits>=1 ){ + + int nNeighUp = 0; + AliHLTTPCCAHit h0 = tracker.GetHit( s.fFirst + ih ); + float y = h0.Y(); + float z = h0.Z(); + AliHLTTPCCAHitArea areaDn, areaUp; + areaUp.Init( tracker, s.fGridUp, s.fFirstUp, y*s.fUpTx, z*s.fUpTx, kAreaSize, kAreaSize ); + areaDn.Init( tracker, s.fGridDn, s.fFirstDn, y*s.fDnTx, z*s.fDnTx, kAreaSize, kAreaSize ); + do{ + AliHLTTPCCAHit h; + Int_t i = areaUp.GetNext( tracker,h ); + if( i<0 ) break; + neighUp[nNeighUp] = (UShort_t) i; + yzUp[nNeighUp] = make_float2( s.fDnDx*(h.Y()-y), s.fDnDx*(h.Z()-z) ); + if( ++nNeighUp>=kMaxN ) break; + }while(1); + + int nNeighDn=0; + if( nNeighUp>0 ){ + + int bestDn=-1, bestUp=-1; + float bestD=1.e10; + + do{ + AliHLTTPCCAHit h; + Int_t i = areaDn.GetNext( tracker,h ); + if( i<0 ) break; + nNeighDn++; + float2 yzdn = make_float2( s.fUpDx*(h.Y()-y), s.fUpDx*(h.Z()-z) ); + + for( int iUp=0; iUpglobal @@ -126,7 +128,7 @@ void AliHLTTPCCAParam::Slice2Global( Float_t x, Float_t y, Float_t z, *Z = z; } -void AliHLTTPCCAParam::Global2Slice( Float_t X, Float_t Y, Float_t Z, +GPUd() void AliHLTTPCCAParam::Global2Slice( Float_t X, Float_t Y, Float_t Z, Float_t *x, Float_t *y, Float_t *z ) const { // conversion of coorinates global->sector @@ -135,17 +137,18 @@ void AliHLTTPCCAParam::Global2Slice( Float_t X, Float_t Y, Float_t Z, *z = Z; } -Float_t AliHLTTPCCAParam::GetClusterError2( Int_t yz, Int_t type, Float_t z, Float_t angle ) +GPUd() Float_t AliHLTTPCCAParam::GetClusterError2( Int_t yz, Int_t type, Float_t z, Float_t angle ) const { //* recalculate the cluster error wih respect to the track slope Float_t angle2 = angle*angle; - Float_t *c = fParamS0Par[yz][type]; + const Float_t *c = fParamS0Par[yz][type]; Float_t v = c[0] + z*(c[1] + c[3]*z) + angle2*(c[2] + angle2*c[4] + c[5]*z ); - return TMath::Abs(v); + return CAMath::Abs(v); } -void AliHLTTPCCAParam::WriteSettings( ostream &out ) +GPUh() void AliHLTTPCCAParam::WriteSettings( std::ostream &out ) const { + // write settings to the file out << fISlice<> fISlice; in >> fNRows; in >> fAlpha; diff --git a/HLT/TPCLib/tracking-ca/AliHLTTPCCAParam.h b/HLT/TPCLib/tracking-ca/AliHLTTPCCAParam.h index a175a5eec98..c7311b81c55 100644 --- a/HLT/TPCLib/tracking-ca/AliHLTTPCCAParam.h +++ b/HLT/TPCLib/tracking-ca/AliHLTTPCCAParam.h @@ -8,8 +8,9 @@ #ifndef ALIHLTTPCCAPARAM_H #define ALIHLTTPCCAPARAM_H -#include "Rtypes.h" -#include "Riostream.h" +#include "AliHLTTPCCADef.h" +#include + /** * @class ALIHLTTPCCAParam @@ -23,54 +24,57 @@ class AliHLTTPCCAParam { public: - AliHLTTPCCAParam(); - virtual ~AliHLTTPCCAParam(){;} +#if !defined(HLTCA_GPUCODE) + GPUd() AliHLTTPCCAParam(); +#endif + + ~AliHLTTPCCAParam(){;} - void Initialize( Int_t iSlice, Int_t nRows, Float_t rowX[], + GPUd() void Initialize( Int_t iSlice, Int_t nRows, Float_t rowX[], Float_t alpha, Float_t dAlpha, Float_t rMin, Float_t rMax, Float_t zMin, Float_t zMax, Float_t padPitch, Float_t zSigma, Float_t bz ); - void Update(); + GPUd() void Update(); - void Slice2Global( Float_t x, Float_t y, Float_t z, + GPUd() void Slice2Global( Float_t x, Float_t y, Float_t z, Float_t *X, Float_t *Y, Float_t *Z ) const; - void Global2Slice( Float_t x, Float_t y, Float_t z, + GPUd() GPUd() void Global2Slice( Float_t x, Float_t y, Float_t z, Float_t *X, Float_t *Y, Float_t *Z ) const; - Int_t &ISlice(){ return fISlice;} - Int_t &NRows(){ return fNRows;} + GPUhd() Int_t &ISlice(){ return fISlice;} + GPUhd() Int_t &NRows(){ return fNRows;} - Float_t &RowX( Int_t iRow ){ return fRowX[iRow]; } + GPUhd() Float_t &RowX( Int_t iRow ){ return fRowX[iRow]; } - Float_t &Alpha(){ return fAlpha;} - Float_t &DAlpha(){ return fDAlpha;} - Float_t &CosAlpha(){ return fCosAlpha;} - Float_t &SinAlpha(){ return fSinAlpha;} - Float_t &AngleMin(){ return fAngleMin;} - Float_t &AngleMax(){ return fAngleMax;} - Float_t &RMin(){ return fRMin;} - Float_t &RMax(){ return fRMax;} - Float_t &ZMin(){ return fZMin;} - Float_t &ZMax(){ return fZMax;} - Float_t &ErrZ(){ return fErrZ;} - Float_t &ErrX(){ return fErrX;} - Float_t &ErrY(){ return fErrY;} - Float_t &Bz(){ return fBz;} - - Float_t &TrackConnectionFactor(){ return fTrackConnectionFactor; } - Float_t &TrackChiCut() { return fTrackChiCut; } - Float_t &TrackChi2Cut(){ return fTrackChi2Cut; } - Int_t &MaxTrackMatchDRow(){ return fMaxTrackMatchDRow; } - Float_t &YErrorCorrection(){ return fYErrorCorrection; } - Float_t &ZErrorCorrection(){ return fZErrorCorrection; } - Float_t &CellConnectionAngleXY(){ return fCellConnectionAngleXY; } - Float_t &CellConnectionAngleXZ(){ return fCellConnectionAngleXZ; } - - Float_t GetClusterError2(Int_t yz, Int_t type, Float_t z, Float_t angle ); - - void WriteSettings( ostream &out ); - void ReadSettings( istream &in ); - - protected: + GPUd() Float_t &Alpha(){ return fAlpha;} + GPUd() Float_t &DAlpha(){ return fDAlpha;} + GPUd() Float_t &CosAlpha(){ return fCosAlpha;} + GPUd() Float_t &SinAlpha(){ return fSinAlpha;} + GPUd() Float_t &AngleMin(){ return fAngleMin;} + GPUd() Float_t &AngleMax(){ return fAngleMax;} + GPUd() Float_t &RMin(){ return fRMin;} + GPUd() GPUd() Float_t &RMax(){ return fRMax;} + GPUd() GPUd() Float_t &ZMin(){ return fZMin;} + GPUd() Float_t &ZMax(){ return fZMax;} + GPUd() Float_t &ErrZ(){ return fErrZ;} + GPUd() Float_t &ErrX(){ return fErrX;} + GPUd() Float_t &ErrY(){ return fErrY;} + GPUd() Float_t &Bz(){ return fBz;} + + GPUd() Float_t &TrackConnectionFactor(){ return fTrackConnectionFactor; } + GPUd() Float_t &TrackChiCut() { return fTrackChiCut; } + GPUd() Float_t &TrackChi2Cut(){ return fTrackChi2Cut; } + GPUd() Int_t &MaxTrackMatchDRow(){ return fMaxTrackMatchDRow; } + GPUd() Float_t &YErrorCorrection(){ return fYErrorCorrection; } + GPUd() Float_t &ZErrorCorrection(){ return fZErrorCorrection; } + GPUd() Float_t &CellConnectionAngleXY(){ return fCellConnectionAngleXY; } + GPUd() Float_t &CellConnectionAngleXZ(){ return fCellConnectionAngleXZ; } + + GPUd() Float_t GetClusterError2(Int_t yz, Int_t type, Float_t z, Float_t angle ) const; + + void WriteSettings( std::ostream &out ) const; + void ReadSettings( std::istream &in ); + + protected: Int_t fISlice; // slice number Int_t fNRows; // number of rows @@ -97,7 +101,6 @@ class AliHLTTPCCAParam Float_t fRowX[200];// X-coordinate of rows Float_t fParamS0Par[2][3][7]; // cluster error parameterization coeficients - ClassDef(AliHLTTPCCAParam,1); }; diff --git a/HLT/TPCLib/tracking-ca/AliHLTTPCCAPerformance.cxx b/HLT/TPCLib/tracking-ca/AliHLTTPCCAPerformance.cxx index 1b79e3928fc..11dee866793 100644 --- a/HLT/TPCLib/tracking-ca/AliHLTTPCCAPerformance.cxx +++ b/HLT/TPCLib/tracking-ca/AliHLTTPCCAPerformance.cxx @@ -35,127 +35,145 @@ #include "TProfile.h" -ClassImp(AliHLTTPCCAPerformance) - - - AliHLTTPCCAPerformance::AliHLTTPCCAPerformance() - : TObject(), - fTracker(0), - fHitLabels(0), - fNHits(0), - fMCTracks(0), - fNMCTracks(0), - fMCPoints(0), - fNMCPoints(0), - fDoClusterPulls(1), - fStatNEvents(0), - fStatNRecTot(0), - fStatNRecOut(0), - fStatNGhost(0), - fStatNMCAll(0), - fStatNRecAll(0), - fStatNClonesAll(0), - fStatNMCRef(0), - fStatNRecRef(0), - fStatNClonesRef(0), - fStatGBNRecTot(0), - fStatGBNRecOut(0), - fStatGBNGhost(0), - fStatGBNMCAll(0), - fStatGBNRecAll(0), - fStatGBNClonesAll(0), - fStatGBNMCRef(0), - fStatGBNRecRef(0), - fStatGBNClonesRef(0), - fHistoDir(0), - fhResY(0), - fhResZ(0), - fhResSinPhi(0), - fhResDzDs(0), - fhResPt(0), - fhPullY(0), - fhPullZ(0), - fhPullSinPhi(0), - fhPullDzDs(0), - fhPullQPt(0), - fhHitErrY(0), - fhHitErrZ(0), - fhHitResY(0), - fhHitResZ(0), - fhHitPullY(0), - fhHitPullZ(0), - fhHitResY1(0), - fhHitResZ1(0), - fhHitPullY1(0), - fhHitPullZ1(0), - fhCellPurity(0), - fhCellNHits(0), - fhCellPurityVsN(0), - fhCellPurityVsPt(0), - fhEffVsP(0), - fhGBEffVsP(0) + : + fTracker(0), + fHitLabels(0), + fNHits(0), + fMCTracks(0), + fNMCTracks(0), + fMCPoints(0), + fNMCPoints(0), + fDoClusterPulls(0), + fStatNEvents(0), + fStatNRecTot(0), + fStatNRecOut(0), + fStatNGhost(0), + fStatNMCAll(0), + fStatNRecAll(0), + fStatNClonesAll(0), + fStatNMCRef(0), + fStatNRecRef(0), + fStatNClonesRef(0), + fStatGBNRecTot(0), + fStatGBNRecOut(0), + fStatGBNGhost(0), + fStatGBNMCAll(0), + fStatGBNRecAll(0), + fStatGBNClonesAll(0), + fStatGBNMCRef(0), + fStatGBNRecRef(0), + fStatGBNClonesRef(0), + fHistoDir(0), + fhResY(0), + fhResZ(0), + fhResSinPhi(0), + fhResDzDs(0), + fhResPt(0), + fhPullY(0), + fhPullZ(0), + fhPullSinPhi(0), + fhPullDzDs(0), + fhPullQPt(0), + fhHitErrY(0), + fhHitErrZ(0), + fhHitResY(0), + fhHitResZ(0), + fhHitPullY(0), + fhHitPullZ(0), + fhHitResY1(0), + fhHitResZ1(0), + fhHitPullY1(0), + fhHitPullZ1(0), + fhCellPurity(0), + fhCellNHits(0), + fhCellPurityVsN(0), + fhCellPurityVsPt(0), + fhEffVsP(0), + fhGBEffVsP(0), + fhNeighQuality(0), + fhNeighEff(0), + fhNeighQualityVsPt(0), + fhNeighEffVsPt(0), + fhNeighDy(0), + fhNeighDz(0), + fhNeighChi(0), + fhNeighDyVsPt(0), + fhNeighDzVsPt(0), + fhNeighChiVsPt(0), + fhNeighNCombVsArea(0) { //* constructor } AliHLTTPCCAPerformance::AliHLTTPCCAPerformance(const AliHLTTPCCAPerformance&) - : TObject(), - fTracker(0), - fHitLabels(0), - fNHits(0), - fMCTracks(0), - fNMCTracks(0), - fMCPoints(0), - fNMCPoints(0), - fDoClusterPulls(1), - fStatNEvents(0), - fStatNRecTot(0), - fStatNRecOut(0), - fStatNGhost(0), - fStatNMCAll(0), - fStatNRecAll(0), - fStatNClonesAll(0), - fStatNMCRef(0), - fStatNRecRef(0), - fStatNClonesRef(0), - fStatGBNRecTot(0), - fStatGBNRecOut(0), - fStatGBNGhost(0), - fStatGBNMCAll(0), - fStatGBNRecAll(0), - fStatGBNClonesAll(0), - fStatGBNMCRef(0), - fStatGBNRecRef(0), - fStatGBNClonesRef(0), - fHistoDir(0), - fhResY(0), - fhResZ(0), - fhResSinPhi(0), - fhResDzDs(0), - fhResPt(0), - fhPullY(0), - fhPullZ(0), - fhPullSinPhi(0), - fhPullDzDs(0), - fhPullQPt(0), - fhHitErrY(0), - fhHitErrZ(0), - fhHitResY(0), - fhHitResZ(0), - fhHitPullY(0), - fhHitPullZ(0), - fhHitResY1(0), - fhHitResZ1(0), - fhHitPullY1(0), - fhHitPullZ1(0), - fhCellPurity(0), - fhCellNHits(0), - fhCellPurityVsN(0), - fhCellPurityVsPt(0), - fhEffVsP(0), - fhGBEffVsP(0) + : + fTracker(0), + fHitLabels(0), + fNHits(0), + fMCTracks(0), + fNMCTracks(0), + fMCPoints(0), + fNMCPoints(0), + fDoClusterPulls(0), + fStatNEvents(0), + fStatNRecTot(0), + fStatNRecOut(0), + fStatNGhost(0), + fStatNMCAll(0), + fStatNRecAll(0), + fStatNClonesAll(0), + fStatNMCRef(0), + fStatNRecRef(0), + fStatNClonesRef(0), + fStatGBNRecTot(0), + fStatGBNRecOut(0), + fStatGBNGhost(0), + fStatGBNMCAll(0), + fStatGBNRecAll(0), + fStatGBNClonesAll(0), + fStatGBNMCRef(0), + fStatGBNRecRef(0), + fStatGBNClonesRef(0), + fHistoDir(0), + fhResY(0), + fhResZ(0), + fhResSinPhi(0), + fhResDzDs(0), + fhResPt(0), + fhPullY(0), + fhPullZ(0), + fhPullSinPhi(0), + fhPullDzDs(0), + fhPullQPt(0), + fhHitErrY(0), + fhHitErrZ(0), + fhHitResY(0), + fhHitResZ(0), + fhHitPullY(0), + fhHitPullZ(0), + fhHitResY1(0), + fhHitResZ1(0), + fhHitPullY1(0), + fhHitPullZ1(0), + fhCellPurity(0), + fhCellNHits(0), + fhCellPurityVsN(0), + fhCellPurityVsPt(0), + fhEffVsP(0), + fhGBEffVsP(0), + fhNeighQuality(0), + fhNeighEff(0), + fhNeighQualityVsPt(0), + fhNeighEffVsPt(0), + fhNeighDy(0), + fhNeighDz(0), + fhNeighChi(0), + fhNeighDyVsPt(0), + fhNeighDzVsPt(0), + fhNeighChiVsPt(0), + fhNeighNCombVsArea(0) { //* dummy } @@ -265,6 +283,26 @@ void AliHLTTPCCAPerformance::CreateHistos() TDirectory *curdir = gDirectory; fHistoDir = gROOT->mkdir("HLTTPCCATrackerPerformance"); fHistoDir->cd(); + + gDirectory->mkdir("Neighbours"); + gDirectory->cd("Neighbours"); + + fhNeighQuality = new TProfile("NeighQuality", "Neighbours Quality vs row", 160, 0., 160.); + fhNeighEff = new TProfile("NeighEff", "Neighbours Efficiency vs row", 160, 0., 160.); + fhNeighQualityVsPt = new TProfile("NeighQualityVsPt", "Neighbours Quality vs Pt", 100, 0., 5.); + fhNeighEffVsPt = new TProfile("NeighEffVsPt", "Neighbours Efficiency vs Pt", 100, 0., 5.); + fhNeighDy = new TH1D("NeighDy","Neighbours dy",100,-10,10); + fhNeighDz = new TH1D("NeighDz","Neighbours dz",100,-10,10); + fhNeighChi = new TH1D("NeighChi","Neighbours chi",100,0,20); + + fhNeighDyVsPt = new TH2D("NeighDyVsPt","NeighDyVsPt", 100,0,5, 100, -20,20); + fhNeighDzVsPt = new TH2D("NeighDzVsPt","NeighDzVsPt", 100,0,5, 100, -20,20); + fhNeighChiVsPt = new TH2D("NeighChiVsPt","NeighChiVsPt", 100,0,5, 100, 0,40); + fhNeighNCombVsArea = new TH2D("NeighNCombVsArea","NeighNCombVsArea", 15,0,3, 40, 0,40); + + gDirectory->cd(".."); + + gDirectory->mkdir("TrackFit"); gDirectory->cd("TrackFit"); @@ -359,62 +397,279 @@ void AliHLTTPCCAPerformance::SlicePerformance( Int_t iSlice, Bool_t PrintFlag ) for( ; endSliceHitNHits(); endSliceHit++){ if( fTracker->Hits()[endSliceHit].ISlice()!=iSlice ) break; } - - { // Cell construction performance + + { // Efficiency and quality of found neighbours +#ifdef XXX + for( Int_t iRow=1; iRowNHits()<Hits()[index].ID()]; //cout<=0 ) lb[nla++]= l.fLab[0]; @@ -485,7 +740,7 @@ void AliHLTTPCCAPerformance::SlicePerformance( Int_t iSlice, Bool_t PrintFlag ) } lmax = 0; for( Int_t ihit=0; ihitHits()[index].ID()]; if( l.fLab[0] == labmax || l.fLab[1] == labmax || l.fLab[2] == labmax ) lmax++; @@ -604,10 +859,12 @@ void AliHLTTPCCAPerformance::Performance() fStatNClonesRef=0; */ fStatNEvents++; + for( Int_t islice=0; isliceNSlices(); islice++){ SlicePerformance(islice,0); } + // global tracker performance { if( !fTracker ) return; @@ -926,9 +1183,9 @@ void AliHLTTPCCAPerformance::Performance() <StatTime(2)/fTracker->StatNEvents()*1.e3<<" " <StatTime(3)/fTracker->StatNEvents()*1.e3<<" " <StatTime(4)/fTracker->StatNEvents()*1.e3<<" " - <StatTime(5)/fTracker->StatNEvents()*1.e3<<"[" - <StatTime(6)/fTracker->StatNEvents()*1.e3<<"/" - <StatTime(7)/fTracker->StatNEvents()*1.e3<<"] " + <StatTime(5)/fTracker->StatNEvents()*1.e3<<" " + <StatTime(6)/fTracker->StatNEvents()*1.e3<<" " + <StatTime(7)/fTracker->StatNEvents()*1.e3<<" " <StatTime(8)/fTracker->StatNEvents()*1.e3<<" " <<" msec/event "< +GPUg() void AliHLTTPCCAProcess() +{ + AliHLTTPCCATracker &tracker = *((AliHLTTPCCATracker*) cTracker); + + GPUshared() typename TProcess::AliHLTTPCCASharedMemory smem; + + TProcess::Thread( gridDim.x, blockDim.x, blockIdx.x, threadIdx.x, 0, smem, tracker ); + +#define GPUPROCESS(iSync) \ + if( TProcess::NThreadSyncPoints()>=iSync ){ \ + GPUsync(); \ + TProcess::Thread( gridDim.x, blockDim.x, blockIdx.x, threadIdx.x, iSync, smem, tracker ); \ + } + + GPUPROCESS(1) + GPUPROCESS(2) + GPUPROCESS(3) + + //for( int iSync=0; iSync<=TProcess::NThreadSyncPoints(); iSync++){ + //__syncthreads(); + //TProcess::ThreadGPU( gridDim.x, blockDim.x, blockIdx.x, threadIdx.x, iSync, smem, tracker ); + //} + +#undef GPUPROCESS +} + +#else + +template +GPUg() void AliHLTTPCCAProcess( Int_t nBlocks, Int_t nThreads, AliHLTTPCCATracker &tracker ) +{ + for( Int_t iB=0; iB +GPUg() void AliHLTTPCCAProcess1() +{ + AliHLTTPCCATracker &tracker = *((AliHLTTPCCATracker*) cTracker); + AliHLTTPCCATrackParam1 tParam; + + GPUshared() typename TProcess::AliHLTTPCCASharedMemory sMem; + + typename TProcess::ThreadMemory rMem; + + for( int iSync=0; iSync<=TProcess::NThreadSyncPoints(); iSync++){ + GPUsync(); + TProcess::Thread( gridDim.x, blockDim.x, blockIdx.x, threadIdx.x, iSync, + sMem, rMem, tracker, tParam ); + } +} + +#else + +template +GPUg() void AliHLTTPCCAProcess1( Int_t nBlocks, Int_t nThreads, AliHLTTPCCATracker &tracker ) +{ + for( Int_t iB=0; iBhits - AliHLTTPCCAEndPoint *fEndPoints; // array of track end points - Int_t fNHits, fNCells, fNEndPoints; // number of hits and cells + Int_t fFirstHit; // index of the first hit in the hit array + Int_t fNHits; // number of hits Float_t fX; // X coordinate of the row Float_t fMaxY; // maximal Y coordinate of the row - Float_t fDeltaY; // allowed Y deviation to the next row - Float_t fDeltaZ; // allowed Z deviation to the next row + AliHLTTPCCAGrid fGrid; // grid of hits + + float fHy0,fHz0, fHstepY,fHstepZ, fHstepYi, fHstepZi; // temporary variables + int fFullSize, fFullOffset, fFullGridOffset,fFullLinkOffset; // temporary variables - ClassDef(AliHLTTPCCARow,1); }; #endif diff --git a/HLT/TPCLib/tracking-ca/AliHLTTPCCAStartHitsFinder.cxx b/HLT/TPCLib/tracking-ca/AliHLTTPCCAStartHitsFinder.cxx new file mode 100644 index 00000000000..09bcb247f4a --- /dev/null +++ b/HLT/TPCLib/tracking-ca/AliHLTTPCCAStartHitsFinder.cxx @@ -0,0 +1,72 @@ +// @(#) $Id: AliHLTTPCCAStartHitsFinder.cxx 27042 2008-07-02 12:06:02Z richterm $ +//*************************************************************************** +// This file is property of and copyright by the ALICE HLT Project * +// ALICE Experiment at CERN, All rights reserved. * +// * +// Primary Authors: Sergey Gorbunov * +// Ivan Kisel * +// for The ALICE HLT Project. * +// * +// Permission to use, copy, modify and distribute this software and its * +// documentation strictly for non-commercial purposes is hereby granted * +// without fee, provided that the above copyright notice appears in all * +// copies and that both the copyright notice and this permission notice * +// appear in the supporting documentation. The authors make no claims * +// about the suitability of this software for any purpose. It is * +// provided "as is" without express or implied warranty. * +//*************************************************************************** + +#include "AliHLTTPCCAStartHitsFinder.h" +#include "AliHLTTPCCATracker.h" +#include "AliHLTTPCCAMath.h" + +GPUd() void AliHLTTPCCAStartHitsFinder::Thread +( Int_t /*nBlocks*/, Int_t nThreads, Int_t iBlock, Int_t iThread, Int_t iSync, + AliHLTTPCCASharedMemory &s, AliHLTTPCCATracker &tracker ) +{ + // find start hits for tracklets + + if( iSync==0 ) + { + if( iThread==0 ){ + Int_t *startHits = tracker.StartHits(); + if( iBlock==0 ){ + CAMath::atomicExch(startHits,0); + } + s.fNRows = tracker.Param().NRows(); + s.fIRow = iBlock+1; + s.fNRowStartHits = 0; + if( s.fIRow <= s.fNRows-4 ){ + int first = tracker.Rows()[s.fIRow].FirstHit(); + s.fNHits = tracker.Rows()[s.fIRow].NHits(); + if( s.fNHits>=1024 ) s.fNHits = 1023; + s.fHitLinkUp = tracker.HitLinkUp() + first; + s.fHitLinkDown = tracker.HitLinkDown() + first; + } else s.fNHits = -1; + } + } + else if( iSync==1 ) + { + for( int ih=iThread; ih=0 ) ){ + int oldNRowStartHits = CAMath::atomicAdd(&s.fNRowStartHits,1); + s.fRowStartHits[oldNRowStartHits] = AliHLTTPCCATracker::IRowIHit2ID(s.fIRow, ih); + } + } + } + else if( iSync==2 ) + { + if( iThread == 0 ){ + Int_t *startHits = tracker.StartHits(); + s.fNOldStartHits = CAMath::atomicAdd(startHits,s.fNRowStartHits); + } + } + else if( iSync==3 ) + { + Int_t *startHits = tracker.StartHits(); + for( int ish=iThread; ishcell pointer array - Int_t fNCells; // number of track cells - Int_t fCellID[3]; // ID of first,middle,last cell - Int_t fPointID[2]; // ID of the track endpoints + Int_t fFirstHitID; // index of the first track cell in the track->cell pointer array + Int_t fNHits; // number of track cells + AliHLTTPCCATrackParam fParam; // track parameters - private: - void Dummy(); // to make rulechecker happy by having something in .cxx file +private: + //void Dummy(); // to make rulechecker happy by having something in .cxx file - //ClassDef(AliHLTTPCCATrack,1); + //ClassDef(AliHLTTPCCATrack,1) }; #endif diff --git a/HLT/TPCLib/tracking-ca/AliHLTTPCCATrackConvertor.cxx b/HLT/TPCLib/tracking-ca/AliHLTTPCCATrackConvertor.cxx index 395eb638dff..0a15d32fa0e 100644 --- a/HLT/TPCLib/tracking-ca/AliHLTTPCCATrackConvertor.cxx +++ b/HLT/TPCLib/tracking-ca/AliHLTTPCCATrackConvertor.cxx @@ -19,7 +19,7 @@ #include "AliHLTTPCCATrackConvertor.h" #include "AliExternalTrackParam.h" #include "AliHLTTPCCATrackParam.h" -#include "TMath.h" +#include "AliHLTTPCCAMath.h" void AliHLTTPCCATrackConvertor::GetExtParam( const AliHLTTPCCATrackParam &T1, AliExternalTrackParam &T2, Double_t alpha, Double_t Bz ) @@ -37,7 +37,7 @@ void AliHLTTPCCATrackConvertor::GetExtParam( const AliHLTTPCCATrackParam &T1, Al { // kappa => 1/pt const Double_t kCLight = 0.000299792458; Double_t c = 1.e4; - if( TMath::Abs(Bz)>1.e-4 ) c = 1./(Bz*kCLight); + if( CAMath::Abs(Bz)>1.e-4 ) c = 1./(Bz*kCLight); par[4] *= c; cov[10]*= c; cov[11]*= c; @@ -56,7 +56,7 @@ void AliHLTTPCCATrackConvertor::GetExtParam( const AliHLTTPCCATrackParam &T1, Al cov[10] = -cov[10]; cov[11] = -cov[11]; } - T2.Set((double)T1.GetX(),alpha,par,cov); + T2.Set(T1.GetX(),alpha,par,cov); } void AliHLTTPCCATrackConvertor::SetExtParam( AliHLTTPCCATrackParam &T1, const AliExternalTrackParam &T2, Double_t Bz ) @@ -68,7 +68,7 @@ void AliHLTTPCCATrackConvertor::SetExtParam( AliHLTTPCCATrackParam &T1, const Al T1.X() = T2.GetX(); if(T1.SinPhi()>.99 ) T1.SinPhi()=.99; if(T1.SinPhi()<-.99 ) T1.SinPhi()=-.99; - T1.CosPhi() = TMath::Sqrt(1.-T1.SinPhi()*T1.SinPhi()); + T1.CosPhi() = CAMath::Sqrt(1.-T1.SinPhi()*T1.SinPhi()); const Double_t kCLight = 0.000299792458; Double_t c = Bz*kCLight; { // 1/pt -> kappa diff --git a/HLT/TPCLib/tracking-ca/AliHLTTPCCATrackConvertor.h b/HLT/TPCLib/tracking-ca/AliHLTTPCCATrackConvertor.h index 0d3d8e27b70..5930a24fe8c 100644 --- a/HLT/TPCLib/tracking-ca/AliHLTTPCCATrackConvertor.h +++ b/HLT/TPCLib/tracking-ca/AliHLTTPCCATrackConvertor.h @@ -7,7 +7,8 @@ #ifndef ALIHLTTPCCATRACKCONVERTOR_H #define ALIHLTTPCCATRACKCONVERTOR_H -#include "Rtypes.h" + +#include "AliHLTTPCCADef.h" class AliExternalTrackParam; class AliHLTTPCCATrackParam; diff --git a/HLT/TPCLib/tracking-ca/AliHLTTPCCATrackParam.cxx b/HLT/TPCLib/tracking-ca/AliHLTTPCCATrackParam.cxx index 6e5bf236687..825a9afaaa0 100644 --- a/HLT/TPCLib/tracking-ca/AliHLTTPCCATrackParam.cxx +++ b/HLT/TPCLib/tracking-ca/AliHLTTPCCATrackParam.cxx @@ -17,10 +17,8 @@ //*************************************************************************** #include "AliHLTTPCCATrackParam.h" -#include "TMath.h" -#include "Riostream.h" +#include "AliHLTTPCCAMath.h" -//ClassImp(AliHLTTPCCATrackParam) // // Circle in XY: @@ -32,7 +30,7 @@ -void AliHLTTPCCATrackParam::ConstructXY3( const Float_t x[3], const Float_t y[3], +GPUd() void AliHLTTPCCATrackParam::ConstructXY3( const Float_t x[3], const Float_t y[3], const Float_t sigmaY2[3], Float_t CosPhi0 ) { //* Construct the track in XY plane by 3 points @@ -49,7 +47,7 @@ void AliHLTTPCCATrackParam::ConstructXY3( const Float_t x[3], const Float_t y[3] Float_t a = 2*(x1*y2 - y1*x2); Float_t lx = a1*y2 - a2*y1; Float_t ly = -a1*x2 + a2*x1; - Float_t l = TMath::Sqrt(lx*lx + ly*ly); + Float_t l = CAMath::Sqrt(lx*lx + ly*ly); Float_t li = 1./l; Float_t li2 = li*li; @@ -129,7 +127,7 @@ void AliHLTTPCCATrackParam::ConstructXY3( const Float_t x[3], const Float_t y[3] } -Float_t AliHLTTPCCATrackParam::GetS( Float_t x, Float_t y ) const +GPUd() Float_t AliHLTTPCCATrackParam::GetS( Float_t x, Float_t y ) const { //* Get XY path length to the given point @@ -139,12 +137,12 @@ Float_t AliHLTTPCCATrackParam::GetS( Float_t x, Float_t y ) const x-= GetX(); y-= GetY(); Float_t dS = x*ex + y*ey; - if( TMath::Abs(k)>1.e-4 ) dS = TMath::ATan2( k*dS, 1+k*(x*ey-y*ex) )/k; + if( CAMath::Abs(k)>1.e-4 ) dS = CAMath::ATan2( k*dS, 1+k*(x*ey-y*ex) )/k; return dS; } -void AliHLTTPCCATrackParam::GetDCAPoint( Float_t x, Float_t y, Float_t z, - Float_t &xp, Float_t &yp, Float_t &zp ) const +GPUd() void AliHLTTPCCATrackParam::GetDCAPoint( Float_t x, Float_t y, Float_t z, + Float_t &xp, Float_t &yp, Float_t &zp ) const { //* Get the track point closest to the (x,y,z) @@ -162,15 +160,15 @@ void AliHLTTPCCATrackParam::GetDCAPoint( Float_t x, Float_t y, Float_t z, yp = y0 + (dy + ex*( (dx*dx+dy*dy)*k - 2*(-dx*ey+dy*ex) )/(a+1) )/a; Float_t s = GetS(x,y); zp = GetZ() + GetDzDs()*s; - if( TMath::Abs(k)>1.e-2 ){ - Float_t dZ = TMath::Abs( GetDzDs()*TMath::TwoPi()/k ); + if( CAMath::Abs(k)>1.e-2 ){ + Float_t dZ = CAMath::Abs( GetDzDs()*CAMath::TwoPi()/k ); if( dZ>.1 ){ - zp+= TMath::Nint((z-zp)/dZ)*dZ; + zp+= CAMath::Nint((z-zp)/dZ)*dZ; } } } -void AliHLTTPCCATrackParam::ConstructXYZ3( const Float_t p0[5], const Float_t p1[5], +GPUd() void AliHLTTPCCATrackParam::ConstructXYZ3( const Float_t p0[5], const Float_t p1[5], const Float_t p2[5], Float_t CosPhi0, Float_t t0[] ) { @@ -187,21 +185,21 @@ void AliHLTTPCCATrackParam::ConstructXYZ3( const Float_t p0[5], const Float_t p1 Float_t pS[3] = { GetS(px[0],py[0]), GetS(px[1],py[1]), GetS(px[2],py[2]) }; Float_t k = Kappa(); - if( TMath::Abs(k)>1.e-2 ){ - Float_t dS = TMath::Abs( TMath::TwoPi()/k ); - pS[1]+= TMath::Nint( (pS[0]-pS[1])/dS )*dS; // not more than half turn - pS[2]+= TMath::Nint( (pS[1]-pS[2])/dS )*dS; + if( CAMath::Abs(k)>1.e-2 ){ + Float_t dS = CAMath::Abs( CAMath::TwoPi()/k ); + pS[1]+= CAMath::Nint( (pS[0]-pS[1])/dS )*dS; // not more than half turn + pS[2]+= CAMath::Nint( (pS[1]-pS[2])/dS )*dS; if( t0 ){ - Float_t dZ = TMath::Abs(t0[3]*dS); - if( TMath::Abs(dZ)>1. ){ + Float_t dZ = CAMath::Abs(t0[3]*dS); + if( CAMath::Abs(dZ)>1. ){ Float_t dsDz = 1./t0[3]; if( kold*k<0 ) dsDz = -dsDz; Float_t s0 = (pz[0]-t0[1])*dsDz; Float_t s1 = (pz[1]-t0[1])*dsDz; Float_t s2 = (pz[2]-t0[1])*dsDz; - pS[0]+= TMath::Nint( (s0-pS[0])/dS )*dS ; - pS[1]+= TMath::Nint( (s1-pS[1])/dS )*dS ; - pS[2]+= TMath::Nint( (s2-pS[2])/dS )*dS ; + pS[0]+= CAMath::Nint( (s0-pS[0])/dS )*dS ; + pS[1]+= CAMath::Nint( (s1-pS[1])/dS )*dS ; + pS[2]+= CAMath::Nint( (s2-pS[2])/dS )*dS ; } } } @@ -224,12 +222,10 @@ void AliHLTTPCCATrackParam::ConstructXYZ3( const Float_t p0[5], const Float_t p1 } -Bool_t AliHLTTPCCATrackParam::TransportToX( Float_t x ) +GPUd() Int_t AliHLTTPCCATrackParam::TransportToX( Float_t x, Float_t maxSinPhi ) { //* Transport the track parameters to X=x - Bool_t ret = 1; - Float_t x0 = X(); //Float_t y0 = Y(); Float_t k = Kappa(); @@ -239,49 +235,37 @@ Bool_t AliHLTTPCCATrackParam::TransportToX( Float_t x ) Float_t ey1 = k*dx + ey; Float_t ex1; - if( TMath::Abs(ey1)>1 ){ // no intersection -> check the border - ey1 = ( ey1>0 ) ?1 :-1; - ex1 = 0; - dx = ( TMath::Abs(k)>1.e-4) ? ( (ey1-ey)/k ) :0; - - Float_t ddx = TMath::Abs(x0+dx - x)*k*k; - Float_t hx[] = {0, -k, 1+ey }; - Float_t sx2 = hx[1]*hx[1]*fC[ 3] + hx[2]*hx[2]*fC[ 5]; - if( ddx*ddx>3.5*3.5*sx2 ) ret = 0; // x not withing the error - ret = 0; // any case - return ret; + if( CAMath::Abs(ey1)>maxSinPhi ){ // no intersection + return 0; }else{ - ex1 = TMath::Sqrt(1 - ey1*ey1); + ex1 = CAMath::Sqrt(1 - ey1*ey1); if( ex<0 ) ex1 = -ex1; } Float_t dx2 = dx*dx; - CosPhi() = ex1; Float_t ss = ey+ey1; Float_t cc = ex+ex1; - Float_t tg = 0; - if( TMath::Abs(cc)>1.e-4 ) tg = ss/cc; // tan((phi1+phi)/2) - else ret = 0; + + if( CAMath::Abs(cc)<1.e-4 || CAMath::Abs(ex)<1.e-4 || CAMath::Abs(ex1)<1.e-4 ) return 0; + + Float_t tg = ss/cc; // tan((phi1+phi)/2) + Float_t dy = dx*tg; - Float_t dl = dx*TMath::Sqrt(1+tg*tg); + Float_t dl = dx*CAMath::Sqrt(1+tg*tg); if( cc<0 ) dl = -dl; Float_t dSin = dl*k/2; if( dSin > 1 ) dSin = 1; if( dSin <-1 ) dSin = -1; - Float_t dS = ( TMath::Abs(k)>1.e-4) ? (2*TMath::ASin(dSin)/k) :dl; + Float_t dS = ( CAMath::Abs(k)>1.e-4) ? (2*CAMath::ASin(dSin)/k) :dl; Float_t dz = dS*DzDs(); - Float_t cci = 0, exi = 0, ex1i = 0; - if( TMath::Abs(cc)>1.e-4 ) cci = 1./cc; - else ret = 0; - if( TMath::Abs(ex)>1.e-4 ) exi = 1./ex; - else ret = 0; - if( TMath::Abs(ex1)>1.e-4 ) ex1i = 1./ex1; - else ret = 0; - - if( !ret ) return ret; - + + Float_t cci = 1./cc; + Float_t exi = 1./ex; + Float_t ex1i = 1./ex1; + + CosPhi() = ex1; X() += dx; fP[0]+= dy; fP[1]+= dz; @@ -336,18 +320,19 @@ Bool_t AliHLTTPCCATrackParam::TransportToX( Float_t x ) fC[13]= c43; fC[14]= c44; - return ret; + return 1; } -Bool_t AliHLTTPCCATrackParam::TransportToXWithMaterial( Float_t x, Float_t Bz ) +GPUd() Bool_t AliHLTTPCCATrackParam::TransportToXWithMaterial( Float_t Xto, Float_t Bz ) { + //* Transport the track parameters to X=Xto AliHLTTPCCATrackFitParam par; CalculateFitParameters( par, Bz ); - return TransportToXWithMaterial(x, par ); + return TransportToXWithMaterial(Xto, par ); } -Bool_t AliHLTTPCCATrackParam::TransportToXWithMaterial( Float_t x, AliHLTTPCCATrackFitParam &par ) +GPUd() Bool_t AliHLTTPCCATrackParam::TransportToXWithMaterial( Float_t x, AliHLTTPCCATrackFitParam &par ) { //* Transport the track parameters to X=x @@ -364,19 +349,19 @@ Bool_t AliHLTTPCCATrackParam::TransportToXWithMaterial( Float_t x, AliHLTTPCCAT Float_t ey1 = k*dx + ey; Float_t ex1; - if( TMath::Abs(ey1)>.99 ){ // no intersection -> check the border + if( CAMath::Abs(ey1)>.99 ){ // no intersection -> check the border ey1 = ( ey1>0 ) ?1 :-1; ex1 = 0; - dx = ( TMath::Abs(k)>1.e-4) ? ( (ey1-ey)/k ) :0; + dx = ( CAMath::Abs(k)>1.e-4) ? ( (ey1-ey)/k ) :0; - Float_t ddx = TMath::Abs(x0+dx - x)*k*k; + Float_t ddx = CAMath::Abs(x0+dx - x)*k*k; Float_t hx[] = {0, -k, 1+ey }; Float_t sx2 = hx[1]*hx[1]*fC[ 3] + hx[2]*hx[2]*fC[ 5]; if( ddx*ddx>3.5*3.5*sx2 ) ret = 0; // x not withing the error ret = 0; // any case return ret; }else{ - ex1 = TMath::Sqrt(1 - ey1*ey1); + ex1 = CAMath::Sqrt(1 - ey1*ey1); if( ex<0 ) ex1 = -ex1; } @@ -385,24 +370,24 @@ Bool_t AliHLTTPCCATrackParam::TransportToXWithMaterial( Float_t x, AliHLTTPCCAT Float_t ss = ey+ey1; Float_t cc = ex+ex1; Float_t tg = 0; - if( TMath::Abs(cc)>1.e-4 ) tg = ss/cc; // tan((phi1+phi)/2) + if( CAMath::Abs(cc)>1.e-4 ) tg = ss/cc; // tan((phi1+phi)/2) else ret = 0; Float_t dy = dx*tg; - Float_t dl = dx*TMath::Sqrt(1+tg*tg); + Float_t dl = dx*CAMath::Sqrt(1+tg*tg); if( cc<0 ) dl = -dl; Float_t dSin = dl*k/2; if( dSin > 1 ) dSin = 1; if( dSin <-1 ) dSin = -1; - Float_t dS = ( TMath::Abs(k)>1.e-4) ? (2*TMath::ASin(dSin)/k) :dl; + Float_t dS = ( CAMath::Abs(k)>1.e-4) ? (2*CAMath::ASin(dSin)/k) :dl; Float_t dz = dS*DzDs(); Float_t cci = 0, exi = 0, ex1i = 0; - if( TMath::Abs(cc)>1.e-4 ) cci = 1./cc; + if( CAMath::Abs(cc)>1.e-4 ) cci = 1./cc; else ret = 0; - if( TMath::Abs(ex)>1.e-4 ) exi = 1./ex; + if( CAMath::Abs(ex)>1.e-4 ) exi = 1./ex; else ret = 0; - if( TMath::Abs(ex1)>1.e-4 ) ex1i = 1./ex1; + if( CAMath::Abs(ex1)>1.e-4 ) ex1i = 1./ex1; else ret = 0; if( !ret ) return ret; @@ -461,7 +446,7 @@ Bool_t AliHLTTPCCATrackParam::TransportToXWithMaterial( Float_t x, AliHLTTPCCAT fC[13]= c43; fC[14]= c44; - Float_t d = TMath::Sqrt(dS*dS + dz*dz ); + Float_t d = CAMath::Sqrt(dS*dS + dz*dz ); if (oldX > GetX() ) d = -d; { @@ -475,7 +460,7 @@ Bool_t AliHLTTPCCATrackParam::TransportToXWithMaterial( Float_t x, AliHLTTPCCAT -Float_t AliHLTTPCCATrackParam::ApproximateBetheBloch( Float_t beta2 ) +GPUd() Float_t AliHLTTPCCATrackParam::ApproximateBetheBloch( Float_t beta2 ) { //------------------------------------------------------------------ // This is an approximation of the Bethe-Bloch formula with @@ -490,7 +475,7 @@ Float_t AliHLTTPCCATrackParam::ApproximateBetheBloch( Float_t beta2 ) } -void AliHLTTPCCATrackParam::CalculateFitParameters( AliHLTTPCCATrackFitParam &par, Float_t Bz, Float_t mass ) +GPUd() void AliHLTTPCCATrackParam::CalculateFitParameters( AliHLTTPCCATrackFitParam &par, Float_t Bz, Float_t mass ) { //*! @@ -499,11 +484,11 @@ void AliHLTTPCCATrackParam::CalculateFitParameters( AliHLTTPCCATrackFitParam &pa Float_t p2 = (1.+ fP[3]*fP[3])*c*c; Float_t k2 = fP[4]*fP[4]; Float_t beta2= p2 / (p2 + mass*mass*k2); - Float_t Bethe = ApproximateBetheBloch(beta2); + Float_t bethe = ApproximateBetheBloch(beta2); - Float_t P2 = (k2>1.e-8) ?p2/k2 :10000; // impuls 2 - par.fBethe = Bethe; - par.fE = TMath::Sqrt( P2 + mass*mass); + Float_t pp2 = (k2>1.e-8) ?p2/k2 :10000; // impuls 2 + par.fBethe = bethe; + par.fE = CAMath::Sqrt( pp2 + mass*mass); par.fTheta2 = 14.1*14.1/(beta2*p2*1e6)*k2; par.fEP2 = par.fE/p2*k2; @@ -520,7 +505,7 @@ void AliHLTTPCCATrackParam::CalculateFitParameters( AliHLTTPCCATrackFitParam &pa } -Bool_t AliHLTTPCCATrackParam::CorrectForMeanMaterial( Float_t xOverX0, Float_t xTimesRho, AliHLTTPCCATrackFitParam &par ) +GPUd() Bool_t AliHLTTPCCATrackParam::CorrectForMeanMaterial( Float_t xOverX0, Float_t xTimesRho, AliHLTTPCCATrackFitParam &par ) { //------------------------------------------------------------------ // This function corrects the track parameters for the crossed material. @@ -539,7 +524,7 @@ Bool_t AliHLTTPCCATrackParam::CorrectForMeanMaterial( Float_t xOverX0, Float_t //Energy losses************************ Float_t dE = par.fBethe*xTimesRho; - if ( TMath::Abs(dE) > 0.3*par.fE ) return 0; //30% energy loss is too much! + if ( CAMath::Abs(dE) > 0.3*par.fE ) return 0; //30% energy loss is too much! Float_t corr = (1.- par.fEP2*dE); if( corr<0.3 ) return 0; fP[4]*= corr; @@ -548,12 +533,12 @@ Bool_t AliHLTTPCCATrackParam::CorrectForMeanMaterial( Float_t xOverX0, Float_t fC42*= corr; fC43*= corr; fC44*= corr*corr; - fC44+= par.fSigmadE2*TMath::Abs(dE); + fC44+= par.fSigmadE2*CAMath::Abs(dE); //Multiple scattering****************** - Float_t theta2 = par.fTheta2*TMath::Abs(xOverX0); + Float_t theta2 = par.fTheta2*CAMath::Abs(xOverX0); fC22 += theta2*par.fK22*(1.- fP[2]*fP[2]); fC33 += theta2*par.fK33; fC43 += theta2*par.fK43; @@ -564,19 +549,17 @@ Bool_t AliHLTTPCCATrackParam::CorrectForMeanMaterial( Float_t xOverX0, Float_t -#include "Riostream.h" - -Bool_t AliHLTTPCCATrackParam::Rotate( Float_t alpha ) +GPUd() Bool_t AliHLTTPCCATrackParam::Rotate( Float_t alpha ) { //* Rotate the coordinate system in XY on the angle alpha - Float_t cA = TMath::Cos( alpha ); - Float_t sA = TMath::Sin( alpha ); + Float_t cA = CAMath::Cos( alpha ); + Float_t sA = CAMath::Sin( alpha ); Float_t x = X(), y= Y(), sP= SinPhi(), cP= CosPhi(); Float_t cosPhi = cP*cA + sP*sA; Float_t sinPhi =-cP*sA + sP*cA; - if( TMath::Abs(sinPhi)>.99 || TMath::Abs(cosPhi)<1.e-2 || TMath::Abs(cP)<1.e-2 ) return 0; + if( CAMath::Abs(sinPhi)>.99 || CAMath::Abs(cosPhi)<1.e-2 || CAMath::Abs(cP)<1.e-2 ) return 0; Float_t j0 = cP/cosPhi; Float_t j2 = cosPhi/cP; @@ -610,7 +593,7 @@ Bool_t AliHLTTPCCATrackParam::Rotate( Float_t alpha ) } -Bool_t AliHLTTPCCATrackParam::Filter2( Float_t y, Float_t z, Float_t err2Y, Float_t err2Z ) +GPUd() Bool_t AliHLTTPCCATrackParam::Filter2( Float_t y, Float_t z, Float_t err2Y, Float_t err2Z, Float_t maxSinPhi ) { //* Add the y,z measurement with the Kalman filter @@ -649,7 +632,7 @@ Bool_t AliHLTTPCCATrackParam::Filter2( Float_t y, Float_t z, Float_t err2Y, Floa k40 = c40*mSi[0] + c41*mSi[1]; k41 = c40*mSi[1] + c41*mSi[2] ; Float_t sinPhi = fP[2] + k20*z0 + k21*z1 ; - if( TMath::Abs(sinPhi)>=0.99 ) return 0; + if( CAMath::Abs(sinPhi)>= maxSinPhi ) return 0; fNDF += 2; fChi2 += ( +(mSi[0]*z0 + mSi[1]*z1 )*z0 @@ -683,14 +666,95 @@ Bool_t AliHLTTPCCATrackParam::Filter2( Float_t y, Float_t z, Float_t err2Y, Floa fC[14]-= k40*c40 + k41*c41 ; if( CosPhi()>=0 ){ - CosPhi() = TMath::Sqrt(1-SinPhi()*SinPhi()); + CosPhi() = CAMath::Sqrt(1-SinPhi()*SinPhi()); }else{ - CosPhi() = -TMath::Sqrt(1-SinPhi()*SinPhi()); + CosPhi() = -CAMath::Sqrt(1-SinPhi()*SinPhi()); } return 1; } -void AliHLTTPCCATrackParam::FilterY( Float_t y, Float_t erry ) +GPUd() Bool_t AliHLTTPCCATrackParam::Filter2v1( Float_t y, Float_t z, Float_t err2Y, Float_t err2Z, Float_t maxSinPhi ) +{ + //* Add the y,z measurement with the Kalman filter + + Float_t + c00 = fC[ 0], + c10 = fC[ 1], c11 = fC[ 2], + c20 = fC[ 3], c21 = fC[ 4], + c30 = fC[ 6], c31 = fC[ 7], + c40 = fC[10], c41 = fC[11]; + + err2Y+=c00; + err2Z+=c11; + + Float_t + z0 = y-fP[0], + z1 = z-fP[1]; + + Float_t det = ( err2Y*err2Z - c10*c10); + if( det < 1.e-8 ) return 0; + + det = 1./det; + + Float_t mS0 = err2Z*det; + Float_t mS1 = -c10*det; + Float_t mS2 = err2Y*det; + + // K = CHtS + + Float_t k00, k01 , k10, k11, k20, k21, k30, k31, k40, k41; + + k00 = c00*mS0 + c10*mS1; k01 = c00*mS1 + c10*mS2; + k10 = c10*mS0 + c11*mS1; k11 = c10*mS1 + c11*mS2; + k20 = c20*mS0 + c21*mS1; k21 = c20*mS1 + c21*mS2; + k30 = c30*mS0 + c31*mS1; k31 = c30*mS1 + c31*mS2; + k40 = c40*mS0 + c41*mS1; k41 = c40*mS1 + c41*mS2; + + Float_t sinPhi = fP[2] + k20*z0 + k21*z1 ; + if( CAMath::Abs(sinPhi)>= maxSinPhi ) return 0; + + fNDF += 2; + fChi2 += (mS0*z0 + mS1*z1 )*z0 + (mS1*z0 + mS2*z1 )*z1 ; + + fP[ 0]+= k00*z0 + k01*z1 ; + fP[ 1]+= k10*z0 + k11*z1 ; + fP[ 2] = sinPhi; + fP[ 3]+= k30*z0 + k31*z1 ; + fP[ 4]+= k40*z0 + k41*z1 ; + + + fC[ 0]-= k00*c00 + k01*c10 ; + + fC[ 1]-= k10*c00 + k11*c10 ; + fC[ 2]-= k10*c10 + k11*c11 ; + + fC[ 3]-= k20*c00 + k21*c10 ; + fC[ 4]-= k20*c10 + k21*c11 ; + fC[ 5]-= k20*c20 + k21*c21 ; + + fC[ 6]-= k30*c00 + k31*c10 ; + fC[ 7]-= k30*c10 + k31*c11 ; + fC[ 8]-= k30*c20 + k31*c21 ; + fC[ 9]-= k30*c30 + k31*c31 ; + + fC[10]-= k40*c00 + k41*c10 ; + fC[11]-= k40*c10 + k41*c11 ; + fC[12]-= k40*c20 + k41*c21 ; + fC[13]-= k40*c30 + k41*c31 ; + fC[14]-= k40*c40 + k41*c41 ; + + if( CosPhi()>=0 ){ + CosPhi() = CAMath::Sqrt(1-sinPhi*sinPhi); + }else{ + CosPhi() = -CAMath::Sqrt(1-sinPhi*sinPhi); + } + return 1; +} + + + + +GPUd() void AliHLTTPCCATrackParam::FilterY( Float_t y, Float_t erry ) { //* Add the y measurement with the Kalman filter @@ -705,7 +769,7 @@ void AliHLTTPCCATrackParam::FilterY( Float_t y, Float_t erry ) z0 = y-fP[0]; Float_t s = { c00+erry*erry }; - if( TMath::Abs(s)<1.e-4 ) return; + if( CAMath::Abs(s)<1.e-4 ) return; Float_t si = 1/s; @@ -723,7 +787,7 @@ void AliHLTTPCCATrackParam::FilterY( Float_t y, Float_t erry ) k4 = c40*si; Float_t sinPhi = fP[2] + k2*z0 ; - if( TMath::Abs(sinPhi)>=0.99 ) return; + if( CAMath::Abs(sinPhi)>=0.99 ) return; fP[ 0]+= k0*z0 ; fP[ 1]+= k1*z0 ; @@ -752,14 +816,14 @@ void AliHLTTPCCATrackParam::FilterY( Float_t y, Float_t erry ) fC[14]-= k4*c40; if( CosPhi()>=0 ){ - CosPhi() = TMath::Sqrt(1-SinPhi()*SinPhi()); + CosPhi() = CAMath::Sqrt(1-SinPhi()*SinPhi()); }else{ - CosPhi() = -TMath::Sqrt(1-SinPhi()*SinPhi()); + CosPhi() = -CAMath::Sqrt(1-SinPhi()*SinPhi()); } } -void AliHLTTPCCATrackParam::FilterZ( Float_t z, Float_t errz ) +GPUd() void AliHLTTPCCATrackParam::FilterZ( Float_t z, Float_t errz ) { //* Add the z measurement with the Kalman filter @@ -774,7 +838,7 @@ void AliHLTTPCCATrackParam::FilterZ( Float_t z, Float_t errz ) z1 = z-fP[1]; Float_t s = c11 + errz*errz; - if( TMath::Abs(s)<1.e-4 ) return; + if( CAMath::Abs(s)<1.e-4 ) return; Float_t si = 1./s; @@ -792,7 +856,7 @@ void AliHLTTPCCATrackParam::FilterZ( Float_t z, Float_t errz ) k4 = 0;//c41*si; Float_t sinPhi = fP[2] + k2*z1 ; - if( TMath::Abs(sinPhi)>=0.99 ) return; + if( CAMath::Abs(sinPhi)>=0.99 ) return; fP[ 0]+= k0*z1 ; fP[ 1]+= k1*z1 ; @@ -822,15 +886,27 @@ void AliHLTTPCCATrackParam::FilterZ( Float_t z, Float_t errz ) fC[14]-= k4*c41 ; if( CosPhi()>=0 ){ - CosPhi() = TMath::Sqrt(1-SinPhi()*SinPhi()); + CosPhi() = CAMath::Sqrt(1-SinPhi()*SinPhi()); }else{ - CosPhi() = -TMath::Sqrt(1-SinPhi()*SinPhi()); + CosPhi() = -CAMath::Sqrt(1-SinPhi()*SinPhi()); } } -void AliHLTTPCCATrackParam::Print() const +#if !defined(HLTCA_GPUCODE) +#if defined( HLTCA_STANDALONE ) +#include +#else +#include "Riostream.h" +#endif +#endif + +GPUd() void AliHLTTPCCATrackParam::Print() const { + //* print parameters + +#if !defined(HLTCA_GPUCODE) cout<<"track: "< * +// Ivan Kisel * +// for The ALICE HLT Project. * +// * +// Permission to use, copy, modify and distribute this software and its * +// documentation strictly for non-commercial purposes is hereby granted * +// without fee, provided that the above copyright notice appears in all * +// copies and that both the copyright notice and this permission notice * +// appear in the supporting documentation. The authors make no claims * +// about the suitability of this software for any purpose. It is * +// provided "as is" without express or implied warranty. * +//*************************************************************************** + +#include "AliHLTTPCCATrackParam1.h" +#include "AliHLTTPCCAMath.h" + + +// +// Circle in XY: +// +// R = 1/TMath::Abs(Kappa); +// Xc = X - sin(Phi)/Kappa; +// Yc = Y + cos(Phi)/Kappa; +// + + + +GPUd() void AliHLTTPCCATrackParam1::ConstructXY3( const Float_t x[3], const Float_t y[3], + const Float_t sigmaY2[3], Float_t CosPhi0 ) +{ + //* Construct the track in XY plane by 3 points + + Float_t x0 = x[0]; + Float_t y0 = y[0]; + Float_t x1 = x[1] - x0; + Float_t y1 = y[1] - y0; + Float_t x2 = x[2] - x0; + Float_t y2 = y[2] - y0; + + Float_t a1 = x1*x1 + y1*y1; + Float_t a2 = x2*x2 + y2*y2; + Float_t a = 2*(x1*y2 - y1*x2); + Float_t lx = a1*y2 - a2*y1; + Float_t ly = -a1*x2 + a2*x1; + Float_t l = CAMath::Sqrt(lx*lx + ly*ly); + + Float_t li = 1./l; + Float_t li2 = li*li; + Float_t li3 = li2*li; + Float_t cosPhi = ly*li; + + Float_t sinPhi = -lx*li; + Float_t kappa = a/l; + + Float_t dlx = a2 - a1; // D lx / D y0 + Float_t dly = -a; // D ly / D y0 + Float_t dA = 2*(x2 - x1); // D a / D y0 + Float_t dl = (lx*dlx + ly*dly)*li; + + // D sinPhi,kappa / D y0 + + Float_t d0[2] = { -(dlx*ly-lx*dly)*ly*li3, (dA*l-a*dl)*li2 }; + + // D sinPhi,kappa / D y1 + + dlx = -a2 + 2*y1*y2; + dly = -2*x2*y1; + dA = -2*x2; + dl = (lx*dlx + ly*dly)*li; + + Float_t d1[2] = { -(dlx*ly-lx*dly)*ly*li3, (dA*l-a*dl)*li2 }; + + // D sinPhi,kappa / D y2 + + dlx = a1 - 2*y1*y2; + dly = -2*x1*y2; + dA = 2*x1; + dl = (lx*dlx + ly*dly)*li; + + Float_t d2[2] = { -(dlx*ly-lx*dly)*ly*li3, (dA*l-a*dl)*li2 }; + + if( CosPhi0*cosPhi <0 ){ + cosPhi = -cosPhi; + sinPhi = -sinPhi; + kappa = -kappa; + d0[0] = -d0[0]; + d0[1] = -d0[1]; + d1[0] = -d1[0]; + d1[1] = -d1[1]; + d2[0] = -d2[0]; + d2[1] = -d2[1]; + } + + X() = x0; + Y() = y0; + SinPhi() = sinPhi; + Kappa() = kappa; + CosPhi() = cosPhi; + + Float_t s0 = sigmaY2[0]; + Float_t s1 = sigmaY2[1]; + Float_t s2 = sigmaY2[2]; + + fC[0] = s0; + fC[1] = 0; + fC[2] = 100.; + + fC[3] = d0[0]*s0; + fC[4] = 0; + fC[5] = d0[0]*d0[0]*s0 + d1[0]*d1[0]*s1 + d2[0]*d2[0]*s2; + + fC[6] = 0; + fC[7] = 0; + fC[8] = 0; + fC[9] = 100.; + + fC[10] = d0[1]*s0; + fC[11] = 0; + fC[12] = d0[0]*d0[1]*s0 + d1[0]*d1[1]*s1 + d2[0]*d2[1]*s2; + fC[13] = 0; + fC[14] = d0[1]*d0[1]*s0 + d1[1]*d1[1]*s1 + d2[1]*d2[1]*s2; +} + + +GPUd() Float_t AliHLTTPCCATrackParam1::GetS( Float_t x, Float_t y ) const +{ + //* Get XY path length to the given point + + Float_t k = GetKappa(); + Float_t ex = GetCosPhi(); + Float_t ey = GetSinPhi(); + x-= GetX(); + y-= GetY(); + Float_t dS = x*ex + y*ey; + if( CAMath::Abs(k)>1.e-4 ) dS = CAMath::ATan2( k*dS, 1+k*(x*ey-y*ex) )/k; + return dS; +} + +GPUd() void AliHLTTPCCATrackParam1::GetDCAPoint( Float_t x, Float_t y, Float_t z, + Float_t &xp, Float_t &yp, Float_t &zp ) const +{ + //* Get the track point closest to the (x,y,z) + + Float_t x0 = GetX(); + Float_t y0 = GetY(); + Float_t k = GetKappa(); + Float_t ex = GetCosPhi(); + Float_t ey = GetSinPhi(); + Float_t dx = x - x0; + Float_t dy = y - y0; + Float_t ax = dx*k+ey; + Float_t ay = dy*k-ex; + Float_t a = sqrt( ax*ax+ay*ay ); + xp = x0 + (dx - ey*( (dx*dx+dy*dy)*k - 2*(-dx*ey+dy*ex) )/(a+1) )/a; + yp = y0 + (dy + ex*( (dx*dx+dy*dy)*k - 2*(-dx*ey+dy*ex) )/(a+1) )/a; + Float_t s = GetS(x,y); + zp = GetZ() + GetDzDs()*s; + if( CAMath::Abs(k)>1.e-2 ){ + Float_t dZ = CAMath::Abs( GetDzDs()*CAMath::TwoPi()/k ); + if( dZ>.1 ){ + zp+= CAMath::Nint((z-zp)/dZ)*dZ; + } + } +} + +GPUd() void AliHLTTPCCATrackParam1::ConstructXYZ3( const Float_t p0[5], const Float_t p1[5], + const Float_t p2[5], + Float_t CosPhi0, Float_t t0[] ) +{ + //* Construct the track in XYZ by 3 points + + Float_t px[3] = { p0[0], p1[0], p2[0] }; + Float_t py[3] = { p0[1], p1[1], p2[1] }; + Float_t pz[3] = { p0[2], p1[2], p2[2] }; + Float_t ps2y[3] = { p0[3]*p0[3], p1[3]*p1[3], p2[3]*p2[3] }; + Float_t ps2z[3] = { p0[4]*p0[4], p1[4]*p1[4], p2[4]*p2[4] }; + + Float_t kold = t0 ?t0[4] :0; + ConstructXY3( px, py, ps2y, CosPhi0 ); + + Float_t pS[3] = { GetS(px[0],py[0]), GetS(px[1],py[1]), GetS(px[2],py[2]) }; + Float_t k = Kappa(); + if( CAMath::Abs(k)>1.e-2 ){ + Float_t dS = CAMath::Abs( CAMath::TwoPi()/k ); + pS[1]+= CAMath::Nint( (pS[0]-pS[1])/dS )*dS; // not more than half turn + pS[2]+= CAMath::Nint( (pS[1]-pS[2])/dS )*dS; + if( t0 ){ + Float_t dZ = CAMath::Abs(t0[3]*dS); + if( CAMath::Abs(dZ)>1. ){ + Float_t dsDz = 1./t0[3]; + if( kold*k<0 ) dsDz = -dsDz; + Float_t s0 = (pz[0]-t0[1])*dsDz; + Float_t s1 = (pz[1]-t0[1])*dsDz; + Float_t s2 = (pz[2]-t0[1])*dsDz; + pS[0]+= CAMath::Nint( (s0-pS[0])/dS )*dS ; + pS[1]+= CAMath::Nint( (s1-pS[1])/dS )*dS ; + pS[2]+= CAMath::Nint( (s2-pS[2])/dS )*dS ; + } + } + } + + Float_t s = pS[0] + pS[1] + pS[2]; + Float_t z = pz[0] + pz[1] + pz[2]; + Float_t sz = pS[0]*pz[0] + pS[1]*pz[1] + pS[2]*pz[2]; + Float_t ss = pS[0]*pS[0] + pS[1]*pS[1] + pS[2]*pS[2]; + + Float_t a = 3*ss-s*s; + Z() = (z*ss-sz*s)/a; // z0 + DzDs() = (3*sz-z*s)/a; // t = dz/ds + + Float_t dz0[3] = {ss - pS[0]*s,ss - pS[1]*s,ss - pS[2]*s }; + Float_t dt [3] = {3*pS[0] - s, 3*pS[1] - s, 3*pS[2] - s }; + + fC[2] = (dz0[0]*dz0[0]*ps2z[0] + dz0[1]*dz0[1]*ps2z[1] + dz0[2]*dz0[2]*ps2z[2])/a/a; + fC[7]= (dz0[0]*dt [0]*ps2z[0] + dz0[1]*dt [1]*ps2z[1] + dz0[2]*dt [2]*ps2z[2])/a/a; + fC[9]= (dt [0]*dt [0]*ps2z[0] + dt [1]*dt [1]*ps2z[1] + dt [2]*dt [2]*ps2z[2])/a/a; +} + + +GPUd() Int_t AliHLTTPCCATrackParam1::TransportToX( Float_t x, Float_t maxSinPhi ) +{ + //* Transport the track parameters to X=x + + Float_t x0 = X(); + //Float_t y0 = Y(); + Float_t k = Kappa(); + Float_t ex = CosPhi(); + Float_t ey = SinPhi(); + Float_t dx = x - x0; + + Float_t ey1 = k*dx + ey; + Float_t ex1; + if( CAMath::Abs(ey1)>maxSinPhi ){ // no intersection + return 0; + }else{ + ex1 = CAMath::Sqrt(1 - ey1*ey1); + if( ex<0 ) ex1 = -ex1; + } + + Float_t dx2 = dx*dx; + Float_t ss = ey+ey1; + Float_t cc = ex+ex1; + + if( CAMath::Abs(cc)<1.e-4 || CAMath::Abs(ex)<1.e-4 || CAMath::Abs(ex1)<1.e-4 ) return 0; + + Float_t tg = ss/cc; // tan((phi1+phi)/2) + + Float_t dy = dx*tg; + Float_t dl = dx*CAMath::Sqrt(1+tg*tg); + + if( cc<0 ) dl = -dl; + Float_t dSin = dl*k/2; + if( dSin > 1 ) dSin = 1; + if( dSin <-1 ) dSin = -1; + Float_t dS = ( CAMath::Abs(k)>1.e-4) ? (2*CAMath::ASin(dSin)/k) :dl; + Float_t dz = dS*DzDs(); + + + Float_t cci = 1./cc; + Float_t exi = 1./ex; + Float_t ex1i = 1./ex1; + + CosPhi() = ex1; + X() += dx; + fP[0]+= dy; + fP[1]+= dz; + fP[2] = ey1; + fP[3] = fP[3]; + fP[4] = fP[4]; + + Float_t h2 = dx*(1+ ex*ex1 + ey*ey1 )*cci*exi*ex1i; + Float_t h4 = dx2*(cc + ss*ey1*ex1i )*cci*cci; + + //Float_t H0[5] = { 1,0, h2, 0, h4 }; + //Float_t H1[5] = { 0, 1, 0, dS, 0 }; + //Float_t H2[5] = { 0, 0, 1, 0, dx }; + //Float_t H3[5] = { 0, 0, 0, 1, 0 }; + //Float_t H4[5] = { 0, 0, 0, 0, 1 }; + + Float_t c00 = fC[0]; + Float_t c11 = fC[2]; + Float_t c20 = fC[3]; + Float_t c22 = fC[5]; + Float_t c31 = fC[7]; + Float_t c33 = fC[9]; + Float_t c40 = fC[10]; + Float_t c42 = fC[12]; + Float_t c44 = fC[14]; + + float dSc33 = dS*c33; + float h2c22 = h2*c22; + float h4c42 = h4*c42; + float h4c44 = h4*c44; + float h2c42 = h2*c42; + + Float_t fC10 = c40 + h2c42 + h4c44; + Float_t fC12 = c42 + dx*c44; + Float_t fC7 = c31 + dSc33; + fC[0]= c00 + h2*h2c22 + h4*h4c44 + 2*( h2*c20 + h4*c40 + h2*h4c42 ) ; + fC[3]= c20 + h2c22 + h4c42 + dx*fC10; + fC[10]= fC10; + fC[12]= fC12; + fC[5]= c22 + dx*( c42 + fC12 ); + + fC[7]= fC7; + fC[2]= c11 + dS*(c31 + fC7); + + return 1; +} + +GPUd() Int_t AliHLTTPCCATrackParam1::TransportToX0( Float_t x, Float_t /**/ ) +{ + //* Transport the track parameters to X=x + + Float_t ex = fCosPhi; + if( CAMath::Abs(ex)<1.e-4 ) return 0; + + Float_t ey = fP[2]; + Float_t dx = x - fX; + Float_t exi = 1./ex; + Float_t dS = dx*exi; + Float_t dy = dS*ey; + Float_t dz = dS*fP[3]; + Float_t h2 = dS*exi*exi; + Float_t h4 = 0.5*dx*h2; + + fX = x; + fP[0]+= dy + h4*fP[4]; + fP[1]+= dz; + fP[2]+= dx*fP[4]; + + //Float_t H0[5] = { 1,0, h2, 0, h4 }; + //Float_t H1[5] = { 0, 1, 0, dS, 0 }; + //Float_t H2[5] = { 0, 0, 1, 0, dx }; + //Float_t H3[5] = { 0, 0, 0, 1, 0 }; + //Float_t H4[5] = { 0, 0, 0, 0, 1 }; + + + Float_t c00 = fC[0]; + Float_t c11 = fC[2]; + Float_t c20 = fC[3]; + Float_t c22 = fC[5]; + Float_t c31 = fC[7]; + Float_t c33 = fC[9]; + Float_t c40 = fC[10]; + Float_t c42 = fC[12]; + Float_t c44 = fC[14]; + + Float_t c40ph4c44 = c40 + h4*c44; + Float_t c20ph4c42 = c20 + h4*c42; + Float_t c20ph2c22ph4c42 = h2*c22 + c20ph4c42; + Float_t fC10 = c40ph4c44 + h2*c42; + Float_t fC12 = c42 + dx*c44; + Float_t fC7 = c31 + dS*c33; + + fC[10]= fC10; + fC[12]= fC12; + fC[7]= fC7; + fC[0]= c00 + h2*( c20ph2c22ph4c42 + c20ph4c42) + h4*(c40ph4c44 + c40 ) ; + fC[3]= c20ph2c22ph4c42 + dx*fC10; + fC[5]= c22 + dx*( c42 + fC12 ); + fC[2]= c11 + dS*(c31 + fC7); + + return 1; +} + +/* +GPUd() Int_t AliHLTTPCCATrackParam1::TransportToX0( Float_t x, Float_t maxSinPhi ) +{ +// Transport the track parameters to X=x + + Float_t x0 = X(); + Float_t ex = CosPhi(); + Float_t ey = SinPhi(); + Float_t dx = x - x0; + + if( CAMath::Abs(ex)<1.e-4 ) return 0; + + Float_t exi = 1./ex; + + Float_t dl = dx*exi; + Float_t dy = dl*ey; + Float_t dS = dl; + Float_t dz = dS*DzDs(); + + CosPhi() = ex; + X() += dx; + fP[0]+= dy; + fP[1]+= dz; + fP[2] = ey; + fP[3] = fP[3]; + fP[4] = fP[4]; + + Float_t h2 = dx*exi*exi*exi; + Float_t h4 = dx*h2/2; + + //Float_t H0[5] = { 1,0, h2, 0, h4 }; + //Float_t H1[5] = { 0, 1, 0, dS, 0 }; + //Float_t H2[5] = { 0, 0, 1, 0, dx }; + //Float_t H3[5] = { 0, 0, 0, 1, 0 }; + //Float_t H4[5] = { 0, 0, 0, 0, 1 }; + + Float_t c00 = fC[0]; + Float_t c11 = fC[2]; + Float_t c20 = fC[3]; + Float_t c22 = fC[5]; + Float_t c31 = fC[7]; + Float_t c33 = fC[9]; + Float_t c40 = fC[10]; + Float_t c42 = fC[12]; + Float_t c44 = fC[14]; + + float dSc33 = dS*c33; + float h2c22 = h2*c22; + float h4c42 = h4*c42; + float h4c44 = h4*c44; + float h2c42 = h2*c42; + + Float_t fC10 = c40 + h2c42 + h4c44; + Float_t fC12 = c42 + dx*c44; + Float_t fC7 = c31 + dSc33; + fC[0]= c00 + h2*h2c22 + h4*h4c44 + 2*( h2*c20 + h4*c40 + h2*h4c42 ) ; + fC[3]= c20 + h2c22 + h4c42 + dx*fC10; + fC[10]= fC10; + fC[12]= fC12; + fC[5]= c22 + dx*( c42 + fC12 ); + + fC[7]= fC7; + fC[2]= c11 + dS*(c31 + fC7); + + return 1; +} +*/ + +GPUd() Int_t AliHLTTPCCATrackParam1::TransportToXMatrix( Float_t x, + Float_t &OutY, Float_t &OutZ, + Float_t &OutErr2Y, Float_t &OutErr2Z, + Float_t &H2, Float_t &H4, + Float_t &dS, Float_t &dx, + Float_t maxSinPhi ) +{ + //* Transport the track parameters to X=x + + Float_t x0 = X(); + //Float_t y0 = Y(); + Float_t k = Kappa(); + Float_t ex = CosPhi(); + Float_t ey = SinPhi(); + dx = x - x0; + + Float_t ey1 = k*dx + ey; + Float_t ex1; + if( CAMath::Abs(ey1)>maxSinPhi ){ // no intersection + return 0; + }else{ + ex1 = CAMath::Sqrt(1 - ey1*ey1); + if( ex<0 ) ex1 = -ex1; + } + + Float_t dx2 = dx*dx; + Float_t ss = ey+ey1; + Float_t cc = ex+ex1; + + if( CAMath::Abs(cc)<1.e-4 || CAMath::Abs(ex)<1.e-4 || CAMath::Abs(ex1)<1.e-4 ) return 0; + + Float_t tg = ss/cc; // tan((phi1+phi)/2) + + Float_t dy = dx*tg; + Float_t dl = dx*CAMath::Sqrt(1+tg*tg); + + if( cc<0 ) dl = -dl; + Float_t dSin = dl*k/2; + if( dSin > 1 ) dSin = 1; + if( dSin <-1 ) dSin = -1; + + dS = ( CAMath::Abs(k)>1.e-4) ? (2*CAMath::ASin(dSin)/k) :dl; + + Float_t dz = dS*DzDs(); + + + Float_t cci = 1./cc; + Float_t exi = 1./ex; + Float_t ex1i = 1./ex1; + + Float_t h2 = dx*(1+ ex*ex1 + ey*ey1 )*cci*exi*ex1i; + Float_t h4 = dx2*(cc + ss*ey1*ex1i )*cci*cci; + + //Float_t H0[5] = { 1,0, h2, 0, h4 }; + //Float_t H1[5] = { 0, 1, 0, dS, 0 }; + //Float_t H2[5] = { 0, 0, 1, 0, dx }; + //Float_t H3[5] = { 0, 0, 0, 1, 0 }; + //Float_t H4[5] = { 0, 0, 0, 0, 1 }; + + Float_t c20 = fC[3]; + Float_t c22 = fC[5]; + Float_t c31 = fC[7]; + Float_t c33 = fC[9]; + Float_t c40 = fC[10]; + Float_t c42 = fC[12]; + Float_t c44 = fC[14]; + + float dSc33 = dS*c33; + float h2c22 = h2*c22; + float h4c42 = h4*c42; + float h4c44 = h4*c44; + + OutY = fP[0] + dy; + OutZ = fP[1] + dz; + OutErr2Y = fC[0] + h2*h2c22 + h4*h4c44 + 2*( h2*c20 + h4*c40 + h2*h4c42 ) ; + OutErr2Z = fC[2] + dS*(c31 + c31 + dSc33); + + H2 = h2; + H4 = h4; + + return 1; +} + + +GPUd() Bool_t AliHLTTPCCATrackParam1::TransportToXWithMaterial( Float_t Xto, Float_t Bz ) +{ + //* Transport the track parameters to X=x + AliHLTTPCCATrackFitParam par; + CalculateFitParameters( par, Bz ); + return TransportToXWithMaterial(Xto, par ); +} + + +GPUd() Bool_t AliHLTTPCCATrackParam1::TransportToXWithMaterial( Float_t x, AliHLTTPCCATrackFitParam &par ) +{ + //* Transport the track parameters to X=x + + Bool_t ret = 1; + + Float_t oldX=GetX(); + + Float_t x0 = X(); + //Float_t y0 = Y(); + Float_t k = Kappa(); + Float_t ex = CosPhi(); + Float_t ey = SinPhi(); + Float_t dx = x - x0; + + Float_t ey1 = k*dx + ey; + Float_t ex1; + if( CAMath::Abs(ey1)>.99 ){ // no intersection -> check the border + ey1 = ( ey1>0 ) ?1 :-1; + ex1 = 0; + dx = ( CAMath::Abs(k)>1.e-4) ? ( (ey1-ey)/k ) :0; + + Float_t ddx = CAMath::Abs(x0+dx - x)*k*k; + Float_t hx[] = {0, -k, 1+ey }; + Float_t sx2 = hx[1]*hx[1]*fC[ 3] + hx[2]*hx[2]*fC[ 5]; + if( ddx*ddx>3.5*3.5*sx2 ) ret = 0; // x not withing the error + ret = 0; // any case + return ret; + }else{ + ex1 = CAMath::Sqrt(1 - ey1*ey1); + if( ex<0 ) ex1 = -ex1; + } + + Float_t dx2 = dx*dx; + CosPhi() = ex1; + Float_t ss = ey+ey1; + Float_t cc = ex+ex1; + Float_t tg = 0; + if( CAMath::Abs(cc)>1.e-4 ) tg = ss/cc; // tan((phi1+phi)/2) + else ret = 0; + Float_t dy = dx*tg; + Float_t dl = dx*CAMath::Sqrt(1+tg*tg); + + if( cc<0 ) dl = -dl; + Float_t dSin = dl*k/2; + if( dSin > 1 ) dSin = 1; + if( dSin <-1 ) dSin = -1; + Float_t dS = ( CAMath::Abs(k)>1.e-4) ? (2*CAMath::ASin(dSin)/k) :dl; + Float_t dz = dS*DzDs(); + + Float_t cci = 0, exi = 0, ex1i = 0; + if( CAMath::Abs(cc)>1.e-4 ) cci = 1./cc; + else ret = 0; + if( CAMath::Abs(ex)>1.e-4 ) exi = 1./ex; + else ret = 0; + if( CAMath::Abs(ex1)>1.e-4 ) ex1i = 1./ex1; + else ret = 0; + + if( !ret ) return ret; + + X() += dx; + fP[0]+= dy; + fP[1]+= dz; + fP[2] = ey1; + fP[3] = fP[3]; + fP[4] = fP[4]; + + Float_t h2 = dx*(1+ ex*ex1 + ey*ey1 )*cci*exi*ex1i; + Float_t h4 = dx2*(cc + ss*ey1*ex1i )*cci*cci; + + Float_t c00 = fC[0]; + Float_t c11 = fC[2]; + Float_t c20 = fC[3]; + Float_t c22 = fC[5]; + Float_t c31 = fC[7]; + Float_t c33 = fC[9]; + Float_t c40 = fC[10]; + Float_t c42 = fC[12]; + Float_t c44 = fC[14]; + + //Float_t H0[5] = { 1,0, h2, 0, h4 }; + //Float_t H1[5] = { 0, 1, 0, dS, 0 }; + //Float_t H2[5] = { 0, 0, 1, 0, dx }; + //Float_t H3[5] = { 0, 0, 0, 1, 0 }; + //Float_t H4[5] = { 0, 0, 0, 0, 1 }; + + + fC[0]=( c00 + h2*h2*c22 + h4*h4*c44 + + 2*( h2*c20 + h4*c40 + h2*h4*c42 ) ); + + fC[2]= c11 + 2*dS*c31 + dS*dS*c33; + + fC[3]= c20 + h2*c22 + h4*c42 + dx*( c40 + h2*c42 + h4*c44); + fC[5]= c22 +2*dx*c42 + dx2*c44; + + fC[7]= c31 + dS*c33; + fC[9]= c33; + + fC[10]= c40 + h2*c42 + h4*c44; + fC[12]= c42 + dx*c44; + fC[14]= c44; + + Float_t d = CAMath::Sqrt(dS*dS + dz*dz ); + + if (oldX > GetX() ) d = -d; + { + Float_t rho=0.9e-3; + Float_t radLen=28.94; + CorrectForMeanMaterial(d*rho/radLen,d*rho,par); + } + + return ret; +} + + + +GPUd() Float_t AliHLTTPCCATrackParam1::ApproximateBetheBloch( Float_t beta2 ) +{ + //------------------------------------------------------------------ + // This is an approximation of the Bethe-Bloch formula with + // the density effect taken into account at beta*gamma > 3.5 + // (the approximation is reasonable only for solid materials) + //------------------------------------------------------------------ + if (beta2 >= 1) return 0; + + if (beta2/(1-beta2)>3.5*3.5) + return 0.153e-3/beta2*( log(3.5*5940)+0.5*log(beta2/(1-beta2)) - beta2); + return 0.153e-3/beta2*(log(5940*beta2/(1-beta2)) - beta2); +} + + +GPUd() void AliHLTTPCCATrackParam1::CalculateFitParameters( AliHLTTPCCATrackFitParam &par, Float_t Bz, Float_t mass ) +{ + //*! + + const Float_t kCLight = 0.000299792458; + Float_t c = Bz*kCLight; + Float_t p2 = (1.+ fP[3]*fP[3])*c*c; + Float_t k2 = fP[4]*fP[4]; + Float_t beta2= p2 / (p2 + mass*mass*k2); + Float_t bethe = ApproximateBetheBloch(beta2); + + Float_t pp2 = (k2>1.e-8) ?p2/k2 :10000; // impuls 2 + par.fBethe = bethe; + par.fE = CAMath::Sqrt( pp2 + mass*mass); + par.fTheta2 = 14.1*14.1/(beta2*p2*1e6)*k2; + par.fEP2 = par.fE/p2*k2; + + // Approximate energy loss fluctuation (M.Ivanov) + + const Float_t knst=0.07; // To be tuned. + par.fSigmadE2 = knst*par.fEP2*fP[4]; + par.fSigmadE2 = par.fSigmadE2 * par.fSigmadE2; + + par.fK22 = (1. + fP[3]*fP[3]); + par.fK33 = par.fK22*par.fK22; + par.fK43 = fP[3]*fP[4]*par.fK22; + par.fK44 = fP[3]*fP[3]*fP[4]*fP[4]; +} + + +GPUd() Bool_t AliHLTTPCCATrackParam1::CorrectForMeanMaterial( Float_t xOverX0, Float_t xTimesRho, AliHLTTPCCATrackFitParam &par ) +{ + //------------------------------------------------------------------ + // This function corrects the track parameters for the crossed material. + // "xOverX0" - X/X0, the thickness in units of the radiation length. + // "xTimesRho" - is the product length*density (g/cm^2). + //------------------------------------------------------------------ + + Float_t &fC22=fC[5]; + Float_t &fC33=fC[9]; + Float_t &fC40=fC[10]; + Float_t &fC42=fC[12]; + Float_t &fC44=fC[14]; + + //Energy losses************************ + + Float_t dE = par.fBethe*xTimesRho; + if ( CAMath::Abs(dE) > 0.3*par.fE ) return 0; //30% energy loss is too much! + Float_t corr = (1.- par.fEP2*dE); + if( corr<0.3 ) return 0; + fP[4]*= corr; + fC40*= corr; + fC42*= corr; + fC44*= corr*corr; + fC44+= par.fSigmadE2*CAMath::Abs(dE); + + + //Multiple scattering****************** + + Float_t theta2 = par.fTheta2*CAMath::Abs(xOverX0); + fC22 += theta2*par.fK22*(1.- fP[2]*fP[2]); + fC33 += theta2*par.fK33; + fC44 += theta2*par.fK44; + + return 1; +} + + + +//#include "Riostream.h" + +GPUd() Bool_t AliHLTTPCCATrackParam1::Rotate( Float_t alpha ) +{ + //* Rotate the coordinate system in XY on the angle alpha + + Float_t cA = CAMath::Cos( alpha ); + Float_t sA = CAMath::Sin( alpha ); + Float_t x = X(), y= Y(), sP= SinPhi(), cP= CosPhi(); + Float_t cosPhi = cP*cA + sP*sA; + Float_t sinPhi =-cP*sA + sP*cA; + + if( CAMath::Abs(sinPhi)>.99 || CAMath::Abs(cosPhi)<1.e-2 || CAMath::Abs(cP)<1.e-2 ) return 0; + + Float_t j0 = cP/cosPhi; + Float_t j2 = cosPhi/cP; + + X() = x*cA + y*sA; + Y() = -x*sA + y*cA; + CosPhi() = cosPhi; + SinPhi() = sinPhi; + + + //Float_t J[5][5] = { { j0, 0, 0, 0, 0 }, // Y + // { 0, 1, 0, 0, 0 }, // Z + // { 0, 0, j2, 0, 0 }, // SinPhi + // { 0, 0, 0, 1, 0 }, // DzDs + // { 0, 0, 0, 0, 1 } }; // Kappa + //cout<<"alpha="<= maxSinPhi ) return 0; + + Float_t cosPhi = CAMath::Sqrt(1-sinPhi*sinPhi); + fNDF += 2; + fChi2 += mS0*z0*z0 + mS2*z1*z1 ; + + fP[ 0]+= k00*z0 ; + fP[ 1]+= k11*z1 ; + fP[ 2] = sinPhi ; + fP[ 3]+= k31*z1 ; + fP[ 4]+= k40*z0 ; + + fC[ 0]-= k00*c00 ; + fC[ 3]-= k20*c00 ; + fC[ 5]-= k20*c20 ; + fC[10]-= k40*c00 ; + fC[12]-= k40*c20 ; + fC[14]-= k40*c40 ; + + fC[ 2]-= k11*c11 ; + fC[ 7]-= k31*c11 ; + fC[ 9]-= k31*c31 ; + + fCosPhi = ( fCosPhi >=0 ) ?cosPhi :-cosPhi; + return 1; +} + + + + +GPUd() void AliHLTTPCCATrackParam1::FilterY( Float_t y, Float_t erry ) +{ + //* Add the y measurement with the Kalman filter + + Float_t + c00 = fC[ 0], + c20 = fC[ 3], + c40 = fC[10]; + + Float_t + z0 = y-fP[0]; + + Float_t s = { c00+erry*erry }; + if( CAMath::Abs(s)<1.e-4 ) return; + + Float_t si = 1/s; + + fNDF += 1; + fChi2 += si*z0*z0; + + // K = CHtS + + Float_t k0, k2, k3, k4; + + k0 = c00*si; + k2 = c20*si; + k3 = 0; + k4 = c40*si; + + Float_t sinPhi = fP[2] + k2*z0 ; + if( CAMath::Abs(sinPhi)>=0.99 ) return; + + fP[ 0]+= k0*z0 ; + fP[ 2] = sinPhi; + fP[ 3]+= k3*z0 ; + fP[ 4]+= k4*z0 ; + + fC[ 0]-= k0*c00; + + fC[ 3]-= k2*c00; + fC[ 5]-= k2*c20; + + fC[10]-= k4*c00; + fC[12]-= k4*c20; + fC[14]-= k4*c40; + + if( CosPhi()>=0 ){ + CosPhi() = CAMath::Sqrt(1-SinPhi()*SinPhi()); + }else{ + CosPhi() = -CAMath::Sqrt(1-SinPhi()*SinPhi()); + } + +} + +GPUd() void AliHLTTPCCATrackParam1::FilterZ( Float_t z, Float_t errz ) +{ + //* Add the z measurement with the Kalman filter + + Float_t + c11 = fC[ 2], + c31 = fC[ 7]; + + Float_t + z1 = z-fP[1]; + + Float_t s = c11 + errz*errz; + if( CAMath::Abs(s)<1.e-4 ) return; + + Float_t si = 1./s; + + fNDF += 1; + fChi2 += si*z1*z1; + + // K = CHtS + + Float_t k1 , k3; + + k1 = c11*si; + k3 = c31*si; + + + fP[ 1]+= k1*z1 ; + fP[ 3]+= k3*z1 ; + + fC[ 2]-= k1*c11 ; + + fC[ 7]-= k3*c11 ; + fC[ 9]-= k3*c31 ; + +} + +#if !defined(HLTCA_GPUCODE) +#if defined( HLTCA_STANDALONE ) +#include +#else +#include "Riostream.h" +#endif +#endif + +GPUd() void AliHLTTPCCATrackParam1::Print() const +{ + //* Print parameters + +#if !defined(HLTCA_GPUCODE) + cout<<"track: "< #include "TStopwatch.h" +#include "AliHLTTPCCAHitArea.h" +#include "AliHLTTPCCANeighboursFinder.h" +#include "AliHLTTPCCANeighboursCleaner.h" +#include "AliHLTTPCCAStartHitsFinder.h" +#include "AliHLTTPCCATrackletConstructor.h" +#include "AliHLTTPCCATrackletSelector.h" +#include "AliHLTTPCCAProcess.h" +#include "AliHLTTPCCALinksWriter.h" +#include "AliHLTTPCCAUsedHitsInitialiser.h" + +#include "AliHLTTPCCATrackParam.h" +#include "AliHLTTPCCATrackParam1.h" + +#if !defined(HLTCA_GPUCODE) +#if defined( HLTCA_STANDALONE ) +#include +#else +#include "Riostream.h" +#endif +#endif //#define DRAW @@ -40,9 +55,35 @@ ClassImp(AliHLTTPCCATracker) +#if !defined(HLTCA_GPUCODE) AliHLTTPCCATracker::AliHLTTPCCATracker() - :fParam(),fRows(0),fOutTrackHits(0),fNOutTrackHits(0),fOutTracks(0),fNOutTracks(0),fNHitsTotal(0),fTracks(0),fNTracks(0),fCellHitPointers(0),fCells(0),fEndPoints(0) + : + fParam(), + fNHitsTotal(0), + fGridSizeTotal(0), + fGrid1SizeTotal(0), + fHits(0), + fHits1(0), + fGridContents(0), + fGrid1Contents(0), + fHitsID(0), + fHitLinkUp(0), + fHitLinkDown(0), + fHitIsUsed(0), + fStartHits(0), + fTracklets(0), + fNTracks(0), + fTracks(0), + fTrackHits(0), + fNOutTracks(0), + fNOutTrackHits(0), + fOutTracks(0), + fOutTrackHits(0), + fEventMemory(0), + fEventMemSize(0), + fTexHitsFullData(0), + fTexHitsFullSize(0) { // constructor //fRows = new AliHLTTPCCARow[fParam.NRows()]; @@ -50,7 +91,32 @@ AliHLTTPCCATracker::AliHLTTPCCATracker() } AliHLTTPCCATracker::AliHLTTPCCATracker( const AliHLTTPCCATracker& ) - :fParam(),fRows(0),fOutTrackHits(0),fNOutTrackHits(0),fOutTracks(0),fNOutTracks(0),fNHitsTotal(0),fTracks(0),fNTracks(0),fCellHitPointers(0),fCells(0),fEndPoints(0) + : + fParam(), + fNHitsTotal(0), + fGridSizeTotal(0), + fGrid1SizeTotal(0), + fHits(0), + fHits1(0), + fGridContents(0), + fGrid1Contents(0), + fHitsID(0), + fHitLinkUp(0), + fHitLinkDown(0), + fHitIsUsed(0), + fStartHits(0), + fTracklets(0), + fNTracks(0), + fTracks(0), + fTrackHits(0), + fNOutTracks(0), + fNOutTrackHits(0), + fOutTracks(0), + fOutTrackHits(0), + fEventMemory(0), + fEventMemSize(0), + fTexHitsFullData(0), + fTexHitsFullSize(0) { // dummy } @@ -58,1092 +124,610 @@ AliHLTTPCCATracker::AliHLTTPCCATracker( const AliHLTTPCCATracker& ) AliHLTTPCCATracker &AliHLTTPCCATracker::operator=( const AliHLTTPCCATracker& ) { // dummy - fRows=0; fOutTrackHits=0; fOutTracks=0; fNOutTracks=0; + fTrackHits = 0; + fEventMemory = 0; return *this; } -AliHLTTPCCATracker::~AliHLTTPCCATracker() +GPUd() AliHLTTPCCATracker::~AliHLTTPCCATracker() { // destructor StartEvent(); - delete[] fRows; } +#endif + + +GPUd() UChar_t AliHLTTPCCATracker::GetGridContent( UInt_t i ) const +{ + //* get grid content +#if defined(HLTCA_GPUSTEP) + return (UChar_t) tex1Dfetch(texGrid,i).x; +#else + return fGridContents[i]; +#endif +} + + +GPUd() AliHLTTPCCAHit AliHLTTPCCATracker::GetHit( UInt_t i ) const +{ + //* get hit +#if defined(HLTCA_USE_GPU) + AliHLTTPCCAHit h; + float2 f = tex1Dfetch(texHits,i); + h.Y() = f.x; + h.Z() = f.y; + return h; +#else + return fHits[i]; +#endif +} + + + // ---------------------------------------------------------------------------------- -void AliHLTTPCCATracker::Initialize( AliHLTTPCCAParam ¶m ) +GPUd() void AliHLTTPCCATracker::Initialize( AliHLTTPCCAParam ¶m ) { // initialisation StartEvent(); - delete[] fRows; - fRows = 0; - fParam = param; - fParam.Update(); - fRows = new AliHLTTPCCARow[fParam.NRows()]; - Float_t xStep = 1; - Float_t deltaY = TMath::Tan(fParam.CellConnectionAngleXY()); - Float_t deltaZ = TMath::Tan(fParam.CellConnectionAngleXZ()); + fParam = param; + fParam.Update(); for( Int_t irow=0; irow Y[j] ) yMin = Y[j]; + if( zMax < Z[j] ) zMax = Z[j]; + if( zMin > Z[j] ) zMin = Z[j]; + } + if( nGrid == 0 ){ + yMin = yMax = zMin = zMax = 0; + nGrid = 1; + } + + row.Grid().Create( yMin, yMax, zMin, zMax, nGrid ); + + float sy = ( CAMath::Abs( row.Grid().StepYInv() ) >1.e-4 ) ?1./row.Grid().StepYInv() :1; + float sz = ( CAMath::Abs( row.Grid().StepZInv() ) >1.e-4 ) ?1./row.Grid().StepZInv() :1; + + //cout<<"grid n = "< ymax ) continue; - if( cont1.Zmax < zmin ) break;// in the grid cell hit Y is decreasing - if( cont1.Zmin > zmax ) continue; - - if( cont1.Ymin < ymin ){ ymin = cont1.Ymin; repeat = 1; } - if( cont1.Ymax > ymax ){ ymax = cont1.Ymax; repeat = 1; } - if( cont1.Zmin < zmin ){ zmin = cont1.Zmin; repeat = 1; } - if( cont1.Zmax > zmax ){ zmax = cont1.Zmax; repeat = 1; } - if( cont1.binYmin < binYmin ){ binYmin = cont1.binYmin; repeat = 1; } - if( cont1.binYmax > binYmax ){ binYmax = cont1.binYmax; repeat = 1; } - if( cont1.binZmin < binZmin ){ binZmin = cont1.binZmin; repeat = 1; } - if( cont1.binZmax > binZmax ){ binZmax = cont1.binZmax; repeat = 1; } - - row.CellHitPointers()[nPointers++] = jh; - cell.NHits()++; - cont1.used = 1; -#ifdef DRAW - //AliHLTTPCCADisplay::Instance().DrawHit( irow, jh, kRed ); - //AliHLTTPCCADisplay::Instance().Ask(); -#endif - } - } - } + grid.Offset() = fGridSizeTotal; + Int_t off= grid.N()+3+10; + fGridSizeTotal+=off; + Int_t n2 = grid.N()/2; + grid.Content2() = c[n2]; + UChar_t *cnew = fGridContents + grid.Offset(); + + for( Int_t i=0; i=256 ){ + cout<<" ERROR!!! "<=256 ){ + cout<<" ERROR 1 !!! "< 63 ) g0n = 63; + if( g1n > 63 ) g1n = 63; + cnew1[i] = (g1n<<26) + (g1<<16) + (g0n<<10) + g0; + } + { + float y0 = row.Grid().YMin(); + float stepY = (row.Grid().YMax() - y0)*(1./65535.); + float z0 = row.Grid().ZMin(); + float stepZ = (row.Grid().ZMax() - z0)*(1./65535.); + float stepYi = 1./stepY; + float stepZi = 1./stepZ; - Float_t deltaY = row1.DeltaY(); - Float_t deltaZ = row1.DeltaZ(); - Float_t xStep = 1; - if( iRow1 < fParam.NRows()-1 ) xStep = fParam.RowX(iRow1+1) - fParam.RowX(iRow1); - Float_t tx = xStep/row1.X(); - - Int_t lastRow2 = iRow1+3; - if( lastRow2>=fParam.NRows() ) lastRow2 = fParam.NRows()-1; - - for (Int_t i1 = 0; i1 yMax ) continue; - if( y2Max < yMin ) continue; - if( z2Min > zMax ) break; - if( z2Max < zMin ) continue; - - // c1 & c2 are neighbours - - found = 1; - - if( c1.Link() ==-1 && c2.Status()==0 ){ - // one-to-one connection - OK - c1.Link() = IRowICell2ID(iRow2,i2); - c2.Status() = 1; - }else{ - // multi-connection - break all links - if( c1.Link()>=0 ) ID2Cell(c1.Link()).Status() = -1; - c1.Link() = -2; - c2.Status() = -1; - } + for( Int_t ih=0; ih=65536 || yy>= 65536 ){ + cout<<"!!!! hit packing error!!! "<0 ) continue; // not a starting cell - - Int_t firstID = IRowICell2ID( iRow, iCell ); - Int_t midID = firstID; - Int_t lastID = firstID; - - AliHLTTPCCATrack &track = fTracks[fNTracks]; - track.Alive() = 1; - track.NCells() = 1; - AliHLTTPCCACell *last = &c; - while( last->Link() >=0 ){ - Int_t nextID = last->Link(); - AliHLTTPCCACell *next = & ID2Cell(nextID); - if(next->Status()!=1 ){ - last->Link() = -1; - break; - } - track.NCells()++; - last = next; - lastID = nextID; + h.x = (UShort_t) xx;//((hh.Y() - y0)*stepYi); + h.y = (UShort_t) yy;//((hh.Z() - z0)*stepZi); } - Int_t nCells05 = (track.NCells()-1)/2; - for( Int_t i=0; iTrackID() = fNTracks; - track.FirstCellID() = firstID; - track.CellID()[0] = firstID; - track.CellID()[1] = midID; - track.CellID()[2] = lastID; - track.PointID()[0] = -1; - track.PointID()[1] = -1; - //cout<<"Track N "<0 ){ - AliHLTTPCCADisplay::Instance().Ask(); - } + AliHLTTPCCADisplay::Instance().SetSliceView(); + AliHLTTPCCADisplay::Instance().DrawSlice( this ); + //for( Int_t iRow=0; iRow0 ) - AliHLTTPCCADisplay::Instance().Ask(); -#endif + fTimers[0] = 0; // find neighbours + fTimers[1] = 0; // construct tracklets + fTimers[2] = 0; // fit tracklets + fTimers[3] = 0; // prolongation of tracklets + fTimers[4] = 0; // selection + fTimers[5] = 0; // write output + fTimers[6] = 0; + fTimers[7] = 0; + + if( fNHitsTotal < 1 ) return; + //if( fParam.ISlice()!=3 ) return; + TStopwatch timer0; + *fNTracks = 0; +#if !defined(HLTCA_GPUCODE) - TStopwatch timer4; + AliHLTTPCCAProcess( Param().NRows(), 1, *this ); + AliHLTTPCCAProcess( Param().NRows()-2, 1, *this ); + AliHLTTPCCAProcess( Param().NRows()-4, 1, *this ); - Bool_t doMerging=1;//SG!!! + Int_t nStartHits = *fStartHits; - Float_t factor2 = fParam.TrackConnectionFactor()*fParam.TrackConnectionFactor(); - Int_t *refEndPoints = new Int_t[fNHitsTotal]; - Int_t nRefEndPoints = 0; - for( Int_t iRow=0; iRow("<(nBlocks, nThreads,*this); - //cout<<"merge neighbours"<(nBlocks, nThreads,*this); + + Int_t nMemThreads = 128; + nThreads = 256;//96; + nBlocks = nStartHits/nThreads + 1; + if( nBlocks<30 ){ + nBlocks = 30; + nThreads = (nStartHits)/30+1; + if( nThreads%32 ) nThreads = (nThreads/32+1)*32; + } - Int_t nRefEndPointsNew = 0; - for( Int_t iRef=0; iRef(nBlocks, nMemThreads+nThreads,*this); - Int_t ipID = IRowICell2ID(iRow,iPoint); - Int_t jpID = ip.Link(); - AliHLTTPCCAEndPoint &jp = ID2Point(jpID); - - if( jp.Link()!=ipID ){ - //cout<<"broken link: jp.Link()!=iID"<("<=0 && jc.Link()>=0 ) ){ - - Int_t currID = jTrack.CellID()[0]; - jTrack.CellID()[0] = jTrack.CellID()[2]; - jTrack.CellID()[2] = currID; - - Int_t pID = jTrack.PointID()[0]; - jTrack.PointID()[0] = jTrack.PointID()[1]; - jTrack.PointID()[1] = pID; - - currID = jTrack.FirstCellID(); - Int_t lastID = -1; - while( currID>=0 ){ - AliHLTTPCCACell &c = ID2Cell( currID ); - Int_t nextID = c.Link(); - c.Link() = lastID; - lastID = currID; - currID = nextID; - } - jTrack.FirstCellID() = lastID; - } - //cout<<"track i "<jTrack - ic.Link() = jcID; - iTrack.PointID()[1] = jTrack.PointID()[1]; - ID2Point(iTrack.PointID()[1]).TrackID() = itr; - if( jTrack.NCells()<3 ){ - refEndPoints[nRefEndPointsNew++] = iTrack.PointID()[1]; - doMerging = 1; - ID2Point(iTrack.PointID()[1]).Param() = ip.Param();// just to set phi direction - } - if( iTrack.NCells()<3 ){ - refEndPoints[nRefEndPointsNew++] = iTrack.PointID()[0]; - doMerging = 1; - ID2Point(iTrack.PointID()[0]).Param() = jp.Param();// just to set phi direction - } - - if( TMath::Abs(ID2Cell(jTrack.CellID()[2]).Z()-ID2Cell(iTrack.CellID()[0]).Z())> - TMath::Abs(ID2Cell(iTrack.CellID()[2]).Z()-ID2Cell(iTrack.CellID()[0]).Z()) ){ - iTrack.CellID()[2] = jTrack.CellID()[2]; - } - }else{ //match jTrack->iTrack - jc.Link() = icID; - iTrack.FirstCellID()=jTrack.FirstCellID(); - iTrack.PointID()[0] = jTrack.PointID()[0]; - ID2Point(iTrack.PointID()[0]).TrackID() = itr; - if( jTrack.NCells()<3 ){ - refEndPoints[nRefEndPointsNew++] = iTrack.PointID()[0]; - doMerging = 1; - ID2Point(iTrack.PointID()[0]).Param() = ip.Param(); // just to set phi direction - } - if( iTrack.NCells()<3 ){ - refEndPoints[nRefEndPointsNew++] = iTrack.PointID()[1]; - doMerging = 1; - ID2Point(iTrack.PointID()[1]).Param() = jp.Param();// just to set phi direction - } - if( TMath::Abs(ID2Cell(jTrack.CellID()[0]).Z()-ID2Cell(iTrack.CellID()[2]).Z())> - TMath::Abs(ID2Cell(iTrack.CellID()[0]).Z()-ID2Cell(iTrack.CellID()[2]).Z()) ){ - iTrack.CellID()[0] = jTrack.CellID()[0]; - } - } - //cout<<"merged ID: "<(nBlocks, nThreads,*this); - if( jTrack.NCells()>iTrack.NCells() ){ - iTrack.CellID()[1] = jTrack.CellID()[1]; - } - - AliHLTTPCCAEndPoint &p0 = ID2Point(iTrack.PointID()[0]); - AliHLTTPCCAEndPoint &p1 = ID2Point(iTrack.PointID()[1]); - - if( p0.Link() == iTrack.PointID()[1] ){ - p0.Link() = -1; - p1.Link() = -1; - } - //cout<<" NCells itr/jtr= "<0 ) AliHLTTPCCADisplay::Instance().Ask(); -#endif + timer0.Stop(); + fTimers[0] = timer0.CpuTime(); + } -#ifdef DRAWXX - AliHLTTPCCADisplay::Instance().ClearView(); - AliHLTTPCCADisplay::Instance().DrawSlice( this ); - for( Int_t iRow=0; iRowNHits(); iHit++ ){ - AliHLTTPCCAHit &h = row.GetCellHit(*ic,iHit); - - // check for wrong hits - { - Float_t ddy = t0.GetY() - h.Y(); - Float_t ddz = t0.GetZ() - h.Z(); - Float_t dd = ddy*ddy+ddz*ddz; - if( dd 3.5*3.5*(/*t0.GetErr2Y() + */h.ErrY()*h.ErrY() ) ) continue;//SG!!! - //if( dz*dz > 3.5*3.5*(/*t0.GetErr2Z() + */h.ErrZ()*h.ErrZ() ) ) continue; - //if( !t0.Filter2( h.Y(), h.Z(), h.ErrY()*h.ErrY(), h.ErrZ()*h.ErrZ() ) ) continue; - - - if( !t.TransportToX( row.X() ) ) continue; - - //* Update the track - - if( first ){ - t.Cov()[ 0] = .5*.5; - t.Cov()[ 1] = 0; - t.Cov()[ 2] = .5*.5; - t.Cov()[ 3] = 0; - t.Cov()[ 4] = 0; - t.Cov()[ 5] = .2*.2; - t.Cov()[ 6] = 0; - t.Cov()[ 7] = 0; - t.Cov()[ 8] = 0; - t.Cov()[ 9] = .2*.2; - t.Cov()[10] = 0; - t.Cov()[11] = 0; - t.Cov()[12] = 0; - t.Cov()[13] = 0; - t.Cov()[14] = .2*.2; - t.Chi2() = 0; - t.NDF() = -5; - } - - if( t.Filter2( h.Y(), h.Z(), h.ErrY()*h.ErrY(), h.ErrZ()*h.ErrZ() ) ) first = 0; - else continue; - - fOutTrackHits[fNOutTrackHits] = h.ID(); + //cout<<12<<" "<fNHitsTotal ){ + //cout<<"ok"<fNHitsTotal*10 ){ cout<<"fNOutTrackHits>fNHitsTotal"< 3.5*3.5*(/*t0.GetErr2Y() + */h.ErrY()*h.ErrY() ) ) continue;//SG!!! + //if( dz*dz > 3.5*3.5*(/*t0.GetErr2Z() + */h.ErrZ()*h.ErrZ() ) ) continue; + } -void AliHLTTPCCATracker::FitTrack( AliHLTTPCCATrack &track, Float_t t0[] ) const -{ - //* Fit the track - - AliHLTTPCCAEndPoint &p0 = ID2Point(track.PointID()[0]); - AliHLTTPCCAEndPoint &p2 = ID2Point(track.PointID()[1]); - AliHLTTPCCACell &c0 = ID2Cell(p0.CellID()); - AliHLTTPCCACell &c1 = ID2Cell(track.CellID()[1]); - AliHLTTPCCACell &c2 = ID2Cell(p2.CellID()); - AliHLTTPCCARow &row0 = ID2Row(p0.CellID()); - AliHLTTPCCARow &row1 = ID2Row(track.CellID()[1]); - AliHLTTPCCARow &row2 = ID2Row(p2.CellID()); + if( !t.TransportToX( row.X() ) ) continue; + //* Update the track + + if( first ){ + t.Cov()[ 0] = .5*.5; + t.Cov()[ 1] = 0; + t.Cov()[ 2] = .5*.5; + t.Cov()[ 3] = 0; + t.Cov()[ 4] = 0; + t.Cov()[ 5] = .2*.2; + t.Cov()[ 6] = 0; + t.Cov()[ 7] = 0; + t.Cov()[ 8] = 0; + t.Cov()[ 9] = .2*.2; + t.Cov()[10] = 0; + t.Cov()[11] = 0; + t.Cov()[12] = 0; + t.Cov()[13] = 0; + t.Cov()[14] = .2*.2; + t.Chi2() = 0; + t.NDF() = -5; + } + Float_t err2Y, err2Z; + GetErrors2( iRow, t, err2Y, err2Z ); + + if( !t.Filter2( h.Y(), h.Z(), err2Y, err2Z ) ) continue; + + first = 0; + } + /* + Float_t cosPhi = iTrack.Param().GetCosPhi(); + p0.Param().TransportToX(ID2Row( iTrack.PointID()[0] ).X()); + p2.Param().TransportToX(ID2Row( iTrack.PointID()[1] ).X()); + if( p0.Param().GetCosPhi()*cosPhi<0 ){ // change direction + Float_t *par = p0.Param().Par(); + Float_t *cov = p0.Param().Cov(); + par[2] = -par[2]; // sin phi + par[3] = -par[3]; // DzDs + par[4] = -par[4]; // kappa + cov[3] = -cov[3]; + cov[4] = -cov[4]; + cov[6] = -cov[6]; + cov[7] = -cov[7]; + cov[10] = -cov[10]; + cov[11] = -cov[11]; + p0.Param().CosPhi() = -p0.Param().GetCosPhi(); + } + */ +#endif +} +GPUh() void AliHLTTPCCATracker::FitTrack( AliHLTTPCCATrack &/*track*/, Float_t */*t0[]*/ ) const +{ + //* Fit the track +#ifdef XXX + AliHLTTPCCAEndPoint &p2 = ID2Point(track.PointID()[1]); + AliHLTTPCCAHit &c0 = ID2Hit(fTrackHits[p0.TrackHitID()].HitID()); + AliHLTTPCCAHit &c1 = ID2Hit(fTrackHits[track.HitID()[1]].HitID()); + AliHLTTPCCAHit &c2 = ID2Hit(fTrackHits[p2.TrackHitID()].HitID()); + AliHLTTPCCARow &row0 = ID2Row(fTrackHits[p0.TrackHitID()].HitID()); + AliHLTTPCCARow &row1 = ID2Row(fTrackHits[track.HitID()[1]].HitID()); + AliHLTTPCCARow &row2 = ID2Row(fTrackHits[p2.TrackHitID()].HitID()); Float_t sp0[5] = {row0.X(), c0.Y(), c0.Z(), c0.ErrY(), c0.ErrZ() }; Float_t sp1[5] = {row1.X(), c1.Y(), c1.Z(), c1.ErrY(), c1.ErrZ() }; Float_t sp2[5] = {row2.X(), c2.Y(), c2.Z(), c2.ErrY(), c2.ErrZ() }; - if( track.NCells()>=3 ){ + //cout<<"Fit track, points ="<=3 ){ p0.Param().ConstructXYZ3(sp0,sp1,sp2,p0.Param().CosPhi(), t0); p2.Param().ConstructXYZ3(sp2,sp1,sp0,p2.Param().CosPhi(), t0); //p2.Param() = p0.Param(); @@ -1162,4 +746,39 @@ void AliHLTTPCCATracker::FitTrack( AliHLTTPCCATrack &track, Float_t t0[] ) const p2.Param().Err2Y() = c2.ErrY()*c2.ErrY(); p2.Param().Err2Z() = c2.ErrZ()*c2.ErrZ(); } +#endif +} + + + +GPUd() void AliHLTTPCCATracker::GetErrors2( Int_t iRow, const AliHLTTPCCATrackParam &t, Float_t &Err2Y, Float_t &Err2Z ) const +{ + // + // Use calibrated cluster error from OCDB + // + + Float_t z = CAMath::Abs((250.-0.275)-CAMath::Abs(t.GetZ())); + Int_t type = (iRow<63) ? 0: (iRow>126) ? 1:2; + Float_t cosPhiInv = CAMath::Abs(t.GetCosPhi())>1.e-2 ?1./t.GetCosPhi() :0; + Float_t angleY = t.GetSinPhi()*cosPhiInv ; + Float_t angleZ = t.GetDzDs()*cosPhiInv ; + + Err2Y = fParam.GetClusterError2(0,type, z,angleY); + Err2Z = fParam.GetClusterError2(1,type, z,angleZ); +} + +GPUd() void AliHLTTPCCATracker::GetErrors2( Int_t iRow, const AliHLTTPCCATrackParam1 &t, Float_t &Err2Y, Float_t &Err2Z ) const +{ + // + // Use calibrated cluster error from OCDB + // + + Float_t z = CAMath::Abs((250.-0.275)-CAMath::Abs(t.GetZ())); + Int_t type = (iRow<63) ? 0: (iRow>126) ? 1:2; + Float_t cosPhiInv = CAMath::Abs(t.GetCosPhi())>1.e-2 ?1./t.GetCosPhi() :0; + Float_t angleY = t.GetSinPhi()*cosPhiInv ; + Float_t angleZ = t.GetDzDs()*cosPhiInv ; + + Err2Y = fParam.GetClusterError2(0,type, z,angleY); + Err2Z = fParam.GetClusterError2(1,type, z,angleZ); } diff --git a/HLT/TPCLib/tracking-ca/AliHLTTPCCATracker.h b/HLT/TPCLib/tracking-ca/AliHLTTPCCATracker.h index f9d50354261..549c88ad3da 100644 --- a/HLT/TPCLib/tracking-ca/AliHLTTPCCATracker.h +++ b/HLT/TPCLib/tracking-ca/AliHLTTPCCATracker.h @@ -9,16 +9,17 @@ #define ALIHLTTPCCATRACKER_H -#include "Rtypes.h" +#include "AliHLTTPCCADef.h" #include "AliHLTTPCCAParam.h" #include "AliHLTTPCCARow.h" +#include "AliHLTTPCCAHit.h" class AliHLTTPCCATrack; -class AliHLTTPCCAHit; -class AliHLTTPCCACell; class AliHLTTPCCAOutTrack; class AliHLTTPCCATrackParam; -class AliHLTTPCCAEndPoint; +class AliHLTTPCCATrackParam1; + + /** * @class AliHLTTPCCATracker @@ -36,79 +37,122 @@ class AliHLTTPCCATracker { public: +#if !defined(HLTCA_GPUCODE) AliHLTTPCCATracker(); AliHLTTPCCATracker( const AliHLTTPCCATracker& ); AliHLTTPCCATracker &operator=( const AliHLTTPCCATracker& ); - virtual ~AliHLTTPCCATracker(); + GPUd() ~AliHLTTPCCATracker(); +#endif - void Initialize( AliHLTTPCCAParam ¶m ); + GPUd() void Initialize( AliHLTTPCCAParam ¶m ); - void StartEvent(); + GPUd() void StartEvent(); - void ReadHitRow( Int_t iRow, AliHLTTPCCAHit *Row, Int_t NHits ); + GPUd() void ReadEvent( Int_t *RowFirstHit, Int_t *RowNHits, Float_t *Y, Float_t *Z, Int_t NHits ); void Reconstruct(); - void FindCells(); - void MergeCells(); - void FindTracks(); + void WriteOutput(); + + GPUd() void GetErrors2( Int_t iRow, const AliHLTTPCCATrackParam &t, Float_t &Err2Y, Float_t &Err2Z ) const; + GPUd() void GetErrors2( Int_t iRow, const AliHLTTPCCATrackParam1 &t, Float_t &Err2Y, Float_t &Err2Z ) const; - AliHLTTPCCAParam &Param(){ return fParam; } - AliHLTTPCCARow *Rows(){ return fRows; } + GPUhd() AliHLTTPCCAParam &Param(){ return fParam; } + GPUhd() AliHLTTPCCARow *Rows(){ return fRows; } + + Int_t * HitsID(){ return fHitsID; } Int_t *OutTrackHits(){ return fOutTrackHits; } Int_t NOutTrackHits() const { return fNOutTrackHits; } - AliHLTTPCCAOutTrack *OutTracks(){ return fOutTracks; } - Int_t NOutTracks() const { return fNOutTracks; } + GPUd() AliHLTTPCCAOutTrack *OutTracks(){ return fOutTracks; } + GPUd() Int_t NOutTracks() const { return fNOutTracks; } + GPUhd() Int_t *TrackHits(){ return fTrackHits; } - AliHLTTPCCATrack *Tracks(){ return fTracks; } - Int_t NTracks() const { return fNTracks; } + GPUhd() AliHLTTPCCATrack *Tracks(){ return fTracks; } + GPUhd() Int_t &NTracks() { return *fNTracks; } Double_t *Timers(){ return fTimers; } - static Int_t IRowICell2ID( Int_t iRow, Int_t iCell ){ - return (iCell<<8)+iRow; + GPUhd() static Int_t IRowIHit2ID( Int_t iRow, Int_t iHit ){ + return (iHit<<8)+iRow; } - static Int_t ID2IRow( Int_t CellID ){ - return ( CellID%256 ); + GPUhd() static Int_t ID2IRow( Int_t HitID ){ + return ( HitID%256 ); } - static Int_t ID2ICell( Int_t CellID ){ - return ( CellID>>8 ); + GPUhd() static Int_t ID2IHit( Int_t HitID ){ + return ( HitID>>8 ); } - AliHLTTPCCACell &ID2Cell( Int_t CellID ) const{ - return fRows[CellID%256].Cells()[CellID>>8]; + + GPUhd() AliHLTTPCCAHit &ID2Hit( Int_t HitID ) { + return fHits[fRows[HitID%256].FirstHit() + (HitID>>8)]; } - AliHLTTPCCARow &ID2Row( Int_t CellID ) const{ - return fRows[CellID%256]; + GPUhd() AliHLTTPCCARow &ID2Row( Int_t HitID ) { + return fRows[HitID%256]; } - AliHLTTPCCAEndPoint &ID2Point( Int_t PointID ) const{ - return fRows[PointID%256].EndPoints()[PointID>>8]; - } - void FitTrack( AliHLTTPCCATrack &track, Float_t *t0 = 0 ) const; + void FitTrackFull( AliHLTTPCCATrack &track, Float_t *t0 = 0 ) const; + GPUhd() void SetPointers(); + + GPUhd() Short_t *HitLinkUp(){ return fHitLinkUp;} + GPUhd() Short_t *HitLinkDown(){ return fHitLinkDown;} + GPUhd() Int_t *StartHits(){ return fStartHits;} + GPUhd() Int_t *Tracklets(){ return fTracklets;} + GPUhd() Int_t *HitIsUsed(){ return fHitIsUsed;} + GPUhd() AliHLTTPCCAHit *Hits(){ return fHits;} + GPUhd() Int_t &NHitsTotal(){ return fNHitsTotal;} + GPUhd() uint4 *&TexHitsFullData(){ return fTexHitsFullData;} + GPUhd() Int_t &TexHitsFullSize(){ return fTexHitsFullSize;} + + GPUd() UChar_t GetGridContent( UInt_t i ) const; + GPUd() AliHLTTPCCAHit GetHit( UInt_t i ) const; + + private: + + // - protected: - AliHLTTPCCAParam fParam; // parameters + AliHLTTPCCARow fRows[200];// array of hit rows + Double_t fTimers[10]; // running CPU time for different parts of the algorithm + + // event + Int_t fNHitsTotal;// total number of hits in event + Int_t fGridSizeTotal; // total grid size + Int_t fGrid1SizeTotal;// total grid1 size - AliHLTTPCCARow *fRows;// array of hit rows + AliHLTTPCCAHit *fHits; // hits + ushort2 *fHits1; // hits1 + + UChar_t *fGridContents; // grid content + UInt_t *fGrid1Contents; // grid1 content + Int_t *fHitsID; // hit ID's + + // temporary information + + Short_t *fHitLinkUp; // array of up links + Short_t *fHitLinkDown;// array of down links + Int_t *fHitIsUsed; // array of used flags + Int_t *fStartHits; // array of start hits + Int_t *fTracklets; // array of tracklets - Int_t *fOutTrackHits; // output array of ID's of the reconstructed hits + Int_t *fNTracks;// number of reconstructed tracks + AliHLTTPCCATrack *fTracks; // reconstructed tracks + Int_t *fTrackHits; // array of track hit numbers + + // output + + Int_t fNOutTracks; // number of tracks in fOutTracks array Int_t fNOutTrackHits; // number of hits in fOutTrackHits array AliHLTTPCCAOutTrack *fOutTracks; // output array of the reconstructed tracks - Int_t fNOutTracks; // number of tracks in fOutTracks array + Int_t *fOutTrackHits; // output array of ID's of the reconstructed hits - Int_t fNHitsTotal;// total number of hits in event - AliHLTTPCCATrack *fTracks; // reconstructed tracks - Int_t fNTracks;// number of reconstructed tracks - Int_t *fCellHitPointers;// global array of cell->hit pointers - AliHLTTPCCACell *fCells;// global array of cells - AliHLTTPCCAEndPoint *fEndPoints;// global array of endpoints - Double_t fTimers[10]; // running CPU time for different parts of the algorithm + char *fEventMemory; // common event memory + UInt_t fEventMemSize; // size of the event memory - ClassDef(AliHLTTPCCATracker,1); + uint4 *fTexHitsFullData; // CUDA texture for hits + Int_t fTexHitsFullSize; // size of the CUDA texture }; + #endif diff --git a/HLT/TPCLib/tracking-ca/AliHLTTPCCATrackerComponent.cxx b/HLT/TPCLib/tracking-ca/AliHLTTPCCATrackerComponent.cxx index 214d49bdbef..ba25d8e100c 100644 --- a/HLT/TPCLib/tracking-ca/AliHLTTPCCATrackerComponent.cxx +++ b/HLT/TPCLib/tracking-ca/AliHLTTPCCATrackerComponent.cxx @@ -26,11 +26,10 @@ using namespace std; #endif -#include +//#include #include "AliHLTTPCCATrackerComponent.h" #include "AliHLTTPCTransform.h" #include "AliHLTTPCCATracker.h" -#include "AliHLTTPCCAHit.h" #include "AliHLTTPCCAOutTrack.h" #include "AliHLTTPCCAParam.h" #include "AliHLTTPCCATrackConvertor.h" @@ -46,6 +45,13 @@ using namespace std; #include "TStopwatch.h" #include "TMath.h" +/** global object for registration + * Matthias 2009-01-13 temporarily using the global object approach again. + * CA cade had to be disabled because of various compilation problems, so + * the global object approach fits better for the moment. + */ +AliHLTTPCCATrackerComponent gAliHLTTPCCATrackerComponent; + /** ROOT macro for the implementation of ROOT specific class methods */ ClassImp(AliHLTTPCCATrackerComponent) @@ -148,10 +154,10 @@ int AliHLTTPCCATrackerComponent::DoInit( int argc, const char** argv ) fFullTime = 0; fRecoTime = 0; - fNEvents = 0; + fNEvents = 0; fTracker = new AliHLTTPCCATracker(); - + // read command line int i = 0; @@ -287,8 +293,8 @@ int AliHLTTPCCATrackerComponent::DoEvent AliHLTUInt32_t& size, vector& outputBlocks ) { - - AliHLTUInt32_t MaxBufferSize = size; + //* process event + AliHLTUInt32_t maxBufferSize = size; size = 0; // output size if(GetFirstInputBlock( kAliHLTDataTypeSOR ) || GetFirstInputBlock( kAliHLTDataTypeEOR )){ @@ -375,7 +381,6 @@ int AliHLTTPCCATrackerComponent::DoEvent // Initialize the tracker - Float_t Bz = fSolenoidBz; { if( !fTracker ) fTracker = new AliHLTTPCCATracker; @@ -397,19 +402,19 @@ int AliHLTTPCCATrackerComponent::DoEvent //TPCZmin = -249.645, ZMax = 249.778 // Float_t rMin = inRmin; // Float_t rMax = outRmax; - Int_t NRows = AliHLTTPCTransform::GetNRows(); + Int_t nRows = AliHLTTPCTransform::GetNRows(); Float_t padPitch = 0.4; Float_t sigmaZ = 0.228808; - Float_t *rowX = new Float_t [NRows]; - for( Int_t irow=0; irowfPadRow != oldRow ){ - if( oldRow>=0 ){ - if( fTracker->Rows()[oldRow].NHits()!=0 ) HLTError("CA: clusters from row %d are readed twice",oldRow); - fTracker->ReadHitRow( oldRow, vHits+firstRowHit, nRowHits ); - } - oldRow = pSP->fPadRow; - firstRowHit = nHits; - nRowHits = 0; - } - AliHLTTPCCAHit &h = vHits[nHits]; - //if( TMath::Abs(pSP->fX- fTracker->Rows()[pSP->fPadRow].X() )>1.e-4 ) cout<<"row "<<(Int_t)pSP->fPadRow<<" "<Rows()[pSP->fPadRow].X()-pSP->fX <fX- fTracker->Rows()[pSP->fPadRow].X() )>1.e-4 ) HLTError( "row %d, %f",(Int_t)pSP->fPadRow, fTracker->Rows()[pSP->fPadRow].X()-pSP->fX ); - - h.Y() = pSP->fY; - h.Z() = pSP->fZ; - if( TMath::Abs(h.Z())>fClusterZCut) continue; - h.ErrY() = TMath::Sqrt(TMath::Abs(pSP->fSigmaY2)); - h.ErrZ() = TMath::Sqrt(TMath::Abs(pSP->fSigmaZ2)); - if( h.ErrY()<.1 ) h.ErrY() = .1; - if( h.ErrZ()<.1 ) h.ErrZ() = .1; - if( h.ErrY()>1. ) h.ErrY() = 1.; - if( h.ErrZ()>1. ) h.ErrZ() = 1.; - h.ID() = nHits; + if( oldRow>=0 && pSP->fPadRow < oldRow ) + HLTError("CA: clusters from row %d are readed twice",oldRow); + + if( TMath::Abs(pSP->fZ)>fClusterZCut) continue; + vHitStoreX[nHits] = pSP->fX; - vHitStoreY[nHits] = h.Y(); - vHitStoreZ[nHits] = h.Z(); + vHitStoreY[nHits] = pSP->fY; + vHitStoreZ[nHits] = pSP->fZ; + vHitStoreIntID[nHits] = nHits; vHitStoreID[nHits] = pSP->fID; vHitRowID[nHits] = pSP->fPadRow; nHits++; - nRowHits++; + rowNHits[pSP->fPadRow]++; } - if( oldRow>=0 ){ - if( fTracker->Rows()[oldRow].NHits()!=0 ) HLTError("CA: clusters from row %d are readed twice",oldRow); - fTracker->ReadHitRow( oldRow, vHits+firstRowHit, nRowHits ); + + Int_t firstRowHit = 0; + for( Int_t ir=0; ir<200; ir++ ){ + rowFirstHits[ir] = firstRowHit; + firstRowHit+=rowNHits[ir]; } + + fTracker->ReadEvent( rowFirstHits, rowNHits, vHitStoreY, vHitStoreZ, nHits ); } + if( vOrigClusters ) delete[] vOrigClusters; // reconstruct the event @@ -578,8 +571,8 @@ int AliHLTTPCCATrackerComponent::DoEvent UInt_t dSize = sizeof(AliHLTTPCTrackSegmentData) + t.NHits()*sizeof(UInt_t); - if( mySize + dSize > MaxBufferSize ){ - Logging( kHLTLogWarning, "HLT::TPCCATracker::DoEvent", "Wrtite output","Output buffer size exceed (buffer size %d, current size %d), %d tracks are not stored", MaxBufferSize, mySize, ntracks-itr+1); + if( mySize + dSize > maxBufferSize ){ + Logging( kHLTLogWarning, "HLT::TPCCATracker::DoEvent", "Wrtite output","Output buffer size exceed (buffer size %d, current size %d), %d tracks are not stored", maxBufferSize, mySize, ntracks-itr+1); ret = -ENOSPC; break; } @@ -599,7 +592,7 @@ int AliHLTTPCCATrackerComponent::DoEvent AliHLTTPCCATrackParam par = t.StartPoint(); - par.TransportToX( vHitStoreX[iFirstHit] ); + par.TransportToX( vHitStoreX[iFirstHit], .99 ); AliExternalTrackParam tp; AliHLTTPCCATrackConvertor::GetExtParam( par, tp, 0, fSolenoidBz ); @@ -607,7 +600,7 @@ int AliHLTTPCCATrackerComponent::DoEvent currOutTracklet->fX = tp.GetX(); currOutTracklet->fY = tp.GetY(); currOutTracklet->fZ = tp.GetZ(); - currOutTracklet->fCharge = -(Int_t ) tp.GetSign(); + currOutTracklet->fCharge = (Int_t ) tp.GetSign(); currOutTracklet->fPt = TMath::Abs(tp.GetSignedPt()); Float_t snp = tp.GetSnp() ; if( snp>.999 ) snp=.999; @@ -623,7 +616,7 @@ int AliHLTTPCCATrackerComponent::DoEvent currOutTracklet->fPsierr = h*h*tp.GetSigmaSnp2(); currOutTracklet->fTglerr = tp.GetSigmaTgl2(); - if( par.TransportToX( vHitStoreX[iLastHit] ) ){ + if( par.TransportToX( vHitStoreX[iLastHit],.99 ) ){ currOutTracklet->fLastX = par.GetX(); currOutTracklet->fLastY = par.GetY(); currOutTracklet->fLastZ = par.GetZ(); @@ -658,10 +651,10 @@ int AliHLTTPCCATrackerComponent::DoEvent outPtr->fTrackletCnt++; } - if( vHits ) delete[] vHits; if( vHitStoreX ) delete[] vHitStoreX; if( vHitStoreY ) delete[] vHitStoreY; if( vHitStoreZ ) delete[] vHitStoreZ; + if( vHitStoreIntID ) delete[] vHitStoreIntID; if( vHitStoreID ) delete[] vHitStoreID; if( vHitRowID ) delete[] vHitRowID; @@ -688,7 +681,6 @@ int AliHLTTPCCATrackerComponent::DoEvent slice, ntracks, nClusters, minPatch, maxPatch, row[0], row[1], hz, hz1 ); return ret; - } diff --git a/HLT/TPCLib/tracking-ca/AliHLTTPCCATrackerComponent.h b/HLT/TPCLib/tracking-ca/AliHLTTPCCATrackerComponent.h index d65b9234ed7..eeb1d729405 100644 --- a/HLT/TPCLib/tracking-ca/AliHLTTPCCATrackerComponent.h +++ b/HLT/TPCLib/tracking-ca/AliHLTTPCCATrackerComponent.h @@ -35,19 +35,19 @@ public: // Public functions to implement AliHLTComponent's interface. // These functions are required for the registration process - /** @see component interface AliHLTComponent::GetComponentID */ + /** @see component interface @ref AliHLTComponent::GetComponentID */ const char* GetComponentID() ; - /** @see component interface AliHLTComponent::GetInputDataTypes */ + /** @see component interface @ref AliHLTComponent::GetInputDataTypes */ void GetInputDataTypes( vector& list) ; - /** @see component interface AliHLTComponent::GetOutputDataType */ + /** @see component interface @ref AliHLTComponent::GetOutputDataType */ AliHLTComponentDataType GetOutputDataType() ; - /** @see component interface AliHLTComponent::GetOutputDataSize */ + /** @see component interface @ref AliHLTComponent::GetOutputDataSize */ virtual void GetOutputDataSize( unsigned long& constBase, double& inputMultiplier ) ; - /** @see component interface AliHLTComponent::Spawn */ + /** @see component interface @ref AliHLTComponent::Spawn */ AliHLTComponent* Spawn() ; protected: @@ -56,10 +56,10 @@ protected: // These functions provide initialization as well as the actual processing // capabilities of the component. - /** @see component interface AliHLTComponent::DoInit */ + /** @see component interface @ref AliHLTComponent::DoInit */ int DoInit( int argc, const char** argv ); - /** @see component interface AliHLTComponent::DoDeinit */ + /** @see component interface @ref AliHLTComponent::DoDeinit */ int DoDeinit(); /** @see component interface @ref AliHLTProcessor::DoEvent */ diff --git a/HLT/TPCLib/tracking-ca/AliHLTTPCCATrackletConstructor.cxx b/HLT/TPCLib/tracking-ca/AliHLTTPCCATrackletConstructor.cxx new file mode 100644 index 00000000000..06a7f68d016 --- /dev/null +++ b/HLT/TPCLib/tracking-ca/AliHLTTPCCATrackletConstructor.cxx @@ -0,0 +1,472 @@ +// @(#) $Id: AliHLTTPCCATrackletConstructor.cxx 27042 2008-07-02 12:06:02Z richterm $ +//*************************************************************************** +// This file is property of and copyright by the ALICE HLT Project * +// ALICE Experiment at CERN, All rights reserved. * +// * +// Primary Authors: Sergey Gorbunov * +// Ivan Kisel * +// for The ALICE HLT Project. * +// * +// Permission to use, copy, modify and distribute this software and its * +// documentation strictly for non-commercial purposes is hereby granted * +// without fee, provided that the above copyright notice appears in all * +// copies and that both the copyright notice and this permission notice * +// appear in the supporting documentation. The authors make no claims * +// about the suitability of this software for any purpose. It is * +// provided "as is" without express or implied warranty. * +//*************************************************************************** + +#include "AliHLTTPCCATracker.h" +#include "AliHLTTPCCATrackParam.h" +#include "AliHLTTPCCATrackParam1.h" +#include "AliHLTTPCCAGrid.h" +#include "AliHLTTPCCAHitArea.h" +#include "AliHLTTPCCAMath.h" +#include "AliHLTTPCCADef.h" +#include "AliHLTTPCCATrackletConstructor.h" + +//GPUd() void myprintf1(int i, int j){ + //printf("fwd: iS=%d, iRow=%d\n",i,j); +//} +//GPUd() void myprintf2(int i, int j){ + //printf("bck: iS=%d, iRow=%d\n",i,j); +//} + + +GPUd() void AliHLTTPCCATrackletConstructor::Step0 +( Int_t nBlocks, Int_t /*nThreads*/, Int_t iBlock, Int_t iThread, Int_t /*iSync*/, + AliHLTTPCCASharedMemory &s, AliHLTTPCCAThreadMemory &r, AliHLTTPCCATracker &tracker, AliHLTTPCCATrackParam1 &/*tParam*/ ) +{ + // reconstruction of tracklets, step 0 + + const Int_t kNMemThreads = 128; + + r.fIsMemThread = ( iThread nTracks ) s.fItr1 = nTracks; + s.fUsedHits = tracker.HitIsUsed(); + s.fMinStartRow = 158; + s.fMaxStartRow = 0; + } + if( iThread<32 ){ + s.fMinStartRow32[iThread] = 158; + s.fMaxStartRow32[iThread] = 0; + } +} + + +GPUd() void AliHLTTPCCATrackletConstructor::Step1 +( Int_t /*nBlocks*/, Int_t /*nThreads*/, Int_t /*iBlock*/, Int_t iThread, Int_t /*iSync*/, + AliHLTTPCCASharedMemory &s, AliHLTTPCCAThreadMemory &r, AliHLTTPCCATracker &tracker, AliHLTTPCCATrackParam1 &tParam ) +{ + // reconstruction of tracklets, step 1 + + const Int_t kNMemThreads = 128; + + r.fItr= s.fItr0 + ( iThread - kNMemThreads ); + r.fGo = (!r.fIsMemThread) && ( r.fItrmaxStartRow ) maxStartRow = s.fMaxStartRow32[i]; + } + s.fMinStartRow = minStartRow; + s.fMaxStartRow = maxStartRow; + } +} + +GPUd() void AliHLTTPCCATrackletConstructor::ReadData +( Int_t iThread, AliHLTTPCCASharedMemory &s, AliHLTTPCCAThreadMemory &r, AliHLTTPCCATracker &tracker, Int_t iRow ) +{ + + // reconstruction of tracklets, read data step + const Int_t kNMemThreads = 128; + if( r.fIsMemThread ){ + AliHLTTPCCARow &row = tracker.Rows()[iRow]; + bool jr = !r.fCurrentData; + Int_t n = row.FullSize(); + uint4* gMem = tracker.TexHitsFullData() + row.FullOffset(); + uint4 *sMem = s.fData[jr]; + for( int i=iThread; i(tmpint4)) + row.FullGridOffset(); + + UInt_t *sGrid = s.fGridContent1; + + for( int i=iThread; i50); + + if( c[0]<=0 || c[2]<=0 || c[5]<=0 || c[9]<=0 || c[14]<=0 ) ok = 0; + + if(!ok){ + r.fNHits = 0; + break; + } + } + }while(0); + + if( !SAVE() ) return; + + int *store = tracker.Tracklets() + r.fTrackStoreOffset; + int *hitstore = tracker.Tracklets() +r.fHitStoreOffset; + store[0] = r.fNHits; + + if( r.fNHits>0 ){ + store[3] = r.fFirstRow; + store[4] = r.fLastRow; + if( CAMath::Abs(tParam.Par()[4])<1.e-8 ) tParam.Par()[4] = 1.e-8; + *((AliHLTTPCCATrackParam1*)(store+5)) = tParam; + int w = (r.fNHits<<16)+r.fItr; + for( int iRow=0; iRow<160; iRow++ ){ + Int_t ih = hitstore[iRow]; + if( ih>=0 ){ + int ihTot = tracker.Rows()[iRow].FirstHit() + ih; + CAMath::atomicMax( tracker.HitIsUsed() + ihTot, w ); + } + } + } +} + +GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet +( Int_t /*nBlocks*/, Int_t /*nThreads*/, Int_t /*iBlock*/, Int_t /*iThread*/, Int_t /*iSync*/, + AliHLTTPCCASharedMemory &s, AliHLTTPCCAThreadMemory &r, AliHLTTPCCATracker &tracker, AliHLTTPCCATrackParam1 &tParam,Int_t iRow ) +{ + // reconstruction of tracklets, tracklets update step + + if( !r.fGo ) return; + + const Int_t kMaxRowGap = 5; + + int *hitstore = tracker.Tracklets() +r.fHitStoreOffset; + + AliHLTTPCCARow &row = tracker.Rows()[iRow]; + + float y0 = row.Grid().YMin(); + float stepY = row.HstepY(); + float z0 = row.Grid().ZMin(); + float stepZ = row.HstepZ(); + float stepYi = row.HstepYi(); + float stepZi = row.HstepZi(); + + if( r.fStage == 0 ){ // fitting part + do{ + + if( iRow(tmpint4)[r.fCurrIH]; + + Int_t oldIH = r.fCurrIH; + r.fCurrIH = reinterpret_cast(tmpint4)[row.FullLinkOffset() + r.fCurrIH]; + + float x = row.X(); + float y = y0 + hh.x*stepY; + float z = z0 + hh.y*stepZ; + + if( iRow==r.fFirstRow ){ + tParam.X() = x; + tParam.Y() = y; + tParam.Z() = z; + float err2Y, err2Z; + tracker.GetErrors2( iRow, tParam, err2Y, err2Z ); + tParam.Cov()[0] = err2Y; + tParam.Cov()[2] = err2Z; + }else{ + if( !tParam.TransportToX0( x, .95 ) ){ + if( SAVE() ) hitstore[iRow] = -1; + break; + } + float err2Y, err2Z; + tracker.GetErrors2( iRow, *((AliHLTTPCCATrackParam*)&tParam), err2Y, err2Z ); + if( !tParam.Filter2( y, z, err2Y, err2Z, .95 ) ) { + if( SAVE() ) hitstore[iRow] = -1; + break; + } + } + if( SAVE() ) hitstore[iRow] = oldIH; + r.fNHits++; + r.fLastRow = iRow; + if( r.fCurrIH<0 ){ + r.fStage = 1; + if( r.fNHits<3 ){ r.fNHits=0; r.fGo = 0;} + } + break; + } while(0); + } + else // forward/backward searching part + { + do{ + if( r.fStage == 2 && iRow>=r.fFirstRow ) break; + if( r.fNMissed>kMaxRowGap ){ + r.fGo = 0; + break; + } + + r.fNMissed++; + + float x = row.X(); + float err2Y, err2Z; + if( !tParam.TransportToX0( x, .95 ) ) break; + uint4 *tmpint4 = s.fData[r.fCurrentData]; + + ushort2 *hits = reinterpret_cast(tmpint4); + UInt_t *gridContent1 = ((UInt_t*)(s.fGridContent1)); + + float fY = tParam.GetY(); + float fZ = tParam.GetZ(); + Int_t best = -1; + + { // search for the closest hit + + Int_t ds; + Int_t fY0 = (Int_t) ((fY - y0)*stepYi); + Int_t fZ0 = (Int_t) ((fZ - z0)*stepZi); + Int_t ds0 = ( ((int)1)<<30); + ds = ds0; + + UInt_t fIndYmin; + UInt_t fHitYfst=1, fHitYlst=0, fHitYfst1=1, fHitYlst1=0; + + fIndYmin = row.Grid().GetBin( (float)(fY-1.), (float)(fZ-1.) ); + UInt_t c = gridContent1[fIndYmin]; + fHitYfst = c & 0x000003FF; + fHitYlst = fHitYfst + ((c & 0x0000FC00)>>10); + fHitYfst1 = ( c & 0x03FF0000 )>>16; + fHitYlst1 = fHitYfst1 + ((c & 0xFC000000)>>26); + + for( UInt_t fIh = fHitYfst; fIh 1. ) sy2 = 1.; + if( sz2 > 1. ) sz2 = 1.; + if( iRow==63 || iRow==64 || iRow==65 ){ + if( sy2 < 4. ) sy2 = 4.; + if( sz2 < 4. ) sz2 = 4.; + } + + + if( CAMath::fmul_rz(dy,dy)>sy2 || CAMath::fmul_rz(dz,dz)>sz2 ) break; + + if( !tParam.Filter2( y, z, err2Y, err2Z, .95 ) ) break; + + if( SAVE() ) hitstore[ iRow ] = best; + r.fNHits++; + r.fNMissed=0; + }while(0); + } +} + + + +GPUd() void AliHLTTPCCATrackletConstructor::Thread +( Int_t nBlocks, Int_t nThreads, Int_t iBlock, Int_t iThread, Int_t iSync, + AliHLTTPCCASharedMemory &s, AliHLTTPCCAThreadMemory &r, AliHLTTPCCATracker &tracker, AliHLTTPCCATrackParam1 &tParam ) +{ + + // reconstruction of tracklets + if( iSync==0 ) + { + Step0( nBlocks, nThreads, iBlock, iThread, iSync, s, r, tracker, tParam ); + } + else if( iSync==1 ) + { + Step1( nBlocks, nThreads, iBlock, iThread, iSync, s, r, tracker, tParam ); + } + else if( iSync==2 ) + { + Step2( nBlocks, nThreads, iBlock, iThread, iSync, s, r, tracker, tParam ); + } + + else if( iSync==3 ) + + { + r.fCurrentData = 1; + ReadData( iThread, s, r, tracker, s.fMinStartRow ); + r.fCurrentData = 0; + r.fNMissed = 0; + } + else if( iSync==3+159*2+1 )//322 + + { + r.fCurrentData = 1; + Int_t nextRow = s.fMaxStartRow-1; + if( nextRow<0 ) nextRow = 0; + ReadData( iThread, s, r, tracker, nextRow ); + r.fCurrentData = 0; + r.fNMissed = 0; + r.fStage = 2; + } + + else if( iSync<=3+159*2+1+159*2 ) + + { + int iRow, nextRow; + if( iSync<=3+159*2 ){ + iRow = (iSync -4)/2; + //if( iBlock==0 && iThread==0 ) myprintf1(iSync,iRow); + if( iRow < s.fMinStartRow ) return; + nextRow = iRow+1; + if( nextRow>158 ) nextRow = 158; + }else{ + iRow = 158 - (iSync - 4-159*2-1)/2; + //if( iBlock==0 && iThread==0 ) myprintf2(iSync,iRow); + if( iRow >= s.fMaxStartRow ) return; + nextRow = iRow-1; + if( nextRow<0 ) nextRow = 0; + } + + if( iSync%2==0 ){ + UnpackGrid( nBlocks, nThreads, iBlock, iThread, iSync, + s, r, tracker, tParam, iRow ); + }else{ + if( r.fIsMemThread ){ + ReadData( iThread, s, r, tracker, nextRow ); + }else{ + UpdateTracklet( nBlocks, nThreads, iBlock, iThread, iSync, + s, r, tracker, tParam, iRow ); + } + r.fCurrentData = !r.fCurrentData; + } + } + + else if( iSync== 4+159*4 +1+1 ) // 642 + + { + StoreTracklet( nBlocks, nThreads, iBlock, iThread, iSync, //SG!!! + s, r, tracker, tParam ); + } +} + diff --git a/HLT/TPCLib/tracking-ca/AliHLTTPCCATrackletConstructor.h b/HLT/TPCLib/tracking-ca/AliHLTTPCCATrackletConstructor.h new file mode 100644 index 00000000000..a2195d5b3f4 --- /dev/null +++ b/HLT/TPCLib/tracking-ca/AliHLTTPCCATrackletConstructor.h @@ -0,0 +1,113 @@ +//-*- Mode: C++ -*- + +//* This file is property of and copyright by the ALICE HLT Project * +//* ALICE Experiment at CERN, All rights reserved. * +//* See cxx source for full Copyright notice * + +#ifndef ALIHLTTPCCATRACKLETCONSTRUCTOR_H +#define ALIHLTTPCCATRACKLETCONSTRUCTOR_H + + +#include "AliHLTTPCCADef.h" + +/** + * @class AliHLTTPCCATrackletConstructor + * + */ +class AliHLTTPCCATrackletConstructor +{ + public: + + class AliHLTTPCCASharedMemory + { + friend class AliHLTTPCCATrackletConstructor; + public: +#if !defined(HLTCA_GPUCODE) + AliHLTTPCCASharedMemory() + : fItr0(0), fItr1(0), fNRows(0), fUsedHits(0), fMinStartRow(0), fMaxStartRow(0) + {} + + AliHLTTPCCASharedMemory( const AliHLTTPCCASharedMemory& /*dummy*/) + : fItr0(0), fItr1(0), fNRows(0), fUsedHits(0), fMinStartRow(0), fMaxStartRow(0) + {} + AliHLTTPCCASharedMemory& operator=(const AliHLTTPCCASharedMemory& /*dummy*/){ return *this; } +#endif + protected: + uint4 fData[2][(500+500+500)/4]; // temp memory + UInt_t fGridContent1[600]; // grid1 content + int fItr0; // start track index + int fItr1; // end track index + int fNRows; // n rows + Int_t *fUsedHits; // array of used hits + int fMinStartRow; // min start row + int fMinStartRow32[32]; // min start row for each thread in warp + int fMaxStartRow; // max start row + int fMaxStartRow32[32];// max start row for each thread in warp + }; + + class AliHLTTPCCAThreadMemory + { + friend class AliHLTTPCCATrackletConstructor; + public: +#if !defined(HLTCA_GPUCODE) + AliHLTTPCCAThreadMemory() + : fItr(0), fFirstRow(0), fLastRow(0), fCurrIH(0), fIsMemThread(0), fGo(0), fSave(0), fCurrentData(0), fStage(0), fNHits(0), fNMissed(0), fTrackStoreOffset(0), fHitStoreOffset(0) + {} + + AliHLTTPCCAThreadMemory( const AliHLTTPCCAThreadMemory& /*dummy*/) + : fItr(0), fFirstRow(0), fLastRow(0), fCurrIH(0), fIsMemThread(0), fGo(0), fSave(0), fCurrentData(0), fStage(0), fNHits(0), fNMissed(0), fTrackStoreOffset(0), fHitStoreOffset(0) + {} + AliHLTTPCCAThreadMemory& operator=(const AliHLTTPCCAThreadMemory& /*dummy*/){ return *this; } +#endif + protected: + Int_t fItr; // track index + Int_t fFirstRow; // first row index + Int_t fLastRow; // last row index + Int_t fCurrIH; // indef of the current hit + Bool_t fIsMemThread; // is the thread used for memory taken + Bool_t fGo; // do fit/searching flag + Bool_t fSave; // save flag + Bool_t fCurrentData; // index of the current memory array + Int_t fStage; // reco stage + Int_t fNHits; // n track hits + Int_t fNMissed; // n missed hits during search + Int_t fTrackStoreOffset; // offset in the global array + Int_t fHitStoreOffset; // offset in the global array + }; + + GPUd() static Int_t NThreadSyncPoints(){ return 4+159*4 +1+1; } + + GPUd() static void Thread( Int_t nBlocks, Int_t nThreads, Int_t iBlock, Int_t iThread, + Int_t iSync, AliHLTTPCCASharedMemory &s, AliHLTTPCCAThreadMemory &r, + AliHLTTPCCATracker &tracker, AliHLTTPCCATrackParam1 &tParam ); + + GPUd() static void Step0 + ( Int_t nBlocks, Int_t nThreads, Int_t iBlock, Int_t iThread, Int_t iSync, + AliHLTTPCCASharedMemory &s, AliHLTTPCCAThreadMemory &r, AliHLTTPCCATracker &tracker, AliHLTTPCCATrackParam1 &tParam ); + GPUd() static void Step1 + ( Int_t nBlocks, Int_t nThreads, Int_t iBlock, Int_t iThread, Int_t iSync, + AliHLTTPCCASharedMemory &s, AliHLTTPCCAThreadMemory &r, AliHLTTPCCATracker &tracker, AliHLTTPCCATrackParam1 &tParam ); + GPUd() static void Step2 + ( Int_t nBlocks, Int_t nThreads, Int_t iBlock, Int_t iThread, Int_t iSync, + AliHLTTPCCASharedMemory &s, AliHLTTPCCAThreadMemory &r, AliHLTTPCCATracker &tracker, AliHLTTPCCATrackParam1 &tParam ); + + GPUd() static void ReadData( Int_t iThread, AliHLTTPCCASharedMemory &s, AliHLTTPCCAThreadMemory &r, AliHLTTPCCATracker &tracker, Int_t iRow ); + + GPUd() static void UpdateTracklet + ( Int_t nBlocks, Int_t nThreads, Int_t iBlock, Int_t iThread, Int_t iSync, + AliHLTTPCCASharedMemory &s, AliHLTTPCCAThreadMemory &r, AliHLTTPCCATracker &tracker, AliHLTTPCCATrackParam1 &tParam, Int_t iRow ); + + GPUd() static void UnpackGrid + ( Int_t nBlocks, Int_t nThreads, Int_t iBlock, Int_t iThread, Int_t iSync, + AliHLTTPCCASharedMemory &s, AliHLTTPCCAThreadMemory &r, AliHLTTPCCATracker &tracker, AliHLTTPCCATrackParam1 &tParam, Int_t iRow ); + + GPUd() static void StoreTracklet + ( Int_t nBlocks, Int_t nThreads, Int_t iBlock, Int_t iThread, Int_t iSync, + AliHLTTPCCASharedMemory &s, AliHLTTPCCAThreadMemory &r, AliHLTTPCCATracker &tracker, AliHLTTPCCATrackParam1 &tParam); + + static Bool_t SAVE(){ return 1; } + +}; + + +#endif diff --git a/HLT/TPCLib/tracking-ca/AliHLTTPCCATrackletSelector.cxx b/HLT/TPCLib/tracking-ca/AliHLTTPCCATrackletSelector.cxx new file mode 100644 index 00000000000..27944864ba3 --- /dev/null +++ b/HLT/TPCLib/tracking-ca/AliHLTTPCCATrackletSelector.cxx @@ -0,0 +1,80 @@ +// @(#) $Id: AliHLTTPCCATrackletSelector.cxx 27042 2008-07-02 12:06:02Z richterm $ +//*************************************************************************** +// This file is property of and copyright by the ALICE HLT Project * +// ALICE Experiment at CERN, All rights reserved. * +// * +// Primary Authors: Sergey Gorbunov * +// Ivan Kisel * +// for The ALICE HLT Project. * +// * +// Permission to use, copy, modify and distribute this software and its * +// documentation strictly for non-commercial purposes is hereby granted * +// without fee, provided that the above copyright notice appears in all * +// copies and that both the copyright notice and this permission notice * +// appear in the supporting documentation. The authors make no claims * +// about the suitability of this software for any purpose. It is * +// provided "as is" without express or implied warranty. * +//*************************************************************************** + +#include "AliHLTTPCCATrackletSelector.h" +#include "AliHLTTPCCATrack.h" +#include "AliHLTTPCCATracker.h" +#include "AliHLTTPCCATrackParam.h" +#include "AliHLTTPCCAMath.h" + +GPUd() void AliHLTTPCCATrackletSelector::Thread +( Int_t nBlocks, Int_t nThreads, Int_t iBlock, Int_t iThread, Int_t iSync, + AliHLTTPCCASharedMemory &s, AliHLTTPCCATracker &tracker ) +{ + // select best tracklets and kill clones + + if( iSync==0 ) + { + if( iThread==0 ){ + if(iBlock==0){ + CAMath::atomicExch(&(tracker.NTracks()),0); + CAMath::atomicExch(tracker.TrackHits(),0); + } + s.fNTracklets = tracker.Tracklets()[0]; + s.fNThreadsTotal = nThreads*nBlocks; + s.fItr0 = nThreads*iBlock; + //if( iBlock==0 ) tracker.StartHits()[0] = 0;//SG!!! + } + } + else if( iSync==1 ) + { + AliHLTTPCCATrack tout; + Int_t trackHits[160]; + + for( Int_t itr= s.fItr0 + iThread; itr w ) continue; + Int_t th = AliHLTTPCCATracker::IRowIHit2ID(irow,ih); + trackHits[tout.NHits()] = th; + tout.NHits()++; + } + if( tout.NHits()<10 ) continue;//SG!!! + Int_t itrout = CAMath::atomicAdd(&(tracker.NTracks()),1); + tout.FirstHitID() = CAMath::atomicAdd( tracker.TrackHits(), tout.NHits() ) + 1; + tout.Param() = *( (AliHLTTPCCATrackParam*)( t+5) ); + tout.Alive() = 1; + tracker.Tracks()[itrout] = tout; + for( Int_t ih=0; ih -#include -#include "AliCluster.h" +#include "TTree.h" +#include "Riostream.h" +//#include "AliCluster.h" #include "AliTPCClustersRow.h" #include "AliTPCParam.h" #include "AliRun.h" #include "AliRunLoader.h" #include "AliStack.h" -#include "AliHLTTPCCATracker.h" -#include "AliHLTTPCCAGBHit.h" #include "AliHLTTPCCAGBTracker.h" +#include "AliHLTTPCCAGBHit.h" #include "AliHLTTPCCAGBTrack.h" -#include "AliHLTTPCCAMCTrack.h" -#include "AliHLTTPCCAOutTrack.h" #include "AliHLTTPCCAPerformance.h" #include "AliHLTTPCCAParam.h" #include "AliHLTTPCCATrackConvertor.h" +#include "AliHLTTPCCATracker.h" #include "TMath.h" #include "AliTPCLoader.h" @@ -43,11 +41,11 @@ #include "AliTPCclusterMI.h" #include "AliTPCTransform.h" #include "AliTPCcalibDB.h" -#include "AliTPCReconstructor.h" #include "AliTPCtrack.h" #include "AliESDtrack.h" #include "AliESDEvent.h" #include "AliTrackReference.h" +#include "TStopwatch.h" //#include @@ -84,9 +82,9 @@ AliTPCtrackerCA::AliTPCtrackerCA(const AliTPCParam *par): AliTracker(),fParam(par), fClusters(0), fNClusters(0), fHLTTracker(0), fHLTPerformance(0),fDoHLTPerformance(0),fDoHLTPerformanceClusters(0),fStatNEvents(0) { //* constructor - + DoHLTPerformance() = 1; - DoHLTPerformanceClusters() = 1; + DoHLTPerformanceClusters() = 0; fHLTTracker = new AliHLTTPCCAGBTracker; fHLTTracker->SetNSlices( fParam->GetNSector()/2 ); @@ -121,19 +119,18 @@ AliTPCtrackerCA::AliTPCtrackerCA(const AliTPCParam *par): Float_t padPitch = 0.4; Float_t sigmaZ = 0.228808; - Int_t NRows = fParam->GetNRowLow()+fParam->GetNRowUp(); - + Int_t nRows = fParam->GetNRowLow()+fParam->GetNRowUp(); Float_t rowX[200]; for( Int_t irow=0; irowGetNRowLow(); irow++){ rowX[irow] = fParam->GetPadRowRadiiLow(irow); } for( Int_t irow=0; irowGetNRowUp(); irow++){ rowX[fParam->GetNRowLow()+irow] = fParam->GetPadRowRadiiUp(irow); - } + } AliHLTTPCCAParam param; - param.Initialize( iSlice, NRows, rowX, alpha, dalpha, + param.Initialize( iSlice, nRows, rowX, alpha, dalpha, inRmin, outRmax, zMin, zMax, padPitch, sigmaZ, bz ); - param.YErrorCorrection() = 1.;//.33; + param.YErrorCorrection() = 1.;// .33; param.ZErrorCorrection() = 1.;//.33; param.MaxTrackMatchDRow() = 5; param.TrackConnectionFactor() = 3.5; @@ -143,8 +140,9 @@ AliTPCtrackerCA::AliTPCtrackerCA(const AliTPCParam *par): -Int_t AliTPCtrackerCA::LoadClusters (TTree * tree) +Int_t AliTPCtrackerCA::LoadClusters (TTree * fromTree) { + // load clusters to the local arrays fNClusters = 0; if( fClusters ) delete[] fClusters; @@ -156,7 +154,7 @@ Int_t AliTPCtrackerCA::LoadClusters (TTree * tree) // load mc tracks while( fDoHLTPerformance ){ if( !gAlice ) break; - AliRunLoader *rl = AliRunLoader::GetRunLoader(); + AliRunLoader *rl = gAlice->GetRunLoader(); if( !rl ) break; rl->LoadKinematics(); AliStack *stack = rl->Stack(); @@ -174,15 +172,15 @@ Int_t AliTPCtrackerCA::LoadClusters (TTree * tree) isTPC = new Bool_t [stack->GetNtrack()]; for( Int_t i=0; iGetNtrack(); i++ ) isTPC[i] = 0; rl->LoadTrackRefs(); - TTree *TR = rl->TreeTR(); - if( !TR ) break; - TBranch *branch=TR->GetBranch("TrackReferences"); + TTree *mcTree = rl->TreeTR(); + if( !mcTree ) break; + TBranch *branch=mcTree->GetBranch("TrackReferences"); if (!branch ) break; TClonesArray tpcdummy("AliTrackReference",1000), *tpcRefs=&tpcdummy; branch->SetAddress(&tpcRefs); - Int_t nr=(Int_t)TR->GetEntries(); + Int_t nr=(Int_t)mcTree->GetEntries(); for (Int_t r=0; rGetEvent(r); + mcTree->GetEvent(r); Int_t nref = tpcRefs->GetEntriesFast(); if (!nref) continue; AliTrackReference *tpcRef= 0x0; @@ -236,7 +234,7 @@ Int_t AliTPCtrackerCA::LoadClusters (TTree * tree) break; } - TBranch * br = tree->GetBranch("Segment"); + TBranch * br = fromTree->GetBranch("Segment"); if( !br ) return 1; AliTPCClustersRow *clrow = new AliTPCClustersRow; @@ -247,10 +245,10 @@ Int_t AliTPCtrackerCA::LoadClusters (TTree * tree) br->SetAddress(&clrow); // - Int_t NEnt=Int_t(tree->GetEntries()); + Int_t nEnt=Int_t(fromTree->GetEntries()); fNClusters = 0; - for (Int_t i=0; iGetEntry(i); Int_t sec,row; fParam->AdjustSectorRow(clrow->GetID(),sec,row); @@ -261,13 +259,13 @@ Int_t AliTPCtrackerCA::LoadClusters (TTree * tree) fHLTTracker->SetNHits( fNClusters ); if( fDoHLTPerformance ) fHLTPerformance->SetNHits( fNClusters ); int ind=0; - for (Int_t i=0; iGetEntry(i); Int_t sec,row; fParam->AdjustSectorRow(clrow->GetID(),sec,row); - int NClu = clrow->GetArray()->GetEntriesFast(); + int nClu = clrow->GetArray()->GetEntriesFast(); Float_t x = fParam->GetPadRowRadii(sec,row); - for (Int_t icl=0; iclGetDetector()}; transform->Transform(xx,id,0,1); //if (!AliTPCReconstructor::GetRecoParam()->GetBYMirror()){ - if (cluster->GetDetector()%36>17){ - xx[1]*=-1; - } + //if (cluster->GetDetector()%36>17){ + //xx[1]*=-1; + //} //} cluster->SetX(xx[0]); @@ -332,9 +330,12 @@ AliCluster * AliTPCtrackerCA::GetCluster(Int_t index) const Int_t AliTPCtrackerCA::Clusters2Tracks( AliESDEvent *event ) { - //cout<<"Start of AliTPCtrackerCA"<FindTracks(); + cout<<"Do performance.."<Performance(); if( 0 ) {// Write Event @@ -347,18 +348,18 @@ Int_t AliTPCtrackerCA::Clusters2Tracks( AliESDEvent *event ) geo.close(); } - fstream eventStream; + fstream hits; char name[255]; sprintf( name,"CAEvents/%i.event.dat",fStatNEvents ); - eventStream.open(name, ios::out); - if( eventStream.is_open() ){ - fHLTTracker->WriteEvent(eventStream); + hits.open(name, ios::out); + if( hits.is_open() ){ + fHLTTracker->WriteEvent(hits); fstream tracks; sprintf( name,"CAEvents/%i.tracks.dat",fStatNEvents ); tracks.open(name, ios::out); fHLTTracker->WriteTracks(tracks); } - eventStream.close(); + hits.close(); if( fDoHLTPerformance ){ fstream mcevent, mcpoints; char mcname[255]; @@ -367,7 +368,7 @@ Int_t AliTPCtrackerCA::Clusters2Tracks( AliESDEvent *event ) if( mcevent.is_open() ){ fHLTPerformance->WriteMCEvent(mcevent); } - if(fDoHLTPerformanceClusters ){ + if(1 && fDoHLTPerformanceClusters ){ sprintf( mcname,"CAEvents/%i.mcpoints.dat",fStatNEvents ); mcpoints.open(mcname, ios::out); if( mcpoints.is_open() ){ @@ -391,17 +392,17 @@ Int_t AliTPCtrackerCA::Clusters2Tracks( AliESDEvent *event ) AliHLTTPCCATrackConvertor::GetExtParam( par, tTPC, tCA.Alpha(), bz ); tTPC.SetMass(0.13957); tTPC.SetdEdx( tCA.DeDx() ); - if( TMath::Abs(tTPC.GetSigned1Pt())>1./0.1 ) continue; + if( TMath::Abs(tTPC.GetSigned1Pt())>1./0.02 ) continue; int nhits = tCA.NHits(); if( nhits>kMaxRow ) nhits = kMaxRow; tTPC.SetNumberOfClusters(nhits); for( Int_t ih=0; ihTrackHits()[tCA.FirstHitRef()+ih]; - Int_t ext_index = fHLTTracker->Hits()[index].ID(); - tTPC.SetClusterIndex(ih, ext_index); + Int_t extIndex = fHLTTracker->Hits()[index].ID(); + tTPC.SetClusterIndex(ih, extIndex); } CookLabel(&tTPC,0.1); - { + if(0){ Double_t xTPC=fParam->GetInnerRadiusLow(); Double_t dAlpha = fParam->GetInnerAngle()/180.*TMath::Pi(); if (tTPC.AliExternalTrackParam::PropagateTo(xTPC,bz)) { @@ -423,6 +424,13 @@ Int_t AliTPCtrackerCA::Clusters2Tracks( AliESDEvent *event ) event->AddTrack(&tESD); } } + timer.Stop(); + static double time=0, time1 = 0; + static int ncalls = 0; + time+=timer.CpuTime(); + time1+=timer.RealTime(); + ncalls++; + cout<<"\n\nCA tracker speed: cpu = "<