X-Git-Url: https://bilbo.iut-bm.univ-fcomte.fr/and/gitweb/book_gpu.git/blobdiff_plain/4eb0d6980c190aa2e92700dd01c5f685405590bd..1eda0e6e69a2160009bfcb2b86237a544898530d:/BookGPU/Chapters/chapter2/ch2.tex diff --git a/BookGPU/Chapters/chapter2/ch2.tex b/BookGPU/Chapters/chapter2/ch2.tex index 501e34e..8222660 100755 --- a/BookGPU/Chapters/chapter2/ch2.tex +++ b/BookGPU/Chapters/chapter2/ch2.tex @@ -1,15 +1,15 @@ \chapterauthor{Raphaël Couturier}{Femto-ST Institute, University of Franche-Comte} -\chapter{Introduction to CUDA} +\chapter{Introduction to Cuda} \label{chapter2} \section{Introduction} \label{ch2:intro} -In this chapter we give some simple examples on CUDA programming. The goal is -not to provide an exhaustive presentation of all the functionalities of CUDA but -rather giving some basic elements. Of course, readers that do not know CUDA are -invited to read other books that are specialized on CUDA programming (for +In this chapter we give some simple examples on Cuda programming. The goal is +not to provide an exhaustive presentation of all the functionalities of Cuda but +rather giving some basic elements. Of course, readers that do not know Cuda are +invited to read other books that are specialized on Cuda programming (for example: \cite{ch2:Sanders:2010:CEI}). @@ -17,53 +17,57 @@ example: \cite{ch2:Sanders:2010:CEI}). \label{ch2:1ex} This first example is intented to show how to build a very simple example with -CUDA. The goal of this example is to performed the sum of two arrays and -putting the result into a third array. A cuda program consists in a C code -which calls CUDA kernels that are executed on a GPU. The listing of this code is +Cuda. The goal of this example is to perform the sum of two arrays and +put the result into a third array. A Cuda program consists in a C code +which calls Cuda kernels that are executed on a GPU. The listing of this code is in Listing~\ref{ch2:lst:ex1}. As GPUs have their own memory, the first step consists in allocating memory on -the GPU. A call to \texttt{cudaMalloc} allows to allocate memory on the GPU. The -first parameter of this function is a pointer on a memory on the device -(i.e. the GPU). In this example, \texttt{d\_} is added on each variable allocated -on the GPU meaning this variable is on the GPU. The second parameter represents -the size of the allocated variables, this size is in bits. +the GPU. A call to \texttt{cudaMalloc}\index{Cuda~functions!cudaMalloc} allows +to allocate memory on the GPU. The first parameter of this function is a pointer +on a memory on the device (i.e. the GPU). In this example, \texttt{d\_} is added +on each variable allocated on the GPU, meaning this variable is on the GPU. The +second parameter represents the size of the allocated variables, this size is in +bits. In this example, we want to compare the execution time of the additions of two arrays in CPU and GPU. So for both these operations, a timer is created to -measure the time. CUDA proposes to manipulate timers quick easily. The first -step is to create the timer, then to start it and at the end to stop it. For -each of these operations a dedicated functions is used. +measure the time. Cuda proposes to manipulate timers quite easily. The first +step is to create the timer\index{Cuda~functions!timer}, then to start it and at +the end to stop it. For each of these operations a dedicated functions is used. -In order to compute the same sum with a GPU, the first step consits in -transferring the data from the CPU (considered as the host with CUDA) to the GPU -(considered as the device with CUDA). A call to \texttt{cudaMalloc} allows to +In order to compute the same sum with a GPU, the first step consists in +transferring the data from the CPU (considered as the host with Cuda) to the GPU +(considered as the device with Cuda). A call to \texttt{cudaMemcpy} allows to copy the content of an array allocated in the host to the device when the fourth -parameter is set to \texttt{cudaMemcpyHostToDevice}. The first parameter of the -function is the destination array, the second is the source array and the third -is the number of elements to copy (exprimed in bytes). - -Now that the GPU contains the data needed to perform the addition. In sequential -such addition is achieved out with a loop on all the elements. With a GPU, it -is possible to perform the addition of all elements of the arrays in parallel +parameter is set +to \texttt{cudaMemcpyHostToDevice}\index{Cuda~functions!cudaMemcpy}. The first +parameter of the function is the destination array, the second is the +source array and the third is the number of elements to copy (exprimed in +bytes). + +Now the GPU contains the data needed to perform the addition. In sequential such +addition is achieved out with a loop on all the elements. With a GPU, it is +possible to perform the addition of all elements of the two arrays in parallel (if the number of blocks and threads per blocks is sufficient). In Listing\ref{ch2:lst:ex1} at the beginning, a simple kernel, called \texttt{addition} is defined to compute in parallel the summation of the -two arrays. With CUDA, a kernel starts with the -keyword \texttt{\_\_global\_\_} \index{CUDA~keywords!\_\_shared\_\_} which +two arrays. With Cuda, a kernel starts with the +keyword \texttt{\_\_global\_\_} \index{Cuda~keywords!\_\_shared\_\_} which indicates that this kernel can be called from the C code. The first instruction in this kernel is used to compute the variable \texttt{tid} which represents the thread index. This thread index\index{thread index} is computed according to -the values of the block index (it is a variable of CUDA -called \texttt{blockIdx}\index{CUDA~keywords!blockIdx}). Blocks of threads can -be decomposed into 1 dimension, 2 dimensions or 3 dimensions. According to the -dimension of data manipulated, the appropriate dimension can be useful. In our -example, only one dimension is used. Then using notation \texttt{.x} we can -access to the first dimension (\texttt{.y} and \texttt{.z} allow respectively to -access to the second and third dimension). The -variable \texttt{blockDim}\index{CUDA~keywords!blockDim} gives the size of each -block. +the values of the block index +(called \texttt{blockIdx} \index{Cuda~keywords!blockIdx} in Cuda) and of the +thread index (called \texttt{blockIdx}\index{Cuda~keywords!threadIdx} in +Cuda). Blocks of threads and thread indexes can be decomposed into 1 dimension, 2 +dimensions or 3 dimensions. According to the dimension of data manipulated, the +appropriate dimension can be useful. In our example, only one dimension is used. +Then using notation \texttt{.x} we can access to the first dimension +(\texttt{.y} and \texttt{.z} allow respectively to access to the second and +third dimension). The variable \texttt{blockDim}\index{Cuda~keywords!blockDim} +gives the size of each block. @@ -72,24 +76,24 @@ block. \section{Second example: using CUBLAS} \label{ch2:2ex} -The Basic Linear Algebra Subprograms (BLAS) allows programmer to use performant -routines that are often used. Those routines are heavily used in many scientific -applications and are very optimized for vector operations, matrix-vector -operations and matrix-matrix +The Basic Linear Algebra Subprograms (BLAS) allows programmers to use efficient +routines that are often required. Those routines are heavily used in many +scientific applications and are very optimized for vector operations, +matrix-vector operations and matrix-matrix operations~\cite{ch2:journals/ijhpca/Dongarra02}. Some of those operations seem -to be easy to implement with CUDA. Nevertheless, as soon as a reduction is -needed, implementing an efficient reduction routines with CUDA is far from being +to be easy to implement with Cuda. Nevertheless, as soon as a reduction is +needed, implementing an efficient reduction routines with Cuda is far from being simple. Roughly speaking, a reduction operation\index{reduction~operation} is an operation which combines all the elements of an array and extract a number computed with all the elements. For example, a sum, a maximum or a dot product -are reduction operations. +are reduction operations. In this second example, we consider that we have two vectors $A$ and $B$. First of all, we want to compute the sum of both vectors in a vector $C$. Then we want to compute the scalar product between $1/C$ and $1/A$. This is just an example -which has no direct interest except to show how to program it with CUDA. +which has no direct interest except to show how to program it with Cuda. -Listing~\ref{ch2:lst:ex2} shows this example with CUDA. The first kernel for the +Listing~\ref{ch2:lst:ex2} shows this example with Cuda. The first kernel for the addition of two arrays is exactly the same as the one described in the previous example. @@ -137,12 +141,60 @@ consider we have a squared matrix of size \texttt{size}. So with a 1D array, \texttt{A[i*size+j]} allows us to access to the element of the $i^{th}$ row and of the $j^{th}$ column. -On C2070M Tesla card, this code take 37.68ms to perform the multiplication. On a -Intel Xeon E31245 at 3.30GHz, it takes 2465ms without any parallelization (using -only one core). Consequently the speed up between the CPU and GPU version is -about 65 which is very good regarding the difficulty of parallelizing this code. +With a sequential programming, the matrix multiplication is performed using +three loops. Supposing that $A$, $B$ represent two square matrices and that the +result of the multiplication of $A \times B$ is $C$. The +element \texttt{C[i*size+j]} is computed as follows: +\begin{equation} +C[i*size+j]=\sum_{k=0}^{size-1} A[i*size+k]*B[k*size+j]; +\end{equation} + +In Listing~\ref{ch2:lst:ex3}, in the CPU computation, this part of code is +performed using 3 loops, one for $i$, one for $j$ and one for $k$. In order to +perform the same computation on a GPU, a naive solution consists in considering +that the matrix $C$ is split into 2 dimensional blocks. The size of each block +must be chosen such as the number of threads per block is inferior to $1,024$. + + +In Listing~\ref{ch2:lst:ex3}, we consider that a block contains 16 threads in +each dimension, the variable \texttt{width} is used for that. The +variable \texttt{nbTh} represents the number of threads per block. So to be able +to compute the matrix-matrix product on a GPU, each block of threads is assigned +to compute the result of the product for the elements of this block. The main +part of the code is quite similar to the previous code. Arrays are allocated in +the CPU and the GPU. Matrices $A$ and $B$ are randomly initialized. Then +arrays are transfered inside the GPU memory with call to \texttt{cudaMemcpy}. +So the first step for each thread of a block is to compute the corresponding row +and column. With a 2 dimensional decomposition, \texttt{int i= +blockIdx.y*blockDim.y+ threadIdx.y;} allows us to compute the corresponding line +and \texttt{int j= blockIdx.x*blockDim.x+ threadIdx.x;} the corresponding +column. Then each thread has to compute the sum of the product of the line of +$A$ per the column of $B$. In order to use a register, the +kernel \texttt{matmul} uses a variable called \texttt{sum} to compute the +sum. Then the result is set into the matrix at the right place. The computation +of CPU matrix-matrix multiplication is performed as described previously. A +timer measures the time. In order to use 2 dimensional blocks, \texttt{dim3 +dimGrid(size/width,size/width);} allows us to create \texttt{size/width} blocks +in each dimension. Likewise, \texttt{dim3 dimBlock(width,width);} is used to +create \texttt{width} thread in each dimension. After that, the kernel for the +matrix multiplication is called. At the end of the listing, the matrix $C$ +computed by the GPU is transfered back in the CPU and we check if both matrices +C computed by the CPU and the GPU are identical with a precision of $10^{-4}$. + + +On C2070M Tesla card, this code take $37.68$ms to perform the multiplication. On +a Intel Xeon E31245 at $3.30$GHz, it takes $2465$ms without any parallelization +(using only one core). Consequently the speed up between the CPU and GPU version +is about $65$ which is very good regarding the difficulty of parallelizing this +code. \lstinputlisting[label=ch2:lst:ex3,caption=simple Matrix-matrix multiplication with cuda]{Chapters/chapter2/ex3.cu} +\section{Conclusion} +In this chapter 3 simple Cuda examples have been presented. Those examples are +quite simple and they cannot present all the possibilities of the Cuda +programming. Interested readers are invited to consult Cuda programming +introduction books if some issues regarding the Cuda programming is not clear. + \putbib[Chapters/chapter2/biblio]