Update from David Rohr:
authorsgorbuno <sgorbuno@f7af4fe6-9843-0410-8265-dc069ae4e863>
Wed, 21 Oct 2009 20:58:33 +0000 (20:58 +0000)
committersgorbuno <sgorbuno@f7af4fe6-9843-0410-8265-dc069ae4e863>
Wed, 21 Oct 2009 20:58:33 +0000 (20:58 +0000)
fixes a bug for the standalone benchmark, improves scheduling for tracklet selector and
fixes the logging for the CATrackerComponent

HLT/TPCLib/tracking-ca/AliHLTTPCCAGPUConfig.h
HLT/TPCLib/tracking-ca/AliHLTTPCCAGPUTrackerNVCC.cu
HLT/TPCLib/tracking-ca/AliHLTTPCCAParam.cxx
HLT/TPCLib/tracking-ca/AliHLTTPCCAParam.h
HLT/TPCLib/tracking-ca/AliHLTTPCCATrackerComponent.cxx

index 47da945..a750a53 100644 (file)
@@ -1,70 +1,79 @@
 #ifndef ALIHLTTPCCAGPUCONFIG_H
 #define ALIHLTTPCCAGPUCONFIG_H
 
+//GPU Run Configuration
 #define HLTCA_GPU_BLOCK_COUNT 30
 #define HLTCA_GPU_THREAD_COUNT 256
 
+//GPU Parameters
 #define HLTCA_GPU_WARP_SIZE 32
 #define HLTCA_GPU_REGS 64
+
+//Detector Parameters
 #define HLTCA_ROW_COUNT 159
 
-#define HLTCA_GPU_ROWALIGNMENT uint4
-#define HLTCA_GPU_ROWCOPY int
-#define HLTCA_GPU_TRACKLET_CONSTRUCTOR_NMEMTHREDS 32
-//#define HLTCA_GPU_PREFETCHDATA
-//#define HLTCA_GPU_PREFETCH_ROWBLOCK_ONLY
+#define HLTCA_GPU_ROWALIGNMENT uint4                                   //Align Row Hits and Grid
+#define HLTCA_GPU_ROWCOPY int                                                  //must not be bigger than row alignment!!!
+#define HLTCA_GPU_TRACKLET_CONSTRUCTOR_NMEMTHREDS 32   //Amound of threads to reserve for memory copy
+//#define HLTCA_GPU_PREFETCHDATA                                               //Fetch Row Data (Hits / Grid) into shared memory during Tracklet Construction
+//#define HLTCA_GPU_PREFETCH_ROWBLOCK_ONLY                             //Prefetch only row parameters for current block during tracklet construction, faster for UpdateTracklet but shared cache can not be used for StoreTracklet then
 
-#define HLTCA_GPU_SCHED_ROW_STEP 32
-#define HLTCA_GPU_SCHED_FIXED_START
-//#define HLTCA_GPU_SCHED_FIXED_SLICE
-#define HLTCA_GPU_RESCHED
+#define HLTCA_GPU_SCHED_ROW_STEP 32                                            //Amount of Rows to process in one step before rescheduling
+#define HLTCA_GPU_SCHED_FIXED_START                                            //Assign each GPU thread a start tracklet to start with instead of using the scheduler to obtain start tracklet
+//#define HLTCA_GPU_SCHED_FIXED_SLICE                                  //Make each Multiprocessor on GPU work only on a single slice during tracklet construction
+#define HLTCA_GPU_RESCHED                                                              //Use dynamic tracklet scheduling
 
-#define HLTCA_GPU_TEXTURE_FETCH
+#define HLTCA_GPU_TEXTURE_FETCH                                                        //Fetch data through texture cache
 
-//#define HLTCA_GPU_TRACKLET_CONSTRUCTOR_DO_PROFILE
-//#define HLTCA_GPU_TIME_PROFILE
+//#define HLTCA_GPU_TRACKLET_CONSTRUCTOR_DO_PROFILE            //Output Profiling Data for Tracklet Constructor Tracklet Scheduling
+//#define HLTCA_GPU_TIME_PROFILE                                               //Output Time Profiling Data for asynchronous DMA transfer
 
 #define HLTCA_GPU_TRACKLET_SELECTOR_HITS_REG_SIZE 12
 #define HLTCA_GPU_TRACKLET_SELECTOR_SLICE_COUNT 3              //Currently must be smaller than avaiable MultiProcessors on GPU or will result in wrong results
 
-#define HLTCA_GPU_SORT_DUMPDATA
+#define HLTCA_GPU_SORT_DUMPDATA                                                        //Sort Start Hits etc before dumping to file
 
