Dachsbacher, C.

Karcher, T.

Tichy, W. F.

Tillmann, M.

KITopen

Application-independentAutotuning for GPUsMartin TILLMANN, Thomas KARCHER, Carsten DACHSBACHER,Walter F. TICHYKarlsruhe Institute of TechnologyAm Fasanengarten 5, Karlsruhe, GermanyAbstract. Autotuning is an established technique for adjusting perfor-mance-critical parameters of applications to their specific run-time envi-ronment. In this paper, we investigate the potential of online autotuningfor general purpose computation on GPUs. Our application-independentautotuner AtuneRT optimizes GPU-specific parameters such as blocksize and loop-unrolling degree. We also discuss the peculiarities of auto-tuning on GPUs. We demonstrate tuning potential using CUDA and byinstrumenting the parallel algorithms library Thrust. We evaluate ouronline autotuning approach with various GPUs and sample applications.Keywords. Autotuning, GPU, GPGPU, CUDA, Thrust.1. IntroductionAutotuning has emerged as an effective and automated technique for optimiz-ing parallel applications. Most work in the autotuning area has concentrated onapplication specific tuning parameters such as the replication factor for tasks orthe number of threads in data parallel operations. On graphic processing units(GPUs), there are various tunable parameters that are application independent.These parameters include the number of warps combined into blocks and the de-gree of loop-unrolling. These parameters must be specified in every GPU call. Thispaper shows how an automatic tuner can optimize block size and loop unrollingand presents the performance gains of four examples on three current GPUs.2. Online AutotuningAutotuning is an automatic optimization technique that adjusts tuning param-eters of an application or its execution environment, with the aim of optimizingsome criterion such as execution time, memory usage or energy consumption. Anonline autotuner operates while the application is running and can therefore ad-just to varying input or environment characteristics. An autotuner drives a feed-back loop that (1) measures one or more sections of a program, and (2) adjuststhe tuning parameter configuration. A search-based algorithm such as Nelder-Mead [1] guides the process. An autotuner may walk into local minima and settleon a suboptimal configuration.Many existing autotuners use application-specific knowledge about tuning pa-rameters to speed up the search. ATLAS [2] exploits specific knowledge about ma-trix sizes and shapes, MATE [3] tunes MPI-specific parameters, and PetaBricks [4]chooses among alternative algorithms. Application-independent autotuners suchas Active Harmony [5] and our own autotuner [6] refrain from using such knowl-edge in exchange for generality.In this work, we adapt our general-purpose run-time autotuner AtuneRT tooptimize GPU kernels. AtuneRT is the successor of Atune [7], an off-line auto-tuner. The following API calls of AtuneRT make applications autotunable:• addParameter(&param, range, default) informs the autotuner about thetunable variable param. The autotuner is allowed to vary param withinthe range given, starting with default. The variable could be, for instance,the number of threads to spawn.• startMeasurement() defines the beginning of a measurement section.Whenever control flow passes this call, the autotuner initializes measure-ment.• stopMeasurement() defines the end of a measurement section. The auto-tuner stops the measurement and calculates a new tuning parameter con-figuration.3. Autotuning on the GPUModern graphics processors provide massive parallelism for a wide range of algo-rithms. To fully utilize GPUs it is important to optimize memory access patterns,balance workloads, and minimize control-flow costs. Optimizing on the GPU is anon-trivial task, given the peculiarities of these multi-threading architectures [8].3.1. Kernel launch configurationIn this work, we employ NVIDIA’s CUDA architecture. CUDA organizes threadsin a hierarchy. A group of 32 threads form a warp, and several warps form ablock. And finally, blocks are organized in a grid. Figure 1 illustrates the threadhierarchy in CUDA.The overall thread number for a given GPU kernel is usually dictated by thenumber of data elements the kernel operates on. However the number of threadsper block, called the block size, is variable and can be chosen at run-time.GPU resources such as temporary registers and shared memory are parti-tioned among the warps and blocks. Due to limited resources available on theGPU not all warps can be active at once.Occupancy is the ratio of active warps to the total number of warps that canrun in parallel. Obviously, it is a good idea to keep as many blocks operating atall times as possible. An important performance aspect is the trade-off betweenoccupancy and the workload per thread. As stated in the ”CUDA C Best PracticesGuide” [9] higher occupancy might not always result in higher performance. ItFigure 1. CUDA Thread Model. Image courtesy of NVIDIA.can be beneficial to run applications at lower occupancy but increased workloadper thread [10].CUDA also allows explicit loop unrolling. The unrolling factor can be specifiedand it is the programmer’s responsibility to ensure best performance.Block size, per-thread workload, and degree of loop unrolling are tuning pa-rameters that apply to any parallel application on CUDA. These parameters arecollectively known as the kernel launch configuration.The optimal values of these parameters are hard to determine as they notonly depend on the algorithm but also the underlying architecture of the respec-tive GPU and the characteristics of the input data. The performance impact ofthese parameters, in particular block size, is well-documented [11]. When tun-ing optimization parameters by hand, a multitude of restrictions and hardwareidiosyncrasies have to be considered. Even though simple heuristics that maxi-mize occupancy exists, the standard practice of finding the optimal kernel launchconfiguration is trial and error.Autotuning removes the need for hand-tuning, because it empirically searchesfor the optimal configuration on any GPU. Furthermore, online autotuning allowsthe application to react to changes that occur at run-time, for example changes inthe data size. As hardware-induced parameters affect the performance of almostall GPU programs, autotuning is applicable to a wide range of software.In order to autotune parameters that need to be known at compile time, suchas CUDA’s preprocessor command for loop unrolling, we used code instantiationvia templates. Code instantiation also benefits from compiler optimizations.block size average time standard deviation128 0.256 0.011256 0.210 0.009512 0.232 0.0281024 0.251 0.022Table 1. Average run-time and standard deviation of an empty kernel. All times are inmilliseconds.3.2. Measuring execution time on the GPUOne important difference to autotuning for CPUs is that measuring execu-tion times is less accurate on GPUs. While CPUs provide exact cycle coun-ters, GPU computation is initiated via driver calls (which might use internalcaching/queuing) and the scheduling on the GPU is a black box for the pro-grammer. On the GPU, one measures the time between the driver call and thereturn of the result. We insert events in the execution pipeline right before andafter the kernel launch. Measuring the time difference of the events occurrencesgives the best approximation of the kernel execution time. The associated, un-predictable overhead makes measurements on the GPU prone to inaccuracies forshort running kernels. Figure 2 illustrates the launch of a GPU kernel and thetiming events.We used an empty kernel to measure the overhead of a kernel launch. Thekernel launch time varies with changing block sizes. Table 1 lists the average timeand standard deviation of 20 kernel launches on the GTX 680 graphics card. Manyalgorithms call multiple kernel which accumulates the overhead and potentialmeasuring errors. It is therefore important to measure time frames that are longenough to dampen differences in the kernel launch time.CPUGPUkernel launch synchronize GPUkernel execution timestart event stop eventFigure 2. Measuring kernel execution time with events.In the next sections we discuss sample applications and show their behaviourduring autotuning. We evaluated them on the following three GPUs:• Geforce GTX680 (GK104, 4096 MBytes memory, GPU clock rate 1058MHz, memory clock rate 1502 MHz, 1536 CUDA cores)• Geforce GTX470 (GF100, 1280 MBytes memory, GPU clock rate 607 MHz,memory clock rate 838 MHz, 448 CUDA cores)• Quadro 6000 (GF100GL, 6144 MBytes memory, GPU clock rate, 574 MHz,memory clock rate 750 MHz, 448 CUDA cores)4. ResultsWe evaluated autotuning on the GPU with four applications that come directlyfrom NVIDIA’s CUDA tool kit and are in most cases already set up for bench-marking. The sample applications documentation discusses performance impli-cations of the kernel launch configurations but in all cases the parameters wereselected by trial and error for one particular GPU.4.1. Merge SortThe merge sort sample implementation uses three different kernels to sort an ar-ray of 222 elements. Each kernel’s block size can be configured independently. Au-totuning these block sizes leads to a 3-dimensional parameter space. A completemeasurement of the parameter space can be seen in figure 3.Figure 3. Merge Sort Parameter Space, block sizes of three kernels.The three tested GPUs react similarly to changes of the kernel launch con-figuration, however the respective optimal values are different. Table 2 lists theachieved speed-up and how many iterations it took for the autotuner to find theoptimal parameters.4.2. N-Body SimulationAn N-body simulation is a simulated dynamic system of mass points that influenceeach others acceleration. An all-pairs approach with a computational complexityof O(N2) is used to compute the position of each mass point in the next timestep. The implemented kernel computes the gravitation for each point by loopingdefault time optimal time speed-up autotuning iterationsQuadro 6000 0.0483 0.0324 32.81% 17GTX 470 0.0457 0.0307 32.85% 20GTX 680 0.0328 0.0234 28.64% 20Table 2. Execution times in seconds. Speed-up of the optimal configuration relative to the defaultparameters.over the other mass points and calculating their gravitational pull. The simulationimplementation yields two parameters of interest:• the block size of the kernel, and• the degree of loop unrolling.The documentation provided by NVIDIA discusses the effects of loop unrollingin relation to the number of simulated points. However, the results presented inthe documentation use trial and error to find good values for a certain graphicscard and a fixed input size. With autotuning it is possible to find the optimaltuning values for different hardware and varying input sizes. We measure theperformance while simulating 131072 mass points.Figures 4 and 5 show the performance in the complete parameter space forthe Quadro 6000 and GTX 680. The GTX 470 has the same architecture as theQuadro 6000, with the main difference being the clock speed. The performancegraphs of those two GPUs therefore have the same shape, with a small offset inthe y-axis.The degree of loop unrolling has a significant impact on the performance ofthe kernel. On the Quadro 6000 the unrolled kernel runs about 23% faster thanthe a kernel without loop unrolling. Important to note is that the loop unrollingparameter and the block size are not independent of each other, as can be seenin figure 4. Tuning the parameters independently will result in a non-optimalconfiguration.The achieved speed-up and the number of autotuning iterations it took tofind the optimum are listed in table 3. Even though the kernel was already handoptimized, AtuneRT was able to find a better configuration on every GPU.default time optimal time speed-up autotuning iterationsQuadro 6000 6.361 6.304 1.0% 26GTX 470 6.044 5.953 1.5% 20GTX 680 2.937 2.824 3.8% 14Table 3. Execution times in seconds. Speed-up of the optimal configuration relative to the defaultparameters.4.3. Thrust ScanThrust1 is a C++ template library for CUDA, similar to the Standard TemplateLibrary (STL), providing various algorithmic building blocks, such as scan, and1http://docs.nvidia.com/cuda/thrust/Figure 4. N-Body simulation – Performance of the Quadro 6000.reduction. Thrust’s algorithms either use hard-coded values or simple heuristicsto determine the block size of the CUDA kernels. We applied autotuning to theseparameters to optimize the kernel launch configuration.We benchmarked 1000 calls to the thrust::inclusive scan-function com-puting a prefix sum over 224 random integers. The results are listed in table 4.The performance for all possible block sizes can be found in figure 6.On the GTX 680 we can achieve a 13% performance gain over the non-tunedThrust function. Our autotuner AtuneRT confirms that 224 is the optimal blocksize for the thrust::inclusive scan-function on the Quadro 6000.4.4. Marching CubesMarching cubes is an algorithm often used in visualization for extracting a trian-gular mesh of an isosurface from three-dimensional scalar fields, e.g. a computer-tomography scan, stored as a 3D regular grid. It processes cells of the 3D gridFigure 5. N-Body simulation – Performance of the GTX 680.Figure 6. Performance of Thrust’s scan function.Figure 7. Marching Cube – Performance.Thrust Scan Marching Cubestime in seconds block size time in seconds block sizeGTX680, no tuner 2.566 224 0.170 160GTX680, with tuner 2.223 128 0.134 128Quadro 6000, no tuner 10.025 224 0.179 160Quadro 6000, with tuner 10.026 224 0.179 224Table 4. Performance measurements for Thrust scan and Marching Cubes.independently and determines the triangles required to represent the part of theisosurface passing through each cell.Figure 7 shows that the best configuration of 5 warps per block for the GTX470 and the Quadro 6000 does not transfer to the GTX 680. The optimal con-figuration for the GTX 470 and the Quadro 6000 perform poorly on the GTX680.The NVIDIA CUDA samples include a marching cube implementation thatuses Thrust. Again we apply AtuneRT to the block size. The results listed intable 4 show that the standard block size gives poor performance on the GTX680. The autotuner achieves a 20% speed-up. A block size of 224 is as good asthe standard block size on the Quadro 6000.5. ConclusionWe showed that autotuning is a feasible tool for optimizing GPU applicationson multiple platforms. Preparing the applications for AtuneRT was easy: ThreeAPI calls sufficed for every example. These three calls could easily be placedinto the kernel API call. In this way, autotuning requires no extra effort forthe programmer. We showed the performance impacts of the tuning parametersand the importance to tune them independently for each GPU. Autotuning alsoprovides a way to configure GPUs without having to deal with, or even know,the hardware specifications. Our evaluation showed that the tuning parameterspace is unintuitive. For every hand-tuned sample application AtuneRT was ableto find the optimal kernel launch configuration on every tested GPU.Due to the heterogeneity and wide variety of GPUs we expect autotuning tobecome an essential part in determining tuning parameters at run-time.References[1] John A. Nelder and Roger Mead. A simplex method for function minimization. Thecomputer journal, 7(4):308–313, 1965.[2] Clint Whaley, Antoine Petitet, and Jack J. Dongarra. Automated empirical optimizationsof software and the atlas project. Parallel Computing, 27(1):3–35, 2001.[3] Anna Morajko, Toma`s Margalef, and Emilio Luque. Design and implementation of adynamic tuning environment. Journal of Parallel and Distributed Computing, 67(4):474–490, 2007.[4] Jason Ansel, Yee Lok Wong, Cy Chan, Marek Olszewski, Alan Edelman, and Saman Ama-rasinghe. Language and compiler support for auto-tuning variable-accuracy algorithms. InProceedings of the 9th Annual IEEE/ACM International Symposium on Code Generationand Optimization, pages 85–96. IEEE Computer Society, April 2011.[5] Cristian T¸a˘pus¸, I-Hsin Chung, Jeffrey K Hollingsworth, et al. Active harmony: towardsautomated performance tuning. In Proceedings of the 2002 ACM/IEEE conference onSupercomputing, pages 1–11. IEEE Computer Society Press, 2002.[6] Thomas Karcher and Victor Pankratius. Run-time automatic performance tuning formulticore applications. In Euro-Par 2011 Parallel Processing, pages 3–14. 2011.[7] Frank Otto, Christoph A. Schaefer, Matthias Dempe, and Walter F. Tichy. A language-based tuning mechanism for task and pipeline parallelism. In Proceedings of the 16thinternational Euro-Par conference on Parallel processing: Part II, Euro-Par’10, pages328–340, Berlin, Heidelberg, 2010.[8] Henry Wong, Misel myrto Papadopoulou, Maryam Sadooghi-alv, and Andreas Moshovos.Demystifying gpu microarchitecture through microbenchmarking. In Performance Anal-ysis of Systems Software (ISPASS), 2010 IEEE International Symposium on, pages 235–246, 2010.[9] CUDA C Best Practices Guide, page 39. October 2012.[10] Vasily Volkov. Better performance at lower occupancy. In Proceedings of the GPU Tech-nology Conference, GTC, volume 10, 2010.[11] Yuri Torres, Arturo Gonzalez-Escribano, and Diego R Llanos. Understanding the impactof cuda tuning techniques for fermi. In High Performance Computing and Simulation(HPCS), 2011 International Conference on, pages 631–639. IEEE, July 2011.

Application-independent Autotuning for GPUs

https://publikationen.bibliothek.kit.edu/1000037784/4204007

Application-independent Autotuning for GPUs

Abstract

Similar works

Full text

Available Versions

KITopen