Skip to content

Commit ba0eeb3

Browse files
f3schf3sch
committed
some alignment
Signed-off-by: Felix Schlepper <felix.schlepper@cern.ch>
1 parent 3ac35d9 commit ba0eeb3

3 files changed

Lines changed: 62 additions & 62 deletions

File tree

‎Detectors/ITSMFT/ITS/tracking/GPU/ITStrackingGPU/LaunchGeometry.h‎

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -11,6 +11,8 @@
1111
///
1212
/// \file LaunchGeometry.h
1313
/// \brief Compile-time launch geometry of the ITS tracking kernels, per GPU family.
14+
/// Poor man's RTC
15+
/// to be removed/reworked entirely once we can use Gabriele's tuner
1416
///
1517

1618
#ifndef ITSTRACKINGGPU_LAUNCHGEOMETRY_H_
@@ -50,7 +52,7 @@ constexpr int WarpSize = 32;
5052
constexpr int ComputeUnits = 68;
5153
constexpr int WarpSize = 32;
5254
#else
53-
#warning "GPU architecture not available setting fallback"
55+
// this is the fallback as we had it before
5456
constexpr int ComputeUnits = 60;
5557
constexpr int WarpSize = 64;
5658
#endif
@@ -131,7 +133,7 @@ constexpr KernelOccupancy ResidentBlocks{
131133
.processNeighboursTrackSeed = 3, // 84 VGPR
132134
.fitTrackSeeds = 4, // 64 VGPR
133135
.fitTrackSeedsExtended = 3, // 84 VGPR
134-
.compileLookupTable = 4, // 8 VGPR, could hold ten; left at the historical grid
136+
.compileLookupTable = 4, // 8 VGPR,
135137
};
136138

137139
#elif defined(__HIPCC__) || defined(__HIP_PLATFORM_AMD__)

‎Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu‎

