Programming Multi-/Many-cores

Using OpenCL
(CS 3006)

Adapted from: see the Acknowledgement slide

Compiled by: Dr. Muhammad Arshad Islam and Dr. Muhammad
Aleem

National University of Computer & Emerging Sciences,

 Lecture
Acknowledgements
Simon McIntosh-Smith University of Bristol
[Link]

Romain Teissyer, Pierre-Franois Lavallée, et al.

[Link]
4

Optimizing OpenCL applications on Intel Xeon Phi

[Link]
Intel-Xeon- [Link]

OpenCL home page

[Link]
[Link]
June/OpenCL-Details-Taiwan_June-[Link]

AMD
[Link]
 OpenCL
• Open Compute Language
• For heterogeneous
parallel-
computing systems
• Cross-platform
− Implementations for
∙ ATI GPUs
∙ NVIDIA GPUs
• Intel MIC
• x86 CPUs
• Many others…..
 OpenCL is Widely Deployed and
Used

[Link]
 Khronos Compute Accleration

[Link]
Industry Standards for Programming Heterogeneous
Platforms

GPUs
CPUs Emerging Increasingly general
Multiple cores driving
Intersection purpose data-parallel
performance increases computing

Graphics
Multi- APIs and
Heterogeneous
processor Shading
programming Computing Languages
–
e.g. OpenMP

OpenCL – Open Computing Language

Open, royalty-free standard for portable, parallel programming of
heterogeneous parallel computing CPUs, GPUs, and other processors
Industry Standards for Programming Heterogeneous
Platforms

OpenCL – Open Computing Language

Open, royalty-free standard for portable, parallel programming of
heterogeneous parallel computing CPUs, GPUs, and other processors
 OpenCL
architecture

[Link]
 OpenCL Platform
Model
……
……
Processing ……
…
Element … … Host
…
…
……
…

Compute Unit OpenCL

Device

∙ One Host and one or more OpenCL Devices

− Each OpenCL Device is composed of one or more
Compute Units
∙ Each Compute Unit is divided into one or more Processing Elements

∙ Memory divided into host memory and device memory

Credits: Trevett and Cyril Zeller, NVIDIA
 OpenCL Program: Host and
Kernel

[Link]
6
 OpenCL Platform
(One node, two CPU sockets, two GPUs)
Example
CPUs: GPUs:
∙ Treated as one OpenCL ∙ Each GPU is a separate
device OpenCL device
− One CU per core
∙ One CU per Streaming
− 1 PE per CU, or if PEs mapped
Multiprocessor
to SIMD lanes, n PEs per CU,
where n matches the SIMD ∙ Can use CPU and all GPU
width devices concurrently through
∙ Remember: OpenCL
− the CPU will also have to be
its own host!

CU = Compute Unit; PE = Processing Element

Heterogeneous Computing – A High level
View

Credits: Andreas Moshovos,

[Link]
[Link] Credits: Andreas Moshovos, [Link]
 OpenCL Memory model Overview

• Private Memory
− Per work-item

• Local Memory
− Shared within a
work-group
• Global /Constant
Memory (Read-only)
Visible to all
work-groups
• Host memory
− On the CPU

Memory management is explicit:

You are responsible for moving data from
 Traditional Vs. OpenCL Parallel
Programming

[Link]
The BIG idea behind OpenCL
• Replace loops with functions (a kernel) executing at each point
in a problem domain
– E.g., process a 1024x1024 image with one kernel invocation per pixel or
1024x1024=1,048,576 kernel executions

Traditional loops OpenCL

void
kernel void
mul(const int n, mul( global const float *a,
const float *a, global const float *b,
const float *b, global float *c)
float *c) {
{ int id = get_global_id(0);
int i; c[id] = a[id] * b[id];
for (i }
= 0; i < n; i++)
over n work-items
c[i] = a[i] * b[i]; // execute
}
 An N-dimensional domain of
•
work-items
Global Dimensions:
– 1024x1024 (whole problem space)
– For example, if we have a 2D problem with dimensions (1024, 1024), then
the
NDRange would be defined as (1024, 1024).

– 128x128 (work-group, executes together)

1024
• Local Dimensions: Synchronization between
items possible only within work-
work- groups: barriers and memory fences
1024

Cannot synchronize among

work- groups

• Choose the dimensions (“best” for your algorithm):

