Skip to content

Commit

Permalink
asdf
Browse files Browse the repository at this point in the history
  • Loading branch information
@mconcas
mconcas committed Dec 6, 2024
1 parent 7842b04 commit 16c5b70
Show file tree
Hide file tree
Showing 5 changed files with 69 additions and 46 deletions.
1 change: 1 addition & 0 deletions Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TimeFrameGPU.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -114,6 +114,7 @@ class TimeFrameGPU : public TimeFrame
Road<nLayers - 2>* getDeviceRoads() { return mRoadsDevice; }
TrackITSExt* getDeviceTrackITSExt() { return mTrackITSExtDevice; }
int* getDeviceNeighboursLUT(const int layer) { return mNeighboursLUTDevice[layer]; }
gsl::span<int*> getDeviceNeighboursLUTs() { return mNeighboursLUTDevice; }
gpuPair<int, int>* getDeviceNeighbourPairs(const int layer) { return mNeighbourPairsDevice[layer]; }
int* getDeviceNeighbours(const int layer) { return mNeighboursDevice[layer]; }
int** getDeviceNeighboursArray() { return mNeighboursDeviceArray; }
Expand Down
6 changes: 3 additions & 3 deletions Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/TrackingKernels.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -180,8 +180,8 @@ int filterCellNeighboursHandler(std::vector<int>&,
int*,
unsigned int);

