Skip to content

Commit 2ed2a8c

Browse files
f3schf3sch
authored
ITS: slab allocator idea (#15678)
* ITS: slab allocator with size estimator Signed-off-by: Felix Schlepper <felix.schlepper@cern.ch> * ITSGPU: count possible cells before Signed-off-by: Felix Schlepper <felix.schlepper@cern.ch> * some alignment Signed-off-by: Felix Schlepper <felix.schlepper@cern.ch> * call fast path & reduce more flots Signed-off-by: Felix Schlepper <felix.schlepper@cern.ch> --------- Signed-off-by: Felix Schlepper <felix.schlepper@cern.ch>
1 parent 0d7617d commit 2ed2a8c

32 files changed

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

0 commit comments

Comments
 (0)