-#define HLTCA_GPU_MAX_TRACKLETS 12288                                  //Should be divisible by 16 at least
-#define HLTCA_GPU_MAX_TRACKS 3072
+#define HLTCA_GPU_MAX_TRACKLETS 12288                                  //Max Number of Tracklets that can be processed by GPU Tracker, Should be divisible by 16 at least
+#define HLTCA_GPU_MAX_TRACKS 3072                                              //Max number of Tracks that can be processd by GPU Tracker
 
-//#define HLTCA_GPU_EMULATION_SINGLE_TRACKLET 1313
+//#define HLTCA_GPU_EMULATION_SINGLE_TRACKLET 1313             //Run Tracklet constructor on on single Tracklet in Device Emulation Mode
 //#define HLTCA_GPU_EMULATION_DEBUG_TRACKLET 1313
 
-#define HLTCA_GPU_TRACKER_CONSTANT_MEM 65000
+#define HLTCA_GPU_TRACKER_CONSTANT_MEM 65000                   //Amount of Constant Memory to reserve
+
+#define HLTCA_GPU_TRACKER_OBJECT_MEMORY 1024 * 1024            //Total amount of Memory to reserve for GPU Tracker Objects
+#define HLTCA_GPU_ROWS_MEMORY 1024 * 1024                              //Total amount of Memory to reserve for GPU Row Parameters
+#define HLTCA_GPU_COMMON_MEMORY 1024 * 1024                            //Total amount of Memory to reserve for CommomMemoryStruct on GPU
+#define HLTCA_GPU_SLICE_DATA_MEMORY 7 * 1024 * 1024            //Amount of Slice Data Memory to reserve per Slice on GPU
+#define HLTCA_GPU_GLOBAL_MEMORY 20 * 1024 * 1024               //Amount of global temporary Memory to reserve per Slice on GPU
+#define HLTCA_GPU_TRACKS_MEMORY 2 * 1024 * 1024                        //Amount of Memory to reserve for Final Tracks per Slice on GPU
 
-#define HLTCA_GPU_TRACKER_OBJECT_MEMORY 1024 * 1024
-#define HLTCA_GPU_ROWS_MEMORY 1024 * 1024
-#define HLTCA_GPU_COMMON_MEMORY 1024 * 1024
-#define HLTCA_GPU_SLICE_DATA_MEMORY 7 * 1024 * 1024
-#define HLTCA_GPU_GLOBAL_MEMORY 20 * 1024 * 1024
-#define HLTCA_GPU_TRACKS_MEMORY 2 * 1024 * 1024
+//Make sure options do not interfere
 
 #ifndef HLTCA_GPUCODE
+//No texture fetch for CPU Tracker
 #ifdef HLTCA_GPU_TEXTURE_FETCH
 #undef HLTCA_GPU_TEXTURE_FETCH
 #endif
-
+//No Shared memory cache for CPU Tracker
 #ifdef HLTCA_GPU_PREFETCHDATA
 #undef HLTCA_GPU_PREFETCHDATA
 #endif
-
+//Do not cache Row Hits during Tracklet selection in Registers for CPU Tracker
 #undef HLTCA_GPU_TRACKLET_SELECTOR_HITS_REG_SIZE
 #define HLTCA_GPU_TRACKLET_SELECTOR_HITS_REG_SIZE 0
-
 #else
+//Sort start hits for GPU tracker
 #define HLTCA_GPU_SORT_STARTHITS
 #endif
 
+//If not using Row Based schreduling or not using shared memory cache do not reserve threads for shared memory copy
 #if !defined(HLTCA_GPU_PREFETCHDATA) | !defined(HLTCA_GPU_RESCHED)
 #undef HLTCA_GPU_TRACKLET_CONSTRUCTOR_NMEMTHREDS
 #define HLTCA_GPU_TRACKLET_CONSTRUCTOR_NMEMTHREDS 0
 #endif
 
+//Error Codes for GPU Tracker
 #define HLTCA_GPU_ERROR_NONE 0
 #define HLTCA_GPU_ERROR_ROWBLOCK_TRACKLET_OVERFLOW 1
 #define HLTCA_GPU_ERROR_TRACKLET_OVERFLOW 2
index c075f10..c9cb853 100644 (file)
@@ -445,7 +445,7 @@ int AliHLTTPCCAGPUTracker::Reconstruct(AliHLTTPCCASliceOutput** pOutput, AliHLTT
        if (fDebugLevel >= 3) HLTInfo("Allocating GPU Tracker memory and initializing constants");
 
 #ifdef HLTCA_GPU_TIME_PROFILE
-       __int64 a, b, c, d;
+       unsigned __int64 a, b, c, d;
        AliHLTTPCCAStandaloneFramework::StandaloneQueryFreq(&c);
        AliHLTTPCCAStandaloneFramework::StandaloneQueryTime(&d);
 #endif
@@ -698,10 +698,12 @@ int AliHLTTPCCAGPUTracker::Reconstruct(AliHLTTPCCASliceOutput** pOutput, AliHLTT
                }
        }
 
