Skip to content

Commit 66676ed

Browse files
f3schf3sch
committed
ITS: slab allocator with size estimator
Signed-off-by: Felix Schlepper <felix.schlepper@cern.ch>
1 parent 4679fed commit 66676ed

31 files changed

Lines changed: 3388 additions & 2665 deletions

‎Detectors/ITSMFT/ITS/tracking/CMakeLists.txt‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -13,6 +13,7 @@ o2_add_library(ITStracking
1313
TARGETVARNAME targetName
1414
SOURCES src/ClusterLines.cxx
1515
src/Cluster.cxx
16+
src/CapacityEstimator.cxx
1617
src/Configuration.cxx
1718
src/FastMultEstConfig.cxx
1819
src/FastMultEst.cxx

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

Lines changed: 192 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,192 @@
1+
// Copyright 2019-2020 CERN and copyright holders of ALICE O2.
2+
// See https://alice-o2.web.cern.ch/copyright for details of the copyright holders.
3+
// All rights not expressly granted are reserved.
4+
//
5+
// This software is distributed under the terms of the GNU General Public
6+
// License v3 (GPL Version 3), copied verbatim in the file "COPYING".
7+
//
8+
// In applying this license CERN does not waive the privileges and immunities
9+
// granted to it by virtue of its status as an Intergovernmental Organization
10+
// or submit itself to any jurisdiction.
11+
///
12+
/// \file LaunchGeometry.h
13+
/// \brief Compile-time launch geometry of the ITS tracking kernels, per GPU family.
14+
///
15+
16+
#ifndef ITSTRACKINGGPU_LAUNCHGEOMETRY_H_
17+
#define ITSTRACKINGGPU_LAUNCHGEOMETRY_H_
18+
19+
namespace o2::its::gpu
20+
{
21+
22+
#if defined(GPUCA_GPUTYPE_VEGA) // gfx906: MI50, Radeon VII
23+
constexpr int ComputeUnits = 60;
24+
constexpr int WarpSize = 64;
25+
#elif defined(GPUCA_GPUTYPE_MI100) // gfx908
26+
constexpr int ComputeUnits = 120;
27+
constexpr int WarpSize = 64;
28+
#elif defined(GPUCA_GPUTYPE_MI210) // gfx90a
29+
constexpr int ComputeUnits = 104;
30+
constexpr int WarpSize = 64;
31+
#elif defined(GPUCA_GPUTYPE_MI300) // gfx942: MI300X (MI300A has 228)
32+
constexpr int ComputeUnits = 304;
33+
constexpr int WarpSize = 64;
34+
#elif defined(GPUCA_GPUTYPE_RDNA) // gfx10xx/11xx consumer parts, wave32
35+
constexpr int ComputeUnits = 60;
36+
constexpr int WarpSize = 32;
37+
#elif defined(GPUCA_GPUTYPE_BLACKWELL) // sm_120: RTX 5080
38+
constexpr int ComputeUnits = 84;
39+
constexpr int WarpSize = 32;
40+
#elif defined(GPUCA_GPUTYPE_HOPPER) // sm_90: H100
41+
constexpr int ComputeUnits = 132;
42+
constexpr int WarpSize = 32;
43+
#elif defined(GPUCA_GPUTYPE_ADA) // sm_89: RTX 4090
44+
constexpr int ComputeUnits = 128;
45+
constexpr int WarpSize = 32;
46+
#elif defined(GPUCA_GPUTYPE_AMPERE) // sm_80/86: A100 has 108, RTX 3090 has 82
47+
constexpr int ComputeUnits = 108;
48+
constexpr int WarpSize = 32;
49+
#elif defined(GPUCA_GPUTYPE_TURING) // sm_75: RTX 2080 Ti
50+
constexpr int ComputeUnits = 68;
51+
constexpr int WarpSize = 32;
52+
#else
53+
#warning "GPU architecture not available setting fallback"
54+
constexpr int ComputeUnits = 60;
55+
constexpr int WarpSize = 64;
56+
#endif
57+
58+
constexpr int GPUThreads = 256;
59+
constexpr int DefaultBlocksPerComputeUnit = 4;
60+
constexpr int MaxBlocksPerComputeUnit = 10;
61+
62+
/// Minimum resident blocks per compute unit to request when no per-kernel measurement exists.
63+
struct KernelOccupancy {
64+
int computeLayerTracklets{1};
65+
int computeLayerCells{1};
66+
int computeLayerCellNeighbours{1};
67+
int processNeighboursCellSeed{1};
68+
int processNeighboursTrackSeed{1};
69+
int fitTrackSeeds{1};
70+
int fitTrackSeedsExtended{1};
71+
int compileLookupTable{1};
72+
73+
/// Return the smallest occupancy value in the table.
74+
constexpr int min() const
75+
{
76+
const int a{computeLayerTracklets < computeLayerCells ? computeLayerTracklets : computeLayerCells};
77+
const int b{computeLayerCellNeighbours < processNeighboursCellSeed ? computeLayerCellNeighbours : processNeighboursCellSeed};
78+
const int c{processNeighboursTrackSeed < fitTrackSeeds ? processNeighboursTrackSeed : fitTrackSeeds};
79+
const int d{fitTrackSeedsExtended < compileLookupTable ? fitTrackSeedsExtended : compileLookupTable};
80+
const int ab{a < b ? a : b};
81+
const int cd{c < d ? c : d};
82+
return ab < cd ? ab : cd;
83+
}
84+
85+
/// Return the largest occupancy value in the table.
86+
constexpr int max() const
87+
{
88+
const int a{computeLayerTracklets > computeLayerCells ? computeLayerTracklets : computeLayerCells};
89+
const int b{computeLayerCellNeighbours > processNeighboursCellSeed ? computeLayerCellNeighbours : processNeighboursCellSeed};
90+
const int c{processNeighboursTrackSeed > fitTrackSeeds ? processNeighboursTrackSeed : fitTrackSeeds};
91+
const int d{fitTrackSeedsExtended > compileLookupTable ? fitTrackSeedsExtended : compileLookupTable};
92+
const int ab{a > b ? a : b};
93+
const int cd{c > d ? c : d};
94+
return ab > cd ? ab : cd;
95+
}
96+
};
97+
98+
/// Use the same occupancy floor for every kernel when no per-kernel measurements are available.
99+
constexpr KernelOccupancy uniformOccupancy(int minBlocks)
100+
{
101+
return {.computeLayerTracklets = minBlocks,
102+
.computeLayerCells = minBlocks,
103+
.computeLayerCellNeighbours = minBlocks,
104+
.processNeighboursCellSeed = minBlocks,
105+
.processNeighboursTrackSeed = minBlocks,
106+
.fitTrackSeeds = minBlocks,
107+
.fitTrackSeedsExtended = minBlocks,
108+
.compileLookupTable = minBlocks};
109+
}
110+
111+
#if defined(GPUCA_GPUTYPE_VEGA) // gfx906: MI50, Radeon VII
112+
113+
/// Per-kernel minimum occupancy floors measured on gfx906.
114+
constexpr KernelOccupancy MinBlocks{
115+
.computeLayerTracklets = 2,
116+
.computeLayerCells = 3,
117+
.computeLayerCellNeighbours = 3,
118+
.processNeighboursCellSeed = 3,
119+
.processNeighboursTrackSeed = 3,
120+
.fitTrackSeeds = 4,
121+
.fitTrackSeedsExtended = 3, // untested: the follower is compiled out of every default iteration
122+
.compileLookupTable = 1,
123+
};
124+
125+
/// Number of blocks per CU used to size the grid for the measured gfx906 kernels.
126+
constexpr KernelOccupancy ResidentBlocks{
127+
.computeLayerTracklets = 4, // 56 VGPR
128+
.computeLayerCells = 3, // 84 VGPR
129+
.computeLayerCellNeighbours = 3, // 84 VGPR
130+
.processNeighboursCellSeed = 3, // 84 VGPR
131+
.processNeighboursTrackSeed = 3, // 84 VGPR
132+
.fitTrackSeeds = 4, // 64 VGPR
133+
.fitTrackSeedsExtended = 3, // 84 VGPR
134+
.compileLookupTable = 4, // 8 VGPR, could hold ten; left at the historical grid
135+
};
136+
137+
#elif defined(__HIPCC__) || defined(__HIP_PLATFORM_AMD__)
138+
/// Other AMD parts: unmeasured.
139+
constexpr KernelOccupancy MinBlocks = uniformOccupancy(3);
140+
constexpr KernelOccupancy ResidentBlocks = uniformOccupancy(DefaultBlocksPerComputeUnit);
141+
#else
142+
/// NVIDIA: unmeasured.
143+
constexpr KernelOccupancy MinBlocks = uniformOccupancy(1);
144+
constexpr KernelOccupancy ResidentBlocks = uniformOccupancy(DefaultBlocksPerComputeUnit);
145+
#endif
146+
147+
/// Number of blocks in a grid whose depth is residentBlocksPerComputeUnit blocks per CU.
148+
constexpr int gridBlocks(int residentBlocksPerComputeUnit)
149+
{
150+
return ComputeUnits * residentBlocksPerComputeUnit;
151+
}
152+
153+
/// Number of threads covered by a grid whose depth is residentBlocksPerComputeUnit blocks per CU.
154+
constexpr int gridThreads(int residentBlocksPerComputeUnit)
155+
{
156+
return gridBlocks(residentBlocksPerComputeUnit) * GPUThreads;
157+
}
158+
159+
static_assert(MinBlocks.min() >= 1,
160+
"an occupancy floor below one resident block is meaningless");
161+
162+
static_assert(MinBlocks.max() <= MaxBlocksPerComputeUnit,
163+
"the occupancy floor cannot exceed the blocks a CU can hold");
164+
165+
static_assert(ResidentBlocks.min() >= 1,
166+
"every kernel must have at least one resident block per CU");
167+
168+
static_assert(ResidentBlocks.max() <= MaxBlocksPerComputeUnit,
169+
"resident blocks per CU cannot exceed what a CU can hold");
170+
171+
/// The grid must provide at least as many blocks per CU as the corresponding occupancy floor.
172+
constexpr bool residentCoversFloor()
173+
{
174+
return ResidentBlocks.computeLayerTracklets >= MinBlocks.computeLayerTracklets &&
175+
ResidentBlocks.computeLayerCells >= MinBlocks.computeLayerCells &&
176+
ResidentBlocks.computeLayerCellNeighbours >= MinBlocks.computeLayerCellNeighbours &&
177+
ResidentBlocks.processNeighboursCellSeed >= MinBlocks.processNeighboursCellSeed &&
178+
ResidentBlocks.processNeighboursTrackSeed >= MinBlocks.processNeighboursTrackSeed &&
179+
ResidentBlocks.fitTrackSeeds >= MinBlocks.fitTrackSeeds &&
180+
ResidentBlocks.fitTrackSeedsExtended >= MinBlocks.fitTrackSeedsExtended &&
181+
ResidentBlocks.compileLookupTable >= MinBlocks.compileLookupTable;
182+
}
183+
184+
static_assert(residentCoversFloor(), "a kernel's grid is narrower than the occupancy its __launch_bounds__ floor demands");
185+
186+
static_assert(GPUThreads % WarpSize == 0, "block size must be a whole number of warps/waves");
187+
188+
static_assert(ComputeUnits > 0 && GPUThreads > 0 && DefaultBlocksPerComputeUnit > 0, "degenerate launch geometry");
189+
190+
} // namespace o2::its::gpu
191+
192+
#endif // ITSTRACKINGGPU_LAUNCHGEOMETRY_H_

0 commit comments

Comments
 (0)