Heterogeneous architectures nowadays are becoming very attractive in the embedded and mobile markets thanks to the possibility to exploit the best computational resource to optimize the performance per Watt figure of merit. Unfortunately, deciding the right resource to use and its operating frequency is a difficult problem that depends on the actual conditions in which the system is operating. This work aims at proposing a runtime controller, integrated in Linux Operating System (OS), for optimizing the power efficiency of a running application deciding the system configuration. Our experimental results over a set of applications from the Polybench suite on the Odroid XU3 board show that our controller is able to obtain a power efficiency of more than 90% of the one achievable via offline profiling.
Heterogeneous System Architectures (HSAs) [
However this increase in heterogeneity comes at the cost of new issues in programmability and runtime management of these resources to achieve the pursued performance/power consumption trade-off. In particular, various kind of processing units imply different type of programming languages and models thus introducing new implementation and integration challenges. Nevertheless, this abundance of resources has to be properly managed in the execution of the workload, since each type of processing unit offers a different level of performance/power efficiency to each single application and part of it.
In 2009, Khronos Group, including Apple, ARM, Samsung
and many other industrial partners, has defined OpenCL [
In this paper, we present a runtime controller integrated within OpenCL applications running on Linux enabling the monitoring of system status and the automated adaptation of the application itself 1. We also propose a novel policy integrated within this controller allowing the application to self-tuning by acting on the mapping and the Dynamic Voltage and Frequency Scaling (DVFS) of the processing units to optimize the performance/power consumption trade-off. Experimental results presented in the paper show the efficiency of the controller to quickly converge to the optimal solution with less than 10% of error.
The rest of the paper is organized as follows. Section 2 briefly discussed the related work. Then Section 3 introduces the working scenario and states the addressed optimization problem. The implementation of the proposed integrated runtime controller and the decision policy are provided in the Sections 4 and 5, respectively. Then, an experimental evaluation of the approach is provided in the subsequent Section 6, and, finally, Section 7 concludes the paper. 2.
Many OpenCL runtime supports have been defined by
vendors for their designed processing units; examples are
Intel for last generations of Pentium, Xeon and HD
Graphics units, NVidia and ARM for GPU devices, and Xilinx
for FPGAs. Moreover, other open source runtime supports,
such as [
CPUs. Finally, OpenCL ICD loaders (e.g., [
Approaches for tuning and optimizing OpenCL
applications on HSAs have been recently investigated by a number
of works, such as [
Further works (e.g., [
This section presents the working scenario considered in this paper, consisting of the target platform and the class of executed applications. Finally, we formulate the optimization problem addressed in the runtime controller we propose. 3.1
In this work we consider an HSA as the Samsung Exynos
5422 [
Regarding the target applications, a set of computational
kernels that is of great interest in the context of
embedded and mobile systems is the family of polyhedral
applications. Polyhedral applications comprise algorithms for video
processing, filters and algebraic transformations which are
at the basis of control, infotainment and augmented reality
software. These applications are characterized by a
computational intensive kernel continuously executed in a loop
on incoming data, such as the frames in a video to be
processed. In this work we considered the OpenCL
implementation of Polybench [
The addressed problem is to identify for a single given application running in our system which is the best operating point in terms of Performance per Watt. In particular, we want to identify at runtime, without previous profiling information, which is the best processing unit to use and the related frequency level.
In a more formal way, let us consider the controlled application A running on the target architecture containing three different processing units P = {BIG, LIT T LE, GP U }; each unit is characterized by a set of possible frequencies (in MHz): fBIG = {200, 300, . . . , 1900, 2000} fLIT T LE = {200, 300, . . . , 1200, 1300} fGP U = {177, 266, 350, 420, 480, 543, 600} The running application will be characterized for each processing unit and supported frequency level, by two direct metrics the throughput, T hr, and the overall power consumption W , and a derived metric called the power efficiency EF F , defined as following:
EF Fp,f = T hrp,f , p ∈ P, f ∈ fp
Wp,f The goal tackled in this work is to find pˆ ∈ P and fˆ ∈ fpˆ such that:
EF Fpˆ,fˆ ≥ EF Fp,f , ∀p ∈ P ∧ ∀f ∈ fp
In order to solve this optimization problem we need to address a set of technical issues related to the monitoring and controllability of the running application on the target system; specifically, we need i) the support for OpenCL for all the processors with the possibility to constrain and move the execution at runtime; ii) to measure the throughput of one iteration of the application and its power consumption; iii) a smart algorithm to explore the power efficiency curves and rapidly identify the best operating point.
The self-adaptive approach proposed in this paper has been implemented in a specific controller C++ class directly instantiated within the application source code. Figure 2 depicts the overall structure of the controller and its integration within the system. Moreover, in order to enable the actuation of the mapping on all OpenCL devices, the application has to be implemented according to a specific template. Listing 1 shows the defined application template and how the controller is instantiated and used. All the details of the controller in Figure 2 are discussed in the following paragraphs.
OpenCL runtime. To enable the support for all the
devices available in the Exynos chip, we have installed both
ARM OpenCL Mali SDK [
OpenCL template. The defined application structure is shown in Listing 1; it slightly enhances the standard OpenCL template. In particular, the latter one requires the programmer to select and configure the desired platform and device to be used to execute the kernel. We extend this template to allow the defined controller to dynamically select at each iteration of the application which device to use.
To enable this capability, in the OpenCL initialization step, all platforms and devices are discovered and set up, as shown in the piece of code reporting the cl_init() function (Lines 8-19). In particular, the function iterates on all platforms and sets up all devices2 and related OpenCL objects, such as the context, the memory objects and the program objects; all these objects are stored in arrays.
Then, a specific variable, curr_device, is used to specify the index of the current device to be used in the execution of the kernel. Thus, the run_kernel() function uses the execution context specified in such a variable to run the kernel (Lines 21-38).
Cgroups actuation. In order to enable cluster-level
mapping on the big.LITTLE CPU, we have exploited OS
facilities for task mapping to force the usage of a subset
of the cores. Linux OS provides two different mechanisms:
sched_set_affinity() and cgroups. sched_set_affinity()
cannot be used in OpenCL applications, since it needs to
know thread IDs; indeed, threads are generated within the
OpenCL runtime and their IDs are not visible externally.
Instead, cgroups offers the possibility to assign a set of cores,
and, more in general, further resources such as CPU quota
and memory amount, by specifying only the application
PID; all the threads spawned by that PID are then
managed automatically. Therefore, as in [
Listing 1: Application template 1 //OpenCL o b j e c t s 2 c l p l a t f o r m i d p l a t f o r m i d s [MAX PLATFORMS ] ; 3 c l d e v i c e i d d e v i c e i d s [ MAX DEVICES ] ; 4 . . . 5 c l u i n t num platforms ; 6 c l u i n t num devices ; 7 8 void c l i n i t ( ) { 9 int i ; 10 // s e t u p OpenCL environment f o r a l l d e v i c e s 11 c l G e t P l a t f o r m I D s ( 0 , NULL, &num platforms ) ; 12 c l G e t P l a t f o r m I D s ( num platforms , p l a t f o r m i d s ,
NULL) ; 13 for ( i =0; i <num platforms ; i ++){ 14 c l G e t D e v i c e I D s ( p l a t f o r m i d s [ i ] ,
CL DEVICE TYPE ALL , 1 , &d e v i c e i d s [ i ] , &num devices ) ; // s e t u p o t h e r OpenCL o b j e c t s f o r d e v i c e [ i ] // i . e . cont ex t , queues , memory , programs . . . 15 16 17 18 } 19 } 20 21 void r u n k e r n e l ( int c u r r D e v i c e ) { 22 // l o a d a p p l i c a t i o n data 23 . . . 24 // s e t u p t h e workgroup s i z e s 25 l o c a l W o r k S i z e [ 0 ] = . . . 26 g l o b a l W o r k S i z e [ 0 ] = . . . 27 // w r i t e memory o b j e c t s 28 c l E n q u e u e W r i t e B u f f e r ( cmdQueue [ c u r r D e v i c e ] ,
mem obj [ c u r r D e v i c e ] , CL TRUE, . . . ) ; . . . // Set t h e arguments o f t h e k e r n e l c l S e t K e r n e l A r g ( c l K e r n e l [ c u r r D e v i c e ] , 0 , s i z e o f ( cl mem ) , ( void ∗)&mem obj [ c u r r D e v i c e ] ) ; . . . // e x e c u t e k e r n e l clEnqueueNDRangeKernel ( cmqQueue [ c u r r D e v i c e ] , c l K e r n e l [ c u r r D e v i c e ] , 2 , NULL, globalWorkSize , l o c a l W o r k S i z e , 0 , NULL,
NULL) ; c l F i n i s h ( cmdQueue [ c u r r D e v i c e ] ) ; // read memory o b j e c t s clEnqueueReadBuffer ( cmdQueue [ c u r r D e v i c e ] ,
mem obj2 [ c u r r D e v i c e ] , CL TRUE, . . . ) ;
Performance monitor. Instruction per Cycle (IPC) or
other classical low-level metrics computed by the OS do not
represent a useful information to the final user to perceive
the actual progress of an application. As an example, IPC
is not able to show if the video application in execution is
providing a minimum Quality of Service (QoS) in terms of
frame/s. Therefore, to enable run-time performance
monitoring, we have integrated in the controller the Heartbeat
mechanism [
Target BIG LGIPTUTLE 2000 1000 Frequency [MHz] 1500
2000 (a) Power consumption 500
1000 Frequency [MHz]
1500 (b) Throughput 0000tt//[(])rrfsaecea1*noPmWW....000010102864 500 SYFrReq21uK0e0n0Ecyff[iMciHenz]1c5y00 TaLBGrIIgPTGeUTt2L0E00 100tt//[()frr]saecea1*onPmWW...075050 500 MFrVeqT1u0e0En0cfyfic[MieHnzc]y1500 (a) LITTLE cluster is the most efficient. (b) big cluster is the most efficient. level information from the application.
The basic idea of the Heartbeat mechanism is to measure the throughput of periodic application by i) measuring the duration of the execution of each single loop and ii) computing the ratio between the amount of processed data and such a duration. Therefore, the controller initialized all necessary data structures (timers and accumulators) during the initialization (Line 47). Then, at the end of the loop (Line 52), the invoked send_heartbeat() function collects the new timestamp and the size of the processed data (directly specified by the programmer) and based on such information computes the current throughput.
Power monitor. The considered Exynos chip integrated various sensors to monitor the status of the hardware platform, such as power consumption and temperature of the various clusters. Such sensors are exposed to the programmer through the virtual file systems of Linux OS.
In order to trace the power consumption of the big and LITTLE clusters and of the GPU, we have implemented an external monitor acting as a separate process daemon and periodically (i.e. every 50ms) collecting power values from the interface provided by the sys virtual file system of Linux OS. The external monitor can be triggered via a message over a named pipe (a Linux interprocess communication mechanism) and instructed to collect power information aggregating them over a period of time. Another message on the same pipe can stop the acquisition returning the average power consumption over the considered period.
Controller. The controller has been implemented in a single class encapsulating all the discussed mechanism and the decision policy. It exposes the following methods: • init(), invoked at Line 47, sets up the environment by initializing the data structures of the Heartbeat, cgroups and of the decision policy; moreover the function connects to the external power monitor by means of the named pipe. • get_curr_config(), invoked at Line 50, analyzes all collected metrics (power and throughput) and executes the decision policy to identify on which device to run the application kernel during the current loop iteration; moreover, it actuates on the cgroup library. • send_heartbeat(), invoked at Line 52, processes current Heartbeat at the end of the main loop to compute the throughput. • destroy(), invoked at Line 55, deallocates all data
structures. 5.
This section describes the controller policy that allows to solve the problem tackled in this paper, defined in Section 3. We will first discuss a preliminary profiling phase carried out on the target applications in order to understand which control strategy should to be adopted, and, then, we will describe the policy itself. 5.1
In a preliminary phase, we carried out an experimental evaluation of the behavior of such applications on the considered heterogeneous platform. We measured the execution time, power consumption and power efficiency of each considered application on each processing unit (big, LITTLE and GPU) at each available frequency level.
During this analysis, we noticed that the power profile is almost the same for all the considered applications presenting a quadratic relation with respect to the frequency, as shown for GEMM application in Figure 3(a). In the same way the execution time of the applications follows the same trend for all the benchmarks where an increase in the frequency level turns into an improvement of the execution time up until a certain value, as shown in Figure 3(b) for the GEMM application.
Combining these two curves we found how the power efficiency of the benchmarks varies (Figure 4). Also in this case,
Listing 2: Selection of next configuration 1 C o n f i g u r a t i o n C o n t r o l l e r : : g e t C o n f i g u r a t i o n ( ) { 2 for ( auto d : d e v i c e s L i s t ) { 3 i f ( ! d−>hasConverged ( ) ) 4 return d−>g e t N e x t C o n f i g u r a t i o n ( ) ; 5 } 6 7 s t d : : v e c t o r <s t d : : p a i r <f l o a t , int>>
b e s t E f f i c i e n c y ; for ( int i =0; i <d e v i c e s L i s t . s i z e ( ) ; i ++){ f l o a t xMax , yMax ; xMax = d e v i c e s L i s t [ i ]−>maxEstimation ; yMax = d e v i c e s L i s t [ i ]−>
getEstimationAtFrequency (xMax) ; b e s t E f f i c i e n c y . push back ( s t d : : p a i r <f l o a t , int >(yMax , i ) ) ; } s t d : : s o r t ( b e s t E f f i c i e n c y . b e g i n ( ) ,
b e s t E f f i c i e n c y . end ( ) ) ; s t d : : p a i r <f l o a t , int> b e s t D e v i c e =
b e s t E f f i c i e n c y [ b e s t E f f i c i e n c y . s i z e ( ) −1]; return d e v i c e s L i s t [ b e s t D e v i c e . second]−>
g e t N e x t C o n f i g u r a t i o n ( ) ;
Listing 3: Code of the get curr config() API 1 unsigned int OpenCl Contr oll er : : g e t c u r r c o n f i g ( ) { 2 this−>c u r r e n t C o n f i g u r a t i o n = g e t C o n f i g u r a t i o n ( ) ; // Set f r e q u e n c y using e i t h e r cpufreq −s e t f o r
CPU or using f i l e s y s t e m f o r GPU // For CPU d e v i c e s a l s o c a l l s t h e proper
cgroups f u n c t i o n s this−>c u r r e n t C o n f i g u r a t i o n . d−> s e t C o n f i g u r a t i o n ( this−> c u r r e n t C o n f i g u r a t i o n . f r e q u e n c y ) ; // Return t h e d e v i c e i d t o use in t h e
a p p l i c a t i o n return this−>c u r r e n t C o n f i g u r a t i o n . d e v i c e −> o p e n C l d e v i c e I d ; the trend of the curves is similar for all the benchmarks. An interesting aspect is that, depending on the actual values of the execution time and power consumption, three different situations can be found where either the big, the LITTLE, or the GPU outperforms the other units (as shown in Figure 4). Finally, all these curves present a maximum which is located around the middle of the frequency range for all architectures. The goal of the policy is then to find at runtime this maximum without any previous profiling information. 5.2
The controller policy is triggered each time the user code requests a configuration to use (i.e. at the beginning of each iteration of the application loop).
Considering the structure of the power efficiency curves, if we fix a device, we can use a ternary search algorithm to solve our optimization problem. This algorithm can be used to find a maximum of a mathematical continuous function F (with a single maximum point), and, actually, the curves characterizing the computational efficiency of the considered applications (refer to Figure 4) present such a shape.
The algorithm needs to find three initial points a, b, c | a < b < c ∧ F (b) ≥ F (a) ∧ F (b) ≥ F (c); to do this we execute the first three iterations of the application loop on a given device at its minimum, maximum and middle frequencies.
Then, at each iteration we acquire a new estimation of the power efficiency of the application at the last operating frequency for a given device. The new a, b, c points will be chosen by sorting the estimation and picking the two frequencies next to the current maximum efficiency found. At this point one operating frequency is chosen in the ranges (a, b) or (b, c), priority is given to the first interval; if no frequency exists in those two ranges, it means that the three points a, b, c are contiguous to each others, and b is the actual maximum we are looking for.
The ternary search algorithm has been implemented in a function called getNextConfiguration() that is used to pick for a given device the frequency to use. This function is then used in the getConfiguration() routine (Listing 2) that decides which device to use to execute the application kernel alongside its frequency. In the getConfiguration() function, the controller checks for all the available devices (devicesList) whether there exists an estimation for the best efficiency. If such estimation does not exist the controller uses the ternary search on the device to find the maximum as explained before (Line 2). When all the estimations are available, the algorithm selects the device with the best power efficiency (Lines 9-14) and asks that device for the next configuration (Line 16).
The whole configuration decision process is invoked when the application code calls the get_curr_config() (Line 50 in Listing 1). The code of this function, which bridges the application code with the controller policy is represented in Listing 3. The API calls the get_curr_config() function and then invokes the function to set the desired configuration: only the frequency for the GPU, both frequency and cgroups configuration for the CPU. Finally, it returns to the user the device id to use when invoking the OpenCL kernel, as explained in the application template section.
This section illustrates the results of our controller.
Experiments have been conducted on an Odroid XU3 board [
In order to test the performance of our controller, we modified each application to execute the computational kernel for a fixed number of 50 iterations to see if the controller was able to converge to the most power efficient solution possible. The best configuration for power efficiency has been identified with an offline profiling phase which performed an exhaustive exploration of all the possible configurations.
Figure 5 illustrates the comparison among the profiled power efficiency and the one found at runtime by our controller. Note that the power measures are subject to a little variability due to different working condition (i.e. background system processes, out of our control). As the figure shows, the runtime power efficiency adheres for all the benchmarks to the best power efficiency found during the profiling phase. In particular, for the first five benchmarks (from 3DCONV to GEMM) we have that the GPU has the best power efficiency and the controller is able to converge and use the same device at runtime. In the last two benchmarks we have that MVT has the best efficiency on the BIG
SYR2K
BICG
Polybench processors while SYR2K3 is optimal on the LITTLE ones; nonetheless the controller found also in these situations a solution near to the optimum. More in details the controller reaches a power efficiency which is at least 90% of the best value found in profiling. The reason behind the 10% error is due to runtime measurements variability. This causes the policy to converge to a frequency that is near to the optimal one; e.g. GEMM converges to 350 instead of 480 MHz.
In all these experiments the time needed to execute our policy is included in the execution time of the kernel iteration and concurs in defining the power efficiency of the application. Since all the controller functionalities are invoked also when the policy has converged, we can state that the overhead introduced by our solution is negligible.
Concerning the convergence time we have that the controller converges by testing less than 20 different configurations; this represents about 50% of the overall design space (that is composed of 38 configurations) as shows in Table 1. At the opposite, the offline profiling strategy has always to explore all the 38 configurations. Furthermore, we have to keep in mind that the offline profiling requires that the working conditions are exactly the same at the moment the application is in execution. At the opposite, runtime adaptation allows to converge also when the working conditions change.
This paper has presented a novel runtime controller integrated within OpenCL applications able to enable its automated adaptation. The controller features a novel policy allowing the application to adapt by acting on the mapping and the DVFS of the processing units to optimize the per3 All the SYR2K power efficiency values have been multiplied by 100 for sake of clarity. formance/power consumption trade-off. Experimental results have demonstrated the efficiency of the controller to quickly converge to the optimal solution with less than 10% of error. Future work deals with the adoption of further actuation knobs for resource usage, such as quota assignment and finer-grained mapping, the improvement of the proposed policy and controller to support the concurrent execution of several applications.
This work has been partially funded by the EU FP7 SAVE project (#610996-SAVE).