Skip to content

Commit a9ae89e

Browse files
asubahasubah
committed
[cuda] Autotuning script for CUDA kernels
1 parent 129aff6 commit a9ae89e

35 files changed

Lines changed: 327 additions & 0 deletions

‎autotuning/README.md‎

Lines changed: 85 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,85 @@
1+
# Autotuning Script
2+
_Requires Python 3.7 or Higher_
3+
4+
A python script that can be used to autotune parameters in the pixeltrack-standalone project. The parameters can be anything as long as there is a n interface in the program that can read the parameters from the configurations folder. Hence, an interface was created to autotune CUDA kernel launch configurations
5+
6+
[`src/cuda/CUDACore/ExecutionConfiguration.h`](src/cuda/CUDACore/ExecutionConfiguration.h)
7+
8+
**Example:**
9+
```c++
10+
#include "CUDACore/ExecutionConfiguration.h"
11+
12+
cms::cuda::ExecutionConfiguration exec;
13+
int threadsPerBlock = exec.configFromFile("RawToDigi_kernel");
14+
15+
int threadsPerBlock = exec.configFromFile("calibDigis");// 256;
16+
int blocks = (std::max(int(wordCounter), int(gpuClustering::MaxNumModules)) + threadsPerBlock - 1) / threadsPerBlock;
17+
18+
gpuCalibPixel::calibDigis<<<blocks, threadsPerBlock, 0, stream>>>(...)
19+
...
20+
```
21+
22+
Other parameters can be tuned also such as number of CPU threads and GPU streams.
23+
24+
## Usage
25+
26+
```bash
27+
usage: tuner.py [-h] -p PROCESS [-c CONFIGURATIONS] [-t TUNABLES] [-r RESULTS] [--validation [VALIDATION]]
28+
[--preheat [PREHEAT]] [--verbose]
29+
30+
Autotuner script.
31+
32+
options:
33+
-h, --help show this help message and exit
34+
-p PROCESS, --process PROCESS
35+
path to the program to be autotuned
36+
-c CONFIGURATIONS, --configurations CONFIGURATIONS
37+
path to save the configurations for the tunable process to read them. Default =
38+
autotuning/kernel_configs/
39+
-t TUNABLES, --tunables TUNABLES
40+
csv file that contains the tunable parameters. Default = autotuning/tunables.csv
41+
-r RESULTS, --results RESULTS
42+
results file, if it exists, results will be appended to the end, if not, it will be created.
43+
Default = autotuning/results.txt
44+
--validation [VALIDATION]
45+
validate the output of the process. Default = False
46+
--preheat [PREHEAT] run the process multiple times before autotuning. Default = False
47+
--verbose, -v
48+
```
49+
50+
Simple run from the project root path:
51+
```bash
52+
python autotuning/tuner.py --process ./cuda
53+
```
54+
55+
You need to provide a `csv` file that contains the tunable parameters as follows:
56+
```
57+
prarmeter_name,lower_bound,upped_bound,incremental_step
58+
```
59+
60+
Example file can be found here:
61+
62+
[`autotuning/tunables.csv`](autotuning/tunables.csv)
63+
64+
Example from results file:
65+
66+
```bash
67+
2022-10-04 14:01:55.184845 1282.89 7.794910e+00 59.5% OK PASSED (128, 128, 64, 128, 128, 64, 64, 128, 128, 64, 256, 256, 128, 64, 64, 128, 256, 128, 192, 256, 256, 128, 64, 256, 192, 64, 192, 128, 64, 128, 256, 12, 4)
68+
2022-10-04 14:02:07.028675 1281.11 7.805726e+00 60.8% OK PASSED (128, 128, 64, 128, 128, 64, 64, 128, 128, 64, 256, 256, 128, 64, 64, 128, 256, 128, 192, 256, 256, 128, 64, 256, 192, 128, 256, 128, 192, 128, 256, 12, 4)
69+
2022-10-04 14:02:23.100261 1279.51 7.815468e+00 60.6% OK PASSED (128, 128, 64, 128, 128, 64, 64, 128, 128, 64, 256, 256, 128, 64, 64, 128, 256, 128, 192, 256, 256, 128, 64, 256, 192, 128, 256, 64, 64, 192, 256, 12, 4)
70+
2022-10-04 14:02:35.089202 1272.24 7.860135e+00 58.4% OK PASSED (128, 128, 64, 128, 128, 64, 64, 128, 128, 64, 256, 256, 128, 64, 64, 128, 256, 128, 192, 256, 256, 128, 64, 256, 192, 64, 192, 256, 64, 192, 192, 12, 4)
71+
2022-10-04 15:09:38.377397 1293.88 7.728671e+00 59.8% OK PASSED (128, 128, 64, 128, 128, 64, 64, 128, 128, 64, 256, 256, 128, 64, 64, 128, 256, 128, 192, 256, 256, 128, 64, 256, 192, 192, 256, 128, 64, 256, 192, 12, 4)
72+
2022-10-04 15:10:02.501392 1284.46 7.785401e+00 60.2% OK PASSED (128, 128, 64, 128, 128, 64, 64, 128, 128, 64, 256, 256, 128, 64, 64, 128, 256, 128, 192, 256, 256, 128, 64, 256, 192, 64, 64, 256, 256, 128, 64, 12, 4)
73+
2022-10-04 15:10:14.319484 1280.77 7.807828e+00 60.1% OK PASSED (128, 128, 64, 128, 128, 64, 64, 128, 128, 64, 256, 256, 128, 64, 64, 128, 256, 128, 192, 256, 256, 128, 64, 256, 192, 256, 256, 256, 192, 128, 128, 12, 4)
74+
2022-10-04 15:20:05.433228 1289.93 7.752346e+00 60.2% OK PASSED (128, 128, 64, 128, 128, 64, 64, 128, 128, 64, 256, 256, 128, 64, 64, 128, 256, 128, 192, 256, 256, 128, 64, 256, 192, 64, 128, 192, 192, 256, 64, 12, 4)
75+
2022-10-04 15:20:17.282590 1287.91 7.764510e+00 60.8% OK PASSED (128, 128, 64, 128, 128, 64, 64, 128, 128, 64, 256, 256, 128, 64, 64, 128, 256, 128, 192, 256, 256, 128, 64, 256, 192, 192, 64, 192, 256, 192, 128, 12, 4)
76+
```
77+
The columns are:
78+
- Date
79+
- Time
80+
- Throughput (Parsed from output)
81+
- Runtime (Parsed from output)
82+
- CPU usage per thread (Parsed from output)
83+
- Tunning process return (OK/ERROR)
84+
- Valudation process return (PASSED/FAILED/ERROR)
85+
- Configuration

‎autotuning/kernel_configs/RawToDigi_kernel‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1 @@
1+
128

‎autotuning/kernel_configs/calibDigis‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1 @@
1+
192

‎autotuning/kernel_configs/clusterChargeCut‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1 @@
1+
128

‎autotuning/kernel_configs/countModules‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1 @@
1+
256

‎autotuning/kernel_configs/finalizeBulk‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1 @@
1+
64

‎autotuning/kernel_configs/findClus‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1 @@
1+
128

‎autotuning/kernel_configs/fishbone‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1 @@
1+
128

‎autotuning/kernel_configs/getHits‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1 @@
1+
64

‎autotuning/kernel_configs/initDoublets‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1 @@
1+
128

0 commit comments

Comments
 (0)