message: add scheduling for up/downward searching concurrently
diff: @@ -444,14 +444,26 @@ GPUdi() int AliHLTTPCCATrackletConstructor::FetchTracklet(AliHLTTPCCATracker &tr
  {
  if (tracker.GPUParameters()->fNextTracklet < nTracklets)
  {
  - const int firstTracklet = CAMath::AtomicAdd(&tracker.GPUParameters()->fNextTracklet, HLTCA_GPU_THREAD_COUNT);
  + const int firstTracklet = /*sMem.fNextTrackletCount != HLTCA_GPU_THREAD_COUNT ? -1 :*/ CAMath::AtomicAdd(&tracker.GPUParameters()->fNextTracklet, sMem.fNextTrackletCount);
  if (firstTracklet < nTracklets) sMem.fNextTrackletFirst = firstTracklet;
  - else sMem.fNextTrackletFirst = -2;
  + else sMem.fNextTrackletFirst = -1;
  }
  else
  + {
  + sMem.fNextTrackletFirst = -1;
  + }
  + }
  + }
  + if (threadIdx.x == 0)
  + {
  + if (sMem.fNextTrackletFirst == -1 && sMem.fNextTrackletCount == HLTCA_GPU_THREAD_COUNT)
  {
  sMem.fNextTrackletFirst = -2;
  + sMem.fNextTrackletCount = HLTCA_GPU_THREAD_COUNT;
  }
  + else
  + {
  + sMem.fNextTrackletCount = 0;
  }
  }
  __syncthreads();
  @@ -485,7 +497,7 @@ GPUdi() void AliHLTTPCCATrackletConstructor::AliHLTTPCCATrackletConstructorGPU(A
  AliHLTTPCCAThreadMemory rMem;
  int tmpTracklet;
  - while (keepTracklet || (tmpTracklet = FetchTracklet(tracker, sMem)) != -2)
  + while ((tmpTracklet = FetchTracklet(tracker, sMem)) != -2)
  {
  if (!keepTracklet)
  {
  @@ -513,6 +525,8 @@ GPUdi() void AliHLTTPCCATrackletConstructor::AliHLTTPCCATrackletConstructorGPU(A
  currentSlice = iSlice;
  }
  + if (iTracklet >= 0)
  + {
  if (keepTracklet) {}
  else if (iTracklet < sMem.fNTracklets)
  {
  @@ -541,8 +555,6 @@ GPUdi() void AliHLTTPCCATrackletConstructor::AliHLTTPCCATrackletConstructorGPU(A
  rMem.fGo = 0;
  rMem.fStartRow = rMem.fEndRow = 0;
  iRowStart = iRowEnd = 0;
  - keepTracklet = 1;////
  - iPhase = 0;////
  }
  for (int j = iRowStart;j != iRowEnd;j+=iRowIncrement)
  @@ -552,11 +564,7 @@ GPUdi() void AliHLTTPCCATrackletConstructor::AliHLTTPCCATrackletConstructorGPU(A
  if (iPhase)
  {
  - if (iTracklet < sMem.fNTracklets)
  - {
  - StoreTracklet( 1, 1, 0, iTracklet, sMem, rMem, tracker, tParam );
  - }
  - keepTracklet = 0;
  + rMem.fGo = 0;
  }
  else
  {
  @@ -569,6 +577,14 @@ GPUdi() void AliHLTTPCCATrackletConstructor::AliHLTTPCCATrackletConstructorGPU(A
  iPhase = 1;
  }
  }
  +
  + if (!rMem.fGo)
  + {
  + if (iTracklet < sMem.fNTracklets && iTracklet >= 0) StoreTracklet( 1, 1, 0, iTracklet, sMem, rMem, tracker, tParam );
  + keepTracklet = 0;
  + iNextLocalTracklet = CAMath::AtomicAdd(&sMem.fNextTrackletCount, 1);
  + }
  + }
  }
  }
  