-The GPU (Graphics Processing Unit) are the future of high performance computing and provides a parallel programming model for general purpose applications thanks to CUDA programming interface. The programming model of GPU architecture is significantly different from the traditional CPU one. This paper presents the advantages of GPU architecture by proposing an algorithm for the linear combination of digital filters in image processing, implemented as a parallel GPU version (a sequential CPU version has been also implemented for comparison). The use of parallel processing CUDA architecture has enabled us to take advantage of the GPU allowing an increase of the performance compared to CPU. Index Terms-Matrices Convolution, Image processing, GPGPU, Parallel Computing, CUDA.
The evolution of computers has allowed to develop more complex calculations and has improved the performance of the simulations in scientific fields. Technology has increased the CPU (Central Processing Unit) frequency with the upper limit of 3 GHz, processing multiple cores: dual core, quad core and octa core in a few decades, unfortunately with the limits due to power dissipation and the increasing temperature.
One of the solutions proposed to push such limits is the
use of the GPU (Graphics Processing Unit). Typically, GPUs
handle the huge amount of data for graphical applications
in three-dimensions with high performances, so the GPGPU
(General-Purpose Computing on GPU) has been introduced.
GPGPU programming has been implemented using two types
of APIs (Application Programming Interface): the complex
OpenGL [
In the recent years, the computer house NVIDIA has developed CUDA (Compute Unified Device Architecture). This platform is much simpler than the previous APIs and and it allows to program using a high level programming language based on C, using a model of parallel computing. NVIDIA CUDA technology has opened a new era for GPGPU computing allowing the design and implementation of parallel GPU-oriented algorithms without needing any knowledge on OpenGL. The computational power of these architectures, nowadays, has received a considerable growth compared to CPUs and due to the GPU ability to perform a huge number of simple operations in parallel. GPUs have cores much simpler compared to CPUs. This allows the realization of Copyright c 2016 held by the authors.
GPU architectures in which the silicon area for control and
management is very reduced compared to that employed in
the CPU and a big increase in the number of cores [
GPUs allow us to compute optimally, in parallel mode, a code with a certain complexity and for the application of digital filters in two-dimension arrays. This paper proposes the use of GPU architecture for the implementation of digital filters with a parallel code that uses a linear combination of digital filters processing images. The aim is to compare the parallel execution speed to the sequential code. The remaining parts of the paper will be dedicated to explain the parallel and serial approaches.
A digital filter performs mathematical operations on discrete-time sampled signals in order to enhance or reduce certain characteristics of the signal. Such filters are implemented by software components, often through mathematical functions or matrices, and loaded inside the processors with programmable hardware. Cost and speed are closely dependent on the used processor.
The application of digital filters has enormous advantages over analog filters as the possible transfer functions are much more flexible for digital filters. The main advantages are: 1) high accuracy due to the absence of physical components; 2) a digital filter is easier to design and implement automatically modifing its frequency response and changing the input; 3) flexibility to change the digital filter parameters, without changing the system hardware; 4) easy simulation and design of the filters being implemented by software with a reduction of system complexity. these properties are essential to the implementation of a high quality filter. As mentioned previously these advantages are accompanied by the limitations of speed and cost, related to the processor used, and frequency. The frequency limit is described by Nyquist theorem, using the following formula: fs > 2B (1) which imposes the filter with a maximum limit of the frequencies in order to reduce the effects of additional aliasing distortion of the signals. B term is the signal band and fs is sampling frequency. The digital filters are widely used in image filtering and image processing. This paper will examine, implement and test an algorithm to filter digital images with a GPU implementation.
Digital filters are an essential tool for image processing. The digital image can be defined as a two dimensional array or matrix and each element represents a pixel of the image. Images are processed by an algorithm. The process of filtering is also known as convolution of a mask with a standard size 3x3, 5x5 or 9x9, which is a matrix. This mask applies different mathematical operators by means of a convolution to the image to achieve digital image processing.
The use of filters to digital image processing allows us to make operations like: 1) the extraction of information from the image, such as the detection of boundaries in images; 2) the exaltation of details, such as the increase of light intensity or color contrast; 3) the elimination and reduction of disturbances in images.
This processing brings a high computational cost especially for large images. For this reason a sequential approach, based on the CPU, is not very efficient, so parallel approaches have been introduced to increase performances. This paper presents an algorithm for the application of a linear combination of filters to local processing based on a model of parallel computation on GPU architectures. For our local processing filters, the image processing method consists of applying a function to each original pixel values and to an appropriate range of pixels, within the radius of the filter matrix. This method is based on the convolution operator for which a brief explanation is needed.
As previously said, the algorithm executes an operation on an input matrix and pre-established filters, in order to provide as a resalt a linear combination of the input. The 2D convolution between two continuous functions, f (x; y) and g(x; y), is defined by the formula:
Z 1 f (x; y) g(x; y) = f ( ; )g(x ; y )d d (2) 1 This expression will be brought to the discrete domain, by assuming f (x; y) and g(x; y) be two discrete arrays of a limited size, as a double sum expressed by:
M 1 N 1 f (x; y) g(x; y) = X X f (m; n)g(x =0 =0 ; y )d d (3) for x = 0; 1; : : : ; M 1 and y = 0; 1; : : : ; N 1.
This convolution operation is defined on discrete-time as a simple operation of “local media” in a range of amplitude defined by the kernel size used obtaining a value for the output matrix for the same position of the source data as shown in figure 1.
This operation is done from left to right, for each of the original matrix element obtaining the matrix of the resulting convolution. Therefore, the convolution algorithm implements the digital filtering operation. Then, it suffices to modify the kernel in order to change the type of filter. We have studied a digital filtering operation of am image, in ASCII format, through two filters of ninth order. The next section describes the convolution operation implemented as a parallel algorithm that can be executed on a CUDA GPU.
This paper presents the convolution algorithm implementation for a bidimensional matrix of input data with a linear combination of filters using a parallel model. In this work two versions of the same algorithm are presented in order to compare the efficiency of the parallel computation and the sequential one, and then demonstrate how the parallel computing support is far more suitable for high performance computing.
In recent decades, high performance computing sector has been having an exponential growth with the widespread use of GPGPU. The use of GPU in this ever-expanding purview can exploit the power of parallel computing for increasingly complex simulations due to the inherent parallel nature of GPUs.
CUDA compatible GPUs are, in fact, based on an architecture made up of a number of MIMD (Multiple Instruction Multiple Data) multiprocessors called Streaming MultiProcessors, whose number depends on the specification and class of GPU performance. This Streaming MultiProcessors are the basic units of the GPU architecture and are implemented as SIMD (single instruction multiple data), and called by NVIDIA as SIMT (Single Instruction Multiple Threads), which has 8 processors said Streaming Processor or CUDA Cores. In this architecture, each Streaming MultiProcessor is able to create, manage, schedule and execute groups of 32 threads called warp. A warp executes one instruction at a time, so as to maximize the efficiency when all 32 threads of a warp agree on their execution path. When one or more blocks of threads are assigned to a multiprocessor to run, they are then partitioned into warps, scheduled by a warp scheduler and executed one at a time. Each of these processors can perform simple mathematical operations (such as addition, subtraction, multiplication, ect.) for integers or floating point numbers. Inside each multiprocessor there is also a shared memory, accessible only by the processors in the same multiprocessor, caches for instructions and for data, and, finally, a unit for decoding the instructions.
Each multiprocessor has access to a global memory shared
among all GPU multiprocessors and called Device
Memory [
Kernel is the model of parallel execution of the code in a device and it is defined as a grid divided in to a certain number of blocks to which it is assigned a multi-processor for each. Inside each block there is a number of fundamental computational units defined as thread.
Kernel, or grid, are sequentially executed between them while blocks and threads are executed in parallel by adopting a SIMT data-parallel model. Each of these thread belongs to a single block and is uniquely identified within the kernel by assigning it an index. In this manner the memory addressing will be simplyfied especially in the case of processing multidimensional data.
In each of the blocks the kernel also has a shared memory accessible only to the threads of the same block. The logical subdivision of a kernel in grid and blocks is a crucial aspect in a code in CUDA for obtaining the parallelization of code. The organization and management of the internal threads allows the implementation of a more efficient code.
The algorithm developed in this work executes the convolution operation between an input matrix A and some preestablished kernels generating an output matrix A~ given by the linear combination of results obtained using this formula A =) A~ = (k1
A) + (k2
A) +
A (4)
While this operation is executed, we want to generate the best parallelized code for the purpose of improving its performances compared to the serial code version. The algorithm shown can be summed up in three basic steps: 1) acquisition of input data matrix; 2) convolution of the input matrix with pre-established kernels; 3) linear combination of the results.
Obviously, each of these steps is performed by both the sequential version of the code and then the parallel one. The following will detail the basic steps previously exposed focusing our attention on the parallel algorithm.
The parallel solution of the algorithm starts in a host by means of code that acquires the input and takes data from a text file in ASCII code and store it in an array on shared memory. Instead of a 2D array for data, complex to manage, we have used a single dimensional array where the rows of the input matrix are reported one after the other, then data are reconstructed by an index. In this way the data will be transfer from host to GPU with the standard function cudaMemcpy(). After a given number of threads and having organized the blocks, the call to kernel is performed, in order to do parallel computations.
Inside the kernel it is assigned a thread to any given input so that we can entrust the execution of the code to a set of parallel threads arranged in blocks and indexed using the following formula: id = threadIdx.x + (blockIdx.x*blockDim.x);
To each thread will be assigned a position in the input array, and the thread will compute the convolution, in parallel to other threads in each block. The input data acquired in the form of array was necessary to implement a mechanism of jump allowing us to move between the various locations in the array in order to perform the products between the elements of the kernel (in our case a simple 3x3 filter), and the components located in the neighborhood of the source location. This problem is solved using the radius of the kernel defined as:
Kernel_radius = (Kernel_order -1)/2;
We can identify all the elements lying in the range of the position of interest. Then kernel radius is subtracted to the input array index and added a variable that increases cyclically through a simple for loop. This let us multiply in parallel the elements of the input with the kernel element lying in the right position around the central element of the first row of the kernel, as shown below: for(i=0;i<N;i++) s[id]+=A[id-z+(i%3)+(i-(i%3)/3)]*k[i];
Then the convolution algorithm performs the operation just described for each row of the kernel: for each row in kernel; for each element in kernel row;
Multiply the element to the location of the Kernel with the corresponding element in the input data matrix.
Sum results, save them on central position.
This algorithm is executed as many times as the kernel number we want to apply to our array of the input data. In our specific case we chose to apply two kernels of ninth order with filter function for the input matrix obtained by ASCII encoding of a two-dimensional image. For a linear combination of the results we have used, finally, a simple program, add, shown below: if(id<((rig-1)*(col-1)))
sum[id]=(a*s[id]+b*s2[id]+g*A[id]); else return;
The add program, called from the host and executed in the device, sums the results of the two performed convolutions loaded in arrays s and s2 with the values of the input data A. Each of these addends is regulated by a default weight which calibrates the effect of a convolution compareted to the other in the final matrix indicated with sum array. An input image having 256 colors has been used as input array, so to ensure the output image had the same range of color variation, two if statements impose this condition to the elements of sum array.
In the following the results and performances of the measures on the execution time in CUDA C and the C versions are compared, hence assessing parallel and sequential versions of the algorithm.
A comparison of the execution time using the function clock belonging library time:h has been performed to the algorithm in the sequential and parallel versions. The execution time measures the time interval between the start and the end time. Matrices of various sizes were used as input data in order to compute the efficiency of the two implementations when varying the amount of input data. Such matrices are matrices representing images, coming from a text file in ASCII code.
The sequential algorithm for processing data runs on a AMD i686 Dual-Core processor with up to 1GHz clock speed, while for the GPU algorithm we have used a NVIDIA GeForce GTX 480 with 480 CUDA cores and 1536 MB GDDR5 video RAM with CUDA 4.2. We developed our GPU code using NVIDIA CUDA API while the CPU code is compiled under nvcc. The GPU results are compared with the CPU results under Linux. For our simulations we used a subdivision of kernel in blocks of 440 threads by applying filters of order 9.
Table I shows the sizes of the matrices used in data input and averages, calculated on 20 samples, the execution times of the sequential and parallel versions.
Note that the execution time of the algorithm on the GPU includes the transfer times of data (HostToDevice and DeviceToHost) and the kernel running time. The experimental results are represented in Table I and Figure 3. 105 106 107
108 dimension matrix
Figure 3 shows the execution times, sequential and parallel, where the parallel code is faster. The result of this comparison shows that the parallel code is more efficient. The comparison between parallel and sequential version has been calculated as the difference between the Sequential Execution T ime e il P arallel Execution T ime, and has been reported using a histogram in which the parallel code presents a significant performance, while increasing the size of the input image, hence more benefits in the use of a parallel CUDA architecture than the typical CPU.
Figure 4 shows the trend of the difference between the Sequential Execution T ime and the P arallel Execution T ime. The trend of this curve is exponential, observing the parallel version it has an execution time smaller than the sequential version, with a greater efficiency and better performances, while increasing the amount of data input. The more calculations the more benefits are gained from CUDA parallelism. data array, have led to the results shown in Table II and plotted in Figure 6.
To test the validity of the results, we were carried out other tests based on the execution of iterated code in order to simulate the processing of a set of images of equal size, then measuring the Execution T ime of the whole process. The results of these tests are shown in Figure 5.
20 40 60 80
100
N umber of Iterations
The advantages of the use of CUDA architecture are directly proportional to the increase in computation. The same tests were performed using a computing architecture more complex and powerful in order to see if an increase in hardware performance could lead to different results. Such tests were performed by loading the source files on a server equipped with Intel Xenon E5- 2630 with 6 cores 2.6 GHz clock and RAM 48 GB, NVIDIA Tesla K10 with 2x1536 Cuda cores and 8 GB GDDR5 video RAM. The two versions of the algorithm, sequential and parallel, varying the size of the input 4096x4096 8192x8192 9000x9000
The numbers give the result of an average of 10 samples of the running times. The results obtained reveal experimentally that by increasing CPU hardware performance and for images having smaller sizes, the sequential version is faster than its parallel version. This is possibly due to CUDA cores of common GeForce having a low clock speed, determing slower performances for a few calculations. On the other hand, the performance of the parallel code is better than the serial version for the larger processed input array. The tests based on the execution of iterated code produce the results shown in Figure 7.
Tests run on the server estimate the trend when increasing the amount of input data. A better hardware has obtained better performance for both versions of the algorithm, producing however, a difference depending on the size of the input data array. In fact, the calculation contained in the algorithm running on the terminal using the GeForce GTX 480 has a better speedup over the sequential version on CPU. As soon as the need grows, it is always better the computational performance with a parallel version. Assuming that the matrices are 9000x9000, the Tesla K10 has a speedup almost 140x with respect to the GeForce GTX 480. So the processing on GPU cards is most often preferable than the computaion on the CPU, the choice of the GPU itself must be weighted on the basis of the knowledge of the computational power needed. The strategy of parallelization brings more efficient algorithms.
In recent years the field of GPGPU has aroused great
interest inspiring many research works in order to include
the best strategies that make it possible to exploit the full
potential of the structure of parallel computing and the CUDA
parallel programming model. In this regard it is possible to
see the work [
A lot of work has been performed for applying the CUDA
model for the image processing field. A paper useful to
understand how to perform image convolution, essential for
image processing, has been produced by Podlozhnyuk [
Finally, the work relating to the performance and application
of parallel programming model is inherent in a huge range of
topics, such as e.g. the work in [
This paper has presented an algorithm for the application of a linear combination of digital filters implemented with parallel programming for image processing on CUDA compatible GPUs, and has compared with its sequential version on the CPU. The solution presented has the aim to exploit the programming model CUDA parallel, more efficient than a sequential processing code with an appropriate management and organization of blocks and thread.
The tests have shown that the implementation of the parallel code on the GPU increases in speed compared to serial implementation of CPU. In general the speed is not less than 40x with peaks higher than 90x. This increase in speed is also evident with the increase of computational power required in processing. The execution time of the parallel code also counts the time of data transfer between the host device and the penalization of some performance of parallel code which remains higher than the serial version. It is possible to improve the result by a more detailed analysis of the strategies of parallelization in order to manage the computing resources provided by the GPU architectures.
For the realization of the C code we have used CUDA programming interface which has made it easier the GPGPU approach, and allowing us to exploit very effectively the potential hidden in GPU computing. The algorithm presented in this article, though designed for the application of filters on digital images, is well suited for many applications including mobile ones that are outside the scope of image processing such as signal analysis or application of spectral masks sampled signals.
Another possible application for our algorithm is that in the industrial field for the detections of the boundary in an image useful for the detection of objects. In this case you could apply the Sobel operator by simply modifying the matrix of the filter. With this work we have shown only a small part the great potential of the GPU, which could gain popularity in the scientific community thanks to the computing power and to its easy availability, which makes GPGPU one of the best choices in high performance computing.