-       for (int iSlice = 0;iSlice < sliceCountLocal;iSlice += HLTCA_GPU_TRACKLET_SELECTOR_SLICE_COUNT)
+       int runSlices = 0;
+       for (int iSlice = 0;iSlice < sliceCountLocal;iSlice += runSlices)
        {
-               if (fDebugLevel >= 3) HLTInfo("Running HLT Tracklet selector (Slice %d to %d)", iSlice, iSlice + HLTCA_GPU_TRACKLET_SELECTOR_SLICE_COUNT);
-               AliHLTTPCCAProcessMulti<AliHLTTPCCATrackletSelector><<<HLTCA_GPU_BLOCK_COUNT, HLTCA_GPU_THREAD_COUNT, 0, cudaStreams[iSlice]>>>(iSlice, CAMath::Min(HLTCA_GPU_TRACKLET_SELECTOR_SLICE_COUNT, sliceCountLocal - iSlice));
+               if (runSlices < HLTCA_GPU_TRACKLET_SELECTOR_SLICE_COUNT) runSlices++;
+               if (fDebugLevel >= 3) HLTInfo("Running HLT Tracklet selector (Slice %d to %d)", iSlice, iSlice + runSlices);
+               AliHLTTPCCAProcessMulti<AliHLTTPCCATrackletSelector><<<HLTCA_GPU_BLOCK_COUNT, HLTCA_GPU_THREAD_COUNT, 0, cudaStreams[iSlice]>>>(iSlice, CAMath::Min(runSlices, sliceCountLocal - iSlice));
        }
        if (CUDASync("Tracklet Selector")) return 1;
        StandalonePerfTime(firstSlice, 9);
index 3e0496b..92f9b19 100644 (file)
@@ -178,6 +178,7 @@ GPUd() void AliHLTTPCCAParam::GetClusterErrors2( int iRow, float z, float sinPhi
   Err2Z = GetClusterError2( 1, type, z, angleZ );
 }
 
+#ifndef HLTCA_GPUCODE
 GPUh() void AliHLTTPCCAParam::WriteSettings( std::ostream &out ) const
 {
   // write settings to the file
@@ -248,3 +249,4 @@ GPUh() void AliHLTTPCCAParam::ReadSettings( std::istream &in )
       for ( int k = 0; k < 7; k++ )
         in >> fParamS0Par[i][j][k];
 }
+#endif
index ab58e61..7a1c463 100644 (file)
 class AliHLTTPCCAParam
 {
   public:
+       AliHLTTPCCAParam();
+    ~AliHLTTPCCAParam() {}
 
 #if !defined(HLTCA_GPUCODE)
-    GPUd() AliHLTTPCCAParam();
-
-    ~AliHLTTPCCAParam() {;}
-
     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 );
     void Update();
-
 #endif
 
        GPUd() void Slice2Global( float x, float y,  float z,
index fc01311..53ce81e 100644 (file)
@@ -651,14 +651,13 @@ int AliHLTTPCCATrackerComponent::DoEvent
 
          if (sliceOutput[islice])
          {
-                 Logging( kHLTLogDebug, "HLT::TPCCATracker::DoEvent", "Reconstruct",
-                                  "%d tracks found for slice %d", sliceOutput[islice]->NOutTracks(), slice );
-
-
                  // write reconstructed tracks
 
                  if ( fOutputTRAKSEGS ) {
 
+                 Logging( kHLTLogDebug, "HLT::TPCCATracker::DoEvent", "Reconstruct",
+                                  "%d tracks found for slice %d", sliceOutput[islice]->NOutTracks(), slice );
+
                        ntracks = sliceOutput[islice]->NOutTracks();
 
                        AliHLTTPCTrackletData* outPtr = ( AliHLTTPCTrackletData* )( outputPtr );
@@ -763,6 +762,10 @@ int AliHLTTPCCATrackerComponent::DoEvent
                        }
 
                  } else { // default output type
+
+                 Logging( kHLTLogDebug, "HLT::TPCCATracker::DoEvent", "Reconstruct",
+                                  "%d tracks found for slice %d", sliceOutput[islice]->NTracks(), slice );
+
                          mySize += sliceOutput[islice]->OutputMemorySize();
                          ntracks += sliceOutput[islice]->NTracks();
                  }