• 1 (one dimension)
• 2 (two dimension)
• 3 (three dimension)
An N-dimensional domain of
work-items

•[Link]
•[Link]
An N-dimensional domain of
work-items

•[Link]
•[Link]
 Context and
• Command-Queues
Context:
– The environment within which kernels
execute and in which
synchronization and memory Device
management is defined.
• The context includes:
– One or more devices Device Memory

– Device memory
– One or more command-queues Queue

• All commands for a device (kernel

execution, synchronization, and
memory operations) are submitted
through a command-queue. Context

• Each command-queue points to a single

device within a context.
 Building Program
Objects
• The program object encapsulates:
OpenCL uses runtime
1. A context compilation … because
2. The program source or binary, in general you don’t
and know the details of the
target device when you
3. List of target devices and build ship the program
• The options
build process to create a
program object:
cl::Program program(context, KernelSource);
kernel void
horizontal_reflect(read_only
write_onlyimage2d_t
image2d_tsrc,
{ dst) Compile for GPU
int x = get_global_id(0); // GPU code
int y = get_global_id(1); x-coord
int width = //
get_image_width(src);
float4 src_val = read_imagef(src,
y-coord sampler,
(int2)(width-1-x, Compile for CPU
y)); write_imagef(dst, (int2)(x, y),
CPU code
} src_val);
 Execution model(kernels)

• OpenCL execution model … define a problem domain

execute an instance of a kernel
and each point in
for the
domain
global float* input,
kernel void times_two(
global float* output)
{ int i = get_global_id(0);
output[i] = 2.0f *
input[i];
}
get_global_id(0)
10
Input 0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25

Output 0 2 4 6 8 10 12 14 16 18 20 22 24 26 28 30 32 34 36 38 40 42 44 46 48 50
 Vector Addition -
Kernel
kernel void
vadd(
global const float *a,
global const float *b,
global float *c)
{
int gid =
get_global_id(0);
c[gid] = a[gid] + b[gid];
}
UNDERSTANDING THE HOST
PROGRAM
 Vector Addition – Host
• The host program is the code that runs on the host to:
– Setup the environment for the OpenCL program
– Create and manage kernels

• 5 simple steps in a basic host program:

1. Define the platform … platform = devices+context+queues
2. Create and Build the program (dynamic library for kernels)
3. Setup memory objects
4. Define the kernel (attach arguments to kernel function)
5. Submit commands … transfer memory objects and execute
kernels
 The C++
Interface
• Khronos has defined a common C++ header file containing
a
high level interface to OpenCL, [Link]

• Key features:
– Uses common defaults for the platform and
command- queue,
– Simplifies the basic API
– Ability to “call” a kernel from the host, like a regular
function
– Error checking can be performed with C++
exceptions
C++ Interface: setting up the host program

• Enable OpenCL API Exceptions. Do this before

#define CL_ENABLE_EXCEPTIONS including the
header files

• Include key header files … both standard and

custom
#include <CL/[Link]> // Khronos C++ Wrapper API
#include <cstdio> // C style IO (e.g. printf)
#include <iostream> // C++ style IO
#include <vector> // C++ vector types
 1. Create a context and
queue
• Grab a context using a device type:
cl::Context context(CL_DEVICE_TYPE_DEFAULT);

Or…CL_DEVICE_TYPE_CPU,
CL_DEVICE_TYPE_GPU,
CL_DEVICE_TYPE_ACCELERATOR,
etc.

• Create a command
the context: queue for the queue(context);
cl::CommandQueue device in
 Commands and Command-Queues
• Commands include:
– Kernel executions
– Memory object management GPU CPU

– Synchronization

Queu Queu
• The only way to submit commands e e
to a device is through a command-
queue.
Conte
• Each command-queue points xt
to a single device within a
context.
Command-Queue execution details
• Command queues can be configured in
different ways to control how commands GPU CPU

execute
Queu Queu
e e
• In-order queues:
– Commands are enqueued and Conte
xt
complete in the order they appear in
the host program (program-order)

• Out-of-order queues:
– Commands are enqueued in
order but can execute (and
program-
complete)
hence in any
order.
 2. Create and Build the
• Define
program
source code for the kernel-program either as a
string literal (for small programs) or read it from a file (for
large applications).

• Create the program object and compile

: “true” tells OpenCL to build
(compile/link) the program object

cl::Program program(context, KernelSource , true );

