update of GPU tracker from David Rohr
authorsgorbuno <sgorbuno@f7af4fe6-9843-0410-8265-dc069ae4e863>
Thu, 23 Jul 2009 13:52:49 +0000 (13:52 +0000)
committersgorbuno <sgorbuno@f7af4fe6-9843-0410-8265-dc069ae4e863>
Thu, 23 Jul 2009 13:52:49 +0000 (13:52 +0000)
39 files changed:
HLT/TPCLib/tracking-ca/AliHLTTPCCAClusterData.cxx
HLT/TPCLib/tracking-ca/AliHLTTPCCAClusterData.h
HLT/TPCLib/tracking-ca/AliHLTTPCCADataCompressor.h
HLT/TPCLib/tracking-ca/AliHLTTPCCADef.h
HLT/TPCLib/tracking-ca/AliHLTTPCCAGPUTracker.cu [new file with mode: 0644]
HLT/TPCLib/tracking-ca/AliHLTTPCCAGPUTracker.h [new file with mode: 0644]
HLT/TPCLib/tracking-ca/AliHLTTPCCAGlobalMergerComponent.cxx
HLT/TPCLib/tracking-ca/AliHLTTPCCAGrid.cxx
HLT/TPCLib/tracking-ca/AliHLTTPCCAGrid.h
HLT/TPCLib/tracking-ca/AliHLTTPCCAHit.h
HLT/TPCLib/tracking-ca/AliHLTTPCCAHitArea.h
HLT/TPCLib/tracking-ca/AliHLTTPCCAHitId.h
HLT/TPCLib/tracking-ca/AliHLTTPCCAMath.h
HLT/TPCLib/tracking-ca/AliHLTTPCCAMerger.cxx
HLT/TPCLib/tracking-ca/AliHLTTPCCANeighboursCleaner.h
HLT/TPCLib/tracking-ca/AliHLTTPCCANeighboursFinder.cxx
HLT/TPCLib/tracking-ca/AliHLTTPCCANeighboursFinder.h
HLT/TPCLib/tracking-ca/AliHLTTPCCAParam.cxx
HLT/TPCLib/tracking-ca/AliHLTTPCCAParam.h
HLT/TPCLib/tracking-ca/AliHLTTPCCAProcess.h
HLT/TPCLib/tracking-ca/AliHLTTPCCARow.h
HLT/TPCLib/tracking-ca/AliHLTTPCCASliceData.cxx
HLT/TPCLib/tracking-ca/AliHLTTPCCASliceData.h
HLT/TPCLib/tracking-ca/AliHLTTPCCASliceOutput.cxx
HLT/TPCLib/tracking-ca/AliHLTTPCCASliceOutput.h
HLT/TPCLib/tracking-ca/AliHLTTPCCAStandaloneFramework.cxx
HLT/TPCLib/tracking-ca/AliHLTTPCCAStandaloneFramework.h
HLT/TPCLib/tracking-ca/AliHLTTPCCAStartHitsFinder.h
HLT/TPCLib/tracking-ca/AliHLTTPCCATrack.h
HLT/TPCLib/tracking-ca/AliHLTTPCCATrackLinearisation.h
HLT/TPCLib/tracking-ca/AliHLTTPCCATrackParam.cxx
HLT/TPCLib/tracking-ca/AliHLTTPCCATrackParam.h
HLT/TPCLib/tracking-ca/AliHLTTPCCATracker.cxx
HLT/TPCLib/tracking-ca/AliHLTTPCCATracker.h
HLT/TPCLib/tracking-ca/AliHLTTPCCATracklet.h
HLT/TPCLib/tracking-ca/AliHLTTPCCATrackletConstructor.cxx
HLT/TPCLib/tracking-ca/AliHLTTPCCATrackletConstructor.h
HLT/TPCLib/tracking-ca/AliHLTTPCCATrackletSelector.h
HLT/TPCLib/tracking-ca/MemoryAssignmentHelpers.h

index 2054e09..cf07955 100644 (file)
@@ -16,6 +16,8 @@
 
 #include "AliHLTTPCCAClusterData.h"
 #include "AliHLTTPCCAMath.h"
+#include <algorithm>
+#include "AliHLTArray.h"
 
 void AliHLTTPCCAClusterData::StartReading( int sliceIndex, int guessForNumberOfClusters )
 {
@@ -59,4 +61,41 @@ void AliHLTTPCCAClusterData::FinishReading()
   fLastRow = row; // the last seen row is the last row in this slice
 }
 
+template <class T> void AliHLTTPCCAClusterData::WriteEventVector(const std::vector<T> &data, std::ostream &out) const
+{
+       AliHLTResizableArray<T> tmpData(data.size());
+       unsigned i;
+       for (i = 0;i < data.size();i++)
+       {
+               tmpData[i] = data[i];
+       }
+       i = data.size();
+       out.write((char*) &i, sizeof(i));
+       out.write((char*) &tmpData[0], i * sizeof(T));
+}
+
+template <class T> void AliHLTTPCCAClusterData::ReadEventVector(std::vector<T> &data, std::istream &in, int MinSize)
+{
+       int i;
+       in.read((char*) &i, sizeof(i));
+       data.reserve(AliHLTTPCCAMath::Max(MinSize, i));
+       data.resize(i);
+       AliHLTResizableArray<T> tmpData(i);
+       in.read((char*) &tmpData[0], i * sizeof(T));
+       for (int j = 0;j < i;j++)
+       {
+               data[j] = tmpData[j];
+       }
+}
+
+void AliHLTTPCCAClusterData::WriteEvent(std::ostream &out) const
+{
+       WriteEventVector<Data>(fData, out);
+}
+
+void AliHLTTPCCAClusterData::ReadEvent(std::istream &in)
+{
+    fData.clear();
+       ReadEventVector<Data>(fData, in, 64);
+}
 
index 12e86e9..758b12f 100644 (file)
@@ -17,6 +17,7 @@
 #ifndef ALIHLTTPCCACLUSTERDATA_H
 #define ALIHLTTPCCACLUSTERDATA_H
 
+#include <iostream>
 #include <vector>
 
 /**
@@ -48,6 +49,13 @@ class AliHLTTPCCAClusterData
      */
     void FinishReading();
 
+    /**
+     * Read/Write Events from/to file
+     */
+    void ReadEvent(std::istream &in);
+    void WriteEvent(std::ostream &out) const;
+       template <class T> void ReadEventVector(std::vector<T> &data, std::istream &in, int MinSize = 0);
+    template <class T> void WriteEventVector(const std::vector<T> &data, std::ostream &out) const;
 
     /**
      * "remove" one cluster and "add" two new ones, keeping history.
@@ -74,7 +82,7 @@ class AliHLTTPCCAClusterData
     /**
      * Return the number of clusters in this slice.
      */
-    int NumberOfClusters() const { return fData.size(); }
+    int NumberOfClusters() const { return (int) fData.size(); }
 
     /**
      * Return the number of clusters in the given row, for this slice.
@@ -94,7 +102,7 @@ class AliHLTTPCCAClusterData
      * for ( int hitIndex = cd.RowOffset( rowIndex ); hitIndex < lastClusterIndex; ++hitIndex )
      * \endcode
      */
