# An Evaluation of Directive-Based Parallelization on the GPU Using a Parboil Benchmark
## 3 OpenACC Programming Model  
OpenACC is an open standard for directive-based parallel programming on heterogeneous systems. It is used for offloading regions of code for execution on accelerator devices, such as GPUs, but also CPUs and FPGA boards. It is implemented on all major parallel platforms, including CPUs and GPUs of different vendors, operating systems, and compiler infrastructure. Supported programming languages include C, C++, and Fortran, which allows easy porting of existing applications via the annotation of source codes implemented in those languages [36]. Similar to OpenMP directives that use sentinel #pragma omp, OpenACC directives for C/C++ use sentinel #pragma acc to annotate the code. The current version of the specification is 3.3. OpenACC offers an implicit, descriptive-based programming model [10]. The programmer annotates code segments that should be executed in parallel on the accelerator device. The programmer is responsible for defining the execution context which include parallel loops and dependencies, as well as the data involved and their movement. Compiler and runtime environment use that information during the compilation and execution time, respectively, to allow for the efficient offloading of the code onto the accelerator. In that sense, parallelization can be seen as assisted, rather than fully automated. Generally, the OpenACC execution model shares similarities with the NVIDIA CUDA low-level programming model. The CPU, which serves as a host, configures and launches computed kernels using a large number of parallel threads. To support different GPU architectures, three levels of parallelism are available in the OpenACC execution model: gang, worker, and vector [37]. Those are mapped to corresponding hardware units depending on the hardware used. This is usually performed by the compiler, but the programmer has some degree of control over the whole process. The execution model assumes that the GPU contains multiple processing units running in parallel that can efficiently perform vector-like (SIMD) operations. In the case of NVIDIA GPUs, the processing unit is a streaming multiprocessor. The gang is mapped to a thread block, the worker represents a warp of threads, and the vector corresponds to a CUDA thread. Two slightly different directives are used to create an OpenACC parallel region. The first one is a parallel construct, which is used to annotate work-sharing loops. The other one is a kernel directive, which allows the assembling of kernels that contain more than one loop [38]. The OpenACC parallel directive is more similar to the same OpenMP construct. It allows more explicit, user-defined parallelism, offering good fine-tuning prospects and control over the parallelization process. As a result, only one kernel will be generated and executed at the launch time with the predetermined number of gangs and workers in the gang. The kernels construct offers more flexibility, as it offers the compiler more space for optimizations. It defines a code region that can be divided into a series of kernels. Typically, execution configuration for each loop transferred into the kernel can differ, as they do not have to use the same number of workers and gangs. Clauses are used to control the level of parallelism in parallel and kernel directives. However, those clauses should be used with caution to avoid possible performance-related issues on different platforms. Like OpenMP for the directive, the loop construct is available to specify which loops are to be parallelized. It is used for work-sharing purposes. Similar to CUDA and OpenCL programming languages, OpenACC does not offer direct synchronization or data sharing among gangs. The main reason for such assumptions is to support the scalable execution of code on architectures with different numbers of processing units. It means that only workers within the same gang are able to share the data, affecting the way parallelism in the code has to be mapped in the kernels [39].