KernelSource is a string … either statically set in the host program or

returned from a function that loads the kernel code from a file.
 3. Setup Memory Objects
• For vector addition we need 3 memory objects: one each
for input vectors A and B, and one for the output vector C

• Create input vectors and assign values on the host:

std::vector<float> h_a(LENGTH), h_b(LENGTH), h_c(LENGTH); for (i =
0; i < LENGTH; i++) {
h_a[i] = rand() /
(float)RAND_MAX; h_b[i] = rand()
/ (float)RAND_MAX;
}
• Define OpenCL device buffers and copy from host buffers:
cl::Buffer d_a(context,
h_a.begin(), h_a.end(), true);
cl::Buffer d_b(context,
h_b.begin(), h_b.end(), true);
cl::Buffer d_c(context, CL_MEM_WRITE_ONLY,

sizeof(float)*LENGTH);
or CL_MEM_READ_ONLY or CL_MEM_READ_WRITE
 What do we put in device memory?
• Memory Objects:

• There are two kinds of memory object

– Buffer object (Always linear):
•Defines a linear collection of bytes.

– Image object:
•Defines a two- or three-dimensional region
of memory.
•Image data can only be accessed with read and
write functions
 Creating and manipulating
buffers
• Buffers are declared on the host as object
cl::Buffer
type:

• Arrays in host memory hold your original host-side

data:
std::vector<float> h_a, h_b;
• Create the device-side buffer (d_a), assign read-only
memory to hold the host array (h_a) and copy it into
device memory:
cl::Buffer d_a(context, h_a.begin(), h_a.end(), true);

Start_iterator and end_iterator for the Stipulates that this is a

container holding host side object read-only buffer

or use function clCreateBuffer ( ) to create a device-side buffer

 Creating and manipulating
buffers
• The last argument sets the read/write access to the Buffer
by the device. true means “read only” while false (the
default) means “read/write”.

• Submit command to copy the device buffer back to host

memory in array “h_c”:
cl::copy(queue, d_c, h_c.begin(), h_c.end());

• Can also copy host memory to device buffers:

cl::copy(queue, h_c.begin(), h_c.end(),
d_c);
Creating and manipulating
buffers
 4. Define the
kernel
• Create a kernel function for the kernels you want to be
able to call in the program:

Must match the pattern of

arguments to the kernel.

cl::make_kernel<cl::Buffer,cl::Buffer,cl::Buffer>
vadd(program, “vadd”);

A previously created The name of the function

“program object” serving as used for the kernel
a dynamic library of kernels

• This means you can ‘call’ the kernel as a

‘function’ in your host code to enqueue the
kernel.
 5. Enqueue
• commands
For kernel launches, specify global and local
dimensions
– cl::NDRange global(1024)
– cl::NDRange local(64)
If you don’t specify a local dimension, it is assumed as
cl::NullRange, and the runtime picks a size for you
 5. Enqueue
commands
Enqueue the kernel for execution (note: returns immediately …
i.e. this is a non-blocking command):
vadd(cl::EnqueueArgs(queue, global), d_a, d_b, d_c);

Read back result (as a blocking operation). We use an in-order

queue to assure the previous commands are completed before
the read can begin

cl::copy(queue, d_c, h_c.begin(), h_c.end());

 OpenCL Vector Addition
#include <iostream>
#include <vector>Example
#include <CL/[Link]>

