in.read((char*) &tmpData[0], i * sizeof(T));
for (int j = 0;j < i;j++)
{
+#ifdef HLTCA_STANDALONE
+ if (tmpData[j].fRow < 0 || tmpData[j].fRow >= HLTCA_ROW_COUNT)
+ {
+ printf("Invalid Row Read %d at Cluster %d\n", tmpData[j].fRow, j);
+ exit(1);
+ }
+#endif
data[j] = tmpData[j];
}
}
texture<signed short, 1, cudaReadModeElementType> gAliTexRefs;
#endif
-#include "AliHLTTPCCAHit.h"
-
//Include CXX Files, GPUd() macro will then produce CUDA device code out of the tracker source code
#include "AliHLTTPCCATrackParam.cxx"
#include "AliHLTTPCCATrack.cxx"
-#include "AliHLTTPCCATrackletSelector.cxx"
-
#include "AliHLTTPCCAHitArea.cxx"
#include "AliHLTTPCCAGrid.cxx"
#include "AliHLTTPCCARow.cxx"
#include "AliHLTTPCCAParam.cxx"
#include "AliHLTTPCCATracker.cxx"
-#include "AliHLTTPCCAOutTrack.cxx"
-
#include "AliHLTTPCCAProcess.h"
+#include "AliHLTTPCCATrackletSelector.cxx"
#include "AliHLTTPCCANeighboursFinder.cxx"
-
#include "AliHLTTPCCANeighboursCleaner.cxx"
#include "AliHLTTPCCAStartHitsFinder.cxx"
#include "AliHLTTPCCAStartHitsSorter.cxx"
{
fSlaveTrackers[i].SetGPUTracker();
fSlaveTrackers[i].SetGPUTrackerCommonMemory((char*) CommonMemory(fHostLockedMemory, i));
- fSlaveTrackers[i].pData()->SetGPUSliceDataMemory(SliceDataMemory(fHostLockedMemory, i), RowMemory(fHostLockedMemory, i));
+ fSlaveTrackers[i].SetGPUSliceDataMemory(SliceDataMemory(fHostLockedMemory, i), RowMemory(fHostLockedMemory, i));
}
fpCudaStreams = malloc(CAMath::Max(3, fSliceCount) * sizeof(cudaStream_t));
//Do one initial run for Benchmark reasons
const int useDebugLevel = fDebugLevel;
fDebugLevel = 0;
- AliHLTTPCCAClusterData tmpCluster;
+ AliHLTTPCCAClusterData* tmpCluster = new AliHLTTPCCAClusterData[sliceCount];
std::ifstream fin;
- fin.open("events/event.0.dump");
- tmpCluster.ReadEvent(fin);
- fin.close();
- AliHLTTPCCASliceOutput *tmpOutput = NULL;
AliHLTTPCCAParam tmpParam;
AliHLTTPCCASliceOutput::outputControlStruct tmpOutputControl;
- fSlaveTrackers[0].SetOutputControl(&tmpOutputControl);
- tmpParam.SetNRows(HLTCA_ROW_COUNT);
- fSlaveTrackers[0].SetParam(tmpParam);
- Reconstruct(&tmpOutput, &tmpCluster, 0, 1);
- free(tmpOutput);
- tmpOutput = NULL;
- fSlaveTrackers[0].SetOutputControl(NULL);
+
+ fin.open("events/settings.dump");
+ int tmpCount;
+ fin >> tmpCount;
+ for (int i = 0;i < sliceCount;i++)
+ {
+ fSlaveTrackers[i].SetOutputControl(&tmpOutputControl);
+ tmpParam.ReadSettings(fin);
+ InitializeSliceParam(i, tmpParam);
+ }
+ fin.close();
+
+ fin.open("eventspbpbc/event.0.dump", std::ifstream::binary);
+ for (int i = 0;i < sliceCount;i++)
+ {
+ tmpCluster[i].StartReading(i, 0);
+ tmpCluster[i].ReadEvent(fin);
+ tmpCluster[i].FinishReading();
+ }
+ fin.close();
+
+ AliHLTTPCCASliceOutput **tmpOutput = new AliHLTTPCCASliceOutput*[sliceCount];
+ memset(tmpOutput, 0, sliceCount * sizeof(AliHLTTPCCASliceOutput*));
+
+ Reconstruct(tmpOutput, tmpCluster, 0, sliceCount);
+ for (int i = 0;i < sliceCount;i++)
+ {
+ free(tmpOutput[i]);
+ tmpOutput[i] = NULL;
+ fSlaveTrackers[i].SetOutputControl(NULL);
+ }
+ delete[] tmpOutput;
+ delete[] tmpCluster;
fDebugLevel = useDebugLevel;
}
#endif
for (int i = 0; i < tracker[iSlice].Param().NRows() / HLTCA_GPU_SCHED_ROW_STEP + 1;i++)
{
*fOutFile << "Rowblock: " << i << ", up " << rowBlockPos[i].y << "/" << rowBlockPos[i].x << ", down " <<
- rowBlockPos[tracker[iSlice].Param().NRows() / HLTCA_GPU_SCHED_ROW_STEP + 1 + i].y << "/" << rowBlockPos[tracker[iSlice].Param().NRows() / HLTCA_GPU_SCHED_ROW_STEP + 1 + i].x << endl << "Phase 1: ";
+ rowBlockPos[tracker[iSlice].Param().NRows() / HLTCA_GPU_SCHED_ROW_STEP + 1 + i].y << "/" << rowBlockPos[tracker[iSlice].Param().NRows() / HLTCA_GPU_SCHED_ROW_STEP + 1 + i].x << std::endl << "Phase 1: ";
for (int j = 0;j < rowBlockPos[i].x;j++)
{
//Use Tracker Object to calculate Offset instead of fGpuTracker, since *fNTracklets of fGpuTracker points to GPU Mem!
HLTError("Error, -1 Tracklet found");
}
}
- *fOutFile << endl << "Phase 2: ";
+ *fOutFile << std::endl << "Phase 2: ";
for (int j = 0;j < rowBlockPos[tracker[iSlice].Param().NRows() / HLTCA_GPU_SCHED_ROW_STEP + 1 + i].x;j++)
{
*fOutFile << rowBlockTracklets[(tracker[iSlice].RowBlockTracklets(1, i) - tracker[iSlice].RowBlockTracklets(0, 0)) + j] << ", ";
}
- *fOutFile << endl;
+ *fOutFile << std::endl;
}
if (check)
{
for (int iSlice = 0;iSlice < sliceCountLocal;iSlice++)
{
- *fOutFile << endl << endl << "Slice: " << fSlaveTrackers[firstSlice + iSlice].Param().ISlice() << endl;
+ *fOutFile << std::endl << std::endl << "Slice: " << fSlaveTrackers[firstSlice + iSlice].Param().ISlice() << std::endl;
}
}
#ifdef HLTCA_GPU_TIME_PROFILE
__int64 a, b, c, d;
- QueryPerformanceFrequency((LARGE_INTEGER*) &c);
- QueryPerformanceCounter((LARGE_INTEGER*) &d);
+ AliHLTTPCCAStandaloneFramework::StandaloneQueryFreq(&c);
+ AliHLTTPCCAStandaloneFramework::StandaloneQueryTime(&d);
#endif
for (int iSlice = 0;iSlice < sliceCountLocal;iSlice++)
//Make this a GPU Tracker
fGpuTracker[iSlice].SetGPUTracker();
fGpuTracker[iSlice].SetGPUTrackerCommonMemory((char*) CommonMemory(fGPUMemory, iSlice));
- fGpuTracker[iSlice].pData()->SetGPUSliceDataMemory(SliceDataMemory(fGPUMemory, iSlice), RowMemory(fGPUMemory, iSlice));
- fGpuTracker[iSlice].pData()->SetPointers(&pClusterData[iSlice], false);
+ fGpuTracker[iSlice].SetGPUSliceDataMemory(SliceDataMemory(fGPUMemory, iSlice), RowMemory(fGPUMemory, iSlice));
+ fGpuTracker[iSlice].SetPointersSliceData(&pClusterData[iSlice], false);
//Set Pointers to GPU Memory
char* tmpMem = (char*) GlobalMemory(fGPUMemory, iSlice);
fGpuTracker[iSlice].GPUParametersConst()->fGPUiSlice = iSlice;
fGpuTracker[iSlice].GPUParametersConst()->fGPUnSlices = sliceCountLocal;
fSlaveTrackers[firstSlice + iSlice].GPUParameters()->fGPUError = 0;
- fGpuTracker[iSlice].pData()->SetGPUTextureBase(fGpuTracker[0].Data().Memory());
+ fGpuTracker[iSlice].SetGPUTextureBase(fGpuTracker[0].Data().Memory());
}
#ifdef HLTCA_GPU_TEXTURE_FETCH
//Initialize GPU Slave Tracker
if (fDebugLevel >= 3) HLTInfo("Creating Slice Data");
- fSlaveTrackers[firstSlice + iSlice].pData()->SetGPUSliceDataMemory(SliceDataMemory(fHostLockedMemory, iSlice), RowMemory(fHostLockedMemory, firstSlice + iSlice));
+ fSlaveTrackers[firstSlice + iSlice].SetGPUSliceDataMemory(SliceDataMemory(fHostLockedMemory, iSlice), RowMemory(fHostLockedMemory, firstSlice + iSlice));
#ifdef HLTCA_GPU_TIME_PROFILE
- QueryPerformanceCounter((LARGE_INTEGER*) &a);
+ AliHLTTPCCAStandaloneFramework::StandaloneQueryTime(&a);
#endif
fSlaveTrackers[firstSlice + iSlice].ReadEvent(&pClusterData[iSlice]);
#ifdef HLTCA_GPU_TIME_PROFILE
- QueryPerformanceCounter((LARGE_INTEGER*) &b);
+ AliHLTTPCCAStandaloneFramework::StandaloneQueryTime(&b);
printf("Read %f %f\n", ((double) b - (double) a) / (double) c, ((double) a - (double) d) / (double) c);
#endif
if (fSlaveTrackers[firstSlice + iSlice].Data().MemorySize() > HLTCA_GPU_SLICE_DATA_MEMORY)
fSlaveTrackers[firstSlice + iSlice].SetOutput(&pOutput[iSlice]);
#ifdef HLTCA_GPU_TIME_PROFILE
- QueryPerformanceCounter((LARGE_INTEGER*) &a);
+ AliHLTTPCCAStandaloneFramework::StandaloneQueryTime(&a);
#endif
fSlaveTrackers[firstSlice + iSlice].WriteOutput();
#ifdef HLTCA_GPU_TIME_PROFILE
- QueryPerformanceCounter((LARGE_INTEGER*) &b);
+ AliHLTTPCCAStandaloneFramework::StandaloneQueryTime(&b);
printf("Write %f %f\n", ((double) b - (double) a) / (double) c, ((double) a - (double) d) / (double) c);
#endif
}
}
-#endif
\ No newline at end of file
+#endif
Err2Z = GetClusterError2( 1, type, z, angleZ );
}
-
+#ifdef HLTCA_GPUCODE
GPUh() void AliHLTTPCCAParam::WriteSettings( std::ostream &out ) const
{
// write settings to the file
for ( int k = 0; k < 7; k++ )
in >> fParamS0Par[i][j][k];
}
+#endif
void AliHLTTPCCATracker::SetGPUTracker()
{
+ //Make this a GPU Tracker
fIsGPUTracker = true;
fData.SetGpuSliceData();
}
-char* AliHLTTPCCATracker::SetGPUTrackerCommonMemory(char* pGPUMemory)
+char* AliHLTTPCCATracker::SetGPUTrackerCommonMemory(char* const pGPUMemory)
{
+ //Set up common Memory Pointer for GPU Tracker
fCommonMem = (commonMemoryStruct*) pGPUMemory;
return(pGPUMemory + sizeof(commonMemoryStruct));
}
char* AliHLTTPCCATracker::SetGPUTrackerHitsMemory(char* pGPUMemory, int MaxNHits)
{
+ //Set up Hits Memory Pointers for GPU Tracker
fHitMemory = (char*) pGPUMemory;
SetPointersHits(MaxNHits);
pGPUMemory += fHitMemorySize;
char* AliHLTTPCCATracker::SetGPUTrackerTrackletsMemory(char* pGPUMemory, int MaxNTracks)
{
+ //Set up Tracklet Memory Pointers for GPU Tracker
fTrackletMemory = (char*) pGPUMemory;
SetPointersTracklets(MaxNTracks);
pGPUMemory += fTrackletMemorySize;
char* AliHLTTPCCATracker::SetGPUTrackerTracksMemory(char* pGPUMemory, int MaxNTracks, int MaxNHits )
{
+ //Set up Tracks Memory Pointer for GPU Tracker
fTrackMemory = (char*) pGPUMemory;
SetPointersTracks(MaxNTracks, MaxNHits);
pGPUMemory += fTrackMemorySize;
void AliHLTTPCCATracker::DumpSliceData(std::ostream &out)
{
- out << "Slice Data:" << endl;
+ //Dump Slice Input Data to File
+ out << "Slice Data:" << std::endl;
for (int i = 0;i < Param().NRows();i++)
{
- out << "Row: " << i << endl;
+ out << "Row: " << i << std::endl;
for (int j = 0;j < Row(i).NHits();j++)
{
if (j && j % 16 == 0) out << std::endl;
out << j << '-' << Data().HitDataY(Row(i), j) << '-' << Data().HitDataZ(Row(i), j) << ", ";
}
- out << endl;
+ out << std::endl;
}
}
void AliHLTTPCCATracker::DumpLinks(std::ostream &out)
{
- out << "Hit Links:" << endl;
+ //Dump Links (after Neighbours Finder / Cleaner) to file
+ out << "Hit Links:" << std::endl;
for (int i = 0;i < Param().NRows();i++)
{
- out << "Row: " << i << endl;
+ out << "Row: " << i << std::endl;
for (int j = 0;j < Row(i).NHits();j++)
{
- if (j && j % 32 == 0) out << endl;
+ if (j && j % 32 == 0) out << std::endl;
out << HitLinkUpData(Row(i), j) << "/" << HitLinkDownData(Row(i), j) << ", ";
}
- out << endl;
+ out << std::endl;
}
}
void AliHLTTPCCATracker::DumpHitWeights(std::ostream &out)
{
- out << "Hit Weights:" << endl;
+ //dump hit weights to file
+ out << "Hit Weights:" << std::endl;
for (int i = 0;i < Param().NRows();i++)
{
- out << "Row: " << i << ":" << endl;
+ out << "Row: " << i << ":" << std::endl;
for (int j = 0;j < Row(i).NHits();j++)
{
- if (j && j % 32 == 0) out << endl;
+ if (j && j % 32 == 0) out << std::endl;
out << HitWeight(Row(i), j) << ", ";
}
- out << endl;
+ out << std::endl;
}
}
int AliHLTTPCCATracker::StarthitSortComparison(const void*a, const void* b)
{
+ //qsort helper function to sort start hits
AliHLTTPCCAHitId* aa = (AliHLTTPCCAHitId*) a;
AliHLTTPCCAHitId* bb = (AliHLTTPCCAHitId*) b;
void AliHLTTPCCATracker::DumpStartHits(std::ostream &out)
{
- out << "Start Hits: (" << *NTracklets() << ")" << endl;
+ //sort start hits and dump to file
+ out << "Start Hits: (" << *NTracklets() << ")" << std::endl;
#ifdef HLTCA_GPU_SORT_DUMPDATA
qsort(TrackletStartHits(), *NTracklets(), sizeof(AliHLTTPCCAHitId), StarthitSortComparison);
#endif
for (int i = 0;i < *NTracklets();i++)
{
- out << TrackletStartHit(i).RowIndex() << "-" << TrackletStartHit(i).HitIndex() << endl;
+ out << TrackletStartHit(i).RowIndex() << "-" << TrackletStartHit(i).HitIndex() << std::endl;
}
- out << endl;
+ out << std::endl;
}
void AliHLTTPCCATracker::DumpTrackHits(std::ostream &out)
{
- out << "Tracks: (" << *NTracks() << ")" << endl;
+ //dump tracks to file
+ out << "Tracks: (" << *NTracks() << ")" << std::endl;
#ifdef HLTCA_GPU_SORT_DUMPDATA
for (int k = 0;k < Param().NRows();k++)
{
{
out << TrackHits()[Tracks()[j].FirstHitID() + i].RowIndex() << "-" << TrackHits()[Tracks()[j].FirstHitID() + i].HitIndex() << ", ";
}
- out << "(Track: " << j << ")" << endl;
+ out << "(Track: " << j << ")" << std::endl;
#ifdef HLTCA_GPU_SORT_DUMPDATA
}
}
void AliHLTTPCCATracker::DumpTrackletHits(std::ostream &out)
{
- out << "Tracklets: (" << *NTracklets() << ")" << endl;
+ //dump tracklets to file
+ out << "Tracklets: (" << *NTracklets() << ")" << std::endl;
#ifdef HLTCA_GPU_SORT_DUMPDATA
AliHLTTPCCAHitId* tmpIds = new AliHLTTPCCAHitId[*NTracklets()];
AliHLTTPCCATracklet* tmpTracklets = new AliHLTTPCCATracklet[*NTracklets()];
#endif
for (int j = 0;j < *NTracklets();j++)
{
- out << "Tracklet " << j << " (Hits: " << setw(3) << Tracklets()[j].NHits() << ", Start: " << setw(3) << TrackletStartHit(j).RowIndex() << "-" << setw(3) << TrackletStartHit(j).HitIndex() << ") ";
+ out << "Tracklet " << j << " (Hits: " << std::setw(3) << Tracklets()[j].NHits() << ", Start: " << std::setw(3) << TrackletStartHit(j).RowIndex() << "-" << std::setw(3) << TrackletStartHit(j).HitIndex() << ") ";
if (Tracklets()[j].NHits() == 0);
else if (Tracklets()[j].LastRow() > Tracklets()[j].FirstRow() && (Tracklets()[j].FirstRow() >= Param().NRows() || Tracklets()[j].LastRow() >= Param().NRows()))
{
#endif
}
}
- out << endl;
+ out << std::endl;
}
}
GPUh() int AliHLTTPCCATracker::CheckEmptySlice() const
{
+ //Check if the Slice is empty, if so set the output apropriate and tell the reconstuct procesdure to terminate
if ( NHitsTotal() < 1 ) {
{
AliHLTTPCCASliceOutput::Allocate(*fOutput, 0, 0, fOutputControl);
void AliHLTTPCCATracker::RunNeighboursFinder()
{
+ //Run the CPU Neighbours Finder
AliHLTTPCCAProcess<AliHLTTPCCANeighboursFinder>( Param().NRows(), 1, *this );
}
void AliHLTTPCCATracker::RunNeighboursCleaner()
{
+ //Run the CPU Neighbours Cleaner
AliHLTTPCCAProcess<AliHLTTPCCANeighboursCleaner>( Param().NRows() - 2, 1, *this );
}
void AliHLTTPCCATracker::RunStartHitsFinder()
{
+ //Run the CPU Start Hits Finder
AliHLTTPCCAProcess<AliHLTTPCCAStartHitsFinder>( Param().NRows() - 4, 1, *this );
}
void AliHLTTPCCATracker::RunTrackletConstructor()
{
+ //Run CPU Tracklet Constructor
AliHLTTPCCATrackletConstructor::AliHLTTPCCATrackletConstructorNewCPU(*this);
}
void AliHLTTPCCATracker::RunTrackletSelector()
{
+ //Run CPU Tracklet Selector
AliHLTTPCCAProcess<AliHLTTPCCATrackletSelector>( 1, fCommonMem->fNTracklets, *this );
}
#ifdef HLTCA_STANDALONE
void AliHLTTPCCATracker::StandalonePerfTime(int i)
{
+ //Query Performance Timer for Standalone Version of Tracker
if (fGPUDebugLevel >= 1)
{
AliHLTTPCCAStandaloneFramework::StandaloneQueryTime(&fPerfTimers[i]);
void RunNeighboursFinder();
void RunNeighboursCleaner();
void RunStartHitsFinder();
- void RunTrackletConstructor() ;
+ void RunTrackletConstructor();
void RunTrackletSelector();
//GPU Tracker Interface
void SetGPUTracker();
void SetGPUDebugLevel(int Level, std::ostream *NewDebugOut = NULL) {fGPUDebugLevel = Level;if (NewDebugOut) fGPUDebugOut = NewDebugOut;}
- char* SetGPUTrackerCommonMemory(char* pGPUMemory);
+ char* SetGPUTrackerCommonMemory(char* const pGPUMemory);
char* SetGPUTrackerHitsMemory(char* pGPUMemory, int MaxNHits);
char* SetGPUTrackerTrackletsMemory(char* pGPUMemory, int MaxNTracklets);
char* SetGPUTrackerTracksMemory(char* pGPUMemory, int MaxNTracks, int MaxNHits );
void SetPointersHits( int MaxNHits );
void SetPointersTracklets ( int MaxNTracklets );
void SetPointersTracks( int MaxNTracks, int MaxNHits );
+ size_t SetPointersSliceData(const AliHLTTPCCAClusterData *data, bool allocate = false) { return(fData.SetPointers(data, allocate)); }
void SetOutput( AliHLTTPCCASliceOutput** out ) { fOutput = out; }
GPUhd() const AliHLTTPCCAParam &Param() const { return fParam; }
GPUhd() void SetParam( const AliHLTTPCCAParam &v ) { fParam = v; }
- GPUhd() const AliHLTTPCCASliceOutput::outputControlStruct* OutputControl() const { return fOutputControl; }
- GPUh() void SetOutputControl( AliHLTTPCCASliceOutput::outputControlStruct* val) { fOutputControl = val; }
-
+ GPUhd() const AliHLTTPCCASliceOutput::outputControlStruct* OutputControl() const { return fOutputControl; }
+ GPUh() void SetOutputControl( AliHLTTPCCASliceOutput::outputControlStruct* const val) { fOutputControl = val; }
+
GPUhd() AliHLTTPCCAClusterData *ClusterData() const { return fClusterData; }
GPUhd() const AliHLTTPCCASliceData &Data() const { return fData; }
GPUhd() AliHLTTPCCASliceData *PData() {return &fData; }
GPUhd() uint2* BlockStartingTracklet() const {return(fBlockStartingTracklet);}
GPUhd() StructGPUParameters* GPUParameters() const {return(&fCommonMem->fGPUParameters);}
GPUhd() StructGPUParametersConst* GPUParametersConst() {return(&fGPUParametersConst);}
-
+
GPUh() unsigned long long int* PerfTimer(unsigned int i) {return &fPerfTimers[i]; }
private:
AliHLTTPCCAHitId *fTrackletTmpStartHits; //Unsorted start hits
AliHLTTPCCATrackletConstructor::AliHLTTPCCAGPUTempMemory *fGPUTrackletTemp; //Temp Memory for GPU Tracklet Constructor
int* fRowBlockTracklets; //Reference which tracklet is processed in which rowblock next
- int4* fRowBlockPos; //x is last tracklet to be processed, y is last tracklet already processed, z is last tracklet to be processed in next iteration, w is initial x value to check if tracklet must be initialized
- uint2* fBlockStartingTracklet; // tracklet which starts the block
-
- StructGPUParametersConst fGPUParametersConst; //
-
+ int4* fRowBlockPos; //x is last tracklet to be processed, y is last tracklet already processed, z is last tracklet to be processed in next iteration, w is initial x value to check if tracklet must be initialized
+ uint2* fBlockStartingTracklet; // First Tracklet that is to be processed by current GPU MP
+
+ StructGPUParametersConst fGPUParametersConst; // Parameters for GPU if this is a GPU tracker
+
// event
commonMemoryStruct *fCommonMem; // common event memory
char *fHitMemory; // event memory for hits
- size_t fHitMemorySize; // size of the event memory [bytes]
-
- char *fTrackletMemory;
- size_t fTrackletMemorySize;
-
+ size_t fHitMemorySize; // size of the event memory for hits [bytes]
+
+ char *fTrackletMemory; //event memory for tracklets
+ size_t fTrackletMemorySize; //size of the event memory for tracklets
+
char *fTrackMemory; // event memory for tracks
- size_t fTrackMemorySize; // size of the event memory [bytes]
-
+ size_t fTrackMemorySize; // size of the event memory for tracks [bytes]
+
AliHLTTPCCAHitId *fTrackletStartHits; // start hits for the tracklets
AliHLTTPCCATracklet *fTracklets; // tracklets
- int *fTrackletRowHits;
-
+ int *fTrackletRowHits; //Hits for each Tracklet in each row
+
//
AliHLTTPCCATrack *fTracks; // reconstructed tracks
AliHLTTPCCAHitId *fTrackHits; // array of track hit numbers
// output
- AliHLTTPCCASliceOutput **fOutput;
+ AliHLTTPCCASliceOutput **fOutput; //address of pointer pointing to SliceOutput Object
// disable copy
AliHLTTPCCATracker( const AliHLTTPCCATracker& );
GPUd() void AliHLTTPCCATrackletConstructor::StoreTracklet
( int /*nBlocks*/, int /*nThreads*/, int /*iBlock*/, int /*iThread*/,
- AliHLTTPCCASharedMemory &/*s*/, AliHLTTPCCAThreadMemory &r, AliHLTTPCCATracker &tracker, AliHLTTPCCATrackParam &tParam )
+ AliHLTTPCCASharedMemory
+#ifdef HLTCA_GPUCODE
+ &s
+#else
+ &/*s*/
+#endif
+ , AliHLTTPCCAThreadMemory &r, AliHLTTPCCATracker &tracker, AliHLTTPCCATrackParam &tParam )
{
// reconstruction of tracklets, tracklet store step
GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
( int /*nBlocks*/, int /*nThreads*/, int /*iBlock*/, int /*iThread*/,
- AliHLTTPCCASharedMemory &/*s*/, AliHLTTPCCAThreadMemory &r, AliHLTTPCCATracker &tracker, AliHLTTPCCATrackParam &tParam, int iRow )
+ AliHLTTPCCASharedMemory
+#ifdef HLTCA_GPUCODE
+ &s
+#else
+ &/*s*/
+#endif
+ , AliHLTTPCCAThreadMemory &r, AliHLTTPCCATracker &tracker, AliHLTTPCCATrackParam &tParam, int iRow )
{
// reconstruction of tracklets, tracklets update step
// hh.y = reinterpret_cast<ushort_v*>( tmpint4 )[NextMultipleOf<sizeof(HLTCA_GPU_ROWALIGNMENT) / sizeof(ushort_v)>(row.NHits()) + r.fCurrIH];
//#else
#if defined(HLTCA_GPU_TEXTURE_FETCH)
- hh = tex1Dfetch(gAliTexRefu2, ((char*) tracker.Data().HitData() - tracker.pData()->GPUTextureBase()) / sizeof(ushort2) + row.HitNumberOffset() + r.fCurrIH);
+ hh = tex1Dfetch(gAliTexRefu2, ((char*) tracker.Data().HitData() - tracker.Data().GPUTextureBase()) / sizeof(ushort2) + row.HitNumberOffset() + r.fCurrIH);
#else
hh = tracker.HitData(row)[r.fCurrIH];
#endif
// r.fCurrIH = reinterpret_cast<short*>( tmpint4 )[2 * NextMultipleOf<sizeof(HLTCA_GPU_ROWALIGNMENT) / sizeof(ushort_v)>(row.NHits()) + r.fCurrIH]; // read from linkup data
//#else
#if defined(HLTCA_GPU_TEXTURE_FETCH)
- r.fCurrIH = tex1Dfetch(gAliTexRefs, ((char*) tracker.Data().HitLinkUpData(row) - tracker.pData()->GPUTextureBase()) / sizeof(unsigned short) + r.fCurrIH);
+ r.fCurrIH = tex1Dfetch(gAliTexRefs, ((char*) tracker.Data().HitLinkUpData(row) - tracker.Data().GPUTextureBase()) / sizeof(unsigned short) + r.fCurrIH);
#else
r.fCurrIH = tracker.HitLinkUpData(row)[r.fCurrIH]; // read from linkup data
#endif
//#endif
#ifdef HLTCA_GPU_TEXTURE_FETCH
- fHitYfst = tex1Dfetch(gAliTexRefu, ((char*) tracker.Data().FirstHitInBin(row) - tracker.pData()->GPUTextureBase()) / sizeof(unsigned short) + fIndYmin);
- fHitYlst = tex1Dfetch(gAliTexRefu, ((char*) tracker.Data().FirstHitInBin(row) - tracker.pData()->GPUTextureBase()) / sizeof(unsigned short) + fIndYmin+2);
- fHitYfst1 = tex1Dfetch(gAliTexRefu, ((char*) tracker.Data().FirstHitInBin(row) - tracker.pData()->GPUTextureBase()) / sizeof(unsigned short) + fIndYmin+nY);
- fHitYlst1 = tex1Dfetch(gAliTexRefu, ((char*) tracker.Data().FirstHitInBin(row) - tracker.pData()->GPUTextureBase()) / sizeof(unsigned short) + fIndYmin+nY+2);
+ fHitYfst = tex1Dfetch(gAliTexRefu, ((char*) tracker.Data().FirstHitInBin(row) - tracker.Data().GPUTextureBase()) / sizeof(unsigned short) + fIndYmin);
+ fHitYlst = tex1Dfetch(gAliTexRefu, ((char*) tracker.Data().FirstHitInBin(row) - tracker.Data().GPUTextureBase()) / sizeof(unsigned short) + fIndYmin+2);
+ fHitYfst1 = tex1Dfetch(gAliTexRefu, ((char*) tracker.Data().FirstHitInBin(row) - tracker.Data().GPUTextureBase()) / sizeof(unsigned short) + fIndYmin+nY);
+ fHitYlst1 = tex1Dfetch(gAliTexRefu, ((char*) tracker.Data().FirstHitInBin(row) - tracker.Data().GPUTextureBase()) / sizeof(unsigned short) + fIndYmin+nY+2);
#else
fHitYfst = sGridP[fIndYmin];
fHitYlst = sGridP[fIndYmin+2];
assert( (signed) fIh < row.NHits() );
ushort2 hh;
#if defined(HLTCA_GPU_TEXTURE_FETCH)
- hh = tex1Dfetch(gAliTexRefu2, ((char*) tracker.Data().HitData() - tracker.pData()->GPUTextureBase()) / sizeof(ushort2) + row.HitNumberOffset() + fIh);
+ hh = tex1Dfetch(gAliTexRefu2, ((char*) tracker.Data().HitData() - tracker.Data().GPUTextureBase()) / sizeof(ushort2) + row.HitNumberOffset() + fIh);
#else
hh = hits[fIh];
#endif
for ( unsigned int fIh = fHitYfst1; fIh < fHitYlst1; fIh++ ) {
ushort2 hh;
#if defined(HLTCA_GPU_TEXTURE_FETCH)
- hh = tex1Dfetch(gAliTexRefu2, ((char*) tracker.Data().HitData() - tracker.pData()->GPUTextureBase()) / sizeof(ushort2) + row.HitNumberOffset() + fIh);
+ hh = tex1Dfetch(gAliTexRefu2, ((char*) tracker.Data().HitData() - tracker.Data().GPUTextureBase()) / sizeof(ushort2) + row.HitNumberOffset() + fIh);
#else
hh = hits[fIh];
#endif
ushort2 hh;
#if defined(HLTCA_GPU_TEXTURE_FETCH)
- hh = tex1Dfetch(gAliTexRefu2, ((char*) tracker.Data().HitData() - tracker.pData()->GPUTextureBase()) / sizeof(ushort2) + row.HitNumberOffset() + best);
+ hh = tex1Dfetch(gAliTexRefu2, ((char*) tracker.Data().HitData() - tracker.Data().GPUTextureBase()) / sizeof(ushort2) + row.HitNumberOffset() + best);
#else
hh = hits[best];
#endif
const int nFillTracks = CAMath::Min(nFetchTracks, nUseTrack + nFetchTracks - (*((volatile int2*) (tracker.RowBlockPos(Reverse, RowBlock)))).x);
if (nFillTracks > 0)
{
- const int nStartFillTrack = CAMath::AtomicAdd(&(*tracker.RowBlockPos(Reverse, RowBlock)).x, nFillTracks);\r
- if (nFillTracks + nStartFillTrack >= HLTCA_GPU_MAX_TRACKLETS)\r
- {\r
- tracker.GPUParameters()->fGPUError = HLTCA_GPU_ERROR_ROWBLOCK_TRACKLET_OVERFLOW;\r
- }\r
- for (int i = 0;i < nFillTracks;i++)\r
- {\r
- tracker.RowBlockTracklets(Reverse, RowBlock)[(nStartFillTrack + i) % HLTCA_GPU_MAX_TRACKLETS] = -3; //Dummy filling track\r
- }\r
+ const int nStartFillTrack = CAMath::AtomicAdd(&(*tracker.RowBlockPos(Reverse, RowBlock)).x, nFillTracks);
+ if (nFillTracks + nStartFillTrack >= HLTCA_GPU_MAX_TRACKLETS)
+ {
+ tracker.GPUParameters()->fGPUError = HLTCA_GPU_ERROR_ROWBLOCK_TRACKLET_OVERFLOW;
+ }
+ for (int i = 0;i < nFillTracks;i++)
+ {
+ tracker.RowBlockTracklets(Reverse, RowBlock)[(nStartFillTrack + i) % HLTCA_GPU_MAX_TRACKLETS] = -3; //Dummy filling track
+ }
}
- sMem.fNextTrackletNoDummy = 0;\r
+ sMem.fNextTrackletNoDummy = 0;
}
}
__syncthreads();
}
else
{
- const int nTrackPos = sMem.fNextTrackletFirst + threadIdx.x - TRACKLET_CONSTRUCTOR_NMEMTHREDS;\r
- mustInit = (nTrackPos < tracker.RowBlockPos(Reverse, RowBlock)->w);\r
- volatile int* const ptrTracklet = &tracker.RowBlockTracklets(Reverse, RowBlock)[nTrackPos % HLTCA_GPU_MAX_TRACKLETS];\r
- int nTracklet;\r
- int nTryCount = 0;\r
- while ((nTracklet = *ptrTracklet) == -1)\r
- {\r
- for (int i = 0;i < 10000;i++)\r
- sMem.fNextTrackletStupidDummy++;\r
- nTryCount++;\r
- if (nTryCount > 20)\r
- {\r
- tracker.GPUParameters()->fGPUError = HLTCA_GPU_ERROR_SCHEDULE_COLLISION;\r
- return(-1);\r
- }\r
- };\r
- return(nTracklet);\r
+ const int nTrackPos = sMem.fNextTrackletFirst + threadIdx.x - TRACKLET_CONSTRUCTOR_NMEMTHREDS;
+ mustInit = (nTrackPos < tracker.RowBlockPos(Reverse, RowBlock)->w);
+ volatile int* const ptrTracklet = &tracker.RowBlockTracklets(Reverse, RowBlock)[nTrackPos % HLTCA_GPU_MAX_TRACKLETS];
+ int nTracklet;
+ int nTryCount = 0;
+ while ((nTracklet = *ptrTracklet) == -1)
+ {
+ for (int i = 0;i < 10000;i++)
+ sMem.fNextTrackletStupidDummy++;
+ nTryCount++;
+ if (nTryCount > 20)
+ {
+ tracker.GPUParameters()->fGPUError = HLTCA_GPU_ERROR_SCHEDULE_COLLISION;
+ return(-1);
+ }
+ };
+ return(nTracklet);
}
}
#else
if (iTracklet >= 0)
{
- for (int j = rMem.fStartRow;j < HLTCA_ROW_COUNT;j++)\r
- {\r
- UpdateTracklet(gridDim.x, blockDim.x, blockIdx.x, threadIdx.x, sMem, rMem, tracker, tParam, j);\r
- if (!rMem.fGo) break;\r
- }\r
-\r
- rMem.fNMissed = 0;\r
- rMem.fStage = 2;\r
- if ( rMem.fGo )\r
- {\r
- if ( !tParam.TransportToX( tracker.Row( rMem.fEndRow ).X(), tracker.Param().ConstBz(), .999 ) ) rMem.fGo = 0;\r
- }\r
-\r
- for (int j = rMem.fEndRow;j >= 0;j--)\r
- {\r
- if (!rMem.fGo) break;\r
- UpdateTracklet( gridDim.x, blockDim.x, blockIdx.x, threadIdx.x, sMem, rMem, tracker, tParam, j);\r
- }\r
-\r
+ for (int j = rMem.fStartRow;j < HLTCA_ROW_COUNT;j++)
+ {
+ UpdateTracklet(gridDim.x, blockDim.x, blockIdx.x, threadIdx.x, sMem, rMem, tracker, tParam, j);
+ if (!rMem.fGo) break;
+ }
+
+ rMem.fNMissed = 0;
+ rMem.fStage = 2;
+ if ( rMem.fGo )
+ {
+ if ( !tParam.TransportToX( tracker.Row( rMem.fEndRow ).X(), tracker.Param().ConstBz(), .999 ) ) rMem.fGo = 0;
+ }
+
+ for (int j = rMem.fEndRow;j >= 0;j--)
+ {
+ if (!rMem.fGo) break;
+ UpdateTracklet( gridDim.x, blockDim.x, blockIdx.x, threadIdx.x, sMem, rMem, tracker, tParam, j);
+ }
+
StoreTracklet( gridDim.x, blockDim.x, blockIdx.x, threadIdx.x, sMem, rMem, tracker, tParam );
}
#endif
rMem.fItr = iTracklet;
rMem.fGo = 1;
- for (int j = rMem.fStartRow;j < tracker.Param().NRows();j++)\r
- {\r
- UpdateTracklet(1, 1, 0, iTracklet, sMem, rMem, tracker, tParam, j);\r
- if (!rMem.fGo) break;\r
- }\r
-\r
- rMem.fNMissed = 0;\r
- rMem.fStage = 2;\r
- if ( rMem.fGo )\r
- {\r
- if ( !tParam.TransportToX( tracker.Row( rMem.fEndRow ).X(), tracker.Param().ConstBz(), .999 ) ) rMem.fGo = 0;\r
- }\r
-\r
- for (int j = rMem.fEndRow;j >= 0;j--)\r
- {\r
- if (!rMem.fGo) break;\r
- UpdateTracklet( 1, 1, 0, iTracklet, sMem, rMem, tracker, tParam, j);\r
- }\r
-\r
+ for (int j = rMem.fStartRow;j < tracker.Param().NRows();j++)
+ {
+ UpdateTracklet(1, 1, 0, iTracklet, sMem, rMem, tracker, tParam, j);
+ if (!rMem.fGo) break;
+ }
+
+ rMem.fNMissed = 0;
+ rMem.fStage = 2;
+ if ( rMem.fGo )
+ {
+ if ( !tParam.TransportToX( tracker.Row( rMem.fEndRow ).X(), tracker.Param().ConstBz(), .999 ) ) rMem.fGo = 0;
+ }
+
+ for (int j = rMem.fEndRow;j >= 0;j--)
+ {
+ if (!rMem.fGo) break;
+ UpdateTracklet( 1, 1, 0, iTracklet, sMem, rMem, tracker, tParam, j);
+ }
+
StoreTracklet( 1, 1, 0, iTracklet, sMem, rMem, tracker, tParam );
}
}