Design and implementation of a high-performance stream-based computing platform on multigenerational GPUs

1 Design and implementation of a highperformance stream-based computing platform on multigenerational GPUs. By Pablo Lamilla Álvarez September 27, 2010 Supervised by: Professor Shinichi Yamagiwa Kochi University of Technology, Japan Professor Francisco José Abad Cerdá Universidad Politécnica de Valencia, Spain 2 Table of Contents Chapter 1 Introduction...................................................................................................... 4 1.1 Overview of the Project.......................................................................................... 4 1.2 Objective of research.............................................................................................. 6 1.3 Expected Outcomes................................................................................................ 7 1.4 Disposition of the document................................................................................... 7 Chapter 2 Background and Definitions............................................................................ 9 2.1 Stream computing................................................................................................... 9 2.2 GPGPU................................................................................................................. 10 2.3 GPGPU computing platforms............................................................................... 13 2.3.1 Caravela Platform.......................................................................................... 14 2.3.2 OpenCL ......................................................................................................... 17 Chapter 3 Stream-based computing platform on multigenerational GPUs.................... 27 3.1 Caravela implementation...................................................................................... 27 3.2 Caravela low-level functions................................................................................ 28 3.3 Flow Model .......................................................................................................... 33 3.3.1 Constant values structure............................................................................... 35 3.4 Flow Model Creator ............................................................................................. 36 3.5 Kernel Structure.................................................................................................... 38 3.6 Swap Mechanism.................................................................................................. 39 3.7 Implementation conclusions................................................................................. 41 Chapter 4 Performance Evaluation................................................................................. 42 4.1 Straightforward application: Matrix multiply....................................................... 42 4.2 Recursive application: IIR filter........................................................................... 46 Chapter 5 Conclusions and future work......................................................................... 50 Bibliography (references)............................................................................................... 52 Apendix I User guide of Flow Model Creator GUI........................................................ 55 1 Introduction ............................................................................................................. 55 1.1 FlowModelCreator ........................................................................................... 55 1.2 Authors............................................................................................................. 55 2 Functionality............................................................................................................ 56 2.1 Basics concepts................................................................................................. 56 2.2 Top Window..................................................................................................... 57 2.3 Set constant values window............................................................................. 61 3 List of Figures Figure 1. A typical system organization with legacy GPU .................................... 10 Figure 2. Graphics processing steps....................................................................... 11 Figure 3 Recent GPU architecture.......................................................................... 12 Figure 4. Structure of the flow-model.................................................................... 14 Figure 5. Resource hierachy in a processing unit................................................... 15 Figure 6. Swap mechanism..................................................................................... 17 Figure 7. Companies that support OpenCL............................................................ 18 Figure 8. Platform Model of OpenCL.................................................................... 18 Figure 9. Example of 2-dimensional NDRange ..................................................... 20 Figure 10. Conceptual OpenCL architecture.......................................................... 21 Figure 11. VectorAdd example in OpenCL kernel language................................. 23 Figure 12. OpenCL host code VectorAdd example. Part one................................ 24 Figure 13. OpenCL host code VectorAdd example. Part two................................ 25 Figure 14. Structure of the Caravela functions....................................................... 28 Figure 15. Flow Model of VectorAdd example ..................................................... 35 Figure 16. Changes in the top window of FlowModelCreator............................... 37 Figure 17. New interface for introducing the constant values................................ 38 Figure 18. Example of kernel header for Caravela................................................. 39 Figure 19. Execution times using matrix multiply................................................. 43 Figure 20 MatrixMultiply Kernels ......................................................................... 44 Figure 21 Execution times using IIR filter............................................................. 46 Figure 22 Kernel code of IIR program................................................................... 47 Figure 23. VectorAdd execution times with 5000 swaps....................................... 48 Figure 24. VectorAdd execution times with 10000 swaps..................................... 49 List of tables Table 1. Main functions of Caravela library........................................................... 16 Table 2. Environment for the evaluation................................................................ 42 4 Chapter 1 Introduction 1.1 Overview of the Project During this decade, high performance computation demand has been increasing more and more, for example in the field of humanities [1]. Scientists and investigators are in need of high speed and performance environments for their research, which need to perform millions of floating points operations per second [2]. One way to achieve this goal is to increase the power of the hardware. Multi-core CPUs and Supercomputers [3] are an example of the evolution of this path. However, even this type of hardware (Supercomputers) has their limits. Supercomputers, besides being very expensive and complicate to build, have a very high electricity consumption [22]. In the meanwhile, another kind of processors, the GPUs, have experimented a great improvement within this decade, and a single GPU IC chip is able to achieve 10 TFLOPS. This performance attracts HPC researchers to the potential computing power of GPUs, so in the last decade the field called GPGPU [17] (General-Purpose computation on Graphics Processing Units) has gained more importance. However programming on the GPU entails a great difficulty, due to its architecture, that is categorized into two types nowadays. One type of GPU architecture has three kinds of processors called, vertex processors, rasterizer and fragment processors that are dedicated respectively to transform the vertices that define the graphics primitives, to transform a primitive into a set of pixels and finally to compute each pixel color, taking into account an illumination equation and maybe a texture. This graphics pipeline can be programmed using graphics libraries such as DirectX 9[7] and OpenGL 2.0 [8]. Also, this architecture is generally called legacy architecture, to distinguish from the newest architectures that have a different structure. These newest architectures integrate tens or hundreds of a standardized general purpose stream processor that can execute any type of shader. When the computational units are unified like this, the architecture is called Unified Shading Architecture. Each 5 processor invokes a program to generate element(s) of output data stream(s) from element(s) of input one(s). Emulating the three-step graphics processes performed by the former architecture, this recent architecture implements the similar operations of the graphics runtimes. In addition, it releases the stream processor resources to general purpose computing via special runtimes such as CUDA [11] and OpenCL [12]. During the era of the legacy GPU architecture the main problem of GPGPU applications was the complete different style of programming compared to CPU programming, due to the fact that legacy GPUs environments were originally designed for graphics processing. Therefore some applications that tried to cover and abstract those differences for the programmers were developed such as Brook [9] and Sh [10]. Brook introduced a programming model for GPUs called stream computing. Stream Computing [4] is a programming paradigm in which the data is processed in a continuous way, as a unique piece of information. This piece is called an element of stream. The point in the stream computing is to manipulate and to operate each element of the stream in a parallel pipeline, i.e. all at the same time. The operations applied to the stream are packed into a function called kernel. For instance, if we want to sum the elements of two vectors (output[i] = inputA[i] + inputB[i]), using the old-sequential style the elements have to be summed one after another inside of a loop. However, using stream computing, each vector is treated as a stream and the sum operation is programmed inside a kernel. Then, this kernel will receive each stream and will apply the sum operation to all the elements of these streams in a parallel way, so the elements of the resulting vector are obtained all at once. Another GPGPU platform, the Caravela [13], was developed to form the true stream-based computing based on the flow-model [14]. This model, used by the Caravela in its execution, is defined by the number of I/O streams, constant values and a program (kernel) invoked on a targeted GPU . The problem with these architectures is that their programming style is not compatible because for programming in the recent architecture one must follow a stream-based computing style, which is not the same style that is followed when programming in the legacy architecture. Therefore algorithms or optimization techniques cannot be shared between architectures. However the Caravela platform has a unified framework to implement stream-based computing using the flow-model, just 6 by defining the number of I/O streams and the target program. This flow-model also can be defined easily with a GUI called FlowModelCreator, which is included in the Caravela package. Thus, the Caravela platform can provide a unified programming style between different GPU generations, because the stream-based computing concept is standardized. This project is focused on the migration of the Caravela platform to OpenCL. Now the Caravela is implemented on the DirectX9 and the OpenGL environments, but the OpenCL support, as a standard and platform-independent language, will allow the Calavera to run into any modern GPU environment. Therefore with this migration process the following problems will be solved: 1) Impossibility to run the Caravela into any GPU environment. 2) Disparities in the programming style of the runtimes of the legacy and recent GPU architectures. This project will be developed in C/C++. C# will also be used in the Graphical User Interface (GUI) that allows the creation of the “Flow Model”, the structure that packs all the information used by Calavera. 1.2 Objective of research There are three main objectives of this research project as listed below: 1) OpenCL support in Caravela platform. This objective will cover the necessity of Caravela platform to run on any GPU environment. To achieve this objective, the relationship between the OpenCL functions and the Caravela functions will be studied and established. 2) OpenCL support in FlowModelCreator GUI. This objective will allow the FlowModelCreator GUI to automatically generate a flow-model XML file with all the new parameters that OpenCL implementation needs. 7 3) Exploiting maximum performance from GPUs in any generation. This objective will focus on preserving the performance of Caravela execution on GPUs of any generation. 1.3 Expected Outcomes There are two outcomes expected by this research project. 1) A software package that includes the new Caravela software with OpenCL support. This package will include the Caravela platform with the new functionalities to support OpenCL environment, as well as the support of the previous environments (OpenGL and DirectX). 2) Multi-Platform “FlowModelCreator” GUI with OpenCL option. This GUI application will generate the flow-model structure for Caravela. With the release of the new version, the changes that allow the support of OpenCL will be implemented. 3) Application examples for the Caravela. Programs samples used for the performance evaluation are also added. This includes the matrix-multiplication and the IIR filter kernel programs, both written in OpenCL, and the main programs using the Caravela API functions to execute the kernels. Thus, with the package of these outcomes, a high-performance stream-based platform on OpenCL will be developed. 1.4 Structure of the document Chapter 2 explains the background of this project’s field, specifically “Stream computing” and GPGPU, as well as a description of the two main environments used in this project, the “Caravela” and the OpenCL. Chapter 3 illustrates the implementation of the new version of Caravela. Technical problems encountered along the migration process and their solutions are explained in detail. 8 Chapter 4 presents the results obtained with the new software. Various performance evaluations are carried out, and their results are commented. The last chapter describes the conclusions for this project and proposes future directions. Finally in the appendix, the User’s Guide of the application and the GUI “FlowModelCreator” are attached at the end of this document. 9 Chapter 2 Background and Definitions 2.1 Stream computing Stream computing (or stream processing) is a computer programming paradigm that allows some applications to easily exploit a limited form of parallel processing [4]. These applications are able to use multiple computational units without the necessity of explicitly managing allocation, synchronization, or communication among those units. Stream computing takes advantage of a SIMD [21] (Single Instruction, Multiple Data) architecture, where the same instruction can be applied to various instances of different data. The main point of stream computing is to use a continuous flow of data called stream as the input and output of the program. A stream is a collection of data which can be operated on in parallel. Each element of a stream is a record of data requiring a similar computation; however it is, in general, independent of the other elements. A series of operations will be applied to each element in the stream all at once. These operations are packed in a function called kernel. Typically, the same kernel function is applied to all elements in the stream (Uniform streaming), but it is not the only option. The benefit of stream computing stems from the highly parallel architecture of GPUs, whereby tens to hundreds of parallel floating points operations are performed with each clock cycle. For this reason, nowadays stream computing is primarily used in the realm of the GPUs, where stream computing can easily take advantage of the large number of parallel processors of GPUs. Stream computing on the GPU was mainly used for graphics purposes, but in the recent years, there has been an increasing interest to use it for general purpose applications [5]. This recently new high performance computing field is called GPGPU [17] and is explained in the next section. 16 Table 1. Main functions of Caravela library With these functions, the programmer can implement target applications in the framework of flow-models just mapping flow-models into one or more shaders. Thus, the programmer does not have the necessity of knowing about graphics runtime environment details, so Caravela can become to a solution to relieve the problem of disparities between graphical environments mentioned in Section 2.2. Moreover, the Caravela platform incorporates optimization functions of the flowmodel execution itself, called swap mechanism [19], which allows executing recursive iterations of a flow-model exchanging the input and the output buffers in the GPU side without copy operations between the host memory and the GPU memory. Avoiding this data transfer between the host memory and the GPU memory, we reduce the execution time of the program considerably. Figure 6 shows this mechanism. 17 Figure 6. Swap mechanism Currently, as mentioned before, the Caravela platform supports the legacy architecture of GPUs, so graphics runtime functions must be used to perform the stream-based computation. Also, recently Caravela has been ported to CUDA, so support for the recent architecture of GPUs has been added. However, because CUDA was developed by NVIDIA, it can be used only on NVIDIA GPUs. Thus, we must give support to Caravela for a runtime that can be used with any GPU regardless of the manufacturer, and the chosen language is OpenCL. The next section gives an overview of this API. 2.3.2 OpenCL OpenCL (Open Computing Language) is an open royalty-free standard, initially proposed by Apple and finally developed by The Khronos Group [20], for general purpose parallel programming across CPUs, GPUs and other processors. Moreover, its most important characteristic compared for example to CUDA, is that OpenCL is crossplatform. It also supports a wide range of applications, from embedded and consumer software to HPC solutions, through a low-level, high-performance, portable abstraction. Besides Khronos Group, many industry-leading companies and institutes have participated and have supported the development of OpenCL. Figure 7 shows many of them. 18 Figure 7. Companies that support OpenCL The objective of OpenCL is to help expert programmers, such as library writers and middleware vendors, to write portable yet efficient code. Therefore OpenCL provides a low-level hardware abstraction as well as a framework to support programming. To explain the main ideas behind this language, we will describe these four models for OpenCL: i) platform, ii) memory, iii) execution and iv) programming models. i) Platform Model The platform model for OpenCL is shown in Figure 8. A host connected to one or more OpenCL devices, such as GPUs, CPUs… Each device is divided into one or more compute units. These compute units are further divided into one or more processing elements which is where the computation occurs. Figure 8. Platform Model of OpenCL 19 The program written in OpenCL runs on the host side submitting commands to execute computations on the processing elements within a device. The processing elements execute a single stream of instructions. ii) Execution Model OpenCL execution can be separated in two parts. First a host program, which defines a context (environment where the kernels execute and the domain in which synchronization and memory management is defined). The context manages the execution of the kernels. Then the kernel functions are executed on one or more OpenCL devices. When the host program submits a kernel for execution, an index space, called NDRange, is defined. For each point of this space an instance of the kernel is executed. This kernel instance is called in OpenCL work-item, which is identified by a global ID in the index space. Each work-item executes the same code but the execution pathway and the data operated can vary per work-item. Moreover, work-items can be organized into work-groups, which provide a more coarse-grained decomposition of the index space. For each work-group a work-group ID is assigned and for each work-item within a work-group a local ID is assigned also. So in summary, each work-item has two IDs (global and local) and each work-group has one ID. The NDRange can be one, two or three dimensional, so each work-item ID will be Ndimensional tuples, where N is the dimension of the NDRange. Figure 9 shows an example of 2-dimensional NDRange. 20 Figure 9. Example of 2-dimensional NDRange In this example each little square is a work-item and each of the nine bigger squares is a work-group. The NDRange is used to identify and to have control of all the workitems that are executed. So for example, if we want to execute a program that does the matrix operation A+B=C, each work-item, when we launch the kernel a NDRange like the above will be created. Each work-item of the NDRange will perform a sum operation to calculate one element of the matrix C, so for example the work-item with the ID (0,0) will perform the operation A[0][0]+B[0][0]=C[0][0]. Because we have the control of all the work-items using the NDRange and the IDs, the sum operation inside the kernel function is just A[x] [y] +B[x] [y] =C[x] [y] where x and y are the IDs of each work-item. As mentioned before, the host program creates a context to manage the executions of the kernel. The context includes the devices, the kernels, the program objects (source and executable that implement the kernels) and memory objects visible to the host and the device. The method to submit orders from the host to the kernel is through a data structure called command-queue, created on the host side. The host program places commands into this structure which are then scheduled onto the device within the context. 21 Examples of commands are the one that executes a kernel, synchronization commands or memory commands which transfer data between memory objects (host->device, device->host or device->device). iii) Memory Model OpenCL divides the memory into four distinct regions: • Global memory. Any work-item has read/write permission to this region. • Constant memory. Is a region of the global memory that remains constant during the execution of a kernel. The host must allocate and initialize the memory objects placed in this region. • Local memory. This region is local to a work-group, which means that is shared only by all the work-items in that work-group. • Private memory. Memory private to a work-item. Any variable defined there is only accessible by this work-item. Figure 10 shows these four regions of memory and how they relate to the platform model. Figure 10. Conceptual OpenCL architecture 22 The host application uses the OpenCL API to create memory objects in global memory, and to enqueue memory commands, to write/read to/from these memory objects. The host and the device memory are almost completely independent of each other, because the host side is defined outside of OpenCL, i.e. everything that is used by OpenCL during the execution (NDRange, memory regions …) is inside the OpenCL device. The host can only communicate with the OpenCL device thorough the OpenCL API functions. However they need to interact in order to pass data between memory objects. This occurs in one of these two ways: by copying data or by mapping and unmapping regions of a memory object. To copy data explicitly, the host has to submit commands to transfer data between a memory object and host memory, which can be blocking or non-blocking. On the other hand, to map/unmap regions of a memory object, the host can map a region from this memory object into its address space. Once a region from the memory object has been mapped, the host is allowed to write/red to this region. When the host finishes all the write/read accesses, the region is unmapped. iv) Programming Model Two types of programming models are supported by OpenCL: data parallel and task parallel models, as well as hybrids of these two. In the data parallel programming model a sequence of instructions (kernel) is applied to multiple elements of a memory object. The index space associated with the OpenCL execution model defines the work-items and how the data maps onto the work-items. Strict one-to-one mapping between the work-item and the element in a memory object is not a requirement. On the other hand, the task parallel programming model defines a model in which a single instance of a kernel is executed independent of any index space. It is equivalent to executing a kernel on a computing unit with an NDRange of 1 work-item and 1 work-group. The programming language of OpenCL kernels is based on C99. Figure 11 shows an example of the vector addition problem written in this language. 23 Figure 11. VectorAdd example in OpenCL kernel language A kernel function must be defined with the __kernel directive. Memory objects created and initialized on the host side are passed to the kernel function as parameters. These parameters can have the directives __global, __local, etc. depending on the memory region that we want to use. In this example we have 3 arrays of floats (2 inputs and 1 output) declared as __global and a single int. The two inputs have also the keyword “const” so they are placed into the Constant Memory region. Inside the kernel function the sum operation is executed for each element of the arrays. Before the operation, we must get the ID to identify the work-item inside the NDRange, stored in the variable iGID, so each work-item will perform one sum operation. This code would be executed on the GPU side, so previously the host side must build and execute this code using the OpenCL API. The following code (Figure 12 and 13) shows the basic steps to prepare the OpenCL environment, execute the kernel function and retrieve the results from it for this VectorAdd example. 24 Figure 12. OpenCL host code VectorAdd example. Part one const unsigned int cnBlockSize = 512; const unsigned int cnBlocks = 3; const unsigned int cnDimension = cnBlocks * cnBlockSize; // create OpenCL device & context cl_context hContext; hContext = clCreateContextFromType(0, CL_DEVICE_TYPE_GPU, 0, 0, 0); // query all devices available to the context size_t nContextDescriptorSize; clGetContextInfo(hContext, CL_CONTEXT_DEVICES, 0, 0, &nContextDescriptorSize); cl_device_id * aDevices = malloc(nContextDescriptorSize); clGetContextInfo(hContext, CL_CONTEXT_DEVICES, nContextDescriptorSize, aDevices, 0); // create a command queue for first device the context reported cl_command_queue hCmdQueue; hCmdQueue = clCreateCommandQueue(hContext, aDevices[0], 0, 0); // create & compile program cl_program hProgram; hProgram = clCreateProgramWithSource(hContext, 1, sProgramSource, 0, 0); clBuildProgram(hProgram, 0, 0, 0, 0, 0); // create kernel cl_kernel hKernel; hKernel = clCreateKernel(hProgram, “vectorAdd”, 0); // allocate host vectors float * pA = new float[cnDimension]; float * pB = new float[cnDimension]; float * pC = new float[cnDimension]; // initialize host memory randomInit(pA, cnDimension); randomInit(pB, cnDimension); // allocate device memory cl_mem hDeviceMemA, hDeviceMemB, hDeviceMemC; hDeviceMemA = clCreateBuffer(hContext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, cnDimension * sizeof(cl_float), pA, 0); hDeviceMemB = clCreateBuffer(hContext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, cnDimension * sizeof(cl_float), pA, 0); hDeviceMemC = clCreateBuffer(hContext, CL_MEM_WRITE_ONLY, cnDimension * sizeof(cl_float), 0, 0); 25 Figure 13. OpenCL host code VectorAdd example. Part two These API functions would be written on the host side to manage the execution of the kernel function. Summarizing, OpenCL provides an API and a C-based kernel language that make an easy and transparent interface for GPGPU. It can be used in most modern GPU due to its platform-independent design. Therefore, the Caravela platform, which has potentially a stream-based computing style, can take advantage of OpenCL characteristics to become almost 100% compatible with any GPU. In order to port the Caravela to OpenCL, we must take care of these issues: • Relationship between the functionality of the Caravela functions and the OpenCL API functions. • Modifications in the flow-model structure in order to pack all the OpenCL information needed. • Modifications in the GUI “FlowModelCreator” in order to accept the new modifications done in the flow-model structure. // setup parameter values clSetKernelArg(hKernel, 0, sizeof(cl_mem), (void *)&hDeviceMemA); clSetKernelArg(hKernel, 1, sizeof(cl_mem), (void *)&hDeviceMemB); clSetKernelArg(hKernel, 2, sizeof(cl_mem), (void *)&hDeviceMemC); // execute kernel clEnqueueNDRangeKernel(hCmdQueue, hKernel, 1, 0, &cnDimension, 0, 0, 0, 0); // copy results from device back to host clEnqueueReadBuffer(hContext, hDeviceMemC, CL_TRUE, 0, cnDimension * sizeof(cl_float), pC, 0, 0, 0); delete[] pA; delete[] pB; delete[] pC; clReleaseMemObj(hDeviceMemA); clReleaseMemObj(hDeviceMemB); clReleaseMemObj(hDeviceMemC); 32 level (the top), so if we return directly an OpenCL buffer, this buffer cannot be initialize without using the OpenCL API function clEnqueueWriteBuffer , and if we force the user to use this function, the way of programming with CalaveraOpenCL would drastically change compared to the previous versions of Caravela. Therefore, the creation of host side buffers is the best solution to preserve the way of programming with Caravela. The array of buffers for the host side is returned in order to initialize them in the main program. 8) _MapInputBufferToShader Here we should transfer the input data from the host buffers to the device buffers. Therefore we use the function clEnqueueWriteBuffer for each input stream to execute the copy operation. 9) _FireShader In this function we should prepare the kernel arguments for the execution and run this kernel. Thus, we use the clSetKernelArg function to assign the input buffers, the output buffers and the constant buffers to the kernel. Then, the function clEnqueueNDRangeKernel is used to execute the kernel. Because the parameters of the clEnqueueNDRangeKernel function change depending on the NDRange dimension selected by the user (for example, if the dimension is one the function expects just an int in the GlobalWorkSize and LocalWorkSize parameters but if the dimension is more than one the function expects an array in those parameters), we have implemented in the code three different calls to the function (one for each dimension) controlled by a conditional operator. 10) _GetOutputBuffer In this function we must return the host side output buffer with the output data after the execution. But before we must transfer the data between the buffers (device to host) so we use the function clEnqueueReadBuffer . 33 11) _FreeOutputBufferFromShader This function frees both the device buffers and the host buffers. Therefore for the device buffers we use the function clReleaseMemObject . We must free also the arrays that contain the buffers. 12) _UnmapInputBufferFromShader The same as the previous function but for the input buffers (device and host). 13) _FinalizeLowlevel Here we release all the OpenCL resources used in the process like the kernel, program, command queue and the context. So the functions clReleaseKernel, clReleaseProgram, clReleaseCommandQueue and clReleaseContext are used. All these low-level functions will be programmed in a new .DLL project called “CaravelaOCL”, and will be used by the Caravela API functions. Next we will explain the FlowModel structure used by Caravela and the changes in this structure required for the OpenCL version. 3.3 Flow Model The flow-model is a structure that packs all the information used by the Caravela in a XML file to execute the program. Using this structure, the programmer only has to worry about writing a flow-model and mapping it to the Caravela environment using the Caravela API functions. Then the Caravela uses the information of this flow-model to perform the execution of the kernel. The values required by the flow-model are: i) NumData. The number of data that will be processed. For example, if we want to add 2 vectors of 1024 elements, NumData will be 1024. ii) DataType. The type of the values of the streams (FLOAT, INT, SHORT). iii) NumInput. Number of input streams. iv) NumOutput. Number of output streams. 34 v) ShaderProgram. Here we write the kernel program that will be executed on the GPU. vi) FunctionName. The name of the function of the kernel program. vii) LangType. Language used to write the shader program. viii) RuntimeType. Runtime that we want to use to execute the program. ix) ShaderVersion. Version of the shader language (not used in OpenCL). x) ConstValues. List of the constant values used in the kernel. This list has a special format that will be explained in the next section. xi) ConstTypes. List of types (FLOAT, INT, SHORT) of the constant values. xii) ConstNames. List of names of the constant values. xiii) NumConstant. Number of constant values. In addition to the previous parameters, OpenCL needs the user to introduce 3 more values to execute the kernel. These are: xiv) Dimension. The dimension (1, 2, 3) of the NDRange. xv) Threads. This list of 3 elements is the GlobalWorkSize in each dimension. xvi) Blocks. This list of 3 elements is the LocalWorkSize in each dimension. The Figure 15 shows an example of an XML flow-model. 35 Figure 15. Flow Model of VectorAdd example Next we will explain the pattern designed for the constant values list. 3.3.1 Constant values structure In the previous version of the Caravela, each constant value had the structure of texture data, i.e. 4 data per pixel (Red, Green, Blue, and Alpha) because it used the graphics runtimes OpenGL and DirectX. However OpenCL does not need that structure, so a new one has been designed. The new structure is described as follows: a list of arrays in which each array is a list of constant values of the same type. Each array is considered one constant value, so if <?xml version="1.0"?> <FlowModelInfo xmlns:xsi="http://www.w3.org/2001/XMLSchema-instance" xmlns:xsd="http://www.w3.org/2001/XMLSchema"> <NumData>8</NumData> <DataType>FLOAT</DataType> <NumInput>2</NumInput> <NumOutput>1</NumOutput> <ShaderProgram> __kernel void VectorAdd(__global const float* input1, __global const float* input2, __global const float* const1, __global float* output){ int iGID = get_global_id(0); if (iGID >= const1[0]) { return ; } output[iGID] = input1[iGID] + input2[iGID]; } </ShaderProgram> <FunctionName>VectorAdd</FunctionName> <LangType>SHADERLANG_OPENCL1.0</LangType> <RuntimeType>RUNTIME_OPENCL</RuntimeType> <ShaderVersion>0</ShaderVersion> <ConstValues>1,256,0,0,0</ConstValues> <ConstTypes>FLOAT4</ConstTypes> <ConstNames>const_1</ConstNames> <NumConstant>1</NumConstant> <Dimension>1</Dimension> <Threads>256,1,1</Threads> <Blocks>256,1,1</Blocks> </FlowModelInfo> 36 we have 3 in the flow-model NumConstant parameter it means that we have 3 arrays of constant values. Each array can have any number of constant values. The ConstValues list follows this pattern: (X, x1,…x(4*X) , Y, y1,…y(4*Y), Z, z1,…z(4*Z)). The first number of each array (X, Y, Z) sets the number of constant values of the array multiplied per 4. Then, the values of the constant values of each array are written. So, for example in this list ( 1,256,0,0,0), we have 1 array of constant values with only one value “256”. The rest three “0” appear in the list because the FlowModelCreator reads a whole row of four values, due to the interface design, so it may be extra non-used values, like in this case. The user can access inside the kernel to every constant value in the arrays. However, the user has not to worry about these implementation details because with the GUI Flow Model Creator, explained in the next section, these issues are hidden. 3.4 Flow Model Creator The FlowModelCreator is a graphical user interface (GUI), programmed in C#, that allows the user to easily create a XML flow-model file. A more detailed explanation of this interface is provided in the appendix A, so in this section we will only explain the changes performed to the GUI in order to create a flow-model for OpenCL. The Figure 16 shows the top window with the new features highlighted. 37 Figure 16. Changes in the top window of FlowModelCreator First of all we modified the combo boxes “Shader Program” and “Runtime type” in order to add a new item for OpenCL. Then new methods for introducing the new OpenCL parameters of the flow-model (dimension, local work size and global work size) were added in the left side of the screen. Finally, a completely new window for introducing the constant values has been developed, as shown in Figure 17. Input Streams Output Stream 38 Figure 17. New interface for introducing the constant values This multi-tab window dynamically generates one tab for each constant array. The name of each array is automatically generated. Then, inside each tab we can introduce the type and values of the constants of the array. After closing the window the changes performed remain until we close the FlowModelCreator. 3.5 Kernel Structure In this section we will explain the structure of the kernel header that must be followed in order to be accepted by the Caravela. As stated before, a flow-model has 2 kinds of inputs (input stream and constant values) and 1 kind of output (output stream). There may be more than one of each kind of these. Therefore the kernel parameters must match this structure of the flow-model. In the CUDA version of the Caravela the solution was to make each parameter an array of arrays like this: function (**input, **constant, **output). With this pattern, each array contains all the streams of one kind (input, output and constants). However this is not possible in OpenCL because the buffers are 1-dimensional, so the parameters in an OpenCL kernel cannot be an array of arrays, i.e. 2-dimensional. Therefore the solution developed is to follow this pattern: function (*input1, *input2 …, *inputN, *const1, *const2 …, *constN, *output1, *output2 …, *outputN). 39 This means that the user has to write the kernel header as follows: first all the input streams, then the constant values arrays and finally the output streams. There must be at least one input stream and one output stream, but the constant value parameter is optional. If this order is not strictly followed the Caravela will not work properly and the results may be unexpected. Also if the number of kernel parameters does not match with the sum of the flow-model parameters “NumInput”, “NumOutput” and “NumConstant”, the Caravela will return an error warning this issue. Figure 18 shows a correct example of kernel header for the Caravela. Figure 18. Example of kernel header for Caravela 3.6 Swap Mechanism As mentioned in section 2.3.1, the swap mechanism of the Caravela allows the user to change (swap) the output and the input streams of the flow-model in order to perform repeated executions of a kernel without copying the results from the device to the host side. Therefore the data remains on the device side until it is required by the CPU. In OpenCL, the implementation of this mechanism has been realized using the clSetKernelArg function. First we defined a new low-level function called _SwapFlowmodelIOOCL that swaps the pointers of the device-side buffers of a previously defined I/O pair with the Caravela API function CARAVELA_CreateSwapIoPair . In order to know which input buffers have been swapped to not make a data transfer on them, we use a global array of bools called “swapped”. So for instance, “swapped [2] == true”, means that the buffer with the index 2 has been swapped. The values of this array are changed in the _SwapFlowmodelIOOCL, and are reinitialized after each execution. The _ SwapFlowmodelIOOCL function also makes “TRUE” a global flag called “swap” in order to warn the Caravela that we are using the swap mechanism so the input buffers do not have to be rewritten during the _ MapInputBufferToShader function. Then during the new fire operation, the swapped buffers are linked again to the kernel using the clSetKernelArg, , , , so the data is never copied back to the host side. __kernel void VectorAdd(__global const float* input1, __global const float* input2, __global const float* const1, __global float* output) 40 The implementation of this mechanism allows the user to rewrite with new data the input buffers that are not swapped. As mentioned before, if we use the swap mechanism the _ MapInputBufferToShader function does not rewrite the input buffers, but the user can change this using the GetInputData API function. If we use this function before the execution, the input buffer that we retrieve will be rewritten in the _ MapInputBufferToShader function even if we are using the swap mechanism. Notice that a swapped buffer will not be rewritten never even if we retrieve it with the GetInputData function. We have implemented this mechanism using a global array of bools called “rewrite”. If rewrite[i] == “true” the input buffer with the index i will be rewritten. The values of the array are changed in the _ GetInputBuffer function and are reinitialized after each execution. Because we swap also the host side buffers during the _SwapFlowmodelIOOCL function, the user does not have to worry about which buffer contains the real output data after the last execution, so he can retrieve the output data normally using the Caravela API function GetOutputData . Therefore with this method we can repeatedly execute a kernel without losing time copying the data to the host side after one execution and transfer it again to the device side before the next execution. 41 3.7 Implementation conclusions According to the implementation showed in this chapter, the flow-model for the OpenCL presents the same structure and behavior implemented on the previous versions of the Caravela. Also the GUI FlowModelCreator keeps the compatibility between different versions of the Caravela, and the user just has to select the runtime and the language in which the kernel program is written. Therefore, just creating a flow-model, the execution framework of the Caravela supports the runtimes with the legacy architecture of the GPUs and the runtimes with the recent architecture. Moreover, since OpenCL is a platform-independent API, the Caravela can be executed in any new GPU independently of the manufacturer. In the next chapter we will discuss the performance aspect on this implementation. 48 However when comparing the two OpenCL versions that do not transfer data back to the host side, we see that the execution time is almost the same in both versions. In the “OCL with copy device to device” version we use the OpenCL API function clEnqueueCopyBuffer to copy the data from the output buffer to the input buffer after each execution of the kernel. On the other hand, in the “OCL with swap” version we do not use this function because we just change the kernel arguments before the executions. For example, if we have a kernel with two arguments “kernel void test(float *input, float *output)” and two read/write buffers “src” and “dst”, in the first execution we use the clSetKernelArg function to set the buffer “src”, initialized with the input data, to the argument 0 (the input) and the buffer “dst” to the argument 1 (the output). Then before the second execution we change the order and we set the buffer “dst” to the argument 0 and the buffer “src” to the argument 1. This way the data stored in “dst” from the first execution is the input data in the second execution without explicitly perform a copy operation between buffers. Because the execution time of these versions is very similar, the function clSetKernelArg might implicitly perform a copy operation using the device bus when setting the buffers to the kernel. Due to this, the swap mechanism implemented in the Calavera, which also uses the clSetKernelArg function to make the swap, may not be a “true” swap between buffer pointers without any kind of copy operation. To clarify this issue, we have performed an extra evaluation. We have measured the execution times of a simple VectorAdd (a [] + b [] = c []) program, swapping the output and the input buffer using the previously mentioned swap methods. The results obtained are showed in figure 23 and 24. VectorAdd 5000 Swaps 115,044 146,46 9,22 7,578 0 50 100 150 200 256K 1MB 16MB Data size Execution Time (secs) Copy method SetKernelArg method Figure 23. VectorAdd execution times with 5000 swaps 49 VectorAdd 10000 Swaps 233,647 292,355 18,455 15,228 0 50 100 150 200 250 300 350 256K 1MB 16MB Data size Execution Time (secs) Copy method SetKernelArg method Figure 24. VectorAdd execution times with 10000 swaps As we see in the above graphs, there is difference in the execution time between these two methods. Using the data of the figure 24, we see that using 16 MB of data size there is a difference of approximately 60 seconds between the methods. Supposing that the cost of SetKernelArg function is 0, a single transference of 16MB of data between buffers on the device side takes approximately 60/10000 = 0,006 seconds. This time still seems very low, but it can be explained if we see our GPU bandwidth, which is approximately 25 GB/s. As a conclusion, difference between the execution times of these methods is small because the power of the hardware of the GPUs nowadays. Finally, as happened in the matrix multiplication evaluation too, the CaravelaOpenCL version is slower than the OpenCL without Caravela version. When we increase the number of swaps, due to the fact that the Calavera needs to access the driver level via several dynamic linked libraries so for each execution we have to access the low-level functions to get the input buffers and fire the kernel, the execution time is affected when compared to the OpenCL without Calavera versions. Therefore, the swap mechanism implemented maintains the performance improvement over methods that make transferences hostdevice each iteration. Also we have demonstrated that the mechanism implemented does not perform any data transference between buffers, even in the device side. 50 Chapter 5 Conclusions and future work In this project we have presented the migration process of a stream-based computing platform, Caravela, to the OpenCL language. As presented in chapter 1, the main objective of this project was to allow the Caravela to take advantage of the new GPU architecture and the new runtimes for GPGPU, without losing performance speed and the compatibility with the legacy runtimes (OpenGL, DirectX). Also another goal of this project was to maintain the structure and the execution style of the flow-model in the new version of the Caravela, as well as all the good features of the Caravela like the swap mechanism. According to the implementation presented in chapter 3, we have succeeded in preserving the flow-model structure and the features of the old version of the Caravela. In terms of performance evaluation, according to the results presented in chapter 4, the comparison using a straightforward application between the pure OpenCL version and the Caravela over OpenCL shows that the performance is only affected by fix degradation due to the process of reading and loading of the flow-model file. Also the swap mechanism, which prevents transferences between host memory and device memory when we execute iteratively a kernel, has been successfully implemented. Therefore we can conclude that the goals proposed in chapter 1 have been achieved. The Caravela has been proven to be a true stream-based computing platform for GPGPU that is able to take advantage of the most recent GPGPU APIs. Therefore it is a good idea to continue with the improvement of this software, such as the followings features: 1. Increase the compatibility of the Caravela porting it to the language DirectCompute. This new environment is part of the Microsoft DirectX collection of APIs and runs in both DirectX 10 and DirectX 11. With the migration of the Calavera to this language, will increase its compatibility with systems with the Windows OS. 2. Meta-pipeline. This mechanism allows a flow-model to be virtually connected to any other flow-model(s). It implements a virtual structure with multiple flow-models, which can be used, for instance, to solve large problems through a set of flow-models executed in different processing units. 51 3. Development of a hardware compiler for the flow-model execution method. This project will aim to develop an environment that generates hardware description from stream-based program. Because the hardware description can be implemented by pipelined components, the hardware design productivity can become better than the naturally implemented application by hand. 52 Bibliography (references) [1] At least nine contributions have emerged in recent years that explore the implications HPC presents for the future of humanities research. See “Links about Digital Humanities and HPC.” Available on-line at https://www.sharcnet.ca/dhhpc/index.php/Links_about_Digital_Humanities_and_HPC [June 29, 2008]. [2] Gustafson, J.; "The program of grand challenge problems: expectations and results," Parallel Algorithms/Architecture Synthesis, 1997. Proceedings. Second Aizu International Symposium , vol., no., pp.2-7, 17-21 Mar 1997 doi: 10.1109/AISPAS.1997.581619 [3] The Jaguar Supercomputer. http://www.nccs.gov/jaguar/. [4] J. Gummaraju and M. Rosenblum. Stream programming on general-purpose processors. In 38th Annual IEEE/ACM International Symposium on Microarchitecture (MICRO), pages 343–354, 2005. [5] GPGPU research projects of ICHEC http://www.ichec.ie/research/gpgpu_projects [6]Kapre, N.; DeHon, A.; "Performance comparison of single-precision SPICE ModelEvaluation on FPGA, GPU, Cell, and multi-core processors," Field Programmable Logic and Applications, 2009. FPL 2009. International Conference on, vol., no., pp.6572, Aug. 31 2009-Sept. 2 2009 doi: 10.1109/FPL.2009.5272548 [7] DirectX homepage. http://www.microsoft.com/directx. [8] OpenGL Architecture Review Board, D. Shreiner, M. Woo, J. Neider, and T. Davis. OpenGL Programming Guide: The Official Guide to Learning OpenGL, Version 2. Addison Wesley, 2005. 53 [9] I. Buck, T. Foley, D. Horn, J. Sugerman, K. Fatahalian, M. Houston, and P. Hanrahan. Brook for GPUs: stream computing on graphics hardware. ACM Trans. Graph., 23(3):777–786, 2004. [10] Sh: A high-level metaprogramming language for modern GPUs. http://libsh.org/. [11] NVIDIA CUDA Zone. Http://www.nvidia.com/cuda. [12] OpenCL. Http://www.khronos.org/opencl/. [13] S.Yamagiwa, L.Sousa. Design and Implementation of a Stream-Based Distributed Computing Platform using Graphics Processing Units. In ACM International Conference on Computing Frontiers, May 2007. [14] Yamagiwa, S.; Sousa, L.; "Caravela: A Novel Stream-Based Distributed Computing Environment," Computer, vol.40, no.5, pp.70-77, May 2007 doi: 10.1109/MC.2007.161 [15] Schatz, M.C., Trapnell, C., Delcher, A.L., Varshney, A. (2007) High-throughput sequence alignment using Graphics Processing Units. BMC Bioinformatics 8:474. [16] S. Yamagiwa, L. Sousa, K. Ferreira, K. Aoki, M. Ono, and K. Wada. Maestro2: Experimental evaluation of communication performance improvement techniques in the link layer. Journal of Interconnection Networks, 7(2):295-318, 2006. [17] GPGPU Homepage. Http://www.gpgpu.org/ [18] P. S. McCormick, J. Inman, J. P. Ahrens, C. Hansen, and G. Roth. Scout: A hardware-Accelerated System for Quantitatively Driven Visualization and Analysis. In VIS ’04: Proceedings of the conference on Visualization ’04, pages 171-178, Washington, DC, Usa, 2004. IEEE Computer Society 54 [19] Shinichi Yamagiwa, Leonel Sousa, Diogo Antao, "Data Buffering optimization methods toward a uniform programming interface for GPU-based applications", Proceeding of ACM Intl' Conference on Computing Frontiers, pp. 205 - 212, May 2007. [20] The Khronos Group. Http://www.khronos.org/. [21] Farrell, C.A.; Kieronska, D.H.; "Formal specification of SIMD execution," Algorithms and Architectures for Parallel Processing, 1996. ICAPP '96. 1996 IEEE Second International Conference on, vol., no., pp.319-325, 11-13 Jun 1996 doi: 10.1109/ICAPP.1996.562891 [22] DeBenedictis, E.P.; "Will Moore's Law Be Sufficient," Supercomputing, 2004. Proceedings of the ACM/IEEE SC2004 Conference, vol., no., pp. 4545, 06-12 Nov. 2004 doi: 10.1109/SC.2004.68 [23] http://en.wikipedia.org/wiki/DirectCompute [24] Thinking Machines Corporation. C* User’s Guide, June 1991. 55 Appendix I User guide of Flow Model Creator GUI. 1 Introduction 1.1 FlowModelCreator The FlowModelCreator is an interface designed for the Caravela platform to help the user to design a flow-model structure. These flow-models can be saved and loaded to/from a XML file. In this guide we will explain the functionality of this application 1.2 Authors The author of the first version of this interface is Professor Shinichi Yamagiwa (SIPS group, INESC-ID Portugal). The changes performed in order to support the OpenCL version of the Caravela, were developed by Pablo Lamilla Álvarez. 56 2 Functionality 2.1 Basics concepts As its name says, the FlowModelCreator creates a structure needed by the Caravela platform called flow-model. The flow-model is just a set of data that includes the information for a number of inputs, a number of outputs, a number of constants, constant values and a shader program. The FlowModelCreator will manage the way of packing all this information into a flow-model XML file. An example of a flow-model is shown in the next figure: In this example, the flow-model has two input data streams, one constant stream and two output streams, and it processes the input streams by the shader program. This program can be written in HLSL (High Level Shader Language) of DirectX, in GLSL (OpenGL Shading Language), or in OpenCL. 57 The flow-model is stored in a XML file that follows the Caravela specification. Because it is very hard for the user to manage manually the file contents and format, this FlowModelCreator provides a very easy interface to the user. 2.2 Top Window The following image illustrates the main window of the FlowModelCreator: 1) Home Tool Bar. From left to right, the buttons of this tool bar contains the followings functionalities: i) Create New Flow Model This button allows the creation of a flow-model from scratch. After clicking this button, you must choose the type of the input and output streams selecting one item from the “Data type” combo box in the Parameters section. After this action, the whole interface is unblocked and ready to use. ii) Open Local Flow Model File Parameters section Home Tool Bar Shader Program Section