int main() {
std::vector<float> a = {1.0f, 2.0f, 3.0f, 4.0f};
std::vector<float> b = {4.0f, 3.0f, 2.0f, 1.0f};
std::vector<float> c([Link]());

try {
// get available OpenCL platforms
std::vector<cl::Platform> platforms;
cl::Platform::get(&platforms);

// choose a platform
cl::Platform platform = platforms[0];

// get available OpenCL devices

std::vector<cl::Device> devices;
[Link](CL_DEVICE_TYPE_ALL, &devices);

// choose a device
cl::Device device = devices[0];

// create an OpenCL context for the device

cl::Context context({device});

// create an OpenCL program from source

cl::Program program(context, "[Link]");
 OpenCL Vector Addition
Example
// build the program for the device
[Link]({device});

// create OpenCL buffers for the input and output vectors

cl::Buffer bufA(context, CL_MEM_READ_ONLY, [Link]() * sizeof(float));
cl::Buffer bufB(context, CL_MEM_READ_ONLY, [Link]() * sizeof(float));
cl::Buffer bufC(context, CL_MEM_WRITE_ONLY, [Link]() * sizeof(float));

// create a command queue for the device

cl::CommandQueue queue(context, device); True:
Readonly
// enqueue data to be transferred to the device
[Link](bufA, CL_TRUE, 0, [Link]() * sizeof(float), [Link]());
[Link](bufB, CL_TRUE, 0, [Link]() * sizeof(float), [Link]());

// create a kernel object for the vector addition kernel

cl::Kernel kernel(program, "vecadd");

// set the arguments of the kernel

[Link](0, bufA); Global NDRange
[Link](1, bufB);
Global offset LocalNDRange
[Link](2, bufC);

// enqueue the kernel for execution

[Link](kernel, cl::NullRange, cl::NDRange([Link]()), cl::NullRange);

// enqueue data to be transferred back to the host

[Link](bufC, CL_TRUE, 0, [Link]() * sizeof(float), [Link]());
 OpenCL Vector Addition Example

// print the
result for (float
x : c) {
std::cout << x << " ";
}

std::cout << std::endl;

} catch (cl::Error& e) {
std::cerr << "OpenCL error: " << [Link]() << " (" << [Link]() << ")"
<< std::endl;
return 1;
}

return 0;
}
INTRODUCTION TO
OPENCL KERNEL
PROGRAMMING
 OpenCL
kernel
• Derived from ISO C99
– A few restrictions: no recursion, function
pointers, functions in C99 standard headers ...
– Preprocessing directives defined by C99
are supported (#include etc.)

• Built-in data types

– Scalar and vector data types, pointers
– Data-type conversion functions:
• convert_type<_sat><_roundingmode>
– Image types: image2d_t, image3d_t and sampler_t
 OpenCL C Language Highlights
• Function qualifiers
– __kernel qualifier declares a function as a kernel
• I.e. makes it visible to host code so it can be enqueued

• Address space qualifiers

– __global, __local, __constant, __private
– Pointer kernel arguments must be declared with
an address space qualifier
 OpenCL C Language Highlights

– uint get_work_dim() … number of dimensions in use (1,2, or 3)

– size_t get_global_id(uint n) … global work-item ID in dim “n”
– size_t get_local_id(uint n) … work-item ID in dim “n” inside work-
group
– size_t get_group_id(uint n) … ID of work-group in dim “n”
– size_t get_global_size(uint n) … num of work-items in dim “n”
– size_t get_local_size(uint n) … num of work-items in work group in
dim “n”
 OpenCL C Language Highlights
• Synchronization functions
– Barriers - all work-items within a work-group must execute
the barrier functionbefore any work-item can continue
– Memory fences - provides ordering between memory
operations
 OpenCL C Language Restrictions
• Pointers to functions are not allowed

• Pointers to pointers allowed within a kernel, but

not as an argument to a kernel invocation

• Variable length arrays and structures are

not supported

• Recursion is not supported (yet!)

 Matrix multiplication: sequential code
We calculate C=AB, dimA = (N x P), dimB=(P x M), dimC=(N x M)

void mat_mul(int Mdim, int Ndim, int Pdim,

float *A, float *B, float *C)
{
int i, j, k;
for (i = 0; i < Ndim; i++) {
for(j = 0; j < Mdim; j++) {
C[i*N+j] = 0.0f;
for (k = 0; k < Pdim; k++) {
// C(i, j) = sum(over k) A(i,k) * B(k,j)
C[i*Ndim+j] += A[i*Ndim+k] * B[k*Pdim+j];
}

} A(i,:)
C(i,j)
} = x B(:,j)
}

Dot product of a row of A and a column of B for each element of

C
 Matrix multiplication: sequential
Wecode
calculate C=AB, where all three matrices are NxN

void mat_mul(int N, float *A, float *B, float *C)

{
Let’s make it
int i, j, k; easier and
for (i = 0; i < N; i++) { specialize to
for(j = 0; j < N; j++) { square matrices
C[i*N+j] = 0.0f;
for (k = 0; k < N; k++) {
// C(i, j) = sum(over k) A(i,k) * B(k,j)
C[i*N+j] += A[i*N+k] * B[k*N+j];
}
}
} A(i,:)
C(i,j)
} = x B(:,j)

Dot product of a row of A and a column of B for each element of

C
 Matrix multiplication
performance
• Serial C code on CPU (single core).
Case MFLOPS
CPU GPU
Sequential C (not OpenCL) 887.2 N/A

Device is Intel® Xeon® CPU, E5649 @ 2.53GHz

using the gcc compiler.

These are not official benchmark results. You

may observe completely different results should
you run these tests on your own system.
Third party names are the property of their owners.
Matrix multiplication: sequential
code
void mat_mul(int N, float *A, float *B, float *C)
{
int i, j, k;
for (i = 0; i <
N; i++) {
for (j = 0; j < N; j++) {
C[i*N+j] = 0.0f;
for (k = 0; k < N; k++) {
// C(i, j) = sum(over k) A(i,k) * B(k,j)
C[i*N+j] += A[i*N+k] * B[k*N+j];
}
}
}
}
Matrix multiplication: sequential code
void mat_mul(int N, float *A, float *B, float *C)

{
int i, j, k;
for (i = 0; i <
N; i++) {
for (j = 0; j < N; j++) {
C[i*N+j] = 0.0f;
for (k = 0; k < N; k++)
{
// C(i, j) = sum(over
k) A(i,k) * B(k,j)
C[i*N+j] += A[i*N+k] * B[k*N+j];
}
}
}
}
Matrix multiplication: OpenCL kernel (1/2)