-    int RowOffset( unsigned int rowIndex ) const { return rowIndex < fRowOffset.size() ? fRowOffset[rowIndex] : fData.size(); }
+    int RowOffset( unsigned int rowIndex ) const { return rowIndex < fRowOffset.size() ? fRowOffset[rowIndex] : (int) fData.size(); }
 
     /**
      * Return the x coordinate of the given cluster.
index ed892d9..9b6800e 100644 (file)
@@ -54,8 +54,8 @@ GPUhd() inline unsigned short AliHLTTPCCADataCompressor::YZ2UShort( float Y, flo
   // compress Y and Z coordinates in range [-3., 3.] to 16 bits
 
   const float kMult = 255. / 6.;
-  Y = ( Y + 3. ) * kMult;
-  Z = ( Z + 3. ) * kMult;
+  Y = ( Y + 3.f ) * kMult;
+  Z = ( Z + 3.f ) * kMult;
   if ( Y < 0. ) Y = 0.;
   else if ( Y > 255. ) Y = 255.;
   if ( Z < 0. ) Z = 0.;
@@ -67,16 +67,16 @@ GPUhd() inline float AliHLTTPCCADataCompressor::UShort2Y( unsigned short iYZ )
 {
   // extract Y coordinate from the compressed 16bits format to [-3.,3.]
 
-  const float kMult = 6. / 255.;
-  return ( iYZ >> 8 )*kMult - 3.;
+  const float kMult = 6.f / 255.f;
+  return ( iYZ >> 8 )*kMult - 3.f;
 }
 
 GPUhd() inline float AliHLTTPCCADataCompressor::UShort2Z( unsigned short iYZ )
 {
   // extract Z coordinate from the compressed 16bits format to [-3.,3.]
 
-  const float kMult = 6. / 255.;
-  return ( iYZ % 256 )*kMult - 3.;
+  const float kMult = 6.f / 255.f;
+  return ( iYZ % 256 )*kMult - 3.f;
 }
 
 #endif
index bbb79c6..88fefc6 100644 (file)
 #define HLTCA_INTERNAL_PERFORMANCE
 
 #ifdef __CUDACC__
-
 #define HLTCA_GPUCODE
+#endif
 
+#ifdef WIN32
+#ifndef R__WIN32
+#define R__WIN32
+#endif
+#endif
+
+#if defined(R__WIN32)
+#ifdef INTEL_RUNTIME\r
+#pragma warning(disable : 1786)\r
+#pragma warning(disable : 1478)\r
+#pragma warning(disable : 161)\r
+#endif\r
+\r
+#ifdef VSNET_RUNTIME\r
+#pragma warning(disable : 4616)\r
+#pragma warning(disable : 4996)\r
+#pragma warning(disable : 1684)\r
+#endif
 #endif
 
 #if defined(HLTCA_STANDALONE) || defined(HLTCA_GPUCODE)
@@ -78,6 +96,17 @@ namespace AliHLTTPCCADefinitions
 
 #endif
 
+#ifdef HLTCA_GPUCODE
+#define ALIHLTTPCCANEIGHBOURS_FINDER_MAX_NNEIGHUP 5
+#define ALIHLTTPCCANEIGHBOURS_FINDER_MAX_FGRIDCONTENTUPDOWN 700
+#define ALIHLTTPCCASTARTHITSFINDER_MAX_FROWSTARTHITS 3500
+#define ALIHLTTPCCATRACKLET_CONSTRUCTOR_TEMP_MEM 650
+#else
+#define ALIHLTTPCCANEIGHBOURS_FINDER_MAX_NNEIGHUP 20
+#define ALIHLTTPCCANEIGHBOURS_FINDER_MAX_FGRIDCONTENTUPDOWN 7000
+#define ALIHLTTPCCASTARTHITSFINDER_MAX_FROWSTARTHITS 10000
+#define ALIHLTTPCCATRACKLET_CONSTRUCTOR_TEMP_MEM 5000
+#endif
 
 #ifdef HLTCA_GPUCODE
 
@@ -106,6 +135,15 @@ struct ushort2 { unsigned short x; unsigned short y; };
 struct uint1 { unsigned int x; };
 struct uint4 { unsigned int x, y, z, w; };
 
+#ifdef R__WIN32
+#include <float.h>
+
+inline bool finite(float x)
+{
+       return(x <= FLT_MAX);
+}
+#endif
+
 #endif
 
 /*
diff --git a/HLT/TPCLib/tracking-ca/AliHLTTPCCAGPUTracker.cu b/HLT/TPCLib/tracking-ca/AliHLTTPCCAGPUTracker.cu
new file mode 100644 (file)
index 0000000..b881a5e
--- /dev/null
@@ -0,0 +1,327 @@
+// **************************************************************************
+// This file is property of and copyright by the ALICE HLT Project          *
+// ALICE Experiment at CERN, All rights reserved.                           *
+//                                                                          *
+// Primary Authors: Sergey Gorbunov <sergey.gorbunov@kip.uni-heidelberg.de> *
+//                  Ivan Kisel <kisel@kip.uni-heidelberg.de>                *
+//                                     David Rohr <drohr@kip.uni-heidelberg.de>                                *
+//                  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 <cutil.h>
+#include <cutil_inline_runtime.h>
+#include <sm_11_atomic_functions.h>
+#include <sm_12_atomic_functions.h>
+
+#include <iostream>
+
+//Disable assertions since they produce errors in GPU Code
+#ifdef assert
+#undef assert
+#endif
+#define assert(param)
+
+#include "AliHLTTPCCAGPUTracker.h"
+
+#ifdef BUILD_GPU
+
+//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 "AliHLTTPCCANeighboursFinder.cxx"
+
+#include "AliHLTTPCCANeighboursCleaner.cxx"
+#include "AliHLTTPCCAStartHitsFinder.cxx"
+#include "AliHLTTPCCATrackletConstructor.cxx"
+#include "AliHLTTPCCASliceOutput.cxx"
+
+#endif
+
+AliHLTTPCCAGPUTracker::AliHLTTPCCAGPUTracker() : gpuTracker(), DebugLevel(0) {}
+AliHLTTPCCAGPUTracker::~AliHLTTPCCAGPUTracker() {}
+
+//Find best CUDA device, initialize and allocate memory
+int AliHLTTPCCAGPUTracker::InitGPU()
+{
+#ifdef BUILD_GPU
+       int cudaDevice = cutGetMaxGflopsDeviceId();
+       cudaSetDevice(cudaDevice);
+
+  cudaDeviceProp prop;
+  cudaGetDeviceProperties(&prop ,0 ); 
+  std::cout<<"CUDA Device Properties: "<<std::endl;
+  std::cout<<"name = "<<prop.name<<std::endl;
+  std::cout<<"totalGlobalMem = "<<prop.totalGlobalMem<<std::endl;
+  std::cout<<"sharedMemPerBlock = "<<prop.sharedMemPerBlock<<std::endl;
+  std::cout<<"regsPerBlock = "<<prop.regsPerBlock<<std::endl;
+  std::cout<<"warpSize = "<<prop.warpSize<<std::endl;
+  std::cout<<"memPitch = "<<prop.memPitch<<std::endl;
+  std::cout<<"maxThreadsPerBlock = "<<prop.maxThreadsPerBlock<<std::endl;
+  std::cout<<"maxThreadsDim = "<<prop.maxThreadsDim[0]<<" "<<prop.maxThreadsDim[1]<<" "<<prop.maxThreadsDim[2]<<std::endl;
+  std::cout<<"maxGridSize = "  <<prop.maxGridSize[0]<<" "<<prop.maxGridSize[1]<<" "<<prop.maxGridSize[2]<<std::endl;
+  std::cout<<"totalConstMem = "<<prop.totalConstMem<<std::endl;
+  std::cout<<"major = "<<prop.major<<std::endl;
+  std::cout<<"minor = "<<prop.minor<<std::endl;
+  std::cout<<"clockRate = "<<prop.clockRate<<std::endl;
+  std::cout<<"textureAlignment = "<<prop.textureAlignment<<std::endl;
+
+  GPUMemSize = (int) prop.totalGlobalMem - 400 * 1024 * 1024;
+  if (CUDA_FAILED_MSG(cudaMalloc(&GPUMemory, (size_t) GPUMemSize)))
+  {
+         std::cout << "CUDA Memory Allocation Error\n";
+         return(1);
+  }
+  std::cout << "CUDA Initialisation successfull\n";
+#endif
+
+       return(0);
+}
+
+//Macro to align Pointers.
+//Will align to start at 1 MB segments, this should be consistent with every alignment in the tracker
+//(As long as every single data structure is <= 1 MB)
+template <class T> inline T* AliHLTTPCCAGPUTracker::alignPointer(T* ptr, int alignment)
+{
+       size_t adr = (size_t) ptr;
+       if (adr % alignment)
+       {
+               adr += alignment - (adr % alignment);
+       }
+       return((T*) adr);
+}
+
+//Check for CUDA Error and in the case of an error display the corresponding error string
+bool AliHLTTPCCAGPUTracker::CUDA_FAILED_MSG(cudaError_t error)
+{
+       if (error == cudaSuccess) return(false);
+       printf("CUDA Error: %d / %s\n", error, cudaGetErrorString(error));
+       return(true);
+}
+
+//Wait for CUDA-Kernel to finish and check for CUDA errors afterwards
+int AliHLTTPCCAGPUTracker::CUDASync()
+{
+       if (DebugLevel == 0) return(0);
+       cudaError cuErr;
+       cuErr = cudaGetLastError();
+       if (cuErr != cudaSuccess)
+       {
+               printf("Cuda Error %s while invoking kernel\n", cudaGetErrorString(cuErr));
+               return(1);
+       }
+       if (CUDA_FAILED_MSG(cudaThreadSynchronize()))
+       {
+               printf("CUDA Error while synchronizing\n");
+               return(1);
+       }
+       if (DebugLevel >= 4) printf("CUDA Sync Done\n");
+       return(0);
+}
+
+void AliHLTTPCCAGPUTracker::SetDebugLevel(int dwLevel, std::ostream *NewOutFile)
+{
+       DebugLevel = dwLevel;
+       if (NewOutFile) OutFile = NewOutFile;
+}
+
+//Primary reconstruction function
+int AliHLTTPCCAGPUTracker::Reconstruct(AliHLTTPCCATracker* tracker)
+{
+    int nThreads;
+    int nBlocks;
+       int size;
+
+       if (tracker->CheckEmptySlice())
+       {
+               if (DebugLevel >= 4) printf("Slice Empty, not running GPU Tracker\n");
+               return(0);
+       }
+
+       if (DebugLevel >= 3)
+       {
+               *OutFile << endl << endl << "Slice: " << tracker->Param().ISlice() << endl;
+       }
+
+       if (DebugLevel >= 4) printf("\n\nInitialising GPU Tracker\n");
+       memcpy(&gpuTracker, tracker, sizeof(AliHLTTPCCATracker));
+       char* tmpMem = alignPointer((char*) GPUMemory, 1024 * 1024);
+       gpuTracker.SetGPUTracker();
+
+       if (DebugLevel >= 4) printf("Initialising GPU Common Memory\n");
+       tmpMem = gpuTracker.SetGPUTrackerCommonMemory(tmpMem);
+       tmpMem = alignPointer(tmpMem, 1024 * 1024);
+
+       if (DebugLevel >= 4) printf("Initialising GPU Hits Memory\n");
+       tmpMem = gpuTracker.SetGPUTrackerHitsMemory(tmpMem, tracker->NHitsTotal());
+       tmpMem = alignPointer(tmpMem, 1024 * 1024);
+
+       if (DebugLevel >= 4) printf("Initialising GPU Slice Data Memory\n");
+       tmpMem = gpuTracker.fData.SetGPUSliceDataMemory(tmpMem, gpuTracker.fClusterData);
+       tmpMem = alignPointer(tmpMem, 1024 * 1024);
+       if (tmpMem - (char*) GPUMemory > GPUMemSize)
+       {
+               printf("Out of CUDA Memory\n");
+               return(1);
+       }
+       
+       CUDA_FAILED_MSG(cudaMemcpy(gpuTracker.fCommonMemory, tracker->fCommonMemory, tracker->fCommonMemorySize, cudaMemcpyHostToDevice));
+       CUDA_FAILED_MSG(cudaMemcpy(gpuTracker.fData.fMemory, tracker->fData.fMemory, tracker->fData.fMemorySize, cudaMemcpyHostToDevice));
+       CUDA_FAILED_MSG(cudaMemcpyToSymbol(gAliHLTTPCCATracker, &gpuTracker, sizeof(AliHLTTPCCATracker)));
+
+       if (DebugLevel >= 4) printf("Running GPU Neighbours Finder\n");
+       AliHLTTPCCAProcess<AliHLTTPCCANeighboursFinder> <<<gpuTracker.Param().NRows(), 256>>>();
+       if (CUDASync()) return 1;
+
+       if (DebugLevel >= 3)
+       {
+               *OutFile << "Neighbours Finder:" << endl;
+               CUDA_FAILED_MSG(cudaMemcpy(tracker->fData.fMemory, gpuTracker.fData.fMemory, tracker->fData.fMemorySize, cudaMemcpyDeviceToHost));
+               tracker->DumpLinks(*OutFile);
+    }
+
+       if (DebugLevel >= 4) printf("Running GPU Neighbours Cleaner\n");
+       AliHLTTPCCAProcess<AliHLTTPCCANeighboursCleaner> <<<gpuTracker.Param().NRows()-2, 256>>>();
+       if (CUDASync()) return 1;
+
+       if (DebugLevel >= 3)
+       {
+               *OutFile << "Neighbours Cleaner:" << endl;
+               CUDA_FAILED_MSG(cudaMemcpy(tracker->fData.fMemory, gpuTracker.fData.fMemory, tracker->fData.fMemorySize, cudaMemcpyDeviceToHost));
+               tracker->DumpLinks(*OutFile);
+    }
+
+       if (DebugLevel >= 4) printf("Running GPU Start Hits Finder\n");
+       AliHLTTPCCAProcess<AliHLTTPCCAStartHitsFinder> <<<gpuTracker.Param().NRows()-4, 256>>>();
+       if (CUDASync()) return 1;
+
+       if (DebugLevel >= 4) printf("Obtaining Number of Start Hits from GPU: ");
+       CUDA_FAILED_MSG(cudaMemcpy(tracker->fCommonMemory, gpuTracker.fCommonMemory, tracker->fCommonMemorySize, cudaMemcpyDeviceToHost));
+       if (DebugLevel >= 4) printf("%d\n", *tracker->NTracklets());
+       else if (DebugLevel >= 2) printf("%3d ", *tracker->NTracklets());
+
+       if (DebugLevel >= 3)
+       {
+               *OutFile << "Start Hits: (" << *tracker->NTracklets() << ")" << endl;
+               CUDA_FAILED_MSG(cudaMemcpy(tracker->fHitMemory, gpuTracker.fHitMemory, tracker->fHitMemorySize, cudaMemcpyDeviceToHost));
+               tracker->DumpStartHits(*OutFile);
+    }
+
+       /*tracker->RunNeighboursFinder();
+       tracker->RunNeighboursCleaner();
+       tracker->RunStartHitsFinder();*/
+
+       if (DebugLevel >= 4) printf("Initialising GPU Track Memory\n");
+       tmpMem = gpuTracker.SetGPUTrackerTracksMemory(tmpMem, *tracker->NTracklets(), tracker->NHitsTotal());
+       tmpMem = alignPointer(tmpMem, 1024 * 1024);
+       if (tmpMem - (char*) GPUMemory > GPUMemSize)
+       {
+               printf("Out of CUDA Memory\n");
+               return(1);
+       }
+
+       tracker->fData.ClearHitWeights();
+       CUDA_FAILED_MSG(cudaMemcpy(gpuTracker.fData.fHitWeights, tracker->fData.fHitWeights, tracker->fData.fNumberOfHits * sizeof(int), cudaMemcpyHostToDevice));
+       CUDA_FAILED_MSG(cudaMemcpyToSymbol(gAliHLTTPCCATracker, &gpuTracker, sizeof(AliHLTTPCCATracker)));
+
+       if (DebugLevel >= 4) printf("Initialising Slice Tracker (CPU) Track Memory\n");
+       tracker->fTrackMemory = reinterpret_cast<char*> ( new uint4 [ gpuTracker.fTrackMemorySize/sizeof( uint4 ) + 100] );
+    tracker->SetPointersTracks( *tracker->NTracklets(), tracker->NHitsTotal() );
+
+/*     tracker->RunTrackletConstructor();
+       if (DebugLevel >= 3)
+       {
+               *OutFile << "Tracklet Hits:" << endl;
+               tracker->DumpTrackletHits(*OutFile);
+       }*/
+
+       int nMemThreads = TRACKLET_CONSTRUCTOR_NMEMTHREDS;
+    nThreads = 256;//96;
+    nBlocks = *tracker->NTracklets()/nThreads + 1;
+    if( nBlocks<30 ){
+               nBlocks = 30;
+               nThreads = (*tracker->NTracklets())/nBlocks+1;
+               if( nThreads%32 ) nThreads = (nThreads/32+1)*32;
+       }
+       if (DebugLevel >= 4) printf("Running GPU Tracklet Constructor\n");
+
+       //AliHLTTPCCAProcess1<AliHLTTPCCATrackletConstructor> <<<nBlocks, nMemThreads+nThreads>>>(); 
+       AliHLTTPCCAProcess1<AliHLTTPCCATrackletConstructor> <<<1, TRACKLET_CONSTRUCTOR_NMEMTHREDS + *tracker->fNTracklets>>>();
+       if (CUDASync()) return 1;
+
+       if (DebugLevel >= 3)
+       {
+               *OutFile << "Tracklet Hits:" << endl;
+               CUDA_FAILED_MSG(cudaMemcpy(tracker->fNTracklets, gpuTracker.fNTracklets, sizeof(int), cudaMemcpyDeviceToHost));
+               CUDA_FAILED_MSG(cudaMemcpy(tracker->fTracklets, gpuTracker.fTracklets, gpuTracker.fTrackMemorySize, cudaMemcpyDeviceToHost));
+               tracker->DumpTrackletHits(*OutFile);
+    }
+
+       //tracker->RunTrackletSelector();
+       
+
+       nThreads = 128;
+       nBlocks = *tracker->NTracklets()/nThreads + 1;
+       if( nBlocks<30 ){
+         nBlocks = 30;  
+         nThreads = *tracker->NTracklets()/nBlocks+1;
+         nThreads = (nThreads/32+1)*32;
+       }
+       if (DebugLevel >= 4) printf("Running GPU Tracklet Selector\n");
+       AliHLTTPCCAProcess<AliHLTTPCCATrackletSelector><<<nBlocks, nThreads>>>();
+       //AliHLTTPCCAProcess<AliHLTTPCCATrackletSelector><<<1, *tracker->fNTracklets>>>();
+       if (CUDASync()) return 1;
+
+       if (DebugLevel >= 4) printf("Transfering Tracks from GPU to Host ");
+       CUDA_FAILED_MSG(cudaMemcpy(tracker->NTracks(), gpuTracker.NTracks(), sizeof(int), cudaMemcpyDeviceToHost));
+       CUDA_FAILED_MSG(cudaMemcpy(tracker->NTrackHits(), gpuTracker.NTrackHits(), sizeof(int), cudaMemcpyDeviceToHost));
+       if (DebugLevel >= 4) printf("%d / %d\n", *tracker->fNTracks, *tracker->fNTrackHits);
+       size = sizeof(AliHLTTPCCATrack) * *tracker->NTracks();
+       CUDA_FAILED_MSG(cudaMemcpy(tracker->Tracks(), gpuTracker.Tracks(), size, cudaMemcpyDeviceToHost));
+       size = sizeof(AliHLTTPCCAHitId) * *tracker->NTrackHits();
+       if (CUDA_FAILED_MSG(cudaMemcpy(tracker->TrackHits(), gpuTracker.TrackHits(), size, cudaMemcpyDeviceToHost)))
+       {
+               printf("CUDA Error during Reconstruction\n");
+               return(1);
+       }
+
+       if (DebugLevel >= 3)
+       {
+               *OutFile << "Track Hits: (" << *tracker->NTracks() << ")" << endl;
+               tracker->DumpTrackHits(*OutFile);
+    }
+
+       if (DebugLevel >= 4) printf("Running WriteOutput\n");
+       tracker->WriteOutput();
+
+       if (DebugLevel >= 4) printf("GPU Reconstruction finished\n");
+       
+       return(0);
+}
+
+int AliHLTTPCCAGPUTracker::ExitGPU()
+{
+       cudaFree(GPUMemory);
+       return(0);
+}
diff --git a/HLT/TPCLib/tracking-ca/AliHLTTPCCAGPUTracker.h b/HLT/TPCLib/tracking-ca/AliHLTTPCCAGPUTracker.h
new file mode 100644 (file)
index 0000000..b44819f
--- /dev/null
@@ -0,0 +1,40 @@
+// ************************************************************************
+// 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                               *
+//                                                                        *
+//*************************************************************************
+
+#include "AliHLTTPCCADef.h"
+#include "AliHLTTPCCATracker.h"
+
+class AliHLTTPCCAGPUTracker
+{
+public:
+       AliHLTTPCCAGPUTracker();
+       ~AliHLTTPCCAGPUTracker();
+
+       int InitGPU();
+       int Reconstruct(AliHLTTPCCATracker* tracker);
+       int ExitGPU();
+
+       void SetDebugLevel(int dwLevel, std::ostream *NewOutFile = NULL);
+
+private:
+       AliHLTTPCCATracker gpuTracker;
+       void* GPUMemory;
+
+       int CUDASync();
+       template <class T> T* alignPointer(T* ptr, int alignment);
+
+       int DebugLevel;
+       std::ostream *OutFile;
+       int GPUMemSize;
+#ifdef HLTCA_GPUCODE
+       bool CUDA_FAILED_MSG(cudaError_t error);
+#endif
+       // disable copy
+       AliHLTTPCCAGPUTracker( const AliHLTTPCCAGPUTracker& );
+       AliHLTTPCCAGPUTracker &operator=( const AliHLTTPCCAGPUTracker& );
+
+};
index 504ff7b..5176a17 100644 (file)
@@ -507,9 +507,9 @@ int AliHLTTPCCAGlobalMergerComponent::DoEvent( const AliHLTComponentEventData &e
   resultData.fSpecification = AliHLTTPCDefinitions::EncodeDataSpecification( 0, 35, 0, 5 );
   outputBlocks.push_back( resultData );
   size = resultData.fSize;
-  */
-
 
+  HLTInfo( "CAGlobalMerger:: output %d tracks", mergerOutput->NTracks() );
+  */
   return iResult;
 }
 
index 89d51e0..eacbaab 100644 (file)
 
 #include "AliHLTTPCCAGrid.h"
 #include "AliHLTTPCCAMath.h"
+
+#ifndef assert
 #include <assert.h>
+#endif
+
 #include <iostream>
 
 GPUd() void AliHLTTPCCAGrid::CreateEmpty()
@@ -65,8 +69,10 @@ GPUd() int AliHLTTPCCAGrid::GetBin( float Y, float Z ) const
   const int yBin = static_cast<int>( CAMath::FMulRZ( Y - fYMin, fStepYInv ) );
   const int zBin = static_cast<int>( CAMath::FMulRZ( Z - fZMin, fStepZInv ) );
   const int bin = CAMath::Mul24( zBin, fNy ) + yBin;