Lines changed: 57 additions & 59 deletions
Original file line numberDiff line numberDiff line change
@@ -1039,29 +1039,28 @@ void TrackingKernels<NLayers>::processNeighboursHandler(const int startLevel,
10391039
gpu::processNeighboursKernel<NLayers, std::remove_pointer_t<decltype(levelSeeds)>><<<gpu::gridBlocks(std::is_same_v<std::remove_pointer_t<decltype(levelSeeds)>, CellSeed>
10401040
? gpu::ResidentBlocks.processNeighboursCellSeed
10411041
: gpu::ResidentBlocks.processNeighboursTrackSeed),
1042-
gpu::GPUThreads>>>(
1043-
topologyId,
1044-
level,
1045-
allCellSeeds,
1046-
levelSeeds,
1047-
levelCellIds,
1048-
levelCellTopologyIds,
1049-
nLevelSeeds,
1050-
thrust::raw_pointer_cast(staged),
1051-
thrust::raw_pointer_cast(stagedCellIds),
1052-
thrust::raw_pointer_cast(stagedCellTopologyIds),
1053-
thrust::raw_pointer_cast(sourceSeeds),
1054-
thrust::raw_pointer_cast(outputCounter),
1055-
out.capacity,
1056-
usedClusters,
1057-
neighbours,
1058-
neighboursDeviceLUTs,
1059-
foundTrackingFrameInfo,
1060-
layerxX0,
1061-
bz,
1062-
maxChi2ClusterAttachment,
1063-
propagator,
1064-
matCorrType);
1042+
gpu::GPUThreads>>>(topologyId,
1043+
level,
1044+
allCellSeeds,
1045+
levelSeeds,
1046+
levelCellIds,
1047+
levelCellTopologyIds,
1048+
nLevelSeeds,
1049+
thrust::raw_pointer_cast(staged),
1050+
thrust::raw_pointer_cast(stagedCellIds),
1051+
thrust::raw_pointer_cast(stagedCellTopologyIds),
1052+
thrust::raw_pointer_cast(sourceSeeds),
1053+
thrust::raw_pointer_cast(outputCounter),
1054+
out.capacity,
1055+
usedClusters,
1056+
neighbours,
1057+
neighboursDeviceLUTs,
1058+
foundTrackingFrameInfo,
1059+
layerxX0,
1060+
bz,
1061+
maxChi2ClusterAttachment,
1062+
propagator,
1063+
matCorrType);
10651064
int wanted{0};
10661065
GPUChkErrS(cudaMemcpyAsync(&wanted, thrust::raw_pointer_cast(outputCounter), sizeof(int), cudaMemcpyDeviceToHost, gpu::Stream::DefaultStream));
10671066
GPUChkErrS(cudaStreamSynchronize(gpu::Stream::DefaultStream));
@@ -1156,42 +1155,41 @@ int TrackingKernels<NLayers>::computeTrackSeedHandler(TrackSeed<NLayers>* trackS
11561155
const auto launchFit = [&](auto extendTracks) {
11571156
gpu::fitTrackSeedsKernel<NLayers, decltype(extendTracks)::value><<<gpu::gridBlocks(decltype(extendTracks)::value ? gpu::ResidentBlocks.fitTrackSeedsExtended
11581157
: gpu::ResidentBlocks.fitTrackSeeds),
1159-
gpu::GPUThreads>>>(
1160-
trackSeeds, // CellSeed*
1161-
foundTrackingFrameInfo, // TrackingFrameInfo**
1162-
unsortedClusters, // Cluster**
1163-
utils, // IndexTableUtils*
1164-
rofMask, // ROFMaskTable::View
1165-
rofOverlaps, // ROFOverlapTable::View
1166-
clusters, // Cluster**
1167-
usedClusters, // unsigned char**
1168-
clustersIndexTables, // int**
1169-
ROFClusters, // int**
1170-
tracks, // TrackITSExt*
1171-
trackSeedIndices, // int*
1172-
outputCounter, // int*
1173-
trackCapacity, // const int
1174-
activeHypotheses, // TrackExtensionHypothesis*
1175-
nextHypotheses, // TrackExtensionHypothesis*
1176-
layerRadii, // const float*
1177-
minPts, // const float*
1178-
layerxX0, // const float*
1179-
nSeeds, // const unsigned int
1180-
bz, // const float
1181-
maxChi2ClusterAttachment, // float
1182-
maxChi2NDF, // float
1183-
reseedIfShorter, // int
1184-
repeatRefitOut, // bool
1185-
shiftRefToCluster, // bool
1186-
nLayers, // int
1187-
phiBins, // int
1188-
maxHypotheses, // int
1189-
extendTop, // bool
1190-
extendBot, // bool
1191-
nSigmaCutPhi, // float
1192-
nSigmaCutZ, // float
1193-
propagator, // const o2::base::Propagator*
1194-
matCorrType); // o2::base::PropagatorF::MatCorrType
1158+
gpu::GPUThreads>>>(trackSeeds, // CellSeed*
1159+
foundTrackingFrameInfo, // TrackingFrameInfo**
1160+
unsortedClusters, // Cluster**
1161+
utils, // IndexTableUtils*
1162+
rofMask, // ROFMaskTable::View
1163+
rofOverlaps, // ROFOverlapTable::View
1164+
clusters, // Cluster**
1165+
usedClusters, // unsigned char**
1166+
clustersIndexTables, // int**
1167+
ROFClusters, // int**
1168+
tracks, // TrackITSExt*
1169+
trackSeedIndices, // int*
1170+
outputCounter, // int*
1171+
trackCapacity, // const int
1172+
activeHypotheses, // TrackExtensionHypothesis*
1173+
nextHypotheses, // TrackExtensionHypothesis*
1174+
layerRadii, // const float*
1175+
minPts, // const float*
1176+
layerxX0, // const float*
1177+
nSeeds, // const unsigned int
1178+
bz, // const float
1179+
maxChi2ClusterAttachment, // float
1180+
maxChi2NDF, // float
1181+
reseedIfShorter, // int
1182+
repeatRefitOut, // bool
1183+
shiftRefToCluster, // bool
1184+
nLayers, // int
1185+
phiBins, // int
1186+
maxHypotheses, // int
1187+
extendTop, // bool
1188+
extendBot, // bool
1189+
nSigmaCutPhi, // float
1190+
nSigmaCutZ, // float
1191+
propagator, // const o2::base::Propagator*
1192+
matCorrType); // o2::base::PropagatorF::MatCorrType
11951193
};
11961194
if (extendTop || extendBot) {
11971195
launchFit(std::true_type{});

‎Detectors/ITSMFT/ITS/tracking/src/CapacityEstimator.cxx‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -31,7 +31,7 @@ struct CapacityEstimator::Impl {
3131
size_t maxEmitted{0};
3232
uint32_t nSamples{0};
3333
uint32_t nLowStreak{0};
34-
uint32_t nOverflows{0}; ///< times the slab was too small and the work had to be redone
34+
uint32_t nOverflows{0};
3535
};
3636

3737
explicit Impl(Config config) : cfg{config} {}

0 commit comments

Comments
 (0)