Skip to content

Commit 129aff6

Browse files
asubahasubah
committed
[cuda] Implement an autotuning interface for CUDA kernels
1 parent 7b601d4 commit 129aff6

7 files changed

Lines changed: 152 additions & 44 deletions

File tree

‎src/cuda/CUDACore/ExecutionConfiguration.h‎

Lines changed: 42 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,42 @@
1+
#ifndef HeterogeneousCore_CUDAUtilities_interface_ExecutionConfiguration_h
2+
#define HeterogeneousCore_CUDAUtilities_interface_ExecutionConfiguration_h
3+
4+
#include <fstream>
5+
6+
namespace cms {
7+
namespace cuda {
8+
9+
class ExecutionConfiguration {
10+
public:
11+
const std::string CONFIG_PATH = "autotuning/kernel_configs/";
12+
13+
ExecutionConfiguration(){};
14+
15+
template <typename T>
16+
void cudaOccCalc(T kernel, int* blockSize, size_t dynamicSMemSize = 0, int blockSizeLimit = 0) {
17+
int minGridSize = 0;
18+
cudaOccupancyMaxPotentialBlockSize(&minGridSize, blockSize, kernel, dynamicSMemSize, blockSizeLimit);
19+
}
20+
21+
size_t configFromFile(std::string filename) {
22+
size_t blockSize;
23+
24+
std::fstream file;
25+
file.open(CONFIG_PATH + filename, std::ios::in);
26+
if (file.is_open())
27+
{
28+
file >> blockSize;
29+
file.close();
30+
} else {
31+
std::cout << "Error in opening file " + filename + "\n";
32+
}
33+
// std::cout << "Filename = " << filename << " Blocksize = " << blockSize << "\n";
34+
35+
return blockSize;
36+
}
37+
};
38+
39+
} // namespace cuda
40+
} // namespace cms
41+
42+
#endif // HeterogeneousCore_CUDAUtilities_interface_ExecutionConfiguration_h

‎src/cuda/CUDACore/HistoContainer.h‎

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -14,6 +14,7 @@
1414
#include "CUDACore/cuda_assert.h"
1515
#include "CUDACore/cudastdAlgorithm.h"
1616
#include "CUDACore/prefixScan.h"
17+
#include "CUDACore/ExecutionConfiguration.h"
1718

