Procesamiento de imágenes cerebrales en GPU

ESCUELA TÉCNICA SUPERIOR DE INGENIERÍA INFORMÁTICA GRADO EN INGENIERÍA INFORMÁTICA PROCESAMIENTO DE IMÁGENES CEREBRALES EN GPU NEUROIMAGE PROCESSING ON GPU USING CUDA Realizado por Javier Cabero Guerra Tutorizado por Manuel Ujaldón Martínez Departamento Arquitectura de Computadores UNIVERSIDAD DE MÁLAGA MÁLAGA, SEPTIEMBRE DE 2015 Fecha defensa: El Secretario del Tribunal Resumen: A lo largo del tiempo, el paradigma de computación de propósito general ha evolucionado, produciendo diferentes arquitecturas hardware cuyas caracteristicas son muy distintas. En este trabajo, trataremos de demostrar, a través de distintas aplicaciones pertenecientes al campo del Procesamiento de Imágenes, la diferencia existente entre tres plataformas hardware de Nvidia: dos de la serie de tarjetas gráficas GeForce, la GTX 480 y la GTX 980 y una plataforma de bajo consumo cuyo propósito es el permitir la ejecución de aplicaciones embebidas a la vez que proporcionar una eficiencia extrema: la Jetson TK1. Respecto a los programas de prueba usaremos cinco ejemplos sacados de los CUDA Samples de Nvidia. Estas aplicaciones tienen una relación directa con el procesamiento de imágenes, dado que los algoritmos implicados en ellas tienen similitudes con los aplicados en el campo del registro de imágenes médico. Tras las pruebas, se mostrará cómo la GTX 980 es tanto el dispositivo con mayor rendimiento como el que mayor consume, se verá que la Jetson TK1 es el dispositivo más eficiente de los tres, se enseñará cómo la GTX 480 es la plataforma que más calor produce y aprenderemos otros efectos producidos por la diferencia entre las arquitecturas que hay entre los dispositivos. Palabras claves: CUDA, GPGPU, Jetson TK1, GTX 480, GTX 980, Rendimiento, Consumo, Eficiencia, Procesamiento Imágenes Cerebrales Abstract: As time has passed, the general purpose programming paradigm has evolved, producing different hardware architectures whose characteristics differ widely. In this work, we are going to demonstrate, through different applications belonging to the field of Image Processing, the existing difference between three Nvidia hardware platforms: two of them belong to the GeForce graphics cards series, the GTX 480 and the GTX 980 and one of the low consumption platforms which purpose is to allow the execution of embedded applications as well as providing an extreme efficiency: the Jetson TK1. With respect to the test applications we will use five examples from Nvidia CUDA Samples. These applications are directly related to Image Processing, as the algorithms they use are similar to those from the field of medical image registration. After the tests, it will be proven that GTX 980 is both the device with the highest computational power and the one that has greater consumption, it will be seen that Jetson TK1 is the most efficient platform, it will be shown that GTX 480 produces more heat than the others and we will learn other effects produced by the existing difference between the architecture of the devices. Keywords: CUDA, GPGPU, Jetson TK1, GTX 480, GTX 980, Performance, Power Drawback, Efficiency, Neuroimaging University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering Contents 1 Introduction 11 1.1 The Testbed .................................. 12 2 The GPGPU movement 15 2.1 The GPU Streaming Processor ....................... 15 2.1.1 Advantages and Drawbacks ..................... 16 2.2 Evolution to a General Purpose Architecture ............... 16 2.2.1 Starting Point ............................. 17 2.2.2 GPGPU First Steps . . . . . . . . . . . . . . . . . . . . . . . . . . 18 2.2.3 The Arrival of CUDA . . . . . . . . . . . . . . . . . . . . . . . . . . 19 2.2.4 OpenCL ................................. 20 2.2.5 Last Years and the Future of GPGPU . . . . . . . . . . . . . . . . 21 3 Programming on Architecture Graphics Using CUDA 23 3.1 CUDA (Compute Unified Device Architecture) .............. 24 3.1.1 Software ................................ 24 3.1.2 Firmware ................................ 24 3.1.3 Hardware ................................ 24 3.2 Programming Model ............................. 25 3.2.1 Processing Levels ........................... 25 3.2.2 Streams ................................. 27 3.2.3 Processing Flow . . . . . . . . . . . . . . . . . . . . . . . . . . . . 27 3.3 Hardware Model ............................... 28 3.4 Evolution of the Architecture by Generations ............... 30 Computer Architecture Dept. 7Javier Cabero GuerraComputer Architecture Dept. 7Javier Cabero GuerraComputer Architecture Dept. 7Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering 3.4.1 The First Generation: Tesla (G80 y GT200) . . . . . . . . . . . . . 30 3.4.2 The Second Generation: Fermi (GF100) . . . . . . . . . . . . . . . 32 3.4.3 The Third Generation: Kepler (GK110 y GK210) . . . . . . . . . . 34 3.4.3.1 Dynamic Parallelism ..................... 36 3.4.3.2 Hyper-Q ............................ 37 3.4.4 The Fourth Generation: Maxwell (GM204) . . . . . . . . . . . . . 38 3.4.4.1 Memory improvement .................... 39 3.4.4.2 Atomic operations ...................... 40 4 GTX 480 vs Jetson TK1 vs GTX 980 43 4.1 Introduction .................................. 43 4.1.1 Dissertation Overview ......................... 43 4.1.2 GeForce GTX 480 ........................... 44 4.1.3 Jetson TK1 ............................... 44 4.1.4 GeForce GTX 980 ........................... 46 4.2 Texture filtering ................................ 48 4.2.1 Description ............................... 48 4.2.2 Performance . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 49 4.2.3 Power Draw and Heat Generation . . . . . . . . . . . . . . . . . . 50 4.2.4 Algorithm Efficiency . . . . . . . . . . . . . . . . . . . . . . . . . . 52 4.3 Bilateral filtering ............................... 54 4.3.1 Description ............................... 54 4.3.2 Performance . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 54 4.3.3 Power Draw and Heat Generation . . . . . . . . . . . . . . . . . . 56 4.3.4 Algorithm Efficiency . . . . . . . . . . . . . . . . . . . . . . . . . . 60 4.4 Box filter .................................... 65 4.4.1 Description ............................... 65 4.4.2 Performance . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 67 4.4.3 Power Draw and Heat Generation .................. 67 4.4.4 Algorithm Efficiency . . . . . . . . . . . . . . . . . . . . . . . . . . 69 4.5 Image denoising ................................ 71 Computer Architecture Dept. 8Javier Cabero GuerraComputer Architecture Dept. 8Javier Cabero GuerraComputer Architecture Dept. 8Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering 4.5.1 Description ............................... 71 4.5.2 Performance . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 71 4.5.3 Power Draw and Heat Generation . . . . . . . . . . . . . . . . . . 72 4.5.4 Algorithm Efficiency . . . . . . . . . . . . . . . . . . . . . . . . . . 74 4.6 Post-process GL ................................ 76 4.6.1 Description ............................... 76 4.6.2 Performance . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 76 4.6.3 Power Draw and Heat Generation . . . . . . . . . . . . . . . . . . 77 4.6.4 Algorithm Efficiency . . . . . . . . . . . . . . . . . . . . . . . . . . 78 Conclusions 81 4.6.5 English ................................. 81 4.6.6 Spanish ................................. 83 5 Bibliography 87 Computer Architecture Dept. 9Javier Cabero GuerraComputer Architecture Dept. 9Javier Cabero GuerraComputer Architecture Dept. 9Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering 1 Introduction New computing technologies put an imperative effort on reducing power consumption. The search of low power platforms derives from an older perspective which focused the increment of computers performance. This idea continued until too many resources were necessary to feed the machine. At this moment, an inflection point occurred in the device targets: instead of computational power they started to concentrate on efficiency. Having more performance is not a forgotten objective but it is now driven by a reasonable power budget. One of the main reasons for improving energy efficiency refers to the mobile market. Battery technologies are stuck and cannot improve their energy capacity using the same size at the same cost [45]. The need of saving the little energy available raises, causing a transition from top performance devices to more efficient ones. On the other hand, we have supercomputers, which have an extraordinary Computer Architecture Dept. 11 Javier Cabero GuerraComputer Architecture Dept. 11 Javier Cabero GuerraComputer Architecture Dept. 11 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering the OpenGL Architecture Review Board. It was also based on C, allowing developers to make cross-platform applications that took advantage from most of the new features of GPUs. It was initially introduced as an extension to OpenGL 1.4 and it was officially included in OpenGL 2.0 in 2004. 2.2.2 GPGPU First Steps At the beginning of the 21st century, GPUs had incredibly increased their programming. However, they had only been used for programming graphics applications up to that moment. The first time they were used as general purpose devices was when some researches from the scientific sphere tried to compute more common algorithms with this platforms. In contrast to the conventional implementation of a CPU algorithm, GPU algorithms need some program layers to restructure incoming data, instructions and operators into geometry such that they behave as rendering graphics information. This way, the problem can be computed by the programmable graphics processors. Unfortunately, developers must check that no side effects or changes occur within the graphics pipeline, as it was not designed for this purpose. These tasks required knowledge of the internal architecture, high skill and previous experience. Algorithms Improvement Particle systems Physic simulations Molecular dynamics 2-3 Database queries Data mining Reduction operations 5-10 Signal processing Volume rendering Image processing Biocomputing 10-20 Raytracing 3D visualization +20 Tabla 2.1: Improvement when executing different kinds of parallel algorithms. Since 2003, it was possible to see codes taking advantage of GPU features. Computer Architecture Dept. 18 Javier Cabero GuerraComputer Architecture Dept. 18 Javier Cabero GuerraComputer Architecture Dept. 18 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering 2008 2015 CUDA GPUs 100.000.000 600.000.000 Supercomputers in top500.org 1 75 University courses 60 840 Scientific articles 4.000 60.000 Tabla 2.2: Evolution of CUDA. These programs made a clear difference between the CPU and the GPU, which would increase in the next years because of the developers that gained experience and improved their algorithms. Table 2.1 shows the differences that were observed. 2.2.3 The Arrival of CUDA In 2003, a team of researchers from outside NVIDIA, led by Ian Buck, announced the first programming model that allowed the development of programs on a GPU using a high level language as if it were a general purpose processor. This not only meant facilities when developing GPU code, but also improved performance. NVIDIA knew his incredibly fast hardware had to be accompanied by a software that was at the cutting edge of technology, so they invited the team to join the company and start developing the next big step in the GPGPU paradigm. As an union of hardware and software, NVIDIA released CUDA in 2006 as the first global solution for general purpose computing on GPUs. Some of the improvements were: •Code readability. •Easy to program and shorter development time. •Easy to debug and optimize code. •Independent code of the GPU. •Complex mathematical operations and accurate results. CUDA computing platform provided developers with a C/C++ based system along with several extensions that allowed programmers to implement parallel applications. It also offered alternatives that gave programmers the ability to express parallelism using other high level languages (Fortran, Python ...) and open standards (such as OpenACC directives). The release of CUDA was widely accepted by scientific, academic and developer communities in general. The new parallel programming paradigm brought a Computer Architecture Dept. 19 Javier Cabero GuerraComputer Architecture Dept. 19 Javier Cabero GuerraComputer Architecture Dept. 19 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering Figure 2.2: The OpenCL model. number of improvements that eliminated all the difficulties encountered in the past. In fact, since its arrival day until the present, the CUDA platform has been used in more than 600.000.000 GPUs and 60.000 research applications (see 2.2). 2.2.4 OpenCL At the end of 2008, OpenCL was released as an open alternative to propietary solutions for GPGPU. OpenCL was the product of many years of development by an open software consortium. It was originally conceived by Apple and developed in conjunction with AMD, IBM, Intel and NVIDIA; then it was given to the Khronos Group, who converted it into an open, royalty-free standard. Unlike CUDA, OpenCL is defined as a general purpose programming standard in heterogeneous systems that can run on different architectures such as CPUs, GPUs and FPGAs. OpenCL provides an API for parallel computing and a programming language based on ISO C99 with extensions for data parallelism. The way OpenCL operates is based on a host machine that distributes the workload between all system devices, which are called computational units. This computational units are then divided into multiple processing elements. Although OpenCL is a valid alternative to CUDA, the distance between both is Computer Architecture Dept. 20 Javier Cabero GuerraComputer Architecture Dept. 20 Javier Cabero GuerraComputer Architecture Dept. 20 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering June 2011 June 2012 June 2013 June 2014 NVIDIA Fermi 12 53 31 18 NVIDIA Kepler 0 0 8 28 Intel Xeon Phi 0 1 11 21 ATI Radeon 2 2 3 3 IBM Cell 5 2 0 0 Hybrid 0 0 1 4 Total 19 58 54 74 Tabla 2.3: Evolution of GPUs in TOP500. sometimes tremendous. If the implementation and distribution of work is perfectly adjusted to the target architecture, OpenCL performance should not be much less than CUDA. However, CUDA has not the portability of OpenCL. 2.2.5 Last Years and the Future of GPGPU The programming of GPUs has evolved a lot in the recent years. However, its evolution needed one more step: the increment in scalability of the GPU itself. To do that, clusters of computers arise and more devices interconnect, operating in groups that act as one graphics device. This led to the emergence of the GPGPU movement to gain momentum in the field of high performance computing. The enhancement was not only limited to the appearance of servers and workstations: it also allowed the raise of the number of heterogeneous supercomputers that incorporated the latest generation of GPUs as coprocessors that were in charge of part of the processing work. Table 2.3 shows the evolution of graphics coprocessors in the TOP500 supercomputers list in the last four years. The change to the GPGPU model is relatively recent, so there is still a long way to go. GPUs offer several orders of magnitude greater performance than the CPU when large amounts of data have to be processed, so they are positioned as an alternative to traditional processors and could be considered as the computing engine for the future. Computer Architecture Dept. 21 Javier Cabero GuerraComputer Architecture Dept. 21 Javier Cabero GuerraComputer Architecture Dept. 21 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering 3 Programming on Architecture Graphics Using CUDA Once seen the increase in the popularity of the graphics programming and the importance of the GPGPU (General-Purpose computation on Graphics Processing Units), we are going to focus on the main model, CUDA, and the hardware platform that executes it. Therefore, this chapter presents the main concepts of the graphical programming with CUDA and the highlight parts of the hardware along with the close relationship between hardware and software. Finally, the evolution of the architecture by generation is explained too. Computer Architecture Dept. 23 Javier Cabero GuerraComputer Architecture Dept. 23 Javier Cabero GuerraComputer Architecture Dept. 23 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering 3.1 CUDA (Compute Unified Device Architecture) CUDA [39] is a parallel computing platform and programming model invented by NVIDIA that allows the programmer to deploy task and data parallelism in three different levels: software, firmware and hardware. 3.1.1 Software The first level, software, has diferent ways of writing code and executing it on the GPU. •Programming language: Although C/C++ is the most usual high-level language for developing code on CUDA, there are also APIs (Application Programming Interface) for other popular languages like Fortran, Java and Python. •Optimized libraries: There are many libraries that allow us to perform GPUaccelerated code with just a few lines of code (cuBLAS, cuFFT, Thrust, etc.). •Compiler directives: Another possibility for accelerating applications is to use standard directives with an open initiative called OpenACC. Programmers identify the data parallelism within the code through simple compiler directives, moving the bulk of the parallelization effort to the compiler. However, this automatic approaches have always a performance payoff. 3.1.2 Firmware NVIDIA offers a driver that is compatible with the one responsible for rendering. This driver has simple APIs for controlling the memory, the device and more. 3.1.3 Hardware Lastly, CUDA provides the programmer with the possibility of using the GPU for general purpose programming by means of a large amount of heterogeneous cores inside multiprocessors which are enveloped by a memory hierarchy. This point is explained in more details in section 3.3. Computer Architecture Dept. 24 Javier Cabero GuerraComputer Architecture Dept. 24 Javier Cabero GuerraComputer Architecture Dept. 24 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering 3.2 Programming Model In the next paragraphs, the CUDA programming model is presented, taking C as the baseline language. CUDA is an extension of the C language which supplies tools for parallel programming on GPU. In this model, the GPU acts like a coprocessor and only executes a minimal fraction of code, the rest is handled by the CPU. This process is transparent for the developer due to the CUDA compiler driver (NVCC) and divides the code in two sections: 1. With the GPU fraction generates PTX 1code files. This code is compatible with different devices so that it is decoupled from hardware implementation. 2. The CPU part is parsed to C compiler code in order to create object files. On Linux, we can use GCC (GNU Compiler Collection). On the other hand, CL (the Microsoft Visual Studio compiler) can be used on Windows platform. Then the linker builds a CPU-GPU executable with the files of both parts. For NVCC to be able to divide the code, it is necessary to introduce new syntax elements are used by the programmer to define kernels. Kernels are C functions that contain code for one thread only, then this code is executed on multiple threads in the graphics device automatically. These threads are very thin and the context switch is immediate. 3.2.1 Processing Levels One of the syntax elements used to define kernels is the __global__ declaration specifier. In addition, the number of threads of each kernel is indicated within «<...»> through two parameters. A thread is identified within the kernel in response to the following hierarchy. 1. The threads are organized in blocks. Each thread has an identifier that is accessible within the kernel by means of the built-in threadIdx variable. 2. Likewise, these blocks are grouped within a grid and, like the threads, to each block is given a unique identifier within the kernel, blockIdx. Both grid and thread blocks can be unidimensional, bidimensional or tridimensional and their size is indicated by the programmer under certain limitations. 1PTX is a low-level Parallel Thread eXecution virtual machine and instruction set architecture (ISA). Computer Architecture Dept. 25 Javier Cabero GuerraComputer Architecture Dept. 25 Javier Cabero GuerraComputer Architecture Dept. 25 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering Block (0, 0) Block (1, 0) Block (2, 0) Block (0, 1) Block (1, 1) Block (2, 1) Grid Thread (0, 0) Thread (1, 0) Thread (2, 0) Thread (3, 0) Thread (0, 1) Thread (1, 1) Thread (2, 1) Thread (3, 1) Thread (0, 2) Thread (1, 2) Thread (2, 2) Thread (3, 2) Block (1,1) Figure 3.1: Graphical representation of a grid with six thread blocks, each one composed of 12 threads. NVIDIA Corporation [39] The dimension of the thread block and the grid are accessible within the kernel through variables blockDim and gridDim respectively. This hierarchy gives to CUDA an important feature: the scalability.2 In addition, threads are grouped within 32 elements groups called warp 3, that is the atomic execution unit, and they are executed in unpredictable order although they could be synchronized if this is necessary. A warp executes one common instruction at a time for all threads, therefore to obtain the maximum efficiency is necessary that all threads within the warp have the same execution path. If due to data-dependence the execution path of a warp is bifurcated, the execution of each branch is serialized disabling the threads that doesn’t participate on each branch. When all paths complete, the threads converge back to the same execution path. This serialization of the execution only occurs within a warp, two differents warps are able to execute distinct paths simultaneously. In the same way blocks are executed in free order too but in contrast they can’t be synchronized. In addition, a thread is able to communicate only with other threads within the same thread block. All those details have to be handled with care by the programmer to guarantee the corretness of the parallel code. 2The code is able to run on any number of cores without recompiling. 3The number of threads per warp could change on future generations of GPUs. Computer Architecture Dept. 26 Javier Cabero GuerraComputer Architecture Dept. 26 Javier Cabero GuerraComputer Architecture Dept. 26 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering 3.2.2 Streams Since the appearance of the second generation of graphics cards, CUDA allows to execute kernels concurrently by means of streams. A stream is a sequence of commands that execute in order. The execution of these commands are out of order with respect to other streams, although CUDA provides functions to synchronize them. By default, all kernels are executed within the same stream. To create a new stream, CUDA C offers a new data type, cudaStream_t, and a new constructor, cudaStreamCreate(). The next code is an example of an array with three streams: cudaStream_t stream[2]; for (int i= 0; i< 2; ++i) cudaStreamCreate(&stream[i]); Kernels are assigned to a stream through the fourth parameter of the kernel launch. The four parameters are: 1. Amount of thread blocks into a grid. 2. Number of threads within a thread blocks. 3. Shared memory allocation size per thread block in bytes. 4. Stream ID. The maximum amount of concurrently streams depends on the generation (16 streams for Fermi and 32 for Kepler). The details of stream concurrence is explained in Section 3.4.3.2 with more detail. 3.2.3 Processing Flow As already mentioned in section 3.2, on CUDA, the GPU (device) acts like a coprocessor of the CPU (host) but with its own memory. Because of that, it is necessary to move the data from host memory to device memory and vice versa. As a result CUDA model has a simple processing flow composed of three steps [17]: 1. Copy the input data from host memory to device memory. 2. Load the program on GPU and run, the data are placed in cache memory to enhance the performance. 3. Move the results from GPU memory to CPU memory. Computer Architecture Dept. 27 Javier Cabero GuerraComputer Architecture Dept. 27 Javier Cabero GuerraComputer Architecture Dept. 27 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering Fermi Memory Hierarchy Thread Shared Memory L1 Cache L2 Cache DRAM Kepler Memory Hierarchy Shared Memory Thread L1 Cache Read-Only Data Cache L2 Cache DRAM Figure 3.7: Fermi and Kepler memory hierarchy. NVIDIA Corporation [31, pg. 19] and [36, pg. 13] accesses. Moreover, this generation incorporates 768KB of L2 cache shared by all stream processors. In the left side of Figure 3.7 the diagram of this hierarchy is visible. 3.4.3 The Third Generation: Kepler (GK110 y GK210) Following the trend introduced by Fermi, Kepler increases the number of cores per SM and reduces the amount of multiprocessors. Even though the GK110 is not the first chip of Kepler architecture, this section is centered in the GK110 and newer ones as they are the most widely used. The quantity of cores per SM is the same in the distinct incremental improvement of the architecture, although the number of stream multiprocessors changes from one to another. Thus, Table 3.1 shows the differents versions and its main features. The Kepler’s SMs (called SMXs) have 192 single precision CUDA cores, and each core has fully pipelined floating-point and integer arithmetic logic units. In addition, these SMs increase the double-precision computation capacity with 64 dedicated units. More over, the GK110 has 32 LD/ST units, doubling the amount of load and store units available in the Fermi architecture. Finally, the SMXs have 32 Special Function Units (SFU). Each SMX has four warp schedulers with two dispatch instruction units each. This allows up to eight warps to be issued and executed concurrently. Computer Architecture Dept. 34 Javier Cabero GuerraComputer Architecture Dept. 34 Javier Cabero GuerraComputer Architecture Dept. 34 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering Memory ControllerMemory Controller Memory Controller Memory ControllerMemory Controller Memory Controller L2 Cache SMX SMX SMX SMX SMX SMX SMX SMX SMXSMXSMX SMX SMXSMXSMX GigaThread Engine PCI Express 3.0 Host Interface Figure 3.8: Kepler GK110 full chip block diagram. NVIDIA Corporation [36, pg. 06] Kepler also follows the memory hierarchy of Fermi, although the texture memory is now accessible for GPGPU as only-read memory of 48KB. In addition, this generation improve all layers of memory: •Register Bank. The amount of 32-bit register per multiprocessor grows until 64K. •Shared Memory and L1 cache. Apart from to the configuration modes of shared memory were seen in Section 3.4.2, a new mode is added in this generation: 32KB for both. •L2 cache. The amount of memory in this layer is doubled to 1536KB. Additionally, the L2 cache on Kepler offers up to 2x of the bandwidth per clock available on Fermi. [36] The GK210 and GK110 have their features explained in Section 3.4.3.1 and Section 3.4.3.2. One as much as the other are Kepler architectures but the GK210 has more resource on-chip than GK110. Thus, both chips share the same amount of core per SMX but the GK210 has 128K register of 32-bit per SMX and 128KB of shared memory/L1 cache with the configurations below: •112KB shared memory + 16KB L1 cache Computer Architecture Dept. 35 Javier Cabero GuerraComputer Architecture Dept. 35 Javier Cabero GuerraComputer Architecture Dept. 35 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering Interconect Network (64 KB Shared Memory / L1 Cache GK110) | (128 KB Shared Memory / L1 Cache GK210) 48 KB Read-Only Data Cache Tex Tex Tex Tex Tex Tex Tex Tex Tex Tex Tex Tex Tex Tex Tex Tex Register File (65,536 x 32-bit GK110) | (131,072 x 32-bit GK210) Core Core Core Core Core CoreDP Unit DP Unit LD/ST SFU Core Core Core Core Core CoreDP Unit DP Unit LD/ST SFU Core Core Core Core Core CoreDP Unit DP Unit LD/ST SFU Core Core Core Core Core CoreDP Unit DP Unit LD/ST SFU Core Core Core Core Core CoreDP Unit DP Unit LD/ST SFU Core Core Core Core Core CoreDP Unit DP Unit LD/ST SFU Core Core Core Core Core CoreDP Unit DP Unit LD/ST SFU Core Core Core Core Core CoreDP Unit DP Unit LD/ST SFU Core Core Core Core Core CoreDP Unit DP Unit LD/ST SFU Core Core Core Core Core CoreDP Unit DP Unit LD/ST SFU Core Core Core Core Core CoreDP Unit DP Unit LD/ST SFU Core Core Core Core Core CoreDP Unit DP Unit LD/ST SFU Core Core Core Core Core CoreDP Unit DP Unit LD/ST SFU Core Core Core Core Core CoreDP Unit DP Unit LD/ST SFU Core Core Core Core Core CoreDP Unit DP Unit LD/ST SFU Core Core Core Core Core CoreDP Unit DP Unit LD/ST SFU Core Core Core Core Core CoreDP Unit DP Unit LD/ST SFU Core Core Core Core Core CoreDP Unit DP Unit LD/ST SFU Core Core Core Core Core CoreDP Unit DP Unit LD/ST SFU Core Core Core Core Core CoreDP Unit DP Unit LD/ST SFU Core Core Core Core Core CoreDP Unit DP Unit LD/ST SFU Core Core Core Core Core CoreDP Unit DP Unit LD/ST SFU Core Core Core Core Core CoreDP Unit DP Unit LD/ST SFU Core Core Core Core Core CoreDP Unit DP Unit LD/ST SFU Core Core Core Core Core CoreDP Unit DP Unit LD/ST SFU Core Core Core Core Core CoreDP Unit DP Unit LD/ST SFU Core Core Core Core Core CoreDP Unit DP Unit LD/ST SFU Core Core Core Core Core CoreDP Unit DP Unit LD/ST SFU Core Core Core Core Core CoreDP Unit DP Unit LD/ST SFU Core Core Core Core Core CoreDP Unit DP Unit LD/ST SFU Core Core Core Core Core CoreDP Unit DP Unit LD/ST SFU Core Core Core Core Core CoreDP Unit DP Unit LD/ST SFU Instruction Cache Warp Scheduler Dispatch Unit Dispatch Unit Warp Scheduler Dispatch Unit Dispatch Unit Warp Scheduler Dispatch Unit Dispatch Unit Warp Scheduler Dispatch Unit Dispatch Unit SMX Figure 3.9: SMX with 192 single-precision CUDA cores, 64 double-precision units, 32 SFU and 32 LD/ST units. NVIDIA Corporation [36, pg. 08] •96KB shared memory + 32KB L1 cache •48KB shared memory + 80KB L1 cache •The anterior amounts reversed. 3.4.3.1 Dynamic Parallelism Until the GK110 was created, the GPU acted like CPU a coprocessor with high speed-up factors, but low autonomy. Now, with dynamic parallelism, the GPU can generate new work for itself. It does not need to interrupt and wait the launch of new kernels in the CPU, create the events and threads required to control dependencies, synchronize the results and control the task scheduling [3]. Figure 3.10 shows an example about how dynamic parallelism behaves releasing work Computer Architecture Dept. 36 Javier Cabero GuerraComputer Architecture Dept. 36 Javier Cabero GuerraComputer Architecture Dept. 36 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering CPU Fermi GPU CPU Kepler GPU GPU Adapts to Data, Dynamically Launches New Threads Dynamic Parallelism Figure 3.10: With Dynamic Parallelism the GPU can generate new work for itself. NVIDIA Corporation [36, pg. 15] from the CPU. This new feature allows the programmer to use recursive techniques in its algorithms. Due to this, the developer is able to make algorithms that were impossible to achieve on FERMI such as quicksort, nested loops with differing amounts of parallelism or even dynamically setting up a grid for a numerical simulation focusing in the interesting zones without an expensive pre-processing. On Fermi, Only the host sends a grid to the CUDA Work Distributor (CWD) and this distributes the blocks among the differents SM. On Kepler, it is necessary a new unit for the management of both device and host grids. This component, called Grid Management Unit (GMU), processes the grids received from CPU and GPU and sends them to CWD. Then, the work distributor, which accepts up to 32 grids, sends the blocks to the SMX. In addition, the GMU can pause the dispatching of new grids due to the two-ways link. In Figure 3.11 it can seen both Fermi and Kepler workflow. 3.4.3.2 Hyper-Q On Fermi until 16 streams could be launched concurrently, but they are implemented underneath using a single queue, only the end of a stream and the start of other could be executed at the same time. On Kepler until 32 streams can be really executed concurrently due to that each stream is managed independently on Computer Architecture Dept. 37 Javier Cabero GuerraComputer Architecture Dept. 37 Javier Cabero GuerraComputer Architecture Dept. 37 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering CUDA-Created Work SMX SMX SMX SMX Work Distributor Actively dispatching grids 32 Active Grids Two-way link allows pausing dispatch 1000’s of pending grids Pending & suspend grids Grid Management Unit Stream Queues Ordered queues of grids Kepler Workflow SM SM SM SM Work Distributor Tracks blocks issued from grids 16 Active Grids One-way Flow Stream Queues Ordered queues of grids Fermi Workflow Figure 3.11: Fermi (left side) and Kepler (right side) workflow. NVIDIA Corporation [36, pg. 19] a different hardware queue. In addition, this allow for executing a streams in parallel that other stream coming from the same or other CUDA program, MPI process or POSIX thread. 3.4.4 The Fourth Generation: Maxwell (GM204) The new generation is focused on maximizing the performance per consumed watt. Thus, NVIDIA has reorganized the internal components of the multiprocessors (SMMs). Now, these are splited in four part. Each CUDA cores processing block contains: 1. 32 int and floating points units (128 per SMM). 2. 1 double precision unit (4 per SMM). 3. 8 Load/Store units (32 per SMM). 4. 8 Special Functions Unit (SFU) (32 per SMM). In addition, each split contains a warp scheduler, which is capable of dispatching two instruction per warp at every clock. This configuration aligns with warp size, making it easier to use efficiently. Computer Architecture Dept. 38 Javier Cabero GuerraComputer Architecture Dept. 38 Javier Cabero GuerraComputer Architecture Dept. 38 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering SMM Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core SFU SFU SFU SFU SFU SFU SFU SFU LD/ST LD/ST LD/ST LD/ST LD/ST LD/ST LD/ST LD/ST Register File (16,384 x 32-bit) Instruction Buffer Warp Scheduler Dispatch Unit Dispatch Unit Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core SFU SFU SFU SFU SFU SFU SFU SFU LD/ST LD/ST LD/ST LD/ST LD/ST LD/ST LD/ST LD/ST Register File (16,384 x 32-bit) Instruction Buffer Warp Scheduler Dispatch Unit Dispatch Unit Texture/L1 Cache Tex Tex Tex Tex Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core SFU SFU SFU SFU SFU SFU SFU SFU LD/ST LD/ST LD/ST LD/ST LD/ST LD/ST LD/ST LD/ST Register File (16,384 x 32-bit) Instruction Buffer Warp Scheduler Dispatch Unit Dispatch Unit Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core Core SFU SFU SFU SFU SFU SFU SFU SFU LD/ST LD/ST LD/ST LD/ST LD/ST LD/ST LD/ST LD/ST Register File (16,384 x 32-bit) Instruction Buffer Warp Scheduler Dispatch Unit Dispatch Unit Texture/L1 Cache Tex Tex Tex Tex PolyMorph Engine 3.0 Vertex Fetch Tessellator Viewport Transform Attribute Setup Stream Output Instruction Buffer 96KB Shared Memory Figure 3.12: GM204 SMM Diagram (GM204 also features 4 DP units per SMM, which are not depicted on this diagram). NVIDIA Corporation [41, pg. 08] 3.4.4.1 Memory improvement The memory hierarchy has changed too, now the shared memory doesn’t share the block with the L1 cache. The L1 caching function is now shared with the texture catching function. The size of shared memory grows to 96KB, although this is limited to 48KB per thread block [23]. Finally, the size of L2 cache is 2MB on GM204. Other improvement which is implemented on Maxwell is the memory compression. To reduce DRAM bandwidth demands, NVIDIA GPUs make use of lossless compression techniques as data is written out to memory. This profit is doubled when clients, such as the Texture Unit, read later the data. Computer Architecture Dept. 39 Javier Cabero GuerraComputer Architecture Dept. 39 Javier Cabero GuerraComputer Architecture Dept. 39 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering Memory ControllerMemory Controller Memory Controller Memory Controller SMM SMM SMM SMM Raster Engine GPC SMM SMM SMM SMM Raster Engine GPC SMM SMM SMM SMM Raster Engine GPC SMM SMM SMM SMM Raster Engine GPC L2 Cache GigaThread Engine PCI Express 3.0 Host Interface Figure 3.13: Maxwell GK204 full chip block diagram.NVIDIA Corporation [41, pg. 06] 3.4.4.2 Atomic operations Maxwell introduces native shared memory atomic operations for 32-bit integers and native shared memory 32-bit and 64-bit compare-and-swap (CAS), which can be used to implement other atomic functions with reduced overhead compared to the Fermi and Kepler methods. This should make it much more efficient to implement things like list and stack data structures shared by the threads of a block [41]. Computer Architecture Dept. 40 Javier Cabero GuerraComputer Architecture Dept. 40 Javier Cabero GuerraComputer Architecture Dept. 40 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering 4 GTX 480 vs Jetson TK1 vs GTX 980 4.1 Introduction 4.1.1 Dissertation Overview NVIDIA Corporation has made a huge step into green computing. Kepler was the first generation that defines itself as an architecture concerned about efficiency. His antecessor, Fermi, was presented as a powerful parallel computing architecture. The predecessor of Kepler is Maxwell generation, which endorses the focus on efficiency. In this section, we ilustrate this transition with a comparison between Fermi, Kepler and Maxwell. However, the devices corresponding to each generation are not only graphics cards. The Kepler one is an embedded platform designed to be extremely efficient. Next sections will introduce them and their features. Computer Architecture Dept. 43 Javier Cabero GuerraComputer Architecture Dept. 43 Javier Cabero GuerraComputer Architecture Dept. 43 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering performance in the expensive algorithms relative to the cheap ones. This makes the Jetson TK1 the device with higher scalability, as the decay in performance on work overload is superior to the GeForce devices for this example. 4.1.11 Texture Filtering Performance (Ubuntu). The same CUDA sample was run in Ubuntu using the GTX 980 (see Figure 4.14.1.11). The results shown a 8% and 6% higher performance for Nearest and Bilinear filters, respectively. Despite of this, the expensive filters do not improve their throughput, having less than 1% lower performance or approximately the same pixels per second. This difference is not enough to question the importance of using the same Operating System to compare the results, as the numbers do not vary widely with the OS but more with the underlying hardware. 4.2.3 Power Draw and Heat Generation In the GTX 980, Bicubic filter has the least power consumption with 115 W. Nearest and Bilinear filters are more stable around 118 W, as well as Fast-Bicubic. Lastly, Catmull-Rom is the highest with 120 W. These differences in wattage are not very important (they are in a range of 5 W wide), meaning that the graphics card operates, in terms of power, more or less the same with each of the filters. However, the small differences reveal how expensive filters consume the most. Bicubic filter is an exception to this rule as it has the lowest power drawback of all. It also has the least pixel processing rate (in GTX 480 and GTX 980 devices). These two indicators prove that the graphics card is not able perfectly fit the work load of the algorithm. In the GTX 480, the first two algorithms have the greatest power draw. Computer Architecture Dept. 50 Javier Cabero GuerraComputer Architecture Dept. 50 Javier Cabero GuerraComputer Architecture Dept. 50 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering 4.1.12 Texture Filtering Power Draw. All the three platforms have small power differences among the different filters. In the GTX 980, the window size is 5 W, the GTX 480 has 7.3 W and the Jetson TK1 0.9 W. The device that required more power is the GTX 980 and the one consuming less is the Jetson TK1. 4.1.13 Texture Filtering Heat Generation. Heat generation in this example showed that filters with less throughtput are the ones that require less power. The harder the algorithm is, the lower the temperature gets. GTX 980 has around 12% less heat generation than the GTX 480. Both devices are cooler when computing the last filters. In general, a strong correlation occurs between the power consumption and the heat generation. Despite of this, the graphics card fan could cause the heat to perform differently, as it dynamically Computer Architecture Dept. 51 Javier Cabero GuerraComputer Architecture Dept. 51 Javier Cabero GuerraComputer Architecture Dept. 51 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering increases its speed depending on the work load. This assumption of the correlation between the heat and the power reflects on GTX 480 and Jetson TK1, but not on the Maxwell graphics card. 4.2.4 Algorithm Efficiency We now compare the efficiency of the devices (see Figure 4.14.1.14). GTX 980 has approximately 220 MPixel/sec per Watt for the first two algorithms while the GTX 480 gets a 180 MPixel/sec per Watt rate. The GTX 980 has higher efficiency also in the expensive algorithms, around 64 MPixel/sec per Watt against the 36 MPixel/sec per Watt of the other device. The conclusion is that, for the cheap algorithms, the GTX 980 is 20% more efficient than the GTX 480 and for the expensive ones, this difference windens to 70%. 4.1.14 Texture Filtering Power Efficiency. Jetson TK1 has the greatest power efficiency in all filters, with a 60% and 50% higher efficiency than the GTX 980 in the first two filters and approximately a 50%, 60% and a 68% in the last three. In the previous section, heat generation for the three devices was shown. The GTX 980 was proven to generate less heat than the GTX 480 and the Jetson TK1 again less than the GTX 980. Jetson TK1 has a very low power consumption, being just 6% and 4% of the power drawn by the GTX 480 and the GTX 980, respectively. However, the heat the device generates doesn’t hold this proportions but much larger ones. Jetson TK1 generates around 53% and 60% of the heat GTX 480 and GTX 980 generates, respectively. The performance increment do not compare to Computer Architecture Dept. 52 Javier Cabero GuerraComputer Architecture Dept. 52 Javier Cabero GuerraComputer Architecture Dept. 52 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering this proportion. Because of this, the heat generation efficiency chart was made to show the heat generated by each performance unit. 4.1.15 Texture Filtering Heat Generation Efficiency. Heat generation efficiency for the Jetson TK1 is now the worst of all devices. It is more efficient in terms of power, but it is fairly warmer proportionally, specially on the last three filters, with 1 degree generated for each 0.11 MPixel/sec that it achieves. Figure 4.14.1.15 shows that the most efficient device is the GTX 980, having 1 degree generated for each 0.01 MPixel/sec. Thus, the GTX 980 has a 9% of the Jetson TK1 heat generation rate (percentage of the Bicubic filter). The most powerful device is the GTX 980, the one that generates more heat is the GTX 480 and the more power efficient is the Jetson TK1. These assertions should be true for all the CUDA samples provided here. The configurations and roofline models will show how hardware systems have evolved to increase their performance and to do a more efficient computation. Further information and examples about this section can be found in [5]. Computer Architecture Dept. 53 Javier Cabero GuerraComputer Architecture Dept. 53 Javier Cabero GuerraComputer Architecture Dept. 53 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering 4.3 Bilateral filtering 4.3.1 Description This is a special type of non-linear filter that smooths out textures but preserves any edges found on it. It is based on four parameters: gaussian delta, euclidean delta, the filter radius and the number of filter iterations. If the value of the second one, the euclidean delta, is high, most parts of the output texture will be filtered away as the edge-preserving nature declines. When this parameter tends to infinite the resulting filter is a gaussian one. Blur effect intensifies with a larger gaussian delta. Having a small euclidean number while incrementing the number of iterations will produce flatter colors without blurring edges, creating a cartoon effect. We use an example of a still life paint, giving the unfiltered and filtered output using the just described technique, to demonstrate the filter result. 4.1.16 Original image. 4.1.17 Filtered image. Note how, in the second image, contours are preserved while all other colors flatten. This is why it is called cartoon effect. The euclidean value for the filtered example is 0.12, a low one to make it very edge-preserving. Gaussian was set to 4 and the filter iteration count and radius size were 5. 4.3.2 Performance To perform the tests, the euclidean value was fixed around 1 and the gaussian at 2. These parameters won’t affect performance as hard as the radius and the iteration count. From 1 to 12, step size 4, values for both radius and iterations were combined into a cartesian product. The performance of the Bilateral filtering is measured using the application framerate or Frames Per Second (FPS). All results Computer Architecture Dept. 54 Javier Cabero GuerraComputer Architecture Dept. 54 Javier Cabero GuerraComputer Architecture Dept. 54 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering can be found in Figures 4.14.1.18,4.14.1.19 and 4.14.1.20, respectively, for GTX 480, Jetson TK1 and GTX 980 (the iterations axis is inverted for clarity). 4.1.18 GTX 480 performance. 4.1.19 Jetson TK1 performance. GTX 480 performance (see Figure 4.14.1.18) drops quickly when increasing the number of iterations. The decay simulates a logarithmic curved surface. GTX Computer Architecture Dept. 55 Javier Cabero GuerraComputer Architecture Dept. 55 Javier Cabero GuerraComputer Architecture Dept. 55 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering 4.1.20 GTX 980 performance. 980 manages to keep up the 60 FPS (framerate limit) further than the GTX 480, with half of the configurations computed at that framerate. The lowest value for this device is the radius 12, 12 iterations configuration which is 15 FPS. In the same configuration, GTX 480 drops below the mark down to 4 FPS. This hardest case makes the GTX 980 to be 3.75x faster than the GTX 480. On Jetson TK1, the 1-1 parameter configuration provides further performance than the limit, as there wasn’t any on the application. It has the least performance and scalability of all the devices, as it drops to less than 10 FPS in most of the configurations. 4.3.3 Power Draw and Heat Generation From previous executions, we obtained power consumption along with generated heat. The depth axis sequence has its standard order and the radius axis is inverted, as power/heat data is better visualized this way. In the GTX 480, the least wattage is 35.4 W and the maximum is 77.6 W. The window size for the power consumption is then 42.2 W. The minimum wattage is close to the idle value and the maximum to the limit achieved in these samples. Because of this, we say that the current sample is complete in terms of power consumption. Computer Architecture Dept. 56 Javier Cabero GuerraComputer Architecture Dept. 56 Javier Cabero GuerraComputer Architecture Dept. 56 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering 4.1.21 GTX 480 Power Draw. 4.1.22 Jetson TK1 Power Draw. Jetson TK1 is, again, the device that consumes the least. Despite the board is connected to a monitor and a keyboard, the power draw never raises higher than 4.5 W. The minimum power consumption for this device in the sample execution is 2.3 W. The window size is then at 2.2 W. Computer Architecture Dept. 57 Javier Cabero GuerraComputer Architecture Dept. 57 Javier Cabero GuerraComputer Architecture Dept. 57 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering 4.1.23 GTX 980 Power Draw. For the GTX 980, the power window size is wider: 90.4 W. The parameter configuration that has the maximum consumption is not the 12-12 but the 8-8. This configuration will have also a great heat generation relative to its position in the configurations set. GTX 980 is having more power draw than any of the devices, not only in the hardests cases but in all the other ones, same as in the previous CUDA sample. The fact that GTX 980 consumes more is because of the hardware quantity: it has 2048 CUDA cores compared to the 480 GTX 480 has. Is it worth it to have higher power consumption? In section 4.3.4 we will discuss this question in more detail. As this CUDA sample is able to get the best of the devices in all the metrics provided here, the temperature will show the greatest values of each of the platforms. Remember that all the tests were performed on summer with almost 30 degrees air temperature. The cooling system for the GeForce graphics cards is the one from a personal computer tower, with one fan in front, another in the back and the CPU one. Jetson TK1 performed all test on its own. GTX 480 has high heat generation (see Figure 4.14.1.24). On the 12-12 configuration, the sensor marked 91 degrees. All the values that have at least one of the parameters set to 1 have considerable less heat generation than the rest, which ranges from 86 to 91 degrees while most configurations of the first have less than 84 degrees. Jetson TK1 degrees go from 40 to 51 (see Figure 4.14.1.25). The window size Computer Architecture Dept. 58 Javier Cabero GuerraComputer Architecture Dept. 58 Javier Cabero GuerraComputer Architecture Dept. 58 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering 4.1.24 GTX 480 Temperature. 4.1.25 Jetson TK1 Temperature. is smaller than the one from GTX 480, just 11 degrees (GTX 480 window had a size of 16). Neither of the GeForce devices have a temperature value less than the maximum Jetson TK1 achieves. For GTX 980, the minimum value is 55 and the maximum 80 (see Figure Computer Architecture Dept. 59 Javier Cabero GuerraComputer Architecture Dept. 59 Javier Cabero GuerraComputer Architecture Dept. 59 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering 4.1.33 Original Lenna. 4.1.34 Radius: 10 Passes: 2. 4.1.35 Radius: 20 Passes: 4. Computer Architecture Dept. 66 Javier Cabero GuerraComputer Architecture Dept. 66 Javier Cabero GuerraComputer Architecture Dept. 66 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering 4.4.2 Performance To test the performance of the algorithm four configuration were used (all with radius size 4) in which the number of passes had the values of 1, 8, 16 and 32. Figure 4.14.1.36 show the graph with the obtained results. 4.1.36 Box Filter Performance. Maxwell architecture has again the best performance related to its predecesors. The lowest performance is the one of the Jetson TK1. When there is only one filter pass the framerate values of the devices are more balanced. Additionaly, the framerate of the GTX 980 is only dropped to 30.2 FPS in the last configuration, which is an acceptable value for a visual application (it is 2.51 times the performance of the GTX 480). The executions performance in the rest of platforms fall down the 30.2 FPS at 8 and 16 iterations. Jetson TK1 will never achieve the 60 FPS limit in this application. 4.4.3 Power Draw and Heat Generation Both power consumption and heat generated in this example are lower than in the previous cases. Jetson TK1 has the same power consumption in the least passes count than in the highest, meaning that the application status doesn’t affect too much to its power draw. GeForce graphics cards increase their consumption with higher number of passes as usual (see Figure 4.14.1.37). In the GTX 480, the power ranges from 30.5 W to 53.2 W. For the GTX 980, it goes from 44.2 W to 72.9 W. We observe how GTX 480 energy consumption is far more stable than the one of the GTX 980. This second device draws more energy Computer Architecture Dept. 67 Javier Cabero GuerraComputer Architecture Dept. 67 Javier Cabero GuerraComputer Architecture Dept. 67 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering 4.1.37 Box Filter Power Draw. than the other platforms in all configurations. GTX 980 has quite small power increment from configuration 16 to 32. Same event occurs on the GTX 480 starting from configuration 8. The point where consumption doesn’t increase as fast as in the first configurations denotes how the device fullfils its main power requirements and doesn’t need to warm up for the rest of the work to be done: the biggest step is the one that makes the machine change from idle to fully working. The temperature values shown in Figure 4.14.1.38 behave similar to those of power drawback from Figure 4.14.1.37. Here, Jetson TK1 does show an increment of temperature when it has to perform more filter passes. The increment is, however, not very significant (0.3 degrees starting from 42.6). 4.1.38 Box Filter Temperature. Computer Architecture Dept. 68 Javier Cabero GuerraComputer Architecture Dept. 68 Javier Cabero GuerraComputer Architecture Dept. 68 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering GTX 480 varies from 83 to 90 degrees. As in the power consumption chart, the greater increment in heat generation is from the 1 pass configuration to the 8 one. The rest are all 1 degree increments. GTX 980 behaves the same way, incrementing from 54 to 63 and then from 63 to 66 in the first configurations, but just to 67 in the last one. 4.4.4 Algorithm Efficiency In Figure 4.14.1.39, similar results to those of the bilateral filter are shown: the best power efficiency is provided by the lighter configurations and the worst by the hardest ones. 4.1.39 Box Filter Power Efficiency. 4.1.40 Box Filter Heat Generation Efficiency. Computer Architecture Dept. 69 Javier Cabero GuerraComputer Architecture Dept. 69 Javier Cabero GuerraComputer Architecture Dept. 69 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering GTX 480 has again higher efficiency than the GTX 980 on the first configuration and lower in the rest. The highest efficiency is provided by the Jetson TK1 in the 1 pass configuration, with 14.6 FPS per Watt, a much higher value than the GeForce graphics cards. About the efficiency on heat generation (see Figure 4.14.1.40), Jetson TK1 has not the lowest in every of the configurations: on the first one, GTX 480 has the highest heat generation. In the rest of configurations, the order from better to worse efficiency is GTX 980, GTX 480 and Jetson TK1. GTX 480 and GTX 980 start in the first configuration with similar efficiency values (1.3 and 0.9, respectively). The difference between them is 0.4, but as the number of passes increases, the different also does. At the end, this value is 5.3 degrees per FPS. Jetson TK1 decreases its heat generation efficiency much faster than the GeForce devices, reaching 22.2 degrees per FPS in the last configuration. Computer Architecture Dept. 70 Javier Cabero GuerraComputer Architecture Dept. 70 Javier Cabero GuerraComputer Architecture Dept. 70 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering 4.5 Image denoising 4.5.1 Description Now we attend at a common problem in image processing known as image denoising. An image that travels through a network can have errors as some pixels may be corrupted during the data transmission (data noise). Hard disks can also make read/write mistakes, but this is less frequent. Situations were these errors tend to occur are those in which the message is not travelling through a wire (e.g. air) and the signal is too weak for the receiver. Denoising algorithms become critical in space programs or military scenarios where perfect information about orders or numbers is required. In the context of our example we will concentrate on the denoising of a picture. These errors manifest in the picture as abnormal color dots because of the radical pixel data change. By looking at it, it is not hard to realise that some of the pixels are incorrect, as they may be very different from their neighbourhood (e.g. red dots in the purple jersey of picture 4.14.1.41). In this case, we could use a box filter to approximate to the original image, as corrupted pixels can be more or less recovered by the information of the healthy surrounding ones. Despite of this, we will use two algorithms called K Nearest Neighbors (KNN) (see Figure 4.14.1.42) and Non Local Means (NLM) (see Figure 4.14.1.43) that fit better in this problem. Again, we won’t go into the details of each of them but it has to be said that the second one is a more complex variation that has greater resource usage. An optimized version of NLM called Quick NLM or NLM2 is also available (see Figure 4.14.1.44). 4.5.2 Performance This example will clearly show the difference between the hardware platforms. The denoising algorithms were executed on all the devices and the obtained data is presented on Figure 4.14.1.45. GTX 480 and GTX 980 have much greater performance than Jetson TK1. On the first denoising algorithm, Jetson TK1 performance is 9% of the GTX 480 one and 5% of the GTX 980. With NLM algorithm, the proportion against the GTX 480 shortens, being Jetson TK1 framerate 15% of the one of that device, but remains around the same in the rate with the GTX 980. Computer Architecture Dept. 71 Javier Cabero GuerraComputer Architecture Dept. 71 Javier Cabero GuerraComputer Architecture Dept. 71 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering 4.1.41 Original Noisy Image. 4.1.42 Applying KNN. 4.1.43 Applying NLM. 4.1.44 Applying Quick NLM (NLM2). 4.5.3 Power Draw and Heat Generation The power draw consumed in this example is not the highest one nor the lowest. In Figure 4.14.1.46, the vertical axis presents a logarithmic scale. The maximum power consumption for GTX 480 is 67.8 W in the KNN filter. NLM and Quick NLM have almost the same power draw. GTX 980 goes higher to more or less 120 W and Jetson TK1 doesn’t go further than 3.9 W. For temperature (see Figure 4.14.1.47), GTX 480 almost reaches the maximum temperature of all the samples with 90 degrees in the NLM filter. Jetson TK1 stops at 48 degrees and GTX 980 at 79 (NLM). Quick NLM has 1 degree less than the NLM in the GeForce devices. The optimized algorithm has much higher framComputer Architecture Dept. 72 Javier Cabero GuerraComputer Architecture Dept. 72 Javier Cabero GuerraComputer Architecture Dept. 72 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering 4.1.45 Image Denoising Performance. 4.1.46 Image Denoising Power Draw. erate and its power drawback and heat generation is a bit lower. As expected, the least temperature values are generated when no algorithm is applied (noisy image) and the higher ones when using NLM denoising algorithm. GTX 980 is the platform that raises higher, changing from 52 degrees to 78 in noisy to KNN swap. The temperature step is 26 degrees. GTX 480 has higher temperature, but as it starts from 77 degrees, the step is smaller. Computer Architecture Dept. 73 Javier Cabero GuerraComputer Architecture Dept. 73 Javier Cabero GuerraComputer Architecture Dept. 73 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering 4.1.47 Image Denoising Heat Generation. 4.5.4 Algorithm Efficiency The efficiency behaves as in the previous case: GTX 480 beats GTX 980 in the first configurations and on harder tasks the GTX 980 is more efficient (see Figure 4.14.1.48). Jetson TK1 is the most efficient platform with more than twice the efficiency in KNN and Quick NLM than the other devices. In NLM, it is more than twice just with respect to GTX 480. 4.1.48 Image Denoising Power Efficiency. Speaking of heat generation efficiency (see Figure 4.14.1.49), GTX 480 has less than GTX 980 but is better than Jetson TK1 in all cases. On NLM, Jetson TK1 Computer Architecture Dept. 74 Javier Cabero GuerraComputer Architecture Dept. 74 Javier Cabero GuerraComputer Architecture Dept. 74 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering generates 0.2 degrees per each FPS it achieves, being the lowest heat generation efficiency in the graph. GTX 980 has again the best scalability in this metric with 0.016 degree per FPS on NLM filter. 4.1.49 Image Denoising Heat Generation Efficiency. Computer Architecture Dept. 75 Javier Cabero GuerraComputer Architecture Dept. 75 Javier Cabero GuerraComputer Architecture Dept. 75 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering Least Greatest Computational Power Jetson TK1 GTX 980 Power Drawback Jetson TK1 GTX 980 Heat Generation Jetson TK1 GTX 480 Power Drawback Efficiency GTX 480 Jetson tK1 Heat Generation Efficiency Jetson TK1 GTX 980 Tabla 4.1: Final Results. there is a high number of devices and the accumulated heat of all together could reduce the lifetime of some system components. Then, a device with higher heat generation efficiency is desired. Whether or not we need performance or efficiency, the different possible platforms will have different features that will make the choice of chosing one or another dependent on the task to be performed, its computational and energetic needs, and the conditions of the target system, both in terms of heat generated and economic cost produce by the time of usage. Computer Architecture Dept. 82 Javier Cabero GuerraComputer Architecture Dept. 82 Javier Cabero GuerraComputer Architecture Dept. 82 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering 4.6.6 Spanish Las pruebas realizadas han demostrado las diferencias existentes entre las distintas plataformas hardware. La diferencia en rendimiento, consumo, generación de calor y eficiencia entre los dispositivos está directamente relacionada con la distancia en el tiempo, la arquitectura y la cantidad de hardware. Gracias a los experimentos, ha sido posible apreciar también cómo la naturaleza de las aplicaciones hace que encajen mejor en un dispositivo o en otro. En resumen, la GTX 980 ha resultado ser el dispositivo que más rendimiendo proporciona en casi todas las configuraciones de todos los ejemplos. No obstante, la GTX 480 ha producido mejores resultados en algunos casos que requerían menor trabajo debido a la mejor adapción del programa al hardware. La generación de calor suele ser directamente proporcional al consumo del dispositivo y cuanto más trabajo se le asigna, más temperatura y consumo se produce. En algunas aplicaciones, la eficiencia de las plataformas podía ser diferente. Generalmente, la Jetson tK1 tiene la eficiencia más alta, aunque en el último ejemplo fue superada por la GTX 980 desde una configuración más primeriza hasta la última. Tanto la eficiencia en generación de calor como en consumo decrementan con tareas más intensas. Aunque la Jetson TK1 tenga la mayor eficiencia en cuanto a consumo gracias a su diseño integrado GPU-CPU, hemos que visto que es la que peor eficiencia tiene en cuanto a eficiencia en la generación de calor. Los números hablan por si mismos, mostrando los puntos fuertes de cada dispositivo (ver Tabla 4.2). Dependiendo de la aplicación, uno puede pensar en usar un dispositivo por su poder computacional, su eficiencia en consumo o generación de calor o simplemente porque genera menos calor, aunque su eficiencia en generación de calor no sea muy buena. Por ejemplo, si no hay necesidad de un alto rendimiento, es mejor usar la GTX 480 dado que consume menos que la GTX 980. Sin embargo, para ejecuciones de mayor importancia la GTX 980 es más eficiente por lo que será un dispositivo más barato ya que, a pesar de consumir más vatios por segundo, finalizará antes. Otro escenario podría ser aquel en el que el calor generado importa: quizás hay un alto número de dispositivos y el calor acumulado de todos ellos puede afectar a la durabilidad de algunos componentes del sistema. Entonces, un dispositivo con más eficiencia en lal generación de calor es deseable. Tanto si se necesita o no rendimiento o eficiencia, las diferentes posibles plataformas tendrán diferentes características que harán que la elección de la plataforma dependa de la tarea que Computer Architecture Dept. 83 Javier Cabero GuerraComputer Architecture Dept. 83 Javier Cabero GuerraComputer Architecture Dept. 83 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering Menor Mayor Poder Computacional Jetson TK1 GTX 980 Consumo Energético Jetson TK1 GTX 980 Generación de Calor Jetson TK1 GTX 480 Eficiencia en Consumo Energético GTX 480 Jetson tK1 Eficiencia en Generación de calor Jetson TK1 GTX 980 Tabla 4.2: Resultados finales. se deba hacer, de sus necesidades computacionales y energéticas, y de las condiciones en las que el sistema destino se encuentre, tanto en lo que respecta al calor como en el coste económico que representa su uso en el tiempo. Computer Architecture Dept. 84 Javier Cabero GuerraComputer Architecture Dept. 84 Javier Cabero GuerraComputer Architecture Dept. 84 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering 5 Bibliography [1] The Green 500 List. URL http://www.green500.org/. [2] Q&A Jetson TK1 FAQ, 2014. URL http://developer.download.nvidia.com/ embedded/jetson/TK1/docs/Jetson_TK1_FAQ_2014May01_V2.pdf. [3] Antonio Ruiz, Manuel Ujaldón. Exploiting Kepler Capabilities on Zernike Moments. 2015. [4] Lisa Gottesfeld Brown. A survey of image registration techniques. ACM Computing Surveys 24:325–376, 1992. [5] C. Tomasi, R. Manduchi. Bilateral Filtering for Gray and Color Images. Proceedings of the 1998 IEEE International. Conference on Computer Vision. Bombay. India, 1998. [6] Lin-Ching Chang, Esam El-Araby, Vinh Q. Dang, and Lam H. Dao. GPU accelComputer Architecture Dept. 87 Javier Cabero GuerraComputer Architecture Dept. 87 Javier Cabero GuerraComputer Architecture Dept. 87 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering eration of nonlinear diffusion tensor estimation using CUDA and MPI. Neurocomputing, pages 328–338, 2014. [7] Chris McClanahan. History and Evolution of GPU Architecture. 2010. [8] Christos Kyrkou. Stream Processors and GPUs: Architectures for High Performance Computing. Unknown. [9] Centre for Medical Image Computing CMIC. NifTK description at the CMIC Software and Open Source Database page, 2015. URL http://cmic.cs.ucl. ac.uk/home/software/. [10] Nvidia Corporation. CUDA Imaging Samples. URL http://docs.nvidia.com/ cuda/cuda-samples/#imaging. [11] Nvidia Corporation. Fermi White Paper, 2009. URL http://www.nvidia.com/ content/PDF/fermi_white_papers/NVIDIA_Fermi_Compute_Architecture_ Whitepaper.pdf. [12] Nvidia Corporation. GeForce GTX 480 Specifications, 2009. URL http://www. nvidia.es/object/product_geforce_gtx_480_es.html. [13] Nvidia Corporation. Nvidia announcements at Consumer Electronics Show (CES), 2014. URL http://www.nvidia.com/object/ces2014.html. [14] Nvidia Corporation. GeForce GTX 980 Whitepaper, 2014. URL http: //international.download.nvidia.com/geforce-com/international/ pdfs/GeForce_GTX_980_Whitepaper_FINAL.PDF. [15] Nvidia Corporation. Meet the Jetson Embedded Platform, 2014. URL https: //developer.nvidia.com/meet-jetson-embedded-platform. [16] Nvidia Corporation. GeForce GTX 980 Specifications, Late 2014. URL http: //www.nvidia.es/object/geforce-gtx-980-es.html#pdpContent=2. [17] Mark Harris. Introduction to CUDA C, 2013. [18] Mark Harris. Jetson TK1: Mobile Embedded Supercomputer Takes CUDA Everywhere, 2014. URL http://devblogs.nvidia.com/parallelforall/ jetson-tk1-mobile-embedded-supercomputer-cuda-everywhere/. [19] Ian Buck. Stream Computing on Graphics Hardware. PhD thesis, September 2006. Computer Architecture Dept. 88 Javier Cabero GuerraComputer Architecture Dept. 88 Javier Cabero GuerraComputer Architecture Dept. 88 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering [20] John Michael McNamee. A Comparison Of Methods For Accurate Summation. ACM SIGSAM Bulletin, 38, March 2004. [21] José A. Lachiondo, Manuel Ujaldón, Regina Berretta, Pablo Moscato. Legendre Moments as High Performance Bone Biomarkers: Computational Methods and GPU Acceleration. Computer Methods in Biomechanics and Biomedical Engineering, 2015. [22] William B. Langdon, Marc Modat, Justyna Petke, and Mark Harman. Improving 3D Medical Image Registration CUDA Software with Genetic Programming. Proceedings of the 2014 conference on Genetic and evolutionary computation (GECCO 2014), pages 951–958, 2014. [23] Mark Harris. Maxwell: The Most Advanced CUDA GPU Ever Made, 2014. URL http://devblogs.nvidia.com/parallelforall/ maxwell-most-advanced-cuda-gpu-ever-made/. [24] Sparsh Mittal and Jeffrey S. Vetter. A survey of methods for analyzing and improving GPU energy efficiency. CoRR, abs/1404.4629, 2014. URL http: //arxiv.org/abs/1404.4629. [25] Marc Modat, Zeike A. Taylor, Josephine Barnes, David J. Hawkes, Nick C. Fox, and Sebastien Ourselin. Fast free-form deformation using the normalised mutual information gradient and graphics processing units. Med Phys, pages 278–284, 2010. [26] Modat, M., Cash, D. M., Daga, P., Winston, G. P., Duncan, J. S., and Ourselin, S. Global image registration using a symmetric block-matching. JOURNAL of Medical Imaging, 1(2):024003–024003, 2014. [27] Modat, M., Ridgway, G. R., Taylor, Z. A., Lehmann, M., Barnes, J., Hawkes, D. J., Fox, N. C., et al. Fast free-form deformation using graphics processing units. Computer Methods And Programs In Biomedicine, 98(3):278–284, 2010. [28] NVIDIA Corporation. NVIDIA GeForce 8800 GPU architecture overview. Technical report, November 2006. [29] NVIDIA Corporation. NVIDIA GeForce GTX 200 GPU architectural overview. Technical report, May 2008. [30] NVIDIA Corporation. NVIDIA’s Next Generation CUDA Compute Architecture: Fermi Whitepaper. 2009. Computer Architecture Dept. 89 Javier Cabero GuerraComputer Architecture Dept. 89 Javier Cabero GuerraComputer Architecture Dept. 89 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering [31] NVIDIA Corporation. NVIDIA GF100 Whitepaper. Technical report, 2010. [32] NVIDIA Corporation. TESLA K10 GPU Accelerator Board Specification. Technical report, November 2012. [33] NVIDIA Corporation. NVIDIA Tesla KSeries Data Sheet, October 2012. [34] NVIDIA Corporation. TESLA K20X GPU Accelerator Board Specification. Technical report, July 2013. [35] NVIDIA Corporation. TESLA K20 GPU Accelerator Board Specification. Technical report, July 2013. [36] NVIDIA Corporation. NVIDIA’s Next Generation CUDA Compute Architecture: Kepler GK110/210 Whitepaper. Technical report, 2014. [37] NVIDIA Corporation. NVIDIA Jetson TK1 Development Kit. Technical report, May 2014. [38] NVIDIA Corporation. NVIDIA Embedded Computing Website, 2015. URL https://developer.nvidia.com/embedded-computing. [39] NVIDIA Corporation. CUDA C Programming Guide, 2015. URL http://docs. nvidia.com/cuda/cuda-c-programming-guide. [40] NVIDIA Corporation. Official Wiki for NVIDIA’s Tegra and Jetson, 2015. URL http://elinux.org/Jetson_TK1. [41] NVIDIA Corporation. NVIDIA GeForce GTX 980 Whitepaper. Technical report, 2015. [42] NVIDIA Corporation. NVIDIA Tegra K1 Whitepaper. Technical report, 2015. [43] NVIDIA Corporation. Maxwell Tuning Guide: 1.4.2.1. Unified L1/Texture Cache, 2015. URL http://docs.nvidia.com/cuda/maxwell-tuning-guide/ #l1-cache. [44] Ourselin, S., Roche, A., Subsol, G., Pennec, X., and Ayache, N. Reconstructing a 3d structure from serial histological sections. Image and Vision Computing, 19(1-2):25–31, 2001. [45] Rana Mohtadi. Magnesium batteries: Current state of the art, issues and future perspectives, 2014. URL http://www.beilstein-journals.org/ bjnano/single/articleFullText.htm?publicId=2190-4286-5-143. Computer Architecture Dept. 90 Javier Cabero GuerraComputer Architecture Dept. 90 Javier Cabero GuerraComputer Architecture Dept. 90 Javier Cabero Guerra University of Malaga School of Computer EngineeringUniversity of Malaga School of Computer EngineeringUniversity of Malaga School of Computer Engineering [46] Rueckert, D., Sonoda, L. I., Hayes, C., Hill, D. L. G., Leach, M. O., and Hawkes, D. J. Nonrigid registration using free-form deformations: Application to breast mr images. IEEE Transactions on Medical Imaging, 18(8):712–721, 1999. [47] Samuel Williams, Andrew Waterman, David Patterson. Roofline: an insightful visual performance model for multicore architectures. Communications of the ACM, 52:65–76, April 2009. doi: 10.1145/1498765.1498785. [48] Stephan Soller. GPGPU origins and GPU hardware architecture. 2011. [49] TechPowerUp. GPU-Z Official Page. URL http://www.techpowerup.com/ gpuz/. [50] UC London Translational Imaging Group. NiftyReg: Open-source software for efficient medical image registration, 2015. URL http://cmictig.cs.ucl.ac. uk/research/software/22-niftyreg. [51] Manuel Ujaldón. VIII Curso Avanzado de GPU: Programación y Rendimiento frente a la CPU, 2014. [52] Vasily Volkov. Better performance at lower occupancy. UC Berkeley Lecture, 2010. URL http://www.cs.berkeley.edu/~volkov/volkov10-GTC.pdf. [53] Barbara Zitová and Jan Flusser. Image registration methods: a survey. Image and Vision Computing, 21:977–1000, 2003. Computer Architecture Dept. 91 Javier Cabero GuerraComputer Architecture Dept. 91 Javier Cabero GuerraComputer Architecture Dept. 91 Javier Cabero Guerra