+#ifndef HLTCA_GPUCODE
   assert( bin >= 0 );
   assert( bin < static_cast<int>( fN ) );
+#endif
   return bin;
 }
 
index 4b15047..7a8af5d 100644 (file)
@@ -43,7 +43,9 @@ class AliHLTTPCCAGrid
     GPUd() float StepYInv() const { return fStepYInv; }
     GPUd() float StepZInv() const { return fStepZInv; }
 
+#ifndef CUDA_DEVICE_EMULATION
   private:
+#endif
 
     unsigned int fNy;        //* N bins in Y
     unsigned int fNz;        //* N bins in Z
index 2e40059..f2279c6 100644 (file)
@@ -29,7 +29,9 @@ class AliHLTTPCCAHit
     GPUhd() void SetY( float v ) { fY = v;    }
     GPUhd() void SetZ( float v ) { fZ = v;    }
 
+#ifndef CUDA_DEVICE_EMULATION
   protected:
+#endif
 
     float fY, fZ;       // Y and Z position of the TPC cluster
 
index 4084a2f..f1b8214 100644 (file)
@@ -59,7 +59,10 @@ class AliHLTTPCCAHitArea
     int  Ny() const { return fNy; }
     int  HitOffset() const { return fHitOffset; }
 
+#ifndef CUDA_DEVICE_EMULATION
   protected:
+#endif
+
     float fY;      // search coordinates
     float fZ;      // search coordinates
     float fMinZ;   // search coordinates
index 7243203..8501283 100644 (file)
 class AliHLTTPCCAHitId
 {
   public:
-    void Set( int row, int hit ) { fId = ( hit << 8 ) | row; }
-    int RowIndex() const { return fId & 0xff; }
-    int HitIndex() const { return fId >> 8; }
+    GPUhd() void Set( int row, int hit ) { fId = ( hit << 8 ) | row; }
+    GPUhd() int RowIndex() const { return fId & 0xff; }
+    GPUhd() int HitIndex() const { return fId >> 8; }
 
+#ifndef CUDA_DEVICE_EMULATION
   private:
+#endif
     int fId;
 };
 
index 66783d5..8d7f825 100644 (file)
@@ -46,6 +46,8 @@ class AliHLTTPCCAMath
     GPUd() static int Nint( float x );
     GPUd() static bool Finite( float x );
 
+       GPUd() static float Log(float x);
+
     GPUd()  static int AtomicExch( int *addr, int val );
     GPUd()  static int AtomicAdd ( int *addr, int val );
     GPUd()  static int AtomicMax ( int *addr, int val );
@@ -115,62 +117,62 @@ GPUd() inline float AliHLTTPCCAMath::Copysign( float x, float y )
 }
 
 
-GPUd() inline float AliHLTTPCCAMath::Sin( float x )
+GPUhd() inline float AliHLTTPCCAMath::Sin( float x )
 {
   return choice( sinf( x ), sin( x ), TMath::Sin( x ) );
 }
 
-GPUd() inline float AliHLTTPCCAMath::Cos( float x )
+GPUhd() inline float AliHLTTPCCAMath::Cos( float x )
 {
   return choice( cosf( x ), cos( x ), TMath::Cos( x ) );
 }
 
-GPUd() inline float AliHLTTPCCAMath::Tan( float x )
+GPUhd() inline float AliHLTTPCCAMath::Tan( float x )
 {
   return choice( tanf( x ), tan( x ), TMath::Tan( x ) );
 }
 
-GPUd() inline float AliHLTTPCCAMath::Min( float x, float y )
+GPUhd() inline float AliHLTTPCCAMath::Min( float x, float y )
 {
   return choice( fminf( x, y ),  ( x < y ? x : y ), TMath::Min( x, y ) );
 }
 
-GPUd() inline float AliHLTTPCCAMath::Max( float x, float y )
+GPUhd() inline float AliHLTTPCCAMath::Max( float x, float y )
 {
   return choice( fmaxf( x, y ),  ( x > y ? x : y ), TMath::Max( x, y ) );
 }
 
-GPUd() inline int AliHLTTPCCAMath::Min( int x, int y )
+GPUhd() inline int AliHLTTPCCAMath::Min( int x, int y )
 {
   return choice( min( x, y ),  ( x < y ? x : y ), TMath::Min( x, y ) );
 }
 
-GPUd() inline int AliHLTTPCCAMath::Max( int x, int y )
+GPUhd() inline int AliHLTTPCCAMath::Max( int x, int y )
 {
   return choice( max( x, y ),  ( x > y ? x : y ), TMath::Max( x, y ) );
 }
 
-GPUd() inline float AliHLTTPCCAMath::Sqrt( float x )
+GPUhd() inline float AliHLTTPCCAMath::Sqrt( float x )
 {
   return choice( sqrtf( x ), sqrt( x ), TMath::Sqrt( x ) );
 }
 
-GPUd() inline float AliHLTTPCCAMath::Abs( float x )
+GPUhd() inline float AliHLTTPCCAMath::Abs( float x )
 {
   return choice( fabsf( x ), fabs( x ), TMath::Abs( x ) );
 }
 
-GPUd() inline double AliHLTTPCCAMath::Abs( double x )
+GPUhd() inline double AliHLTTPCCAMath::Abs( double x )
 {
   return choice( fabs( x ), fabs( x ), TMath::Abs( x ) );
 }
 
-GPUd() inline int AliHLTTPCCAMath::Abs( int x )
+GPUhd() inline int AliHLTTPCCAMath::Abs( int x )
 {
   return choice( abs( x ), ( x >= 0 ? x : -x ), TMath::Abs( x ) );
 }
 
-GPUd() inline float AliHLTTPCCAMath::ASin( float x )
+GPUhd() inline float AliHLTTPCCAMath::ASin( float x )
 {
   return choice( asinf( x ), asin( x ), TMath::ASin( x ) );
 }
@@ -186,6 +188,11 @@ GPUd() inline float AliHLTTPCCAMath::FMulRZ( float a, float b )
   return choice( __fmul_rz( a, b ), a*b, a*b );
 }
 
+GPUhd() inline float AliHLTTPCCAMath::Log(float x)
+{
+       return choice( Log(x), Log(x), TMath::Log(x));
+}
+
 
 GPUd()  inline int AliHLTTPCCAMath::AtomicExch( int *addr, int val )
 {
index 6d3b2de..102dc66 100644 (file)
@@ -17,7 +17,7 @@
 //                                                                          *
 //***************************************************************************
 
-
+#include <stdio.h>
 #include "AliHLTTPCCASliceTrack.h"
 #include "AliHLTTPCCATracker.h"
 #include "AliHLTTPCCATrackParam.h"
@@ -404,7 +404,7 @@ float AliHLTTPCCAMerger::GetChi2( float x1, float y1, float a00, float a10, floa
 
   float mS[3] = { mSi[2], -mSi[1], mSi[0] };
 
-  return TMath::Abs( ( ( mS[0]*d[0] + mS[1]*d[1] )*d[0]
+  return AliHLTTPCCAMath::Abs( ( ( mS[0]*d[0] + mS[1]*d[1] )*d[0]
                        + ( mS[1]*d[0] + mS[2]*d[1] )*d[1] ) / s / 2 );
 
 }
@@ -835,6 +835,9 @@ void AliHLTTPCCAMerger::Merging()
   }
 
   fOutput->SetNTracks( nOutTracks );
+#ifdef HLTCA_STANDALONE
+  printf("Tracks Output: %d\n", nOutTracks);
+#endif
   fOutput->SetNTrackClusters( nOutTrackClusters );
   fOutput->SetPointers();
 
index 1fe362b..4987d8e 100644 (file)
@@ -32,7 +32,11 @@ class AliHLTTPCCANeighboursCleaner
             : fIRow( 0 ), fIRowUp( 0 ), fIRowDn( 0 ), fNRows( 0 ), fNHits( 0 ) {}
         AliHLTTPCCASharedMemory& operator=( const AliHLTTPCCASharedMemory& /*dummy*/ ) { return *this; }
 #endif
+
+#ifndef CUDA_DEVICE_EMULATION
       protected:
+#endif
+
         int fIRow; // current row index
         int fIRowUp; // current row index
         int fIRowDn; // current row index
index b16af8a..7a8df2d 100644 (file)
@@ -95,7 +95,6 @@ GPUd() void AliHLTTPCCANeighboursFinder::Thread
   } else if ( iSync == 2 ) {
     if ( ( s.fIRow <= 1 ) || ( s.fIRow >= s.fNRows - 2 ) ) return;
 
-    //const float kAreaSize = 3;
     float chi2Cut = 3.*3.*4 * ( s.fUpDx * s.fUpDx + s.fDnDx * s.fDnDx );
     const float kAreaSize = 3;
     //float chi2Cut = 3.*3.*(s.fUpDx*s.fUpDx + s.fDnDx*s.fDnDx ); //SG
index d6ae85d..3dd21de 100644 (file)
@@ -33,7 +33,10 @@ class AliHLTTPCCANeighboursFinder
             : fGridUp(), fGridDn(), fNHits( 0 ), fUpNHits( 0 ), fDnNHits( 0 ), fUpDx( 0 ), fDnDx( 0 ), fUpTx( 0 ), fDnTx( 0 ), fIRow( 0 ), fIRowUp( 0 ), fIRowDn( 0 ), fNRows( 0 ) {}
         AliHLTTPCCASharedMemory& operator=( const AliHLTTPCCASharedMemory& /*dummy*/ ) { return *this; }
 #endif
+
+#ifndef CUDA_DEVICE_EMULATION
       protected:
+#endif
         AliHLTTPCCAGrid fGridUp; // grid for the next row
         AliHLTTPCCAGrid fGridDn; // grid for the previous row
         int fNHits; // n hits
@@ -47,10 +50,10 @@ class AliHLTTPCCANeighboursFinder
         int fIRowUp; // next row number
         int fIRowDn;// previous row number
         int fNRows; // number of rows
-        float2 fA[256][20]; // temp memory
-        unsigned short fB[256][20]; // temp memory
-        unsigned short fGridContentUp[7000]; // grid content for the next row
-        unsigned short fGridContentDn[7000];// grid content for the previous row
+        float2 fA[256][ALIHLTTPCCANEIGHBOURS_FINDER_MAX_NNEIGHUP]; // temp memory
+        unsigned short fB[256][ALIHLTTPCCANEIGHBOURS_FINDER_MAX_NNEIGHUP]; // temp memory
+        unsigned short fGridContentUp[ALIHLTTPCCANEIGHBOURS_FINDER_MAX_FGRIDCONTENTUPDOWN]; // grid content for the next row
+        unsigned short fGridContentDn[ALIHLTTPCCANEIGHBOURS_FINDER_MAX_FGRIDCONTENTUPDOWN];// grid content for the previous row
     };
 
     GPUd() static int NThreadSyncPoints() { return 2; }
index 6d14c5f..d7e77a1 100644 (file)
@@ -78,9 +78,8 @@ GPUd() AliHLTTPCCAParam::AliHLTTPCCAParam()
 
   Update();
 }
-#endif
 
