Multi-threaded Programming

Introduction

OpenCL is a standard which defines a framework, an API and a programming language for parallel computation on heterogeneous systems like client computer systems, high- performance computing servers as well as hand-held devices.

The standard is maintained by the Khronos Group and supported by a large consortium of industry leaders including Apple, Intel, AMD, NVIDIA and ARM. Influenced by NVIDIA’s CUDA from the GPU side and by OpenMP which originates from the classical CPU side, the open OpenCL standard is characterized by a formulation which is abstract enough to support both CPU and GPU computing resources.

This is an ambitious goal, since providing an abstract interface together with a peak performance is a challenging task. OpenCL employs a strict isolation of the computation work into fundamental units, the kernels. These kernels can be developed in the OpenCL C programming language, a subset of the C99 language, with some additional OpenCL specific keywords.

In general, these kernels are hardware independent and compiled by the OpenCL runtime when they are loaded. To be able to fully exploit the parallel execution of the kernel code, several kernel instances, the work items, are started to process a set of input values. The actual number of concurrently running work items is determined by the OpenCL system. How a concrete algorithm can be partitioned into work items has to be decided by the programmer.

Reference Material

• Khronos Group OpenCL

  http://www.khronos.org/opencl/

• OpenCL 1.2 Quick Reference Card

  http://www.khronos.org/files/opencl-1-2-quick-reference-card.pdf

• OpenCL 1.2 Full Documentation

  http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/

• Intel SDK for OpenCL

  http://software.intel.com/en-us/articles/vcsource-tools-opencl-sdk/

• AMD OpenCL Zone

  http://developer.amd.com/zones/OpenCLZone/

• NVIDIA OpenCL

  http://www.nvidia.com/object/cuda_opencl_1.html

Programming Environment and Libraries

c++ interface / blah

Project: Boostraping OpenCL and Vector Addition

TODO: give overview of what to do

Compiling and running the test program

Open the folder project_vectoradd, create the build files using CMake and compile the application.

$ cd project_vectoradd/
[hauth@vdt-corei7avx project_vectoradd]$ cmake . 
[hauth@vdt-corei7avx project_vectoradd]$ make 
[100%] Built target vectoradd
[hauth@vdt-corei7avx project_vectoradd]$ ./vectoradd 

Once the application was compiled successully run it. The output should be along the following lines:

$ ./vectoradd
Testing Platform : Intel(R) OpenCL
> Selected Compute Device :        Intel(R) Core(TM) i7-3930K CPU @ 3.20GHz
Transferring data to device memory took 2e-06 s
Running vectorAdd kernel took 2e-06 s
Transferring data to host memory took 0 s
All done

Task: Understand the existing code

The first task is to read and understand the existing source code. To do so, you have to edit the file vectoradd.cpp in your favorite text editor, use nano if you are not sure which tool to use.

$ nano vectoradd.cpp

Take your time to familiarize yourself with the sourcecode which is already in the file. Some of the intial steps of setting up the OpenCL system are already provided:

• OpenCL compute context

  An OpenCL platform is automaticaly selected, depending on the required device type (CPU/GPU) and a compute context is created. You can change the required device type by modifying the constant devType:

  // Desired Device type.
  // can be CL_DEVICE_TYPE_GPU or CL_DEVICE_TYPE_CPU in this example
  const cl_device_type devType = CL_DEVICE_TYPE_CPU;

• OpenCL command queue

  A command queue is created on the context

• Kernel compile

  The variable contains kernelSourceAdd the source code of the OpenCL kernel. This string ist used to compile and register the kernel with the OpenCL runtime system. If you want to extend the kernel code, you have to add to the string contained in the variable kernelSourceAdd.

• Host data buffers

  The variables hostVectorA, hostVectorB and hostVectorC hold buffers which will be used later to transfer data to and from the OpenCL kernel. The underlying C++ type is std::vector which can be conviniently used to acquire a contingious space of memory in host’s RAM.

  FloatingPoint vp = 23.0f;
  FloatingPointVector hostVectorA(itemCount, vp);
  FloatingPointVector hostVectorB(itemCount, vp);
  FloatingPointVector hostVectorC(itemCount, vp);

• Various timing measurements

  To quantify the runtime of the OpenCL operations, various timing objects are present. For the time measurements to be correct, only the operations indicated by the TODO comments must be between the creation of the Timing object and the call to EndWithReport() of the respective timing object.

  Timing t_transfer_input("Transferring data to device memory");
  // TODO: transfer input buffer here
  t_transfer_input.EndWithReport();

Task: Running the OpenCL Kernel

The OpenCL kernel is already created and associated with the compiled program:

cl::Kernel kernelAdd(pProgram, "vectorAdd");

Use the enqueueNDRangeKernel(…) method on the command queue object to add a kernel instance to the queue <ref>http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clEnqueueNDRangeKernel.html </ref>

Call the finish() method on the command queue object to ensure the kernel is done processing before the host program continues execution. <ref>http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/clFinish.html </ref>

Modifications to play around

• switch to double , how does the runtime change for CPU/GPU ?
• switch to float4 vector type. can you perform the same addition operations ?

Project: N-Body Simulation

<references />