kernel void mat_mul(const int N, global float *A,

global float *B, global float
{
int i, j, k;
for (i = 0; i < N; i++) {
for (j = 0; j < N; j++) {
C[i*N+j] = 0.0f;
for (k = 0; k < N; k++) {
// C(i, j) = sum(over k) A(i,k) * B(k,j)
C[i*N+j] += A[i*N+k] * B[k*N+j];
}
}
} Mark as a kernel function and
} specify memory qualifiers
Matrix multiplication: OpenCL kernel (2/2)

kernel void mat_mul(const int N, global float *A,

global float *B, global float
{
int i, j, k;
i = get_global_id(0);
j = get_global_id(1);
C[i*N+j] = 0.0f;
for (k = 0; k < N; k++) {
// C(i, j) = sum(over k) A(i,k) * B(k,j)
C[i*N+j] += A[i*N+k] * B[k*N+j];
}
}
} Replace loops with the
} work item’s global id
Matrix multiplication: kernel
cleaned-up
Rearrange and use a local scalar for intermediate C element
values (a common optimization in Matrix Multiplication functions)

{
kernel voidmmul( int k;
const int N, int i = get_global_id(0);
global float *A, int j = get_global_id(1);
global float *B,
float tmp = 0.0f;
global float *C) for (k = 0; k < N; k++)
tmp += A[i*N+k]*B[k*N+j];

C[i*N+j] = tmp;
}
 Matrix multiplication
performance
• Matrices are stored in global memory.

Case MFLOPS
CPU GPU
Sequential C (not OpenCL) 887.2 N/A
C(i,j) per work-item, all global 3,926. 3,720.9
1

Device is Tesla® M2090 GPU from NVIDIA® with a max of 16 compute units, 512 PEs
Device is Intel® Xeon® CPU, E5649 @ 2.53GHz

These are not official benchmark results. You may

observe completely different results should you run
these tests on your own system.
Third party names are the property of their owners.
Another Test & Optimization
 OpenCL and CUDA
∙ Many OpenCL features have a one to one mapping to CUDA
features

∙ OpenCL
− More complex platform and device management
− More complex kernel launch
− More lower level than CUDA
 OpenCL and CUDA

• Compute Unit (CU) correspond to

− CUDA streaming multiprocessors
(SMs)
− CPU core
− etc.
• Processing Element correspond to
− CUDA streaming processor (SP)
− CPU ALU
 OpenCL and CUDA
• Work Item (CUDA thread) – executes kernel code

• Index Space (CUDA grid) – defines work items and how data is mapped to
them

• Work Group (CUDA block) – work items in a work group can synchronize
 References
Optimizing OpenCL applications on Intel Xeon Phi
[Link]
Intel-
[Link]

OpenCL home page

[Link]

One OpenCL to Rule Them All

Romain Dolbeau, Francois Bodin, Guillame Colin de Verdière
[Link]
them-
[Link]

Ramses Project
Romain Teissyer, Pierre-Franois Lavallée, et al.
[Link]