-GPUd() void AliHLTTPCCAParam::Initialize( int iSlice,
+void AliHLTTPCCAParam::Initialize( int iSlice,
     int nRows, float rowX[],
     float alpha, float dAlpha,
     float rMin, float rMax,
@@ -110,7 +109,7 @@ GPUd() void AliHLTTPCCAParam::Initialize( int iSlice,
   Update();
 }
 
-GPUd() void AliHLTTPCCAParam::Update()
+void AliHLTTPCCAParam::Update()
 {
   // update of calculated values
 
@@ -126,12 +125,15 @@ GPUd() void AliHLTTPCCAParam::Update()
 
   fCosAlpha = CAMath::Cos( fAlpha );
   fSinAlpha = CAMath::Sin( fAlpha );
-  fAngleMin = fAlpha - fDAlpha / 2.;
-  fAngleMax = fAlpha + fDAlpha / 2.;
+  fAngleMin = fAlpha - fDAlpha / 2.f;
+  fAngleMax = fAlpha + fDAlpha / 2.f;
   fErrX = fPadPitch / CAMath::Sqrt( 12. );
   fTrackChi2Cut = fTrackChiCut * fTrackChiCut;
 }
 
+#endif
+
+
 GPUd() void AliHLTTPCCAParam::Slice2Global( float x, float y,  float z,
     float *X, float *Y,  float *Z ) const
 {
index 832f0a8..3d8e381 100644 (file)
@@ -31,17 +31,18 @@ class AliHLTTPCCAParam
 
 #if !defined(HLTCA_GPUCODE)
     GPUd() AliHLTTPCCAParam();
-#endif
 
     ~AliHLTTPCCAParam() {;}
 
-    GPUd() void Initialize( int iSlice, int nRows, float rowX[],
+    void Initialize( int iSlice, int nRows, float rowX[],
                             float alpha, float dAlpha,
                             float rMin, float rMax, float zMin, float zMax,
                             float padPitch, float zSigma, float bz );
-    GPUd() void Update();
+    void Update();
+
+#endif
 
-    GPUd() void Slice2Global( float x, float y,  float z,
+       GPUd() void Slice2Global( float x, float y,  float z,
                               float *X, float *Y,  float *Z ) const;
 
     GPUd() void Global2Slice( float x, float y,  float z,
@@ -117,7 +118,9 @@ class AliHLTTPCCAParam
     GPUd() float GetBz( float x, float y, float z ) const;
     GPUd()  float GetBz( const AliHLTTPCCATrackParam &t ) const;
 
+#ifndef CUDA_DEVICE_EMULATION
   protected:
+#endif
 
     int fISlice; // slice number
     int fNRows; // number of rows
index 0f7a058..3d71214 100644 (file)
@@ -25,7 +25,7 @@ class AliHLTTPCCATracker;
 template<class TProcess>
 GPUg() void AliHLTTPCCAProcess()
 {
-  AliHLTTPCCATracker &tracker = *( ( AliHLTTPCCATracker* ) cTracker );
+  AliHLTTPCCATracker &tracker = *( ( AliHLTTPCCATracker* ) gAliHLTTPCCATracker );
 
   GPUshared() typename TProcess::AliHLTTPCCASharedMemory smem;
 
@@ -72,7 +72,7 @@ GPUg() void AliHLTTPCCAProcess( int nBlocks, int nThreads, AliHLTTPCCATracker &t
 template<typename TProcess>
 GPUg() void AliHLTTPCCAProcess1()
 {
-  AliHLTTPCCATracker &tracker = *( ( AliHLTTPCCATracker* ) cTracker );
+  AliHLTTPCCATracker &tracker = *( ( AliHLTTPCCATracker* ) gAliHLTTPCCATracker );
   AliHLTTPCCATrackParam tParam;
 
   GPUshared() typename TProcess::AliHLTTPCCASharedMemory sMem;
index d04a579..a809048 100644 (file)
@@ -44,7 +44,10 @@ class AliHLTTPCCARow
     GPUhd() int   HitNumberOffset() const { return fHitNumberOffset; }
     GPUhd() unsigned int FirstHitInBinOffset() const { return fFirstHitInBinOffset; }
 
+#ifndef CUDA_DEVICE_EMULATION
   private:
+#endif
+
     int fNHits;            // number of hits
     float fX;              // X coordinate of the row
     float fMaxY;           // maximal Y coordinate of the row
index 972b66a..b6012d0 100644 (file)
@@ -114,23 +114,15 @@ void AliHLTTPCCASliceData::InitializeRows( const AliHLTTPCCAParam &p )
   }
 }
 
-void AliHLTTPCCASliceData::InitFromClusterData( const AliHLTTPCCAClusterData &data )
+GPUh() char* AliHLTTPCCASliceData::SetGPUSliceDataMemory(char* pGPUMemory, const AliHLTTPCCAClusterData *data)
 {
-  // initialisation from cluster data
-
-  ////////////////////////////////////
-  // 1. prepare arrays
-  ////////////////////////////////////
-
-  fNumberOfHits = data.NumberOfClusters();
+       fMemory = (char*) pGPUMemory;
+       return(pGPUMemory + SetPointers(data, false));
+}
 
-  /* TODO Vectorization
-  for ( int rowIndex = data.FirstRow(); rowIndex <= data.LastRow(); ++rowIndex ) {
-    int NumberOfClusters( int rowIndex ) const;
-  }
-  const int memorySize = fNumberOfHits * sizeof( short_v::Type )
-  */
-  const int numberOfRows = data.LastRow() - data.FirstRow();
+size_t AliHLTTPCCASliceData::SetPointers(const AliHLTTPCCAClusterData *data, bool allocate)
+{
+  const int numberOfRows = data->LastRow() - data->FirstRow();
   enum { kVectorAlignment = sizeof( int ) };
   const int numberOfHitsPlusAlignment = NextMultipleOf < kVectorAlignment / sizeof( int ) > ( fNumberOfHits );
   const int memorySize =
@@ -142,9 +134,12 @@ void AliHLTTPCCASliceData::InitFromClusterData( const AliHLTTPCCAClusterData &da
     numberOfHitsPlusAlignment * 2 * sizeof( int );
 
   if ( fMemorySize < memorySize ) {
-    fMemorySize = memorySize;
-    delete[] fMemory;
-    fMemory = new char[fMemorySize + 4];// kVectorAlignment];
+       fMemorySize = memorySize;
+       if (allocate)
+       {
+         delete[] fMemory;
+         fMemory = new char[fMemorySize + 4];// kVectorAlignment];
+       }
   }
 
   char *mem = fMemory;
@@ -155,6 +150,27 @@ void AliHLTTPCCASliceData::InitFromClusterData( const AliHLTTPCCAClusterData &da
   AssignMemory( fFirstHitInBin,  mem, 23 * numberOfRows + 4 * fNumberOfHits + 3 );
   AssignMemory( fHitWeights,   mem, numberOfHitsPlusAlignment );
   AssignMemory( fClusterDataIndex, mem, numberOfHitsPlusAlignment );
+  return(mem - fMemory);
+}
+
+void AliHLTTPCCASliceData::InitFromClusterData( const AliHLTTPCCAClusterData &data )
+{
+  // initialisation from cluster data
+
+  ////////////////////////////////////
+  // 1. prepare arrays
+  ////////////////////////////////////
+
+  const int numberOfRows = data.LastRow() - data.FirstRow();
+  fNumberOfHits = data.NumberOfClusters();
+
+  /* TODO Vectorization
+  for ( int rowIndex = data.FirstRow(); rowIndex <= data.LastRow(); ++rowIndex ) {
+    int NumberOfClusters( int rowIndex ) const;
+  }
+  const int memorySize = fNumberOfHits * sizeof( short_v::Type )
+  */
+  SetPointers(&data, true);
 
   ////////////////////////////////////
   // 2. fill HitData and FirstHitInBin
index e680bb2..ebdd8c6 100644 (file)
 #ifndef ALIHLTTPCCASLICEDATA_H
 #define ALIHLTTPCCASLICEDATA_H
 
+#include "AliHLTTPCCADef.h"
 #include "AliHLTTPCCARow.h"
 #include "AliHLTTPCCAMath.h"
+#include "AliHLTArray.h"
 
 typedef int int_v;
 typedef unsigned int uint_v;
@@ -39,6 +41,7 @@ class AliHLTTPCCAParam;
  */
 class AliHLTTPCCASliceData
 {
+       friend class AliHLTTPCCAGPUTracker;
   public:
     AliHLTTPCCASliceData()
         : fNumberOfHits( 0 ), fMemorySize( 0 ), fMemory( 0 ), fLinkUpData( 0 ),
@@ -51,6 +54,9 @@ class AliHLTTPCCASliceData
      * (Re)Create the data that is tuned for optimal performance of the algorithm from the cluster
      * data.
      */
+
+       char* SetGPUSliceDataMemory(char* pGPUMemory, const AliHLTTPCCAClusterData *data);
+       size_t SetPointers(const AliHLTTPCCAClusterData *data, bool allocate = false);
     void InitFromClusterData( const AliHLTTPCCAClusterData &data );
 
     /**
@@ -61,7 +67,7 @@ class AliHLTTPCCASliceData
     /**
      * Return the number of hits in this slice.
      */
-    int NumberOfHits() const { return fNumberOfHits; }
+    GPUhd() int NumberOfHits() const { return fNumberOfHits; }
 
     /**
      * Access to the hit links.
@@ -120,7 +126,9 @@ class AliHLTTPCCASliceData
      */
     const AliHLTTPCCARow &Row( int rowIndex ) const;
 
+#ifndef CUDA_DEVICE_EMULATION
   private:
+#endif
 
     AliHLTTPCCASliceData( const AliHLTTPCCASliceData & )
         : fNumberOfHits( 0 ), fMemorySize( 0 ), fMemory( 0 ), fLinkUpData( 0 ),
@@ -158,57 +166,57 @@ class AliHLTTPCCASliceData
 
 };
 
-inline short_v AliHLTTPCCASliceData::HitLinkUpData  ( const AliHLTTPCCARow &row, const short_v &hitIndex ) const
+GPUd() inline short_v AliHLTTPCCASliceData::HitLinkUpData  ( const AliHLTTPCCARow &row, const short_v &hitIndex ) const
 {
   return fLinkUpData[row.fHitNumberOffset + hitIndex];
 }
 
-inline short_v AliHLTTPCCASliceData::HitLinkDownData( const AliHLTTPCCARow &row, const short_v &hitIndex ) const
+GPUd() inline short_v AliHLTTPCCASliceData::HitLinkDownData( const AliHLTTPCCARow &row, const short_v &hitIndex ) const
 {
   return fLinkDownData[row.fHitNumberOffset + hitIndex];
 }
 
-inline void AliHLTTPCCASliceData::SetHitLinkUpData  ( const AliHLTTPCCARow &row, const short_v &hitIndex, const short_v &value )
+GPUd() inline void AliHLTTPCCASliceData::SetHitLinkUpData  ( const AliHLTTPCCARow &row, const short_v &hitIndex, const short_v &value )
 {
   fLinkUpData[row.fHitNumberOffset + hitIndex] = value;
 }
 
-inline void AliHLTTPCCASliceData::SetHitLinkDownData( const AliHLTTPCCARow &row, const short_v &hitIndex, const short_v &value )
+GPUd() inline void AliHLTTPCCASliceData::SetHitLinkDownData( const AliHLTTPCCARow &row, const short_v &hitIndex, const short_v &value )
 {
   fLinkDownData[row.fHitNumberOffset + hitIndex] = value;
 }
 
-inline short_v AliHLTTPCCASliceData::HitDataY( const AliHLTTPCCARow &row, const uint_v &hitIndex ) const
+GPUd() inline short_v AliHLTTPCCASliceData::HitDataY( const AliHLTTPCCARow &row, const uint_v &hitIndex ) const
 {
   return fHitDataY[row.fHitNumberOffset + hitIndex];
 }
 
-inline short_v AliHLTTPCCASliceData::HitDataZ( const AliHLTTPCCARow &row, const uint_v &hitIndex ) const
+GPUd() inline short_v AliHLTTPCCASliceData::HitDataZ( const AliHLTTPCCARow &row, const uint_v &hitIndex ) const
 {
   return fHitDataZ[row.fHitNumberOffset + hitIndex];
 }
 
-inline ushort_v AliHLTTPCCASliceData::FirstHitInBin( const AliHLTTPCCARow &row, ushort_v binIndexes ) const
+GPUd() inline ushort_v AliHLTTPCCASliceData::FirstHitInBin( const AliHLTTPCCARow &row, ushort_v binIndexes ) const
 {
   return fFirstHitInBin[row.fFirstHitInBinOffset + binIndexes];
 }
 
-inline int_v AliHLTTPCCASliceData::ClusterDataIndex( const AliHLTTPCCARow &row, uint_v hitIndex ) const
+GPUhd() inline int_v AliHLTTPCCASliceData::ClusterDataIndex( const AliHLTTPCCARow &row, uint_v hitIndex ) const
 {
   return fClusterDataIndex[row.fHitNumberOffset + hitIndex];
 }
 
-inline const AliHLTTPCCARow &AliHLTTPCCASliceData::Row( int rowIndex ) const
+GPUhd() inline const AliHLTTPCCARow &AliHLTTPCCASliceData::Row( int rowIndex ) const
 {
   return fRows[rowIndex];
 }
 
-inline void AliHLTTPCCASliceData::MaximizeHitWeight( const AliHLTTPCCARow &row, uint_v hitIndex, int_v weight )
+GPUd() inline void AliHLTTPCCASliceData::MaximizeHitWeight( const AliHLTTPCCARow &row, uint_v hitIndex, int_v weight )
 {
   CAMath::AtomicMax( &fHitWeights[row.fHitNumberOffset + hitIndex], weight );
 }
 
-inline int_v AliHLTTPCCASliceData::HitWeight( const AliHLTTPCCARow &row, uint_v hitIndex ) const
+GPUd() inline int_v AliHLTTPCCASliceData::HitWeight( const AliHLTTPCCARow &row, uint_v hitIndex ) const
 {
   return fHitWeights[row.fHitNumberOffset + hitIndex];
 }
index aad494a..842661d 100644 (file)
@@ -28,6 +28,7 @@ GPUhd() int AliHLTTPCCASliceOutput::EstimateSize( int nOfTracks, int nOfTrackClu
   return sizeof( AliHLTTPCCASliceOutput ) + sizeof( AliHLTTPCCASliceTrack )*nOfTracks + kClusterDataSize*nOfTrackClusters;
 }
 
+#ifndef HLTCA_GPUCODE
 template<typename T> inline void AssignNoAlignment( T *&dst, char *&mem, int count )
 {
   // assign memory to the pointer dst
@@ -49,4 +50,4 @@ GPUhd() void AliHLTTPCCASliceOutput::SetPointers()
   AssignNoAlignment( fClusterPackedAmp,  mem, fNTrackClusters );
 
 }
-
+#endif
index aaa03ee..d1e83fc 100644 (file)
@@ -41,7 +41,9 @@ class AliHLTTPCCASliceOutput
     GPUhd() float    ClusterUnpackedX  ( int i )  const { return fClusterUnpackedX[i]; }
 
     GPUhd() static int EstimateSize( int nOfTracks, int nOfTrackClusters );
+#ifndef HLTCA_GPUCODE
     GPUhd() void SetPointers();
+#endif
 
     GPUhd() void SetNTracks       ( int v )  { fNTracks = v;        }
     GPUhd() void SetNTrackClusters( int v )  { fNTrackClusters = v; }
index 34f8134..d76a026 100644 (file)
 #include "AliHLTTPCCAClusterData.h"
 #include "TStopwatch.h"
 
+//If not building GPU Code then build dummy functions to link against
+#ifndef BUILD_GPU
+AliHLTTPCCAGPUTracker::AliHLTTPCCAGPUTracker() : gpuTracker(), DebugLevel(0) {}
+AliHLTTPCCAGPUTracker::~AliHLTTPCCAGPUTracker() {}
+int AliHLTTPCCAGPUTracker::InitGPU() {return(0);}
+//template <class T> inline T* AliHLTTPCCAGPUTracker::alignPointer(T* ptr, int alignment) {return(NULL);}
+//bool AliHLTTPCCAGPUTracker::CUDA_FAILED_MSG(cudaError_t error) {return(true);}
+//int AliHLTTPCCAGPUTracker::CUDASync() {return(1);}
+//void AliHLTTPCCAGPUTracker::SetDebugLevel(int dwLevel, std::ostream *NewOutFile) {};
+int AliHLTTPCCAGPUTracker::Reconstruct(AliHLTTPCCATracker* tracker) {return(1);}
+int AliHLTTPCCAGPUTracker::ExitGPU() {return(0);}
+#endif
 
 AliHLTTPCCAStandaloneFramework &AliHLTTPCCAStandaloneFramework::Instance()
 {
@@ -35,7 +47,7 @@ AliHLTTPCCAStandaloneFramework &AliHLTTPCCAStandaloneFramework::Instance()
 }
 
 AliHLTTPCCAStandaloneFramework::AliHLTTPCCAStandaloneFramework()
-    : fMerger(), fStatNEvents( 0 )
+    : fMerger(), fStatNEvents( 0 ), fUseGPUTracker(false), fGPUDebugLevel(0)
 {
   //* constructor
 
@@ -78,12 +90,43 @@ void AliHLTTPCCAStandaloneFramework::FinishDataReading()
 {
   // finish reading of the event
 
+  /*static int event_number = 0;
+  char filename[256];
+
+  sprintf(filename, "events/event.%d.dump", event_number);
+  printf("Dumping event into file %s\n", filename);
+  std::ofstream outfile(filename, std::ofstream::binary);
+  if (outfile.fail())
+  {
+    printf("Error opening event dump file\n");
+    exit(1);
+  }
+  WriteEvent(outfile);
+  if (outfile.fail())
+  {
+    printf("Error writing event dump file\n");
+    exit(1);
+  }
+  outfile.close();
+
+  sprintf(filename, "events/settings.%d.dump", event_number);
+  outfile.open(filename);
+  WriteSettings(outfile);
+  outfile.close();
+
+  event_number++;
+  
+  std::ifstream infile(filename, std::ifstream::binary);
+  ReadEvent(infile);
+  infile.close();*/
+
   for ( int i = 0; i < fgkNSlices; i++ ) {
     fClusterData[i].FinishReading();
   }
 }
 
 
+//int
 void AliHLTTPCCAStandaloneFramework::ProcessEvent()
 {
   // perform the event reconstruction
@@ -93,9 +136,26 @@ void AliHLTTPCCAStandaloneFramework::ProcessEvent()
   TStopwatch timer0;
   TStopwatch timer1;
 
-  for ( int iSlice = 0; iSlice < fgkNSlices; iSlice++ ) {
-    fSliceTrackers[iSlice].ReadEvent( &( fClusterData[iSlice] ) );
-    fSliceTrackers[iSlice].Reconstruct();
+  if (!fUseGPUTracker || fGPUDebugLevel >= 3)
+  {
+       for ( int iSlice = 0; iSlice < fgkNSlices; iSlice++ ) {
+         fSliceTrackers[iSlice].ReadEvent( &( fClusterData[iSlice] ) );
+      fSliceTrackers[iSlice].Reconstruct();
+       }
+       if (fGPUDebugLevel >= 2) printf("\n");
+  }
+
+  if (fUseGPUTracker)
+  {
+         for ( int iSlice = 0; iSlice < fgkNSlices; iSlice++ ) {
+           fSliceTrackers[iSlice].ReadEvent( &( fClusterData[iSlice] ) );
+               if (fGPUTracker.Reconstruct(&fSliceTrackers[iSlice]))
+               {
+                       printf("Error during GPU Reconstruction!!!\n");
+                       //return(1);
+               }
+         }
+         if (fGPUDebugLevel >= 2) printf("\n");
   }
 
   timer1.Stop();
@@ -118,6 +178,8 @@ void AliHLTTPCCAStandaloneFramework::ProcessEvent()
   fLastTime[2] = timer2.CpuTime();
 
   for ( int i = 0; i < 3; i++ ) fStatTime[i] += fLastTime[i];
+
+  //return(0);
 }
 
 
@@ -142,20 +204,19 @@ void AliHLTTPCCAStandaloneFramework::ReadSettings( std::istream &in )
   }
 }
 
-void AliHLTTPCCAStandaloneFramework::WriteEvent( std::ostream &/*out*/ ) const
+void AliHLTTPCCAStandaloneFramework::WriteEvent( std::ostream &out ) const
 {
   // write event to the file
   for ( int iSlice = 0; iSlice < fgkNSlices; iSlice++ ) {
-    //fClusterData[i].WriteEvent( out );
+    fClusterData[iSlice].WriteEvent( out );
   }
 }
 
-void AliHLTTPCCAStandaloneFramework::ReadEvent( std::istream &/*in*/ ) const
+void AliHLTTPCCAStandaloneFramework::ReadEvent( std::istream &in )
 {
   //* Read event from file
-
   for ( int iSlice = 0; iSlice < fgkNSlices; iSlice++ ) {
-    //fClusterData[i].ReadEvent( in );
+    fClusterData[iSlice].ReadEvent( in );
   }
 }
 
@@ -177,3 +238,27 @@ void AliHLTTPCCAStandaloneFramework::ReadTracks( std::istream &in )
   }
   //fMerger.Output()->Read( in );
 }
+
+int AliHLTTPCCAStandaloneFramework::InitGPU()
+{
+       if (fUseGPUTracker) return(1);
+       int retVal = fGPUTracker.InitGPU();
+       fUseGPUTracker = retVal == 0;
+       return(retVal);
+}
+
+int AliHLTTPCCAStandaloneFramework::ExitGPU()
+{
+       if (!fUseGPUTracker) return(1);
+       return(fGPUTracker.ExitGPU());
+}
+
+void AliHLTTPCCAStandaloneFramework::SetGPUDebugLevel(int Level, std::ostream *OutFile, std::ostream *GPUOutFile)
+{
+       fGPUTracker.SetDebugLevel(Level, GPUOutFile);
+       fGPUDebugLevel = Level;
+       for (int i = 0;i < fgkNSlices;i++)
+       {
+               fSliceTrackers[i].SetGPUDebugLevel(Level, OutFile);
+       }
+}
index 1eaa05c..175a41f 100644 (file)
@@ -13,7 +13,9 @@
 #include "AliHLTTPCCATracker.h"
 #include "AliHLTTPCCAMerger.h"
 #include "AliHLTTPCCAClusterData.h"
+#include "AliHLTTPCCAGPUTracker.h"
 #include <iostream>
+#include <fstream>
 
 /**
  * @class AliHLTTPCCAStandaloneFramework
@@ -71,9 +73,13 @@ class AliHLTTPCCAStandaloneFramework
     void WriteTracks( std::ostream &out ) const;
 
     void ReadSettings( std::istream &in );
-    void ReadEvent( std::istream &in ) const;
+    void ReadEvent( std::istream &in );
     void ReadTracks( std::istream &in );
 
+       int InitGPU();
+       int ExitGPU();
+       void SetGPUDebugLevel(int Level, std::ostream *OutFile = NULL, std::ostream *GPUOutFile = NULL);
+
   private:
 
     static const int fgkNSlices = 36;       //* N slices
@@ -85,9 +91,14 @@ class AliHLTTPCCAStandaloneFramework
     AliHLTTPCCAMerger fMerger;  //* global merger
     AliHLTTPCCAClusterData fClusterData[fgkNSlices];
 
+  AliHLTTPCCAGPUTracker fGPUTracker;
+
     double fLastTime[20]; //* timers
     double fStatTime[20]; //* timers
     int fStatNEvents;    //* n events proceed
+
+  bool fUseGPUTracker; // use the GPU tracker 
+  int fGPUDebugLevel;  // debug level for the GPU code
 };
 
 #endif
index 3501ae2..709278d 100644 (file)
@@ -33,11 +33,15 @@ class AliHLTTPCCAStartHitsFinder
             : fIRow( 0 ), fNRows( 0 ), fNHits( 0 ), fNOldStartHits( 0 ), fNRowStartHits( 0 ) {}
         AliHLTTPCCASharedMemory& operator=( const AliHLTTPCCASharedMemory& /*dummy*/ ) { return *this; }
 #endif
+
+#ifndef CUDA_DEVICE_EMULATION
       protected:
+#endif
+
         int fIRow; // row index
         int fNRows; // n rows
         int fNHits; // n hits in the row
-        AliHLTTPCCAHitId fRowStartHits[10240]; // temp. array for the start hits
+        AliHLTTPCCAHitId fRowStartHits[ALIHLTTPCCASTARTHITSFINDER_MAX_FROWSTARTHITS]; // temp. array for the start hits
         int fNOldStartHits; // n start hits from other jobs
         int fNRowStartHits; // n start hits for this row
     };
index 0839048..714e974 100644 (file)
@@ -38,7 +38,9 @@ class AliHLTTPCCATrack
     GPUhd() void SetFirstHitID( int v )          { fFirstHitID = v; }
     GPUhd() void SetParam( AliHLTTPCCATrackParam v ) { fParam = v; };
 
+#ifndef CUDA_DEVICE_EMULATION
   private:
+#endif
 
     bool fAlive;       // flag for mark tracks used by the track merger
     int  fFirstHitID; // index of the first track cell in the track->cell pointer array
index e00ce0f..1656f8c 100644 (file)
@@ -28,6 +28,7 @@
  */
 class AliHLTTPCCATrackLinearisation
 {
+
   public:
 
     AliHLTTPCCATrackLinearisation()
@@ -56,7 +57,9 @@ class AliHLTTPCCATrackLinearisation
     GPUd() void SetDzDs( float v )  {  fDzDs   = v; }
     GPUd() void SetQPt( float v )   {  fQPt = v; }
 
+#ifndef CUDA_DEVICE_EMULATION
   private:
+#endif
 
     float fSinPhi; // SinPhi
     float fCosPhi; // CosPhi
@@ -65,7 +68,7 @@ class AliHLTTPCCATrackLinearisation
 };
 
 
-inline AliHLTTPCCATrackLinearisation::AliHLTTPCCATrackLinearisation( const AliHLTTPCCATrackParam &t )
+GPUd() inline AliHLTTPCCATrackLinearisation::AliHLTTPCCATrackLinearisation( const AliHLTTPCCATrackParam &t )
     : fSinPhi( t.SinPhi() ), fCosPhi( 0 ), fDzDs( t.DzDs() ), fQPt( t.QPt() )
 {
   if ( fSinPhi > .999 ) fSinPhi = .999;
index 2ca1103..090a15d 100644 (file)
@@ -378,8 +378,8 @@ float AliHLTTPCCATrackParam::BetheBlochGeant( float bg2,
 
   //*** Density effect
   float d2 = 0.;
-  const float x = 0.5 * TMath::Log( bg2 );
-  const float lhwI = TMath::Log( 28.816 * 1e-9 * TMath::Sqrt( rho * mZA ) / mI );
+  const float x = 0.5 * AliHLTTPCCAMath::Log( bg2 );
+  const float lhwI = AliHLTTPCCAMath::Log( 28.816 * 1e-9 * AliHLTTPCCAMath::Sqrt( rho * mZA ) / mI );
   if ( x > x1 ) {
     d2 = lhwI + x - 0.5;
   } else if ( x > x0 ) {
@@ -387,7 +387,7 @@ float AliHLTTPCCATrackParam::BetheBlochGeant( float bg2,
     d2 = lhwI + x - 0.5 + ( 0.5 - lhwI - x0 ) * r * r * r;
   }
 
-  return mK*mZA*( 1 + bg2 ) / bg2*( 0.5*TMath::Log( 2*me*bg2*maxT / ( mI*mI ) ) - bg2 / ( 1 + bg2 ) - d2 );
+  return mK*mZA*( 1 + bg2 ) / bg2*( 0.5*AliHLTTPCCAMath::Log( 2*me*bg2*maxT / ( mI*mI ) ) - bg2 / ( 1 + bg2 ) - d2 );
 }
 
 float AliHLTTPCCATrackParam::BetheBlochSolid( float bg )
@@ -669,11 +669,11 @@ GPUd() bool AliHLTTPCCATrackParam::CheckNumericalQuality() const
 {
   //* Check that the track parameters and covariance matrix are reasonable
 
-  bool ok = finite( fX ) && finite( fSignCosPhi ) && finite( fChi2 ) && finite( fNDF );
+  bool ok = AliHLTTPCCAMath::Finite( fX ) && AliHLTTPCCAMath::Finite( fSignCosPhi ) && AliHLTTPCCAMath::Finite( fChi2 ) && AliHLTTPCCAMath::Finite( fNDF );
 
   const float *c = Cov();
-  for ( int i = 0; i < 15; i++ ) ok = ok && finite( c[i] );
-  for ( int i = 0; i < 5; i++ ) ok = ok && finite( Par()[i] );
+  for ( int i = 0; i < 15; i++ ) ok = ok && AliHLTTPCCAMath::Finite( c[i] );
+  for ( int i = 0; i < 5; i++ ) ok = ok && AliHLTTPCCAMath::Finite( Par()[i] );
 
   if ( c[0] <= 0 || c[2] <= 0 || c[5] <= 0 || c[9] <= 0 || c[14] <= 0 ) ok = 0;
   if ( c[0] > 5. || c[2] > 5. || c[5] > 2. || c[9] > 2 || ( CAMath::Abs( QPt() ) > 1.e-4 && c[14] > 2. ) ) ok = 0;
index d103c12..74f5bbe 100644 (file)
@@ -136,7 +136,9 @@ class AliHLTTPCCATrackParam
 
     GPUd() void Print() const;
 
+#ifndef CUDA_DEVICE_EMULATION
   private:
+#endif
 
     float fX;      // x position
     float fSignCosPhi; // sign of cosPhi
index e452f45..6acc920 100644 (file)
 
 ClassImp( AliHLTTPCCATracker )
 
-#if !defined(HLTCA_GPUCODE)
-
-AliHLTTPCCATracker::AliHLTTPCCATracker()
-    :
-    fParam(),
-    fClusterData( 0 ),
-    fData(),
-    fCommonMemory( 0 ),
-    fCommonMemorySize( 0 ),
-    fHitMemory( 0 ),
-    fHitMemorySize( 0 ),
-    fTrackMemory( 0 ),
-    fTrackMemorySize( 0 ),
-    fNTracklets( 0 ),
-    fTrackletStartHits( 0 ),
-    fTracklets( 0 ),
-    fNTracks( 0 ),
-    fTracks( 0 ),
-    fNTrackHits( 0 ),
-    fTrackHits( 0 ),
-    fOutput( 0 ),
-    fNOutTracks( 0 ),
-    fOutTracks( 0 ),
-    fNOutTrackHits( 0 ),
-    fOutTrackHits( 0 )
-{
-  // constructor
-}
-
 GPUd() AliHLTTPCCATracker::~AliHLTTPCCATracker()
 {
   // destructor
-  delete[] fCommonMemory;
-  delete[] fHitMemory;
-  delete[] fTrackMemory;
+       if (!fIsGPUTracker)
+       {
+               delete[] fCommonMemory;
+               delete[] fHitMemory;
+               delete[] fTrackMemory;
+       }
 }
-#endif
 
+#if !defined(HLTCA_GPUCODE)
 
 
 // ----------------------------------------------------------------------------------
-GPUd() void AliHLTTPCCATracker::Initialize( const AliHLTTPCCAParam &param )
+void AliHLTTPCCATracker::Initialize( const AliHLTTPCCAParam &param )
 {
   // initialisation
   fParam = param;
@@ -110,29 +84,151 @@ GPUd() void AliHLTTPCCATracker::Initialize( const AliHLTTPCCAParam &param )
   StartEvent();
 }
 
-GPUd() void AliHLTTPCCATracker::StartEvent()
+void AliHLTTPCCATracker::StartEvent()
 {
   // start new event and fresh the memory
 
   SetupCommonMemory();
-  *fNTrackHits = 0;
 }
 
+void AliHLTTPCCATracker::SetGPUTracker()
+{
+       fIsGPUTracker = true;
+}
+
+char* AliHLTTPCCATracker::SetGPUTrackerCommonMemory(char* pGPUMemory)
+{
+       fCommonMemory = (char*) pGPUMemory;
+       SetPointersCommon();
+       return(pGPUMemory + fCommonMemorySize);
+}
+
+
+char* AliHLTTPCCATracker::SetGPUTrackerHitsMemory(char* pGPUMemory, int MaxNHits )
+{
+       fHitMemory = (char*) pGPUMemory;
+       SetPointersHits(MaxNHits);
+       return(pGPUMemory + fHitMemorySize);
+}
+
+
+char* AliHLTTPCCATracker::SetGPUTrackerTracksMemory(char* pGPUMemory, int MaxNTracks, int MaxNHits )
+{
+       fTrackMemory = (char*) pGPUMemory;
+       SetPointersTracks(MaxNTracks, MaxNHits);
+       return(pGPUMemory + fTrackMemorySize);
+}
+
+void AliHLTTPCCATracker::DumpLinks(std::ostream &out)
+{
+       for (int i = 0;i < Param().NRows();i++)
+       {
+               out << "Row: " << i << endl;
+               for (int j = 0;j < Row(i).NHits();j++)
+               {
+                       out << HitLinkUpData(Row(i), j) << ", ";
+               }
+               out << endl;
+       }
+}
+
+void AliHLTTPCCATracker::DumpStartHits(std::ostream &out)
+{
+       for (int j = 0;j < Param().NRows();j++)
+       {
+               for (int i = 0;i < *NTracklets();i++)
+               {
+                       if (TrackletStartHit(i).RowIndex() == j)
+                               out << TrackletStartHit(i).RowIndex() << "-" << TrackletStartHit(i).HitIndex() << endl;
+               }
+       }
+       out << endl;
+}
+
+void AliHLTTPCCATracker::DumpTrackHits(std::ostream &out)
+{
+       for (int k = 0;k < Param().NRows();k++)
+       {
+               for (int j = 0;j < *NTracks();j++)
+               {
+                       if (Tracks()[j].NHits() == 0 || !Tracks()[j].Alive()) continue;
+                       if (TrackHits()[Tracks()[j].FirstHitID()].RowIndex() == k)
+                       {
+                               for (int i = 0;i < Tracks()[j].NHits();i++)
+                               {
+                                       out << TrackHits()[Tracks()[j].FirstHitID() + i].RowIndex() << "-" << TrackHits()[Tracks()[j].FirstHitID() + i].HitIndex() << ", ";
+                               }
+                               out << "(Track: " << j << ")" << endl;
+                       }
+               }
+       }
+}
+
+int trackletSortComparison(const void* a, const void* b)
+{
+       const AliHLTTPCCATracklet* aa = (AliHLTTPCCATracklet*) a;
+       const AliHLTTPCCATracklet* bb = (AliHLTTPCCATracklet*) b;
+       if (aa->NHits() == 0) return(-1);
+       if (bb->NHits() == 0) return(1);
+       if (aa->FirstRow() != bb->FirstRow())
+       {
+               return(aa->FirstRow() - bb->FirstRow());
+       }
+       for (int i = aa->FirstRow();i <= aa->LastRow();i++)
+       {
+               if (i >= bb->LastRow()) return(-1);
+               if (aa->RowHit(i) != bb->RowHit(i))
+               {
+                       return(aa->RowHit(i) - bb->RowHit(i));
+               }
+       }
+       return(0);
+}
+
+void AliHLTTPCCATracker::DumpTrackletHits(std::ostream &out)
+{
+       qsort(Tracklets(), *NTracklets(), sizeof(AliHLTTPCCATracklet), trackletSortComparison);
+       for (int k = 0;k < Param().NRows();k++)
+       {
+               for (int j = 0;j < *NTracklets();j++)
+               {
+                       if (Tracklets()[j].NHits() == 0) continue;
+                       if (Tracklets()[j].LastRow() > Tracklets()[j].FirstRow() && (Tracklets()[j].FirstRow() >= Param().NRows() || Tracklets()[j].LastRow() >= Param().NRows()))
+                       {
+                               printf("\nError: First %d Last %d Num %d", Tracklets()[j].FirstRow(), Tracklets()[j].LastRow(), Tracklets()[j].NHits());
+                       }
+                       else if (Tracklets()[j].NHits() && Tracklets()[j].FirstRow() == k && Tracklets()[j].LastRow() > Tracklets()[j].FirstRow())
+                       {
+                               for (int i = Tracklets()[j].FirstRow();i <= Tracklets()[j].LastRow();i++)
+                               {
+                                       if (Tracklets()[j].RowHit(i) != -1)
+                                               out << i << "-" << Tracklets()[j].RowHit(i) << ", ";
+                               }
+                               out << endl;
+                       }
+               }
+       }
+}
+
+
 void  AliHLTTPCCATracker::SetupCommonMemory()
 {
   // set up common memory
 
-  if ( !fCommonMemory ) {
-    SetPointersCommon(); // just to calculate the size
-    // the 1600 extra bytes are not used unless fCommonMemorySize increases with a later event
-    fCommonMemory = reinterpret_cast<char*> ( new uint4 [ fCommonMemorySize/sizeof( uint4 ) + 100] );
-    SetPointersCommon();// set pointers
-  }
+  if (!fIsGPUTracker)
+  {
+    if ( !fCommonMemory ) {
+      SetPointersCommon(); // just to calculate the size
+      // the 1600 extra bytes are not used unless fCommonMemorySize increases with a later event
+      fCommonMemory = reinterpret_cast<char*> ( new uint4 [ fCommonMemorySize/sizeof( uint4 ) + 100] );
+      SetPointersCommon();// set pointers
+    }
 
-  delete[] fHitMemory;
-  fHitMemory = 0;
-  delete[] fTrackMemory;
-  fTrackMemory = 0;
+    delete[] fHitMemory;
+    delete[] fTrackMemory;
+    fHitMemory = 0;
+    fTrackMemory = 0;
+  }
 
   fData.Clear();
   *fNTracklets = 0;
@@ -142,6 +238,29 @@ void  AliHLTTPCCATracker::SetupCommonMemory()
   *fNOutTrackHits = 0;
 }
 
+void AliHLTTPCCATracker::ReadEvent( AliHLTTPCCAClusterData *clusterData )
+{
+  // read event
+
+  fClusterData = clusterData;
+
+  StartEvent();
+
+  //* Convert input hits, create grids, etc.
+  fData.InitFromClusterData( *clusterData );
+
+  {
+    SetPointersHits( fData.NumberOfHits() ); // to calculate the size
+    fHitMemory = reinterpret_cast<char*> ( new uint4 [ fHitMemorySize/sizeof( uint4 ) + 100] );
+    SetPointersHits( fData.NumberOfHits() ); // set pointers for hits
+    *fNTracklets = 0;
+    *fNTracks = 0 ;
+    *fNOutTracks = 0;
+    *fNOutTrackHits = 0;
+  }
+}
+
+
 GPUhd() void  AliHLTTPCCATracker::SetPointersCommon()
 {
   // set all pointers to the event memory
@@ -210,27 +329,45 @@ GPUhd() void  AliHLTTPCCATracker::SetPointersTracks( int MaxNTracks, int MaxNHit
   fTrackMemorySize = mem - fTrackMemory;
 }
 
+GPUh() int AliHLTTPCCATracker::CheckEmptySlice()
+{
+  if ( NHitsTotal() < 1 ) {
+    {
+      SetPointersTracks( 1, 1 ); // to calculate the size
+      fTrackMemory = reinterpret_cast<char*> ( new uint4 [ fTrackMemorySize/sizeof( uint4 ) + 100] );
+      SetPointersTracks( 1, 1 ); // set pointers for tracks
+      fOutput->SetNTracks( 0 );
+      fOutput->SetNTrackClusters( 0 );
+    }
+
+    return 1;
+  }
+  return 0;
+}
 
-void AliHLTTPCCATracker::ReadEvent( AliHLTTPCCAClusterData *clusterData )
+void AliHLTTPCCATracker::RunNeighboursFinder()
 {
-  // read event
+       AliHLTTPCCAProcess<AliHLTTPCCANeighboursFinder>( Param().NRows(), 1, *this );
+}
 
-  fClusterData = clusterData;
+void AliHLTTPCCATracker::RunNeighboursCleaner()
+{
+       AliHLTTPCCAProcess<AliHLTTPCCANeighboursCleaner>( Param().NRows() - 2, 1, *this );
+}
 
-  StartEvent();
+void AliHLTTPCCATracker::RunStartHitsFinder()
+{
+       AliHLTTPCCAProcess<AliHLTTPCCAStartHitsFinder>( Param().NRows() - 4, 1, *this );
+}
 
-  //* Convert input hits, create grids, etc.
-  fData.InitFromClusterData( *clusterData );
+void AliHLTTPCCATracker::RunTrackletConstructor()
+{
+  AliHLTTPCCAProcess1<AliHLTTPCCATrackletConstructor>( 1, TRACKLET_CONSTRUCTOR_NMEMTHREDS + *fNTracklets, *this );
+}
 
-  {
-    SetPointersHits( fData.NumberOfHits() ); // to calculate the size
-    fHitMemory = reinterpret_cast<char*> ( new uint4 [ fHitMemorySize/sizeof( uint4 ) + 100] );
-    SetPointersHits( fData.NumberOfHits() ); // set pointers for hits
-    *fNTracklets = 0;
-    *fNTracks = 0 ;
-    *fNOutTracks = 0;
-    *fNOutTrackHits = 0;
-  }
+void AliHLTTPCCATracker::RunTrackletSelector()
+{
+  AliHLTTPCCAProcess<AliHLTTPCCATrackletSelector>( 1, *fNTracklets, *this );
 }
 
 GPUh() void AliHLTTPCCATracker::Reconstruct()
@@ -250,18 +387,8 @@ GPUh() void AliHLTTPCCATracker::Reconstruct()
   //if( fParam.ISlice()<1 ) return; //SG!!!
 
   TStopwatch timer0;
+  if (CheckEmptySlice()) return;
 
-  if ( NHitsTotal() < 1 ) {
-    {
-      SetPointersTracks( 1, 1 ); // to calculate the size
-      fTrackMemory = reinterpret_cast<char*> ( new uint4 [ fTrackMemorySize/sizeof( uint4 ) + 100] );
-      SetPointersTracks( 1, 1 ); // set pointers for tracks
-      fOutput->SetNTracks( 0 );
-      fOutput->SetNTrackClusters( 0 );
-    }
-
-    return;
-  }
 #ifdef DRAW1
   //if( fParam.ISlice()==15){
   AliHLTTPCCADisplay::Instance().ClearView();
@@ -280,8 +407,18 @@ GPUh() void AliHLTTPCCATracker::Reconstruct()
 
 #if !defined(HLTCA_GPUCODE)
 
-  AliHLTTPCCAProcess<AliHLTTPCCANeighboursFinder>( Param().NRows(), 1, *this );
+  if (fGPUDebugLevel >= 3)
+  {
+         *fGPUDebugOut << endl << endl << "Slice: " << Param().ISlice() << endl;
+  }
 
+  RunNeighboursFinder();
+
+  if (fGPUDebugLevel >= 3)
+  {
+         *fGPUDebugOut << "Neighbours Finder:" << endl;
+         DumpLinks(*fGPUDebugOut);
+  }
 #ifdef HLTCA_INTERNAL_PERFORMANCE
   //if( Param().ISlice()<=2 )
   //AliHLTTPCCAPerformance::Instance().LinkPerformance( Param().ISlice() );
@@ -295,75 +432,54 @@ GPUh() void AliHLTTPCCATracker::Reconstruct()
   }
 #endif
 
+  RunNeighboursCleaner();
 
-  AliHLTTPCCAProcess<AliHLTTPCCANeighboursCleaner>( Param().NRows() - 2, 1, *this );
-  AliHLTTPCCAProcess<AliHLTTPCCAStartHitsFinder>( Param().NRows() - 4, 1, *this );
-
-  int nStartHits = *fNTracklets;
-
-  int nThreads = 128;
-  int nBlocks = NHitsTotal() / nThreads + 1;
-  if ( nBlocks < 12 ) {
-    nBlocks = 12;
-    nThreads = NHitsTotal() / 12 + 1;
-    if ( nThreads % 32 ) nThreads = ( nThreads / 32 + 1 ) * 32;
+  if (fGPUDebugLevel >= 3)
+  {
+         *fGPUDebugOut << "Neighbours Cleaner:" << endl;
+         DumpLinks(*fGPUDebugOut);
   }
 
-  nThreads = NHitsTotal();
-  nBlocks = 1;
-
-  fData.ClearHitWeights();
-  //AliHLTTPCCAProcess<AliHLTTPCCAUsedHitsInitialiser>( nBlocks, nThreads, *this );
-
+  RunStartHitsFinder();
 
+  if (fGPUDebugLevel >= 3)
   {
-    SetPointersTracks( nStartHits*2, NHitsTotal() ); // to calculate the size
-    fTrackMemory = reinterpret_cast<char*> ( new uint4 [ fTrackMemorySize/sizeof( uint4 ) + 100] );
-    SetPointersTracks( nStartHits*2, NHitsTotal() ); // set pointers for hits
-  }
-
-  int nMemThreads = AliHLTTPCCATrackletConstructor::NMemThreads();
-  nThreads = 256;//96;
-  nBlocks = nStartHits / nThreads + 1;
-  if ( nBlocks < 30 ) {
-    nBlocks = 30;
-    nThreads = ( nStartHits ) / 30 + 1;
-    if ( nThreads % 32 ) nThreads = ( nThreads / 32 + 1 ) * 32;
+         *fGPUDebugOut << "Start Hits: (" << *fNTracklets << ")" << endl;
+         DumpStartHits(*fGPUDebugOut);
   }
+  
+  if (fGPUDebugLevel >= 2) printf("%3d ", *fNTracklets);
 
-  nThreads = nStartHits;
-  nBlocks = 1;
+  fData.ClearHitWeights();
 
-  AliHLTTPCCAProcess1<AliHLTTPCCATrackletConstructor>( nBlocks, nMemThreads + nThreads, *this );
+  SetPointersTracks( *fNTracklets * 2, NHitsTotal() ); // to calculate the size
+  fTrackMemory = reinterpret_cast<char*> ( new uint4 [ fTrackMemorySize/sizeof( uint4 ) + 100] );
+  SetPointersTracks( *fNTracklets * 2, NHitsTotal() ); // set pointers for hits
 
-  //std::cout<<"Slice "<<Param().ISlice()<<": NHits="<<NHitsTotal()<<", NTracklets="<<*NTracklets()<<std::endl;
+  RunTrackletConstructor();
 
+  if (fGPUDebugLevel >= 3)
   {
-    nThreads = 128;
-    nBlocks = nStartHits / nThreads + 1;
-    if ( nBlocks < 12 ) {
-      nBlocks = 12;
-      nThreads = nStartHits / 12 + 1;
-      nThreads = ( nThreads / 32 + 1 ) * 32;
-    }
-
-    *fNTrackHits = 0;
+         *fGPUDebugOut << "Tracklet Hits:" << endl;
+         DumpTrackletHits(*fGPUDebugOut);
+  }
 
-    nThreads = nStartHits;
-    nBlocks = 1;
+  //std::cout<<"Slice "<<Param().ISlice()<<": NHits="<<NHitsTotal()<<", NTracklets="<<*NTracklets()<<std::endl;
 
+  RunTrackletSelector();
 
-    AliHLTTPCCAProcess<AliHLTTPCCATrackletSelector>( nBlocks, nThreads, *this );
+  //std::cout<<"Slice "<<Param().ISlice()<<": N start hits/tracklets/tracks = "<<nStartHits<<" "<<nStartHits<<" "<<*fNTracks<<std::endl;
 
-    //std::cout<<"Slice "<<Param().ISlice()<<": N start hits/tracklets/tracks = "<<nStartHits<<" "<<nStartHits<<" "<<*fNTracks<<std::endl;
+  if (fGPUDebugLevel >= 3)
+  {
+         *fGPUDebugOut << "Track Hits: (" << *NTracks() << ")" << endl;
+         DumpTrackHits(*fGPUDebugOut);
   }
 
   //std::cout<<"Memory used for slice "<<fParam.ISlice()<<" : "<<fCommonMemorySize/1024./1024.<<" + "<<fHitMemorySize/1024./1024.<<" + "<<fTrackMemorySize/1024./1024.<<" = "<<( fCommonMemorySize+fHitMemorySize+fTrackMemorySize )/1024./1024.<<" Mb "<<std::endl;
 
-
   WriteOutput();
 
-
 #endif
 
 #ifdef DRAW
@@ -397,9 +513,6 @@ GPUh() void AliHLTTPCCATracker::Reconstruct()
 
 }
 
-
-
-
 GPUh() void AliHLTTPCCATracker::WriteOutput()
 {
   // write output
@@ -535,6 +648,8 @@ GPUh() void AliHLTTPCCATracker::WriteOutput()
   fTimers[5] += timer.CpuTime();
 }
 
+#endif
+
 GPUh() void AliHLTTPCCATracker::FitTrackFull( const AliHLTTPCCATrack &/**/, float * /**/ ) const
 {
   // fit track with material
index a0931de..b6ec633 100644 (file)
@@ -41,22 +41,71 @@ class AliHLTTPCCASliceOutput;
  */
 class AliHLTTPCCATracker
 {
+       friend class AliHLTTPCCAGPUTracker;
   public:
 
-#if !defined(HLTCA_GPUCODE)
-    AliHLTTPCCATracker();
-
+       AliHLTTPCCATracker()
+               :
+               fParam(),
+               fClusterData( 0 ),
+               fData(),
+               fIsGPUTracker( false ),
+               fGPUDebugLevel( 0 ),
+               fGPUDebugOut( 0 ),
+               fCommonMemory( 0 ),
+               fCommonMemorySize( 0 ),
+               fHitMemory( 0 ),
+               fHitMemorySize( 0 ),
+               fTrackMemory( 0 ),
+               fTrackMemorySize( 0 ),
+               fNTracklets( 0 ),
+               fTrackletStartHits( 0 ),
+               fTracklets( 0 ),
+               fNTracks( 0 ),
+               fTracks( 0 ),
+               fNTrackHits( 0 ),
+               fTrackHits( 0 ),
+               fOutput( 0 ),
+               fNOutTracks( 0 ),
+               fOutTracks( 0 ),
+               fNOutTrackHits( 0 ),
+               fOutTrackHits( 0 )
+       {
+         // constructor
+       }
     GPUd() ~AliHLTTPCCATracker();
-#endif
 
-    GPUd() void Initialize( const AliHLTTPCCAParam &param );
+    void Initialize( const AliHLTTPCCAParam &param );
 
-    GPUd() void StartEvent();
+    void StartEvent();
 
-    void ReadEvent( AliHLTTPCCAClusterData *clusterData );
+       int CheckEmptySlice();
+       void WriteOutput();
 
+#if !defined(HLTCA_GPUCODE)
     void Reconstruct();
-    void WriteOutput();
+#endif
+
+       //Make Reconstruction steps directly callable (Used for GPU debugging)
+       void RunNeighboursFinder();
+       void RunNeighboursCleaner();
+       void RunStartHitsFinder();
+       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* SetGPUTrackerHitsMemory(char* pGPUMemory, int MaxNHits );
+       char* SetGPUTrackerTracksMemory(char* pGPUMemory, int MaxNTracks, int MaxNHits );
+
+       //Debugging Stuff
+       void DumpLinks(std::ostream &out);              //Dump all links to file (for comparison after NeighboursFinder/Cleaner)
+       void DumpStartHits(std::ostream &out);  //Same for Start Hits
+       void DumpTrackHits(std::ostream &out);  //Same for Track Hits
+       void DumpTrackletHits(std::ostream &out);       //Same for Track Hits
 
     GPUd() void GetErrors2( int iRow,  const AliHLTTPCCATrackParam &t, float &Err2Y, float &Err2Z ) const;
     GPUd() void GetErrors2( int iRow, float z, float sinPhi, float cosPhi, float DzDs, float &Err2Y, float &Err2Z ) const;
@@ -64,11 +113,13 @@ class AliHLTTPCCATracker
     void FitTrack( const AliHLTTPCCATrack &track, float *t0 = 0 ) const;
     void FitTrackFull( const AliHLTTPCCATrack &track, float *t0 = 0 ) const;
 
-    GPUhd() void SetPointersCommon();
-    GPUhd() void SetPointersHits( int MaxNHits );
-    GPUhd() void SetPointersTracks( int MaxNTracks, int MaxNHits );
+       void SetPointersCommon();
+    void SetPointersHits( int MaxNHits );
+    void SetPointersTracks( int MaxNTracks, int MaxNHits );
 
 #if !defined(HLTCA_GPUCODE)
+    void ReadEvent( AliHLTTPCCAClusterData *clusterData );
+
     GPUh() void WriteEvent( std::ostream &out );
     GPUh() void WriteTracks( std::ostream &out ) ;
     GPUh() void ReadTracks( std::istream &in );
@@ -80,28 +131,28 @@ class AliHLTTPCCATracker
     GPUhd() const AliHLTTPCCAClusterData *ClusterData() const { return fClusterData; }
     GPUhd() const AliHLTTPCCASliceData &Data() const { return fData; }
     GPUhd() const AliHLTTPCCARow &Row( int rowIndex ) const { return fData.Row( rowIndex ); }
-    GPUhd() const AliHLTTPCCARow &Row( const AliHLTTPCCAHitId &HitId ) const { return fData.Row( HitId.RowIndex() ); }
+    GPUh() const AliHLTTPCCARow &Row( const AliHLTTPCCAHitId &HitId ) const { return fData.Row( HitId.RowIndex() ); }
 
     GPUhd() double Timer( int i ) const { return fTimers[i]; }
     GPUhd() void SetTimer( int i, double v ) { fTimers[i] = v; }
 
     GPUhd() int NHitsTotal() const { return fData.NumberOfHits(); }
 
-    void SetHitLinkUpData( const AliHLTTPCCARow &row, int hitIndex, short v ) { fData.SetHitLinkUpData( row, hitIndex, v ); }
-    void SetHitLinkDownData( const AliHLTTPCCARow &row, int hitIndex, short v ) { fData.SetHitLinkDownData( row, hitIndex, v ); }
-    short HitLinkUpData( const AliHLTTPCCARow &row, int hitIndex ) const { return fData.HitLinkUpData( row, hitIndex ); }
-    short HitLinkDownData( const AliHLTTPCCARow &row, int hitIndex ) const { return fData.HitLinkDownData( row, hitIndex ); }
+    GPUd() void SetHitLinkUpData( const AliHLTTPCCARow &row, int hitIndex, short v ) { fData.SetHitLinkUpData( row, hitIndex, v ); }
+    GPUd() void SetHitLinkDownData( const AliHLTTPCCARow &row, int hitIndex, short v ) { fData.SetHitLinkDownData( row, hitIndex, v ); }
+    GPUd() short HitLinkUpData( const AliHLTTPCCARow &row, int hitIndex ) const { return fData.HitLinkUpData( row, hitIndex ); }
+    GPUd() short HitLinkDownData( const AliHLTTPCCARow &row, int hitIndex ) const { return fData.HitLinkDownData( row, hitIndex ); }
 
-    int FirstHitInBin( const AliHLTTPCCARow &row, int binIndex ) const { return fData.FirstHitInBin( row, binIndex ); }
+    GPUd() int FirstHitInBin( const AliHLTTPCCARow &row, int binIndex ) const { return fData.FirstHitInBin( row, binIndex ); }
 
-    unsigned short HitDataY( const AliHLTTPCCARow &row, int hitIndex ) const {
+    GPUd() unsigned short HitDataY( const AliHLTTPCCARow &row, int hitIndex ) const {
       return fData.HitDataY( row, hitIndex );
     }
-    unsigned short HitDataZ( const AliHLTTPCCARow &row, int hitIndex ) const {
+    GPUd() unsigned short HitDataZ( const AliHLTTPCCARow &row, int hitIndex ) const {
       return fData.HitDataZ( row, hitIndex );
     }
 
-    int HitInputID( const AliHLTTPCCARow &row, int hitIndex ) const { return fData.ClusterDataIndex( row, hitIndex ); }
+    GPUhd() int HitInputID( const AliHLTTPCCARow &row, int hitIndex ) const { return fData.ClusterDataIndex( row, hitIndex ); }
 
     /**
      * The hit weight is used to determine whether a hit belongs to a certain tracklet or another one
@@ -114,10 +165,10 @@ class AliHLTTPCCATracker
     static int CalculateHitWeight( int NHits, int unique ) {
       return ( NHits << 16 ) + unique;
     }
-    void MaximizeHitWeight( const AliHLTTPCCARow &row, int hitIndex, int weight ) {
+    GPUd() void MaximizeHitWeight( const AliHLTTPCCARow &row, int hitIndex, int weight ) {
       fData.MaximizeHitWeight( row, hitIndex, weight );
     }
-    int HitWeight( const AliHLTTPCCARow &row, int hitIndex ) const {
+    GPUd() int HitWeight( const AliHLTTPCCARow &row, int hitIndex ) const {
       return fData.HitWeight( row, hitIndex );
     }
 
@@ -143,8 +194,10 @@ class AliHLTTPCCATracker
     GPUhd()  int *OutTrackHits() const { return  fOutTrackHits; }
     GPUhd()  int OutTrackHit( int i ) const { return  fOutTrackHits[i]; }
 
-
+#ifndef CUDA_DEVICE_EMULATION
   private:
+#endif
+
     void SetupCommonMemory();
 
     AliHLTTPCCAParam fParam; // parameters
@@ -155,16 +208,21 @@ class AliHLTTPCCATracker
     AliHLTTPCCAClusterData *fClusterData; // ^
     AliHLTTPCCASliceData fData; // The SliceData object. It is used to encapsulate the storage in memory from the access
 
+  //Will this tracker run on GPU?
+  bool fIsGPUTracker; // is it GPU tracker
+  int fGPUDebugLevel; // debug level
+  std::ostream *fGPUDebugOut; // debug stream
+
     // event
 
     char *fCommonMemory; // common event memory
-    int   fCommonMemorySize; // size of the event memory [bytes]
+    size_t   fCommonMemorySize; // size of the event memory [bytes]
 
     char *fHitMemory; // event memory for hits
-    int   fHitMemorySize; // size of the event memory [bytes]
+    size_t   fHitMemorySize; // size of the event memory [bytes]
 
     char *fTrackMemory; // event memory for tracks
-    int   fTrackMemorySize; // size of the event memory [bytes]
+    size_t   fTrackMemorySize; // size of the event memory [bytes]
 
 
     int *fNTracklets;     // number of tracklets
index 561e804..5871c81 100644 (file)
@@ -42,8 +42,9 @@ class AliHLTTPCCATracklet
     GPUhd() void SetParam( const AliHLTTPCCATrackParam &v ) { fParam = v;      }
     GPUhd() void SetRowHit( int irow, int ih )  { fRowHits[irow] = ih;    }
 
-
+#ifndef CUDA_DEVICE_EMULATION
   private:
+#endif
 
     int fStartHitID;            // ID of the starting hit
     int fNHits;                 // N hits
index 99dd71d..710de90 100644 (file)
@@ -17,7 +17,6 @@
 //                                                                          *
 //***************************************************************************
 
-
 #include "AliHLTTPCCATracker.h"
 #include "AliHLTTPCCATrackParam.h"
 #include "AliHLTTPCCATrackParam.h"
@@ -43,7 +42,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::Step0
 {
   // reconstruction of tracklets, step 0
 
-  r.fIsMemThread = ( iThread < NMemThreads() );
+  r.fIsMemThread = ( iThread < TRACKLET_CONSTRUCTOR_NMEMTHREDS );
   if ( iThread == 0 ) {
     int nTracks = *tracker.NTracklets();
     int nTrPerBlock = nTracks / nBlocks + 1;
@@ -66,7 +65,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::Step1
 {
   // reconstruction of tracklets, step 1
 
-  r.fItr = s.fItr0 + ( iThread - NMemThreads() );
+  r.fItr = s.fItr0 + ( iThread - TRACKLET_CONSTRUCTOR_NMEMTHREDS );
   r.fGo = ( !r.fIsMemThread ) && ( r.fItr < s.fItr1 );
   r.fSave = r.fGo;
   r.fNHits = 0;
@@ -143,18 +142,18 @@ GPUd() void AliHLTTPCCATrackletConstructor::ReadData
     // FIXME: inefficient copy
     const int numberOfHits = row.NHits();
     ushort2 *sMem1 = reinterpret_cast<ushort2 *>( s.fData[jr] );
-    for ( int i = iThread; i < numberOfHits; i += NMemThreads() ) {
+    for ( int i = iThread; i < numberOfHits; i += TRACKLET_CONSTRUCTOR_NMEMTHREDS ) {
       sMem1[i].x = tracker.HitDataY( row, i );
       sMem1[i].y = tracker.HitDataZ( row, i );
     }
     short *sMem2 = reinterpret_cast<short *>( s.fData[jr] ) + 2 * numberOfHits;
-    for ( int i = iThread; i < numberOfHits; i += NMemThreads() ) {
+    for ( int i = iThread; i < numberOfHits; i += TRACKLET_CONSTRUCTOR_NMEMTHREDS ) {
       sMem2[i] = tracker.HitLinkUpData( row, i );
     }
 
     unsigned short *sMem3 = reinterpret_cast<unsigned short *>( s.fData[jr] ) + 3 * numberOfHits;
     const int n = row.FullSize(); // + grid content size
-    for ( int i = iThread; i < n; i += NMemThreads() ) {
+    for ( int i = iThread; i < n; i += TRACKLET_CONSTRUCTOR_NMEMTHREDS ) {
       sMem3[i] = tracker.FirstHitInBin( row, i );
     }
   }
@@ -173,7 +172,7 @@ GPUd() void AliHLTTPCCATrackletConstructor::StoreTracklet
 
   do {
     {
-      //std::cout<<"tracklet to store: "<<r.fItr<<", nhits = "<<r.fNHits<<std::endl;
+       //std::cout<<"tracklet to store: "<<r.fItr<<", nhits = "<<r.fNHits<<std::endl;
     }
 
     if ( r.fNHits < 5 ) {
@@ -275,7 +274,9 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
       float x = row.X();
       float y = y0 + hh.x * stepY;
       float z = z0 + hh.y * stepZ;
+#ifdef DRAW
       if ( drawFit ) std::cout << " fit tracklet: new hit " << oldIH << ", xyz=" << x << " " << y << " " << z << std::endl;
+#endif
 
       if ( iRow == r.fStartRow ) {
         tParam.SetX( x );
@@ -283,9 +284,9 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
         tParam.SetZ( z );
         r.fLastY = y;
         r.fLastZ = z;
-        //#ifdef DRAW
+        #ifdef DRAW
         if ( drawFit ) std::cout << " fit tracklet " << r.fItr << ", row " << iRow << " first row" << std::endl;
-        //#endif
+        #endif
       } else {
 
         float err2Y, err2Z;
@@ -307,11 +308,11 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
           tParam.SetCov( 2, err2Z );
         }
         if ( drawFit ) {
-          //#ifdef DRAW
+          #ifdef DRAW
           std::cout << " fit tracklet " << r.fItr << ", row " << iRow << " transporting.." << std::endl;
           std::cout << " params before transport=" << std::endl;
           tParam.Print();
-          //#endif
+          #endif
         }
         float sinPhi, cosPhi;
         if ( r.fNHits >= 10 && CAMath::Abs( tParam.SinPhi() ) < .99 ) {
@@ -321,13 +322,13 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
           sinPhi = dy * ri;
           cosPhi = dx * ri;
         }
-        //#ifdef DRAW
+        #ifdef DRAW
         if ( drawFit ) std::cout << "sinPhi0 = " << sinPhi << ", cosPhi0 = " << cosPhi << std::endl;
-        //#endif
+        #endif
         if ( !tParam.TransportToX( x, sinPhi, cosPhi, tracker.Param().ConstBz(), -1 ) ) {
-          //#ifdef DRAW
+          #ifdef DRAW
           if ( drawFit ) std::cout << " tracklet " << r.fItr << ", row " << iRow << ": can not transport!!" << std::endl;
-//#endif
+                 #endif
           if ( SAVE() ) tracklet.SetRowHit( iRow, -1 );
           break;
         }
@@ -337,34 +338,30 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
         //std::cout<<"mark2"<<std::endl;
 
         if ( drawFit ) {
-          //#ifdef DRAW
+          #ifdef DRAW
           std::cout << " params after transport=" << std::endl;
           tParam.Print();
           std::cout << "fit tracklet before filter: " << r.fItr << ", row " << iRow << " errs=" << err2Y << " " << err2Z << std::endl;
-          //#endif
-#ifdef DRAW
           AliHLTTPCCADisplay::Instance().DrawTracklet( tParam, hitstore, kBlue, 2., 1 );
           AliHLTTPCCADisplay::Instance().Ask();
-#endif
+                 #endif
         }
         if ( !tParam.Filter( y, z, err2Y, err2Z, .99 ) ) {
-          //#ifdef DRAW
+          #ifdef DRAW
           if ( drawFit ) std::cout << " tracklet " << r.fItr << ", row " << iRow << ": can not filter!!" << std::endl;
-          //#endif
+          #endif
           if ( SAVE() ) tracklet.SetRowHit( iRow, -1 );
           break;
         }
       }
       if ( SAVE() ) tracklet.SetRowHit( iRow, oldIH );
       if ( drawFit ) {
-        //#ifdef DRAW
+        #ifdef DRAW
         std::cout << "fit tracklet after filter " << r.fItr << ", row " << iRow << std::endl;
         tParam.Print();
-        //#endif
-#ifdef DRAW
         AliHLTTPCCADisplay::Instance().DrawTracklet( tParam, hitstore, kGreen, 2. );
         AliHLTTPCCADisplay::Instance().Ask();
-#endif
+               #endif
       }
       r.fNHits++;
       r.fLastRow = iRow;
@@ -373,28 +370,27 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
     } while ( 0 );
 
     if ( r.fCurrIH < 0 ) {
-      //#ifdef DRAW
+      #ifdef DRAW
       if ( drawFitted ) std::cout << "fitted tracklet " << r.fItr << ", nhits=" << r.fNHits << std::endl;
-      //#endif
+      #endif
       r.fStage = 1;
       //AliHLTTPCCAPerformance::Instance().HNHitsPerSeed()->Fill(r.fNHits);
       if ( r.fNHits < 3 ) { r.fNHits = 0; r.fGo = 0;}//SG!!!
       if ( CAMath::Abs( tParam.SinPhi() ) > .999 ) {
-        //#ifdef DRAW
+        #ifdef DRAW
         if ( drawFitted ) std::cout << " fitted tracklet  error: sinPhi=" << tParam.SinPhi() << std::endl;
-        //#endif
+        #endif
         r.fNHits = 0; r.fGo = 0;
       } else {
         //tParam.SetCosPhi( CAMath::Sqrt(1-tParam.SinPhi()*tParam.SinPhi()) );
       }
       if ( drawFitted ) {
-        //#ifdef DRAW
+        #ifdef DRAW
         std::cout << "fitted tracklet " << r.fItr << " miss=" << r.fNMissed << " go=" << r.fGo << std::endl;
         tParam.Print();
-#ifdef DRAW
         AliHLTTPCCADisplay::Instance().DrawTracklet( tParam, hitstore, kBlue );
         AliHLTTPCCADisplay::Instance().Ask();
-#endif
+               #endif
       }
       if ( r.fGo ) {
         CAMath::AtomicMax( &s.fMaxEndRow, r.fEndRow - 1 );
@@ -403,9 +399,9 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
   } else { // forward/backward searching part
     do {
       if ( drawSearch ) {
-        //#ifdef DRAW
+        #ifdef DRAW
         std::cout << "search tracklet " << r.fItr << " row " << iRow << " miss=" << r.fNMissed << " go=" << r.fGo << " stage=" << r.fStage << std::endl;
-        //#endif
+        #endif
       }
 
       if ( r.fStage == 2 && ( ( iRow >= r.fEndRow ) ||
@@ -420,15 +416,15 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
       float x = row.X();
       float err2Y, err2Z;
       if ( drawSearch ) {
-        //#ifdef DRAW
+        #ifdef DRAW
         std::cout << "tracklet " << r.fItr << " before transport to row " << iRow << " : " << std::endl;
         tParam.Print();
-        //#endif
+        #endif
       }
       if ( !tParam.TransportToX( x, tParam.SinPhi(), tParam.GetCosPhi(), tracker.Param().ConstBz(), .99 ) ) {
-        //#ifdef DRAW
+        #ifdef DRAW
         if ( drawSearch ) std::cout << " tracklet " << r.fItr << ", row " << iRow << ": can not transport!!" << std::endl;
-        //#endif
+        #endif
         break;
       }
       if ( row.NHits() < 1 ) {
@@ -436,13 +432,12 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
         break;
       }
       if ( drawSearch ) {
-        //#ifdef DRAW
+               #ifdef DRAW
         std::cout << "tracklet " << r.fItr << " after transport to row " << iRow << " : " << std::endl;
         tParam.Print();
-#ifdef DRAW
         AliHLTTPCCADisplay::Instance().DrawTracklet( tParam, hitstore, kBlue, 2., 1 );
         AliHLTTPCCADisplay::Instance().Ask();
-#endif
+               #endif
       }
       uint4 *tmpint4 = s.fData[r.fCurrentData];
 
@@ -507,10 +502,10 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
           }
         }
         if ( drawSearch ) {
-          //#ifdef DRAW
+          #ifdef DRAW
           std::cout << " tracklet " << r.fItr << ", row " << iRow << ", yz= " << fY << "," << fZ << ": search hits=" << fHitYfst << " " << fHitYlst << " / " << fHitYfst1 << " " << fHitYlst1 << std::endl;
           std::cout << " hit search :" << std::endl;
-          //#endif
+          #endif
         }
         for ( unsigned int fIh = fHitYfst; fIh < fHitYlst; fIh++ ) {
           assert( fIh < row.NHits() );
@@ -519,9 +514,9 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
           int ddz = ( int )( hh.y ) - fZ0;
           int dds = CAMath::Abs( ddy ) + CAMath::Abs( ddz );
           if ( drawSearch ) {
-            //#ifdef DRAW
+            #ifdef DRAW
             std::cout << fIh << ": hityz= " << hh.x << " " << hh.y << "(" << hh.x*stepY << " " << hh.y*stepZ << "), trackyz=" << fY0 << " " << fZ0 << "(" << fY0*stepY << " " << fZ0*stepZ << "), dy,dz,ds= " << ddy << " " << ddz << " " << dds << "(" << ddy*stepY << " " << ddz*stepZ << std::endl;
-            //#endif
+            #endif
           }
           if ( dds < ds ) {
             ds = dds;
@@ -535,9 +530,9 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
           int ddz = ( int )( hh.y ) - fZ0;
           int dds = CAMath::Abs( ddy ) + CAMath::Abs( ddz );
           if ( drawSearch ) {
-            //#ifdef DRAW
+            #ifdef DRAW
             std::cout << fIh << ": hityz= " << hh.x << " " << hh.y << "(" << hh.x*stepY << " " << hh.y*stepZ << "), trackyz=" << fY0 << " " << fZ0 << "(" << fY0*stepY << " " << fZ0*stepZ << "), dy,dz,ds= " << ddy << " " << ddz << " " << dds << "(" << ddy*stepY << " " << ddz*stepZ << std::endl;
-            //#endif
+            #endif
           }
           if ( dds < ds ) {
             ds = dds;
@@ -548,14 +543,13 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
 
       if ( best < 0 ) break;
       if ( drawSearch ) {
-        //#ifdef DRAW
+        #ifdef DRAW
         std::cout << "hit search " << r.fItr << ", row " << iRow << " hit " << best << " found" << std::endl;
-#ifdef DRAW
         AliHLTTPCCADisplay::Instance().DrawSliceHit( iRow, best, kRed, 1. );
         AliHLTTPCCADisplay::Instance().Ask();
         AliHLTTPCCADisplay::Instance().DrawSliceHit( iRow, best, kWhite, 1 );
         AliHLTTPCCADisplay::Instance().DrawSliceHit( iRow, best );
-#endif
+               #endif
       }
 
       ushort2 hh = hits[best];
@@ -577,17 +571,16 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
       if ( sz2 > 2. ) sz2 = 2.;
 
       if ( drawSearch ) {
-        //#ifdef DRAW
+        #ifdef DRAW
         std::cout << "dy,sy= " << dy << " " << CAMath::Sqrt( sy2 ) << ", dz,sz= " << dz << " " << CAMath::Sqrt( sz2 ) << std::endl;
         std::cout << "dy,dz= " << dy << " " << dz << ", sy,sz= " << CAMath::Sqrt( sy2 ) << " " << CAMath::Sqrt( sz2 ) << ", sy,sz= " << CAMath::Sqrt( kFactor*( tParam.GetErr2Y() +  err2Y ) ) << " " << CAMath::Sqrt( kFactor*( tParam.GetErr2Z() +  err2Z ) ) << std::endl;
-        //#endif
+        #endif
       }
       if ( CAMath::FMulRZ( dy, dy ) > sy2 || CAMath::FMulRZ( dz, dz ) > sz2  ) {
         if ( drawSearch ) {
-          //#ifdef DRAW
-
+          #ifdef DRAW
           std::cout << "found hit is out of the chi2 window\n " << std::endl;
-          //#endif
+          #endif
         }
         break;
       }
@@ -599,21 +592,20 @@ GPUd() void AliHLTTPCCATrackletConstructor::UpdateTracklet
 #endif
       if ( !tParam.Filter( y, z, err2Y, err2Z, .99 ) ) {
         if ( drawSearch ) {
-          //#ifdef DRAW
+          #ifdef DRAW
           std::cout << "tracklet " << r.fItr << " at row " << iRow << " : can not filter!!!! " << std::endl;
-          //#endif
+          #endif
         }
         break;
       }
       if ( SAVE() ) tracklet.SetRowHit( iRow, best );
       if ( drawSearch ) {
-        //#ifdef DRAW
+        #ifdef DRAW
         std::cout << "tracklet " << r.fItr << " after filter at row " << iRow << " : " << std::endl;
         tParam.Print();
-#ifdef DRAW
         AliHLTTPCCADisplay::Instance().DrawTracklet( tParam, hitstore, kRed );
         AliHLTTPCCADisplay::Instance().Ask();
-#endif
+               #endif
       }
       r.fNHits++;
       r.fNMissed = 0;
index 4d74cf0..2df55b2 100644 (file)
@@ -32,8 +32,12 @@ class AliHLTTPCCATrackletConstructor
             : fItr0( 0 ), fItr1( 0 ), fNRows( 0 ), fMinStartRow( 0 ), fMaxEndRow( 0 ) {}
         AliHLTTPCCASharedMemory& operator=( const AliHLTTPCCASharedMemory& /*dummy*/ ) { return *this; }
 #endif
+
+#ifndef CUDA_DEVICE_EMULATION
       protected:
-        uint4 fData[2][( 5000+5000+5000 )/4]; // temp memory
+#endif
+
+        uint4 fData[2][( ALIHLTTPCCATRACKLET_CONSTRUCTOR_TEMP_MEM+ALIHLTTPCCATRACKLET_CONSTRUCTOR_TEMP_MEM+ALIHLTTPCCATRACKLET_CONSTRUCTOR_TEMP_MEM )/4]; // temp memory
         int fItr0; // start track index
         int fItr1; // end track index
         int fNRows; // n rows
@@ -54,7 +58,10 @@ class AliHLTTPCCATrackletConstructor
             : fItr( 0 ), fFirstRow( 0 ), fLastRow( 0 ), fStartRow( 0 ), fEndRow( 0 ), fCurrIH( 0 ), fIsMemThread( 0 ), fGo( 0 ), fSave( 0 ), fCurrentData( 0 ), fStage( 0 ), fNHits( 0 ), fNMissed( 0 ), fLastY( 0 ), fLastZ( 0 ) {}
         AliHLTTPCCAThreadMemory& operator=( const AliHLTTPCCAThreadMemory& /*dummy*/ ) { return *this; }
 #endif
+
+#ifndef CUDA_DEVICE_EMULATION
       protected:
+#endif
         int fItr; // track index
         int fFirstRow;  // first row index
         int fLastRow; // last row index
@@ -101,9 +108,11 @@ class AliHLTTPCCATrackletConstructor
     GPUd() static bool SAVE() { return 1; }
 
 #if defined(HLTCA_GPUCODE)
-    GPUhd() static int NMemThreads() { return 128; }
+    //GPUhd() inline int NMemThreads() { return 128; }
+#define TRACKLET_CONSTRUCTOR_NMEMTHREDS 128
 #else
-    GPUhd() static int NMemThreads() { return 1; }
+    //GPUhd() inline int NMemThreads() { return 1; }
+#define TRACKLET_CONSTRUCTOR_NMEMTHREDS 1
 #endif
 
 };
index f46faaa..7fb238b 100644 (file)
@@ -23,7 +23,11 @@ class AliHLTTPCCATrackletSelector
     class AliHLTTPCCASharedMemory
     {
         friend class AliHLTTPCCATrackletSelector;
+#ifndef CUDA_DEVICE_EMULATION
       protected:
+#else
+         public:
+#endif
         int fItr0; // index of the first track in the block
         int fNThreadsTotal; // total n threads
         int fNTracklets; // n of tracklets
index 4d162c5..8cf3e41 100644 (file)
 #ifndef MEMORYASSIGNMENTHELPERS_H
 #define MEMORYASSIGNMENTHELPERS_H
 
+#ifndef assert
 #include <assert.h>
+#endif
 
 template<unsigned int X>
-static inline void AlignTo( char *&mem )
+GPUhd() static inline void AlignTo( char *&mem )
 {
   STATIC_ASSERT( ( X & ( X - 1 ) ) == 0, X_needs_to_be_a_multiple_of_2 );
-  const int offset = reinterpret_cast<unsigned long>( mem ) & ( X - 1 );
+  const int offset = reinterpret_cast<unsigned long long>( mem ) & ( X - 1 );
   if ( offset > 0 ) {
     mem += ( X - offset );
   }
@@ -56,7 +58,7 @@ template<typename T, unsigned int Alignment> static inline T *AssignMemory( char
   return AssignMemory<T, Alignment>( mem, stride * count );
 }
 
-template<typename T, unsigned int Alignment> static T *_assignMemory( char *&mem, unsigned int size )
+template<typename T, unsigned int Alignment> GPUhd() static T *_assignMemory( char *&mem, unsigned int size )
 {
   STATIC_ASSERT( ( Alignment & ( Alignment - 1 ) ) == 0, Alignment_needs_to_be_a_multiple_of_2 );
   AlignTo<Alignment>( mem );
@@ -65,7 +67,7 @@ template<typename T, unsigned int Alignment> static T *_assignMemory( char *&mem
   return r;
 }
 
-template<typename T> static inline void AssignMemory( T *&dst, char *&mem, int count )
+template<typename T> GPUhd() static inline void AssignMemory( T *&dst, char *&mem, int count )
 {
   dst = _assignMemory < T, ( sizeof( T ) & ( sizeof( T ) - 1 ) ) == 0 && sizeof( T ) <= 16 ? sizeof( T ) : sizeof( void * ) > ( mem, count );
 }