void processNeighboursHandler(const int iteration,
const int startLayer,
template <int nLayers = 7>
void processNeighboursHandler(const int startLayer,
const int startLevel,
CellSeed** allCellSeeds,
CellSeed* currentCellSeeds,
Expand All @@ -192,7 +192,7 @@ void processNeighboursHandler(const int iteration,
// int* updatedCellsIds,
const unsigned char** usedClusters, // Used clusters
int* neighbours,
int* neighboursLUT,
gsl::span<int*> neighboursDeviceLUTs,
const TrackingFrameInfo** foundTrackingFrameInfo,
const float bz,
const float MaxChi2ClusterAttachment,
Expand Down
5 changes: 2 additions & 3 deletions Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackerTraitsGPU.cxx
View file
Original file line number Diff line number Diff line change
Expand Up @@ -279,15 +279,14 @@ void TrackerTraitsGPU<nLayers>::findRoads(const int iteration)
const int minimumLayer{startLevel - 1};
std::vector<CellSeed> trackSeeds;
for (int startLayer{mTrkParams[iteration].CellsPerRoad() - 1}; startLayer >= minimumLayer; --startLayer) {
processNeighboursHandler(iteration,
startLayer,
processNeighboursHandler(startLayer,
startLevel,
mTimeFrameGPU->getDeviceArrayCells(),
mTimeFrameGPU->getDeviceCells()[startLayer],
mTimeFrameGPU->getNCells()[startLayer],
mTimeFrameGPU->getDeviceArrayUsedClusters(),
mTimeFrameGPU->getDeviceNeighbours(startLayer - 1),
mTimeFrameGPU->getDeviceNeighboursLUT(startLayer - 1),
mTimeFrameGPU->getDeviceNeighboursLUTs(),
mTimeFrameGPU->getDeviceArrayTrackingFrameInfo(),
mBz,
mTrkParams[0].MaxChi2ClusterAttachment, // float maxChi2ClusterAttachment
Expand Down
89 changes: 56 additions & 33 deletions Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -606,11 +606,11 @@ GPUg() void processNeighboursKernel(const int layer,
for (unsigned int iCurrentCell = blockIdx.x * blockDim.x + threadIdx.x; iCurrentCell < nCurrentCells; iCurrentCell += blockDim.x * gridDim.x) {
int foundSeeds{0};
const auto& currentCell{currentCellSeeds[iCurrentCell]};
if constexpr (debug) {
if (threadIdx.x == 0 && currentCellIds != nullptr) {
currentCellSeeds[iCurrentCell].printCell();
}
}
// if constexpr (debug) {
// if (threadIdx.x == 0 && currentCellIds != nullptr) {
// currentCellSeeds[iCurrentCell].printCell();
// }
// }
if (currentCell.getLevel() != level) {
continue;
}
Expand Down Expand Up @@ -1176,6 +1176,7 @@ int filterCellNeighboursHandler(std::vector<int>& neighHost, // TODO: eventually
return trimmedSize;
}
template <int nLayers>
void processNeighboursHandler(const int iteration,
const int startLayer,
const int startLevel,
Expand All @@ -1188,7 +1189,7 @@ void processNeighboursHandler(const int iteration,
// int* updatedCellsIds,
const unsigned char** usedClusters, // Used clusters
int* neighbours,
int* neighboursLUT,
gsl::span<int*> neighboursDeviceLUTs,
const TrackingFrameInfo** foundTrackingFrameInfo,
const float bz,
const float MaxChi2ClusterAttachment,
Expand All @@ -1202,14 +1203,14 @@ void processNeighboursHandler(const int iteration,
startLevel,
allCellSeeds,
currentCellSeeds,
nullptr, // currentCellIds,
nullptr, // currentCellIds,
nCurrentCells,
nullptr, // updatedCellSeeds,
nullptr, // updatedCellsIds,
thrust::raw_pointer_cast(&foundSeedsTable[0]), // auxiliary only in GPU code to compute the number of cells per iteration
usedClusters, // Used clusters
neighbours,
neighboursLUT,
neighboursDeviceLUTs[startLayer - 1],
foundTrackingFrameInfo,
bz,
MaxChi2ClusterAttachment,
Expand Down Expand Up @@ -1238,14 +1239,14 @@ void processNeighboursHandler(const int iteration,
startLevel,
allCellSeeds,
currentCellSeeds,
nullptr, // currentCellIds,
nullptr, // currentCellIds,
nCurrentCells,
thrust::raw_pointer_cast(&updatedCellSeeds[0]), // updatedCellSeeds
thrust::raw_pointer_cast(&updatedCellIds[0]), // updatedCellsIds
thrust::raw_pointer_cast(&foundSeedsTable[0]), // auxiliary only in GPU code to compute the number of cells per iteration
usedClusters, // Used clusters
neighbours,
neighboursLUT,
neighboursDeviceLUTs[startLayer - 1],
foundTrackingFrameInfo,
bz,
MaxChi2ClusterAttachment,
Expand All @@ -1254,33 +1255,35 @@ void processNeighboursHandler(const int iteration,
LOGP(info, "1: updatedCellIds {} - updatedCellSeeds {}", updatedCellIds.size(), updatedCellSeeds.size());
// gpu::printCellSeeds<<<1, 1>>>(thrust::raw_pointer_cast(&updatedCellSeeds[0]), foundSeedsTable.back());
// gpu::printBufferLayerOnThread<<<1, 1>>>(0, thrust::raw_pointer_cast(&updatedCellIds[0]), updatedCellIds.size());
// gpu::printCellSeeds<<<1, 1>>>(thrust::raw_pointer_cast(&lastCellSeeds[0]), lastCellSeeds.size());
int level = startLevel;
for (int iLayer{startLayer - 1}; iLayer > 0 && level > 2; --iLayer) {
LOGP(info, "iLayer: {} level {}", iLayer, level);
--level;
// lastCellSeeds.swap(updatedCellSeeds);
// lastCellIds.swap(updatedCellIds);
// // gpu::printCellSeeds<<<1, 1>>>(thrust::raw_pointer_cast(&lastCellSeeds[0]), lastCellSeeds.size());
// foundSeedsTable.resize(lastCellSeeds.size() + 1);
// thrust::fill(foundSeedsTable.begin(), foundSeedsTable.end(), 0);
// gpu::processNeighboursKernel<true, true><<<1, 1>>>(iLayer,
// level,
// allCellSeeds,
// thrust::raw_pointer_cast(&lastCellSeeds[0]),
// thrust::raw_pointer_cast(&lastCellIds[0]), // currentCellIds,
// lastCellSeeds.size(),
// nullptr, // updatedCellSeeds,
// nullptr, // updatedCellsIds,
// thrust::raw_pointer_cast(&foundSeedsTable[0]), // auxiliary only in GPU code to compute the number of cells per iteration
// usedClusters, // Used clusters
// neighbours,
// neighboursLUT,
// foundTrackingFrameInfo,
// bz,
// MaxChi2ClusterAttachment,
// propagator,
// matCorrType);
// gpu::printBufferLayerOnThread<<<1, 1>>>(0, thrust::raw_pointer_cast(&foundSeedsTable[0]), foundSeedsTable.size());
lastCellSeeds.swap(updatedCellSeeds);
lastCellIds.swap(updatedCellIds);
foundSeedsTable.resize(lastCellSeeds.size() + 1);
thrust::fill(foundSeedsTable.begin(), foundSeedsTable.end(), 0);
gpu::processNeighboursKernel<true, true><<<1, 1>>>(iLayer,
level,
allCellSeeds,
thrust::raw_pointer_cast(&lastCellSeeds[0]),
thrust::raw_pointer_cast(&lastCellIds[0]), // currentCellIds,
lastCellSeeds.size(),
nullptr, // updatedCellSeeds,
nullptr, // updatedCellsIds,
thrust::raw_pointer_cast(&foundSeedsTable[0]), // auxiliary only in GPU code to compute the number of cells per iteration
usedClusters, // Used clusters
neighbours,
neighboursDeviceLUTs[iLayer - 1],
foundTrackingFrameInfo,
bz,
MaxChi2ClusterAttachment,
propagator,
matCorrType);
gpu::printBufferLayerOnThread<<<1, 1>>>(0, thrust::raw_pointer_cast(&foundSeedsTable[0]), foundSeedsTable.size());
gpu::printBufferLayerOnThread<<<1, 1>>>(0, thrust::raw_pointer_cast(&lastCellIds[0]), lastCellIds.size());
break;
// gpuCheckError(cudaPeekAtLastError());
// gpuCheckError(cudaDeviceSynchronize());
// gpuCheckError(cub::DeviceScan::ExclusiveSum(d_temp_storage_2, // d_temp_storage
Expand Down Expand Up @@ -1411,4 +1414,24 @@ template void computeTrackletsInROFsHandler<7>(const IndexTableUtils* utils,
std::vector<float>& mulScatAng,
const int nBlocks,
const int nThreads);
template void processNeighboursHandler<7>(const int startLayer,
const int startLevel,
CellSeed** allCellSeeds,
CellSeed* currentCellSeeds,
const unsigned int nCurrentCells,
// const int* currentCellIds,
// const unsigned int nCurrentCellsIds,
// CellSeed* updatedCellSeeds,
// int* updatedCellsIds,
const unsigned char** usedClusters, // Used clusters
int* neighbours,
gsl::span<int*> neighboursDeviceLUTs,
const TrackingFrameInfo** foundTrackingFrameInfo,
const float bz,
const float MaxChi2ClusterAttachment,
const o2::base::Propagator* propagator,
const o2::base::PropagatorF::MatCorrType matCorrType,
const int nBlocks,
const int nThreads);
} // namespace o2::its
14 changes: 7 additions & 7 deletions Detectors/ITSMFT/ITS/tracking/src/TrackerTraits.cxx
View file
Original file line number Diff line number Diff line change
Expand Up @@ -585,25 +585,25 @@ void TrackerTraits::findRoads(const int iteration)
std::vector<int> lastCellId, updatedCellId;
std::vector<CellSeed> lastCellSeed, updatedCellSeed;

printf("\n");
processNeighbours(startLayer, startLevel, mTimeFrame->getCells()[startLayer], lastCellId, updatedCellSeed, updatedCellId);
LOGP(info, "1: updatedCellIds {} - updatedCellSeeds {}", updatedCellId.size(), updatedCellSeed.size());
// for (auto& c : updatedCellSeed) {
// c.printCell();
// }
// for (int i{0}; i < updatedCellId.size(); ++i) {
// if (!(i % 150)) {
// printf("\n layer %d: ===> %d/%d\t", 0, i, (int)updatedCellId.size());
// }
// printf("%d\t", updatedCellId[i]);
// }
// printf("\n");
int level = startLevel;
for (int iLayer{startLayer - 1}; iLayer > 0 && level > 2; --iLayer) {
LOGP(info, "iLayer: {} level {}", iLayer, level);
lastCellSeed.swap(updatedCellSeed);
lastCellId.swap(updatedCellId);
std::vector<CellSeed>().swap(updatedCellSeed); /// tame the memory peaks
updatedCellId.clear();
for (int i{0}; i < lastCellId.size(); ++i) {
if (!(i % 150)) {
printf("\n layer %d: ===> %d/%d\t", 0, i, (int)lastCellId.size());
}
printf("%d\t", lastCellId[i]);
}
processNeighbours(iLayer, --level, lastCellSeed, lastCellId, updatedCellSeed, updatedCellId);
// LOGP(info, " -> 2: updatedCellIds {} - updatedCellSeeds {}", updatedCellId.size(), updatedCellSeed.size());
}
Expand Down

0 comments on commit 16c5b70

Please sign in to comment.