1819
namespace cms {
1920
namespace cuda {
@@ -77,6 +78,7 @@ namespace cms {
7778
#ifdef __CUDACC__
7879
uint32_t *poff = (uint32_t *)((char *)(h) + offsetof(Histo, off));
7980
int32_t *ppsws = (int32_t *)((char *)(h) + offsetof(Histo, psws));
81+
cms::cuda::ExecutionConfiguration exec;
8082
auto nthreads = 1024;
8183
auto nblocks = (Histo::totbins() + nthreads - 1) / nthreads;
8284
multiBlockPrefixScan<<<nblocks, nthreads, sizeof(int32_t) * nblocks, stream>>>(

‎src/cuda/plugin-PixelTriplets/BrokenLineFitOnGPU.cu‎

Lines changed: 26 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -7,9 +7,6 @@ void HelixFitOnGPU::launchBrokenLineKernels(HitsView const *hv,
77
cudaStream_t stream) {
88
assert(tuples_d);
99

10-
auto blockSize = 64;
11-
auto numberOfBlocks = (maxNumberOfConcurrentFits_ + blockSize - 1) / blockSize;
12-
1310
// Fit internals
1411
auto hitsGPU_ = cms::cuda::make_device_unique<double[]>(
1512
maxNumberOfConcurrentFits_ * sizeof(Rfit::Matrix3xNd<4>) / sizeof(double), stream);
@@ -18,13 +15,31 @@ void HelixFitOnGPU::launchBrokenLineKernels(HitsView const *hv,
1815
auto fast_fit_resultsGPU_ = cms::cuda::make_device_unique<double[]>(
1916
maxNumberOfConcurrentFits_ * sizeof(Rfit::Vector4d) / sizeof(double), stream);
2017

18+
cms::cuda::ExecutionConfiguration exec;
19+
auto blockSize_ff3 = exec.configFromFile("kernelFastFit3");
20+
auto numberOfBlocks_ff3 = (maxNumberOfConcurrentFits_ + blockSize_ff3 - 1) / blockSize_ff3;
21+
22+
auto blockSize_blf3 = exec.configFromFile("kernelLineFit3");
23+
auto numberOfBlocks_blf3 = (maxNumberOfConcurrentFits_ + blockSize_blf3 - 1) / blockSize_blf3;
24+
25+
auto blockSize_ff4 = exec.configFromFile("kernelFastFit4");
26+
auto numberOfBlocks_ff4 = (maxNumberOfConcurrentFits_ + blockSize_ff4 - 1) / blockSize_ff4;
27+
28+
auto blockSize_blf4 = exec.configFromFile("kernelLineFit4");
29+
auto numberOfBlocks_blf4 = (maxNumberOfConcurrentFits_ + blockSize_blf4 - 1) / blockSize_blf4;
30+
31+
auto blockSize_ff5 = exec.configFromFile("kernelFastFit5");
32+
auto numberOfBlocks_ff5 = (maxNumberOfConcurrentFits_ + blockSize_ff5 - 1) / blockSize_ff5;
33+
34+
auto blockSize_blf5 = exec.configFromFile("kernelLineFit5");
35+
auto numberOfBlocks_blf5 = (maxNumberOfConcurrentFits_ + blockSize_blf5 - 1) / blockSize_blf5;
2136
for (uint32_t offset = 0; offset < maxNumberOfTuples; offset += maxNumberOfConcurrentFits_) {
2237
// fit triplets
23-
kernelBLFastFit<3><<<numberOfBlocks, blockSize, 0, stream>>>(
38+
kernelBLFastFit<3><<<numberOfBlocks_ff3, blockSize_ff3, 0, stream>>>(
2439
tuples_d, tupleMultiplicity_d, hv, hitsGPU_.get(), hits_geGPU_.get(), fast_fit_resultsGPU_.get(), 3, offset);
2540
cudaCheck(cudaGetLastError());
2641

27-
kernelBLFit<3><<<numberOfBlocks, blockSize, 0, stream>>>(tupleMultiplicity_d,
42+
kernelBLFit<3><<<numberOfBlocks_blf3, blockSize_blf3, 0, stream>>>(tupleMultiplicity_d,
2843
bField_,
2944
outputSoa_d,
3045
hitsGPU_.get(),
@@ -35,11 +50,11 @@ void HelixFitOnGPU::launchBrokenLineKernels(HitsView const *hv,
3550
cudaCheck(cudaGetLastError());
3651

3752
// fit quads
38-
kernelBLFastFit<4><<<numberOfBlocks / 4, blockSize, 0, stream>>>(
53+
kernelBLFastFit<4><<<numberOfBlocks_ff4 / 4, blockSize_ff4, 0, stream>>>(
3954
tuples_d, tupleMultiplicity_d, hv, hitsGPU_.get(), hits_geGPU_.get(), fast_fit_resultsGPU_.get(), 4, offset);
4055
cudaCheck(cudaGetLastError());
4156

42-
kernelBLFit<4><<<numberOfBlocks / 4, blockSize, 0, stream>>>(tupleMultiplicity_d,
57+
kernelBLFit<4><<<numberOfBlocks_blf4 / 4, blockSize_blf4, 0, stream>>>(tupleMultiplicity_d,
4358
bField_,
4459
outputSoa_d,
4560
hitsGPU_.get(),
@@ -51,11 +66,11 @@ void HelixFitOnGPU::launchBrokenLineKernels(HitsView const *hv,
5166

5267
if (fit5as4_) {
5368
// fit penta (only first 4)
54-
kernelBLFastFit<4><<<numberOfBlocks / 4, blockSize, 0, stream>>>(
69+
kernelBLFastFit<4><<<numberOfBlocks_ff4 / 4, blockSize_ff4, 0, stream>>>(
5570
tuples_d, tupleMultiplicity_d, hv, hitsGPU_.get(), hits_geGPU_.get(), fast_fit_resultsGPU_.get(), 5, offset);
5671
cudaCheck(cudaGetLastError());
5772

58-
kernelBLFit<4><<<numberOfBlocks / 4, blockSize, 0, stream>>>(tupleMultiplicity_d,
73+
kernelBLFit<4><<<numberOfBlocks_blf4 / 4, blockSize_blf4, 0, stream>>>(tupleMultiplicity_d,
5974
bField_,
6075
outputSoa_d,
6176
hitsGPU_.get(),
@@ -66,11 +81,11 @@ void HelixFitOnGPU::launchBrokenLineKernels(HitsView const *hv,
6681
cudaCheck(cudaGetLastError());
6782
} else {
6883
// fit penta (all 5)
69-
kernelBLFastFit<5><<<numberOfBlocks / 4, blockSize, 0, stream>>>(
84+
kernelBLFastFit<5><<<numberOfBlocks_ff5 / 4, blockSize_ff5, 0, stream>>>(
7085
tuples_d, tupleMultiplicity_d, hv, hitsGPU_.get(), hits_geGPU_.get(), fast_fit_resultsGPU_.get(), 5, offset);
7186
cudaCheck(cudaGetLastError());
7287

73-
kernelBLFit<5><<<numberOfBlocks / 4, blockSize, 0, stream>>>(tupleMultiplicity_d,
88+
kernelBLFit<5><<<numberOfBlocks_blf5 / 4, blockSize_blf5, 0, stream>>>(tupleMultiplicity_d,
7489
bField_,
7590
outputSoa_d,
7691
hitsGPU_.get(),

‎src/cuda/plugin-PixelTriplets/CAHitNtupletGeneratorKernels.cu‎

Lines changed: 30 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -1,8 +1,10 @@
11
#include "CAHitNtupletGeneratorKernelsImpl.h"
2+
#include "CUDACore/ExecutionConfiguration.h"
23

34
template <>
45
void CAHitNtupletGeneratorKernelsGPU::fillHitDetIndices(HitsView const *hv, TkSoA *tracks_d, cudaStream_t cudaStream) {
5-
auto blockSize = 128;
6+
cms::cuda::ExecutionConfiguration exec;
7+
auto blockSize = exec.configFromFile("kernel_fillHitDetIndices");
68
auto numberOfBlocks = (HitContainer::capacity() + blockSize - 1) / blockSize;
79

810
kernel_fillHitDetIndices<<<numberOfBlocks, blockSize, 0, cudaStream>>>(
@@ -33,7 +35,8 @@ void CAHitNtupletGeneratorKernelsGPU::launchKernels(HitsOnCPU const &hh, TkSoA *
3335
// applying conbinatoric cleaning such as fishbone at this stage is too expensive
3436
//
3537

36-
auto nthTot = 64;
38+
cms::cuda::ExecutionConfiguration exec;
39+
auto nthTot = exec.configFromFile("kernel_connect");
3740
auto stride = 4;
3841
auto blockSize = nthTot / stride;
3942
auto numberOfBlocks = (3 * m_params.maxNumberOfDoublets_ / 4 + blockSize - 1) / blockSize;
@@ -62,7 +65,7 @@ void CAHitNtupletGeneratorKernelsGPU::launchKernels(HitsOnCPU const &hh, TkSoA *
6265
cudaCheck(cudaGetLastError());
6366

6467
if (nhits > 1 && m_params.earlyFishbone_) {
65-
auto nthTot = 128;
68+
auto nthTot = exec.configFromFile("fishbone");
6669
auto stride = 16;
6770
auto blockSize = nthTot / stride;
6871
auto numberOfBlocks = (nhits + blockSize - 1) / blockSize;
@@ -73,7 +76,7 @@ void CAHitNtupletGeneratorKernelsGPU::launchKernels(HitsOnCPU const &hh, TkSoA *
7376
cudaCheck(cudaGetLastError());
7477
}
7578

76-
blockSize = 64;
79+
blockSize = exec.configFromFile("kernel_find_ntuplets");
7780
numberOfBlocks = (3 * m_params.maxNumberOfDoublets_ / 4 + blockSize - 1) / blockSize;
7881
kernel_find_ntuplets<<<numberOfBlocks, blockSize, 0, cudaStream>>>(hh.view(),
7982
device_theCells_.get(),
@@ -85,36 +88,41 @@ void CAHitNtupletGeneratorKernelsGPU::launchKernels(HitsOnCPU const &hh, TkSoA *
8588
m_params.minHitsPerNtuplet_);
8689
cudaCheck(cudaGetLastError());
8790

88-
if (m_params.doStats_)
91+
if (m_params.doStats_) {
92+
blockSize = exec.configFromFile("kernel_mark_used");
93+
numberOfBlocks = (3 * m_params.maxNumberOfDoublets_ / 4 + blockSize - 1) / blockSize;
8994
kernel_mark_used<<<numberOfBlocks, blockSize, 0, cudaStream>>>(hh.view(), device_theCells_.get(), device_nCells_);
90-
cudaCheck(cudaGetLastError());
91-
95+
cudaCheck(cudaGetLastError());
96+
}
9297
#ifdef GPU_DEBUG
9398
cudaDeviceSynchronize();
9499
cudaCheck(cudaGetLastError());
95100
#endif
96101

97-
blockSize = 128;
102+
blockSize = exec.configFromFile("finalizeBulk");
98103
numberOfBlocks = (HitContainer::totbins() + blockSize - 1) / blockSize;
99104
cms::cuda::finalizeBulk<<<numberOfBlocks, blockSize, 0, cudaStream>>>(device_hitTuple_apc_, tuples_d);
100105

101106
// remove duplicates (tracks that share a doublet)
107+
blockSize = exec.configFromFile("kernel_earlyDuplicateRemover");
102108
numberOfBlocks = (3 * m_params.maxNumberOfDoublets_ / 4 + blockSize - 1) / blockSize;
103109
kernel_earlyDuplicateRemover<<<numberOfBlocks, blockSize, 0, cudaStream>>>(
104110
device_theCells_.get(), device_nCells_, tuples_d, quality_d);
105111
cudaCheck(cudaGetLastError());
106112

107-
blockSize = 128;
113+
blockSize = exec.configFromFile("kernel_countMultiplicity");
108114
numberOfBlocks = (3 * CAConstants::maxTuples() / 4 + blockSize - 1) / blockSize;
109115
kernel_countMultiplicity<<<numberOfBlocks, blockSize, 0, cudaStream>>>(
110116
tuples_d, quality_d, device_tupleMultiplicity_.get());
111117
cms::cuda::launchFinalize(device_tupleMultiplicity_.get(), cudaStream);
118+
blockSize = exec.configFromFile("kernel_fillMultiplicity");
119+
numberOfBlocks = (3 * CAConstants::maxTuples() / 4 + blockSize - 1) / blockSize;
112120
kernel_fillMultiplicity<<<numberOfBlocks, blockSize, 0, cudaStream>>>(
113121
tuples_d, quality_d, device_tupleMultiplicity_.get());
114122
cudaCheck(cudaGetLastError());
115123

116124
if (nhits > 1 && m_params.lateFishbone_) {
117-
auto nthTot = 128;
125+
auto nthTot = exec.configFromFile("fishbone");
118126
auto stride = 16;
119127
auto blockSize = nthTot / stride;
120128
auto numberOfBlocks = (nhits + blockSize - 1) / blockSize;
@@ -126,6 +134,7 @@ void CAHitNtupletGeneratorKernelsGPU::launchKernels(HitsOnCPU const &hh, TkSoA *
126134
}
127135

128136
if (m_params.doStats_) {
137+
blockSize = exec.configFromFile("kernel_checkOverflows");
129138
numberOfBlocks = (std::max(nhits, m_params.maxNumberOfDoublets_) + blockSize - 1) / blockSize;
130139
kernel_checkOverflows<<<numberOfBlocks, blockSize, 0, cudaStream>>>(tuples_d,
131140
device_tupleMultiplicity_.get(),
@@ -151,6 +160,7 @@ void CAHitNtupletGeneratorKernelsGPU::launchKernels(HitsOnCPU const &hh, TkSoA *
151160

152161
template <>
153162
void CAHitNtupletGeneratorKernelsGPU::buildDoublets(HitsOnCPU const &hh, cudaStream_t stream) {
163+
cms::cuda::ExecutionConfiguration exec;
154164
auto nhits = hh.nHits();
155165

156166
#ifdef NTUPLE_DEBUG
@@ -176,7 +186,7 @@ void CAHitNtupletGeneratorKernelsGPU::buildDoublets(HitsOnCPU const &hh, cudaStr
176186
CAConstants::maxNumOfActiveDoublets() * sizeof(GPUCACell::CellNeighbors));
177187

178188
{
179-
int threadsPerBlock = 128;
189+
int threadsPerBlock = exec.configFromFile("initDoublets");
180190
// at least one block!
181191
int blocks = (std::max(1U, nhits) + threadsPerBlock - 1) / threadsPerBlock;
182192
gpuPixelDoublets::initDoublets<<<blocks, threadsPerBlock, 0, stream>>>(device_isOuterHitOfCell_.get(),
@@ -238,40 +248,46 @@ void CAHitNtupletGeneratorKernelsGPU::classifyTuples(HitsOnCPU const &hh, TkSoA
238248
auto const *tuples_d = &tracks_d->hitIndices;
239249
auto *quality_d = (Quality *)(&tracks_d->m_quality);
240250

241-
auto blockSize = 64;
242-
243251
// classify tracks based on kinematics
252+
cms::cuda::ExecutionConfiguration exec;
253+
auto blockSize = exec.configFromFile("kernel_classifyTracks");
244254
auto numberOfBlocks = (3 * CAConstants::maxNumberOfQuadruplets() / 4 + blockSize - 1) / blockSize;
245255
kernel_classifyTracks<<<numberOfBlocks, blockSize, 0, cudaStream>>>(tuples_d, tracks_d, m_params.cuts_, quality_d);
246256
cudaCheck(cudaGetLastError());
247257

248258
if (m_params.lateFishbone_) {
249259
// apply fishbone cleaning to good tracks
260+
blockSize = exec.configFromFile("kernel_fishboneCleaner");
250261
numberOfBlocks = (3 * m_params.maxNumberOfDoublets_ / 4 + blockSize - 1) / blockSize;
251262
kernel_fishboneCleaner<<<numberOfBlocks, blockSize, 0, cudaStream>>>(
252263
device_theCells_.get(), device_nCells_, quality_d);
253264
cudaCheck(cudaGetLastError());
254265
}
255266

256267
// remove duplicates (tracks that share a doublet)
268+
blockSize = exec.configFromFile("kernel_fastDuplicateRemover");
257269
numberOfBlocks = (3 * m_params.maxNumberOfDoublets_ / 4 + blockSize - 1) / blockSize;
258270
kernel_fastDuplicateRemover<<<numberOfBlocks, blockSize, 0, cudaStream>>>(
259271
device_theCells_.get(), device_nCells_, tuples_d, tracks_d);
260272
cudaCheck(cudaGetLastError());
261273

262274
if (m_params.minHitsPerNtuplet_ < 4 || m_params.doStats_) {
263275
// fill hit->track "map"
276+
blockSize = exec.configFromFile("kernel_countHitInTracks");
264277
numberOfBlocks = (3 * CAConstants::maxNumberOfQuadruplets() / 4 + blockSize - 1) / blockSize;
265278
kernel_countHitInTracks<<<numberOfBlocks, blockSize, 0, cudaStream>>>(
266279
tuples_d, quality_d, device_hitToTuple_.get());
267280
cudaCheck(cudaGetLastError());
268281
cms::cuda::launchFinalize(device_hitToTuple_.get(), cudaStream);
269282
cudaCheck(cudaGetLastError());
283+
blockSize = exec.configFromFile("kernel_fillHitInTracks");
284+
numberOfBlocks = (3 * CAConstants::maxNumberOfQuadruplets() / 4 + blockSize - 1) / blockSize;
270285
kernel_fillHitInTracks<<<numberOfBlocks, blockSize, 0, cudaStream>>>(tuples_d, quality_d, device_hitToTuple_.get());
271286
cudaCheck(cudaGetLastError());
272287
}
273288
if (m_params.minHitsPerNtuplet_ < 4) {
274289
// remove duplicates (tracks that share a hit)
290+
blockSize = exec.configFromFile("kernel_tripletCleaner");
275291
numberOfBlocks = (HitToTuple::capacity() + blockSize - 1) / blockSize;
276292
kernel_tripletCleaner<<<numberOfBlocks, blockSize, 0, cudaStream>>>(
277293
hh.view(), tuples_d, tracks_d, quality_d, device_hitToTuple_.get());
@@ -280,6 +296,7 @@ void CAHitNtupletGeneratorKernelsGPU::classifyTuples(HitsOnCPU const &hh, TkSoA
280296

281297
if (m_params.doStats_) {
282298
// counters (add flag???)
299+
blockSize = 64;
283300
numberOfBlocks = (HitToTuple::capacity() + blockSize - 1) / blockSize;
284301
kernel_doStatsForHitInTracks<<<numberOfBlocks, blockSize, 0, cudaStream>>>(device_hitToTuple_.get(), counters_);
285302
cudaCheck(cudaGetLastError());

0 commit comments

Comments
 (0)