Tag Archives: c

Collapsing two arrays into one in C and C++

A quick tip – if you have two arrays and you want to merge/collapse/concatenate/append/combine (whatever you want to call it) them into one here is one way to do it in C and C++. The code also prints the combined array to check it’s correct.

Note the subtle differences between C and C++. One of the (innumerable!) problems with C++ is that you can write as much C as you want in it and leave the next developer wondering why you were writing C instead of C++ and whether you had specific justifiable reasons to do so in each instance. I personally also consider it bad practice. This is one of the reasons I dislike hybrid languages – they allow you to blend two languages together in arbitrary proportions making the end result unreadable and unmaintainable and just raising more questions than the problems that it solves.

C

#include <stdio.h>
#include <stdlib.h>
#include <string.h>

int main () {

   double a[4] = {1,2,3,4};
   double b[4] = {5,6,7,8};

   double* c = malloc(8 * sizeof(double));

   memcpy(c,     a, 4 * sizeof(double));
   memcpy(c + 4, b, 4 * sizeof(double));

   for (int i = 0; i < 8; i++) {
      printf("%d %fn", i, c[i]);
   }

   free(c);

}

C++

#include <iostream>
#include <cstring>
#include <iomanip>
       
using namespace std;
       
int main () {
       
   double a[4] = {1,2,3,4};
   double b[4] = {5,6,7,8};
       
   double* c = new double[8];
       
   copy(a, a + 4, c); 
   copy(b, b + 4, c + 4); 
       
   for (int i = 0; i < 8; i++) {
      cout << i << " " << c[i] << endl;
   }
       
   delete[] c;
       
}   

Note that in the C++ program:

  • I’m not using includes from C
  • I’m not using memcpy from C
  • I’m not using printf from C
  • I’m not using malloc from C
  • I’m not using sizeof from C
  • I’m not using free from C

OpenCL Cookbook: Compiling OpenCL with Ubuntu 12.10, Unity, AMD 12.11 beta drivers & AMD APP SDK 2.7

Continuing on in the OpenCL cookbook series here I present a post not about code but about environmental setup further diversifying the scope of the cookbook. It can be a real challenge for the uninitiated to install all the above and compile an opencl c or c++ program on linux. Here’s a short guide. First download and install ubuntu (duh!).

Install ubuntu build tools and linux kernel extras

Then install the following packages which are a prerequisite to the amd installers and the subsequent c/c++ compilation.

sudo apt-get update
sudo apt-get install build-essential
sudo apt-get install linux-source
sudo apt-get install linux-headers-generic

Then download AMD 12.11 beta drivers (amd-driver-installer-catalyst-12.11-beta-x86.x86_64.zip) and AMD APP SDK 2.7 (AMD-APP-SDK-v2.7-lnx64.tgz). Obviously download either 32bit or 64bit based on what your system supports.

AMD 12.11 beta drivers installation

Once you’ve done that install the AMD 12.11 beta drivers as root first. Installation is as simple as extracting the tarball, marking the script inside as executable and running the script as root. Reboot. After the reboot unity should start using the new AMD 12.11 beta driver and you’ll know it’s the beta because you’ll see a watermark at the bottom left of the screen saying ‘AMD Testing use only’. Note that the reason we’re using the beta here is because unity does not work with earlier versions of the driver. You get a problem where you see the desktop background and a mouse pointer but there’s no toolbar or status bar. But the 12.11 beta driver works which is great.

AMD APP SDK 2.7 installation

Then install the AMD APP SDK 2.7 also as root. Again installation is very simple and exactly the same as for the beta driver above. The AMD beta drivers install a video driver and the OpenCL runtime. The AMD APP SDK install the SDK and also OpenCL and OpenGL runtimes. However if you’ve already installed the video driver first you’ll already have the OpenCL runtime on your system in /usr/lib/libamdocl64.so so the APP SDK won’t install another copy in its location of /opt/AMDAPP/lib/x86_64/libOpenCL.so. You’ll see some messages during installation that it’s skipping the opencl runtime and that’s absolutely fine for now.

Test your OpenCL environment

Now you should test your OpenCL environment by compiling and running an example c opencl program. Get my C file to list all devices on your system as an example calling it devices.c and compile as follows.

gcc -L/usr/lib -I/opt/AMDAPP/include devices.c -lamdocl64 -o devices.o # for c
g++ -L/usr/lib -I/opt/AMDAPP/include devices.c -lamdocl64 -o devices.o # for c++

Once compiled run the output file (devices.o) and if it works then you should output similar to that below.

1. Device: Tahiti
 1.1 Hardware version: OpenCL 1.2 AMD-APP (923.1)
 1.2 Software version: CAL 1.4.1741 (VM)
 1.3 OpenCL C version: OpenCL C 1.2 
 1.4 Parallel compute units: 32
2. Device: Intel(R) Xeon(R) CPU E5-2687W 0 @ 3.10GHz
 2.1 Hardware version: OpenCL 1.2 AMD-APP (923.1)
 2.2 Software version: 2.0 (sse2,avx)
 2.3 OpenCL C version: OpenCL C 1.2 
 2.4 Parallel compute units: 32

Enabling multiple gpus for OpenCL

You may find that you are only seeing one gpu in your opencl programs. There are two things you need to do to enable multiple gpus in the OpenCL runtime. The first is to disable all crossfire. You can do this either in the amd catalyst control centre > performance which you start by running amdcccle or you can do it using the awesome amdconfig tool by running amdconfig --crossfire=off. See my post on amdconfig to find out more about this incredibly powerful tool.

The second thing you may or may not need to do is to enable COMPUTE mode as follows.

export COMPUTE=:0

Once you’ve done the above you should see program output from the program above similar to below.

dhruba@debian:~$ ./source/devices.o 
1. Device: Tahiti
 1.1 Hardware version: OpenCL 1.2 AMD-APP (1084.2)
 1.2 Software version: 1084.2 (VM)
 1.3 OpenCL C version: OpenCL C 1.2 
 1.4 Parallel compute units: 32
2. Device: Tahiti
 2.1 Hardware version: OpenCL 1.2 AMD-APP (1084.2)
 2.2 Software version: 1084.2 (VM)
 2.3 OpenCL C version: OpenCL C 1.2 
 2.4 Parallel compute units: 32
3. Device: Tahiti
 3.1 Hardware version: OpenCL 1.2 AMD-APP (1084.2)
 3.2 Software version: 1084.2 (VM)
 3.3 OpenCL C version: OpenCL C 1.2 
 3.4 Parallel compute units: 32
4. Device: Tahiti
 4.1 Hardware version: OpenCL 1.2 AMD-APP (1084.2)
 4.2 Software version: 1084.2 (VM)
 4.3 OpenCL C version: OpenCL C 1.2 
 4.4 Parallel compute units: 32
5. Device: Tahiti
 5.1 Hardware version: OpenCL 1.2 AMD-APP (1084.2)
 5.2 Software version: 1084.2 (VM)
 5.3 OpenCL C version: OpenCL C 1.2 
 5.4 Parallel compute units: 32
6. Device: Tahiti
 6.1 Hardware version: OpenCL 1.2 AMD-APP (1084.2)
 6.2 Software version: 1084.2 (VM)
 6.3 OpenCL C version: OpenCL C 1.2 
 6.4 Parallel compute units: 32
7. Device: Intel(R) Xeon(R) CPU E5-2687W 0 @ 3.10GHz
 7.1 Hardware version: OpenCL 1.2 AMD-APP (1084.2)
 7.2 Software version: 1084.2 (sse2,avx)
 7.3 OpenCL C version: OpenCL C 1.2 
 7.4 Parallel compute units: 32

Standardising the OpenCL runtime library path

Now – it may be that you wish for the OpenCL runtime library to be installed in the standard AMD APP SDK location of /opt/AMDAPP/lib/x86_64/libOpenCL.so as opposed to the non-standard location of /usr/lib/libamdocl64.so which is where the beta driver installation puts it. The proper way to do this would probably be to install the AMD APP SDK first and then the video driver or simply skip the video driver installation (I haven’t tried either of these options so they may need verification).

However, I used a little trick to make this easier since I’d already installed the video driver followed by the APP SDK. I renamed /usr/lib/libamdocl64.so to /usr/lib/libamdocl64.so.x and reinstalled the APP SDK. This time it detected that the runtime wasn’t present and installed another runtime in /opt/AMDAPP/lib/x86_64/libOpenCL.so – the standard SDK runtime path. With the new APP SDK OpenCL runtime in place I was able to compile the same program using the new runtime as below depending on whether you want the c or c++ compiler.

gcc -L/opt/AMDAPP/lib/x86_64/ -I/opt/AMDAPP/include devices.c -lOpenCL -o devices.o # for c
g++ -L/opt/AMDAPP/lib/x86_64/ -I/opt/AMDAPP/include devices.c -lOpenCL -o devices.o # for c++

Summary

And there you have it – an opencl compiler working on ubuntu 12.10 using the AMD 12.11 beta drivers and the AMD APP 2.7 SDK. Sometimes you just need someone else to have done it first and written a guide and I hope this serves to help someone out there.

C++ error LNK2001: unresolved external symbol

A quick Visual C++ tip to help those who get stuck on this problem like I did. If you find yourself getting errors like this:

foo.obj : error LNK2001: unresolved external symbol "public: __thiscall MyMatrix::~MyMatrix(void)" (??1MyMatrix@@QAE@XZ)
foo.obj : error LNK2001: unresolved external symbol "public: __thiscall MyMatrix::MyMatrix(class MyMatrix const &)" (??0MyMatrix@@QAE@ABV0@@Z)
fooData.obj : error LNK2001: unresolved external symbol "public: __thiscall MyMatrix::MyMatrix(unsigned int,unsigned int)" (??0MyMatrix@@QAE@II@Z)
bar.obj : error LNK2001: unresolved external symbol "public: class MyMatrix & __thiscall MyMatrix::resize(unsigned int,unsigned int)" (?resize@MyMatrix@@QAEAAV1@II@Z)
.Release2MCProject-32bit-noxlw.exe : fatal error LNK1120: 4 unresolved externals

it could mean that certain methods are declared but not defined. In other words there may be methods in the header files that have no implementations in cpp files. Above we can see that a destructor, two constructors and a resize call are what the compiler’s complaining about. This was because those methods had not been implemented but were in the header file.

Though the really odd thing is that this project was compiling in 32 bit mode and when I tried to port it to 64 bit mode suddenly these errors crept up. I don’t know why these errors didn’t occur in 32 bit mode. Must be yet another compiling intricacy that I’m not aware of.

C++ can be a nightmare compared to higher level languages.

msvcprtd.lib(MSVCP100D.dll) : fatal error LNK1112: module machine type ‘X86’ conflicts with target machine type ‘x64’

If you get the above error when trying to build a project in 64 bits or port a 32 bits project to 64 bits in Visual Studio 2010 go to Project Property Pages > Configuration Properties > VC++ Directories > Library Directories and make sure you have the appropriate 64 bit directories in there and not the 32 bit equivalents.

I had to change my Library Directories entry from:

$(VCInstallDir)lib;$(VCInstallDir)atlmfclib;$(WindowsSdkDir)lib;$(FrameworkSDKDir)lib;

to:

$(VCInstallDir)libamd64;$(VCInstallDir)atlmfclibamd64;$(WindowsSdkDir)libx64;

In fact check all entries under Project Property Pages > Configuration Properties > VC++ Directories to see they’re 64 bit paths. If the above doesn’t work try creating a new project in VS, converting it from 32bit to 64bit and if it builds use it as a reference and compare with the real project you’re trying to convert to 64bit. Sync the differences and see where you get to. That’s what I did.

This has wasted so much of my time and has been so difficult to track down that I just had to blog it in case it saved other hours, days or even weeks. There are many variants of the above error each of which could be due to different causes. This post only relates to the error prefixed with msvcprtd.lib(MSVCP100D.dll).

Did this help you out? Let me know in the comments.

OpenCL Cookbook: How to leverage multiple devices in OpenCL

So far, in the OpenCL Cookbook series, we’ve only looked at utilising a single device for computation. But what happens when you install more than one card in your host machine? How do you scale your computation across multiple GPUs? Will your code automatically scale to multiple devices or does it require you to consciously think about how to distribute the load of the computation across all available devices and change your code to apply that strategy? Here I look at answers to these questions.

Decide on how you want to use the host binding to support multiple devices

There are two ways in which a given host binding can support multiple devices.

  • A single context across all device and one command queue per device.
  • One context and command queue per device

Let’s look at these in more detail with skeletal implementations in C.

Creating a single context across all devices and one command queue per device

For this particular way of the binding supporting multiple devices we create only one context and share it across one command queue per device. So if we have say two devices we’ll have one context and two command queues each of which share that one context.

#include <iostream>
#include <CL/cl.hpp>
#include <CL/opencl.h>

int main () {

    cl_int err;
    
    // get first platform
    cl_platform_id platform;
    err = clGetPlatformIDs(1, &platform, NULL);
    
    // get device count
    cl_uint deviceCount;
    err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 0, NULL, &deviceCount);
    
    // get all devices
    cl_device_id* devices;
    devices = new cl_device_id[deviceCount];
    err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, deviceCount, devices, NULL);
    
    // create a single context for all devices
    cl_context context = clCreateContext(NULL, deviceCount, devices, NULL, NULL, &err);
    
    // for each device create a separate queue
    cl_command_queue* queues = new cl_command_queue[deviceCount];
    for (int i = 0; i < deviceCount; i++) {
        queues[i] = clCreateCommandQueue(context, devices[i], 0, &err);
    }

    /*
     * Here you have one context across all devices and one command queue per device.
     * You can choose to send your tasks to any of these queues depending on which
     * device you want to execute the task on.
     */

    // cleanup
    for(int i = 0; i < deviceCount; i++) {
        clReleaseDevice(devices[i]);
        clReleaseCommandQueue(queues[i]);
    }
    
    clReleaseContext(context);

    delete[] devices;
    delete[] queues;
    
    return 0;
    
}

Creating one context and one command queue per device

Here I create one context and one command queue per device each of which have their own context rather than sharing one.

#include <iostream>
#include <CL/cl.hpp>
#include <CL/opencl.h>

int main () {

    cl_int err;
    
    // get first platform
    cl_platform_id platform;
    err = clGetPlatformIDs(1, &platform, NULL);
    
    // get device count
    cl_uint deviceCount;
    err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 0, NULL, &deviceCount);
    
    // get all devices
    cl_device_id* devices;
    devices = new cl_device_id[deviceCount];
    err = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, deviceCount, devices, NULL);
    
    // for each device create a separate context AND queue
    cl_context* contexts = new cl_context[deviceCount];
    cl_command_queue* queues = new cl_command_queue[deviceCount];
    for (int i = 0; i < deviceCount; i++) {
        contexts[i] = clCreateContext(NULL, deviceCount, devices, NULL, NULL, &err);
        queues[i] = clCreateCommandQueue(contexts[i], devices[i], 0, &err);
    }

    /*
     * Here you have one context and one command queue per device.
     * You can choose to send your tasks to any of these queues.
     */

    // cleanup
    for(int i = 0; i < deviceCount; i++) {
        clReleaseDevice(devices[i]);
        clReleaseContext(contexts[i]);
        clReleaseCommandQueue(queues[i]);
    }
    
    delete[] devices;
    delete[] contexts;
    delete[] queues;
    
    return 0;

}

How do you scale your computation across multiple devices?

The process of utilising multiple devices for your computation is not done automatically by the binding when new devices are detected sadly. Nor is it possible for it do so. Doing this requires active thought from the host programmer. When using a single device you send all your kernel invocations to the command queue associated with that device. In order to use multiple devices you must have one command queue per device either sharing a context or each queue having its own context. Then you must decide how to distribute your kernel calls across all available queues. It may be as simple as a round robin strategy across all queues for all your computations or it may be more complex.

Bear in mind that if your computation entails reading back a result synchronously then a round robin strategy across queues won’t work. This is because each current call will block and complete prior to you sending to the next queue which will essentially make the process of distributing across queues serial. Obviously this defeats the whole purpose of having multiple devices operating in parallel. What you really need is one host thread per device each sending computations to its own command queue. That way each queue is receiving and processing computations in parallel with other queues. Then you effectively achieve true hardware parallelism.

Which of the two ways should you use?

It depends. I would try the single context option first as it’s likely to use less memory and be faster. If you encounter instability or problems I would switch to the multiple context method. That’s the general rule. There is, however, another reason you may opt for a multiple context method. If you are using multiple threads which all require access to a context it is preferable for each thread to have its own context as the opencl host binding is not guaranteed to be thread safe. If you try to access a single context across multiple threads you may get serious system crashes and reboots so always have thread confined opencl structures.

Using a single context across multiple host threads

You may want to use one thread per device to send tasks to the command queue associated with each device. In this case you will have multiple host threads. But here have to be careful. In my experience it has not been safe to use a single context across multiple host threads. The last time I tried this was in C# using the Cloo host binding. Using a single context across multiple host threads resulted in a Windows 7 blue screen, Windows dumping memory to a file and then rebooting after which Windows failed to come back up until physically rebooted once more from the machine. The solution is to use the multi context option outlined above. Have thread confined separation for opencl resources and you’ll be fine.

OpenCL Cookbook: Hello World using C# Cloo host binding

So far I’ve used the C and C++ bindings in the OpenCL Cookbook series. This time I provide a quick and simple example of how to use Cloo – the C# OpenCL host binding. However, since Cloo, for whatever reason, didn’t work as expected with a char array I will use an integer array instead. In other words – instead of sending a “Hello World!” message to the kernel I will send five integers instead. My guess is that there is some sort of bug with Cloo and char arrays.

Device code using Cloo’s variant of the OpenCL language

kernel void helloWorld(global read_only int* message, int messageSize) {
	for (int i = 0; i < messageSize; i++) {
		printf("%d", message[i]);
	}
}

The kernel above is merely illustrative in that it simply receives an integer array and its size and prints the array.

Note that the OpenCL syntax here is not the same as in C/C++. It has additional keywords to say whether the arguments are read only or write or read write and the kernel keyword is not prefixed with two underscores. The Cloo author must have decided that the original OpenCL syntax was for whatever reason unsuitable for adoption which IMO was a mistake. The OpenCL language syntax should be standard for portability, reusability and also so that there is only a single learning curve.

Host code using Cloo API

using System;
using System.Collections.Concurrent;
using System.Threading.Tasks;
using System.IO;
using Cloo;

namespace test
{
    class Program
    {
        static void Main(string[] args)
        {
            // pick first platform
            ComputePlatform platform = ComputePlatform.Platforms[0];

            // create context with all gpu devices
            ComputeContext context = new ComputeContext(ComputeDeviceTypes.Gpu,
                new ComputeContextPropertyList(platform), null, IntPtr.Zero);

            // create a command queue with first gpu found
            ComputeCommandQueue queue = new ComputeCommandQueue(context,
                context.Devices[0], ComputeCommandQueueFlags.None);

            // load opencl source
            StreamReader streamReader = new StreamReader("..\..\kernels.cl");
            string clSource = streamReader.ReadToEnd();
            streamReader.Close();

            // create program with opencl source
            ComputeProgram program = new ComputeProgram(context, clSource);

            // compile opencl source
            program.Build(null, null, null, IntPtr.Zero);

            // load chosen kernel from program
            ComputeKernel kernel = program.CreateKernel("helloWorld");

            // create a ten integer array and its length
            int[] message = new int[] { 1, 2, 3, 4, 5 };
            int messageSize = message.Length;

            // allocate a memory buffer with the message (the int array)
            ComputeBuffer<int> messageBuffer = new ComputeBuffer<int>(context,
                ComputeMemoryFlags.ReadOnly | ComputeMemoryFlags.UseHostPointer, message);

            kernel.SetMemoryArgument(0, messageBuffer); // set the integer array
            kernel.SetValueArgument(1, messageSize); // set the array size

            // execute kernel
            queue.ExecuteTask(kernel, null);

            // wait for completion
            queue.Finish();
        }
    }
}

The C# program above uses the Cloo object oriented api to interface with the underlying low level opencl implementation. It’s pretty self explanatory if you’ve been following the series so far. The output of the program is 12345.

How to use core affinity to pin a process to a core on Windows using C#

Previously I wrote about how to use core affinity to pin a process to a core on Windows using C/C++. This is just a quick note on how to do so in C#. It’s actually a one liner so much easier than C/C++ unsurprisingly.

System.Diagnostics.Process.GetCurrentProcess().ProcessorAffinity = (System.IntPtr)(1 << coreId);

Above the coreId is a zero indexed number of the core you’d like to pin to. If you pass a coreId that’s incorrect this line of code will fail as below so there’s no need for explicit error checks.

Unhandled Exception: System.ComponentModel.Win32Exception: The parameter is incorrect

In case you’re wondering why anyone would want to limit themselves to one core in a multicore world it’s useful for checking how well single threaded processes perform when running one process per core and how well they scale as the number of processes goes up.

For example you may have 16 cores and because your process is single threaded you may want to run 16 processes each pinned to its respective core. As you deploy more and more processes however you may experience a degradation in how long each takes to perform a set amount of work. This will usually be due to cache overflow and reaching memory bandwidth limitations.

How to use core affinity to pin a process to a core on Windows using C/C++

Here’s how you can use core affinity to pin a particular process to any given core in C/C++ on Windows. The program below works by receiving the core number to pin the process to as the first argument to the executable. So for the first core you’d pass 0, for core 16 you’d pass 15 and so on.

#include <windows.h>
#include <iostream>
#include <algorithm>

using namespace std;

int main (int argc, char **argv) {

    // pin process to a core requested by incoming argument
    BOOL result = SetProcessAffinityMask(GetCurrentProcess(), 1 << atoi(argv[1]));
    if (result == 0) { cout << "SetProcessAffinityMask failed" << endl; return -1; }
    
    // perform long running cpu computation
    for (int i = 0; i < 10; i++) {
    
        // create a large array
        int sampleSize = 100000000;
        int* randoms = new int[sampleSize];
        for (int i = 0; i < sampleSize; i++) {
            randoms[i] = rand();
        }
        
        // sort it to take up some cpu time
        sort(randoms, randoms + sizeof randoms / sizeof randoms[0]);
        
        // cleanup
        delete[] randoms;
        
    }
    
    return 0;
    
}

The line of interest that actually applies core affinity is below.

BOOL result = SetProcessAffinityMask(GetCurrentProcess(), 1 << atoi(argv[1]));

There’s one thing you should watch out for when using core affinity. If you run a long running computation like the one above and you’re checking which cores get spiked on task manager > performance like I was you may see that more than one core gets spiked and you may initially think that the core affinity is not working. However this is a red herring as I explain below.

You’ll notice that other cores will suffer small spikes but only one core (the one you are pinning to) will sustain load the entire duration of the computation. So what’s actually happening is that in the initial period when your program is being set up to run other cores are getting involved but for the execution of your program only the requested core is being used. So always make sure that the computation runs long enough for you to see which one runs under sustained load.

Concurrent producer consumer pattern using C# 4.0, BlockingCollection & Tasks

Here is a very simple illustrative and annotated producer consumer pattern example using C#, .NET 4.0, BlockingCollection and Tasks. The example sets up one producer thread on which it produces 100 integers and two consumer threads which each read that data off a common concurrent blocking queue. Although I use a BlockingCollection here it is backed internally by a concurrent queue.

using System;
using System.Collections.Concurrent;
using System.Threading.Tasks;

namespace test
{

    class Program
    {

        static void Main(string[] args)
        {

            Console.WriteLine();

            // declare blocking collection backed by concurrent queue
            BlockingCollection<int> b = new BlockingCollection<int>();

            // start consumer 1 which waits for data until available
            var consumer1 = Task.Factory.StartNew(() =>
            {
                foreach (int data in b.GetConsumingEnumerable())
                {
                    Console.Write("c1=" + data + ", ");
                }
            });

            // start consumer 2 which waits for data until available
            var consumer2 = Task.Factory.StartNew(() =>
            {
                foreach (int data in b.GetConsumingEnumerable())
                {
                    Console.Write("c2=" + data + ", ");
                }
            });

            // produce 100 integers on a producer thread
            var producer = Task.Factory.StartNew(() =>
            {
                for (int i = 0; i < 100; i++)
                {
                    b.Add(i);
                }
                b.CompleteAdding();
            });

            // wait for producer to finish producing
            producer.Wait();
            
            // wait for all consumers to finish consuming
            Task.WaitAll(consumer1, consumer2);

            Console.WriteLine();
            Console.WriteLine();

        }

    }

}

The output from running the program above is as follows. It’s basically consumer 1 and 2 interleaving as they read integers off the queue as they’re being produced.

c1=0, c1=2, c1=3, c1=4, c1=5, c1=6, c1=7, c1=8, c1=9, c1=10, c1=11, c1=12, c1=13, c1=14, c1=15, c1=16, c1=17, c1=18, c1=19, c2=1, c2=21, c2=22, c2=23, c2=24, c2
=25, c2=26, c2=27, c2=28, c2=29, c2=30, c2=31, c2=32, c2=33, c2=34, c2=35, c2=36, c2=37, c2=38, c2=39, c2=40, c2=41, c2=42, c2=43, c2=44, c2=45, c2=46, c2=47, c
2=48, c2=49, c2=50, c2=51, c2=52, c2=53, c2=54, c2=55, c2=56, c2=57, c2=58, c2=59, c2=60, c2=61, c2=62, c2=63, c2=64, c2=65, c2=66, c2=67, c2=68, c2=69, c2=70,
c2=71, c2=72, c2=73, c2=74, c2=75, c2=76, c2=77, c2=78, c2=79, c2=80, c2=81, c2=82, c2=83, c2=84, c2=85, c2=86, c2=87, c2=88, c2=89, c2=90, c2=91, c2=92, c2=93,
 c2=94, c2=95, c2=96, c2=97, c2=98, c2=99, c1=20,

It’s somewhat sad and disheartening to see just how much further ahead C# is of Java in terms of the expressiveness and richness of the language. I personally really enjoyed using the lambdas above and I hope that with Java 8 the SDK libraries will grow to develop a more fluid and expressive style.

Credit: [1, 2]

OpenCL Cookbook: Parallelise your host loops using OpenCL

Continuing on in our series – this time we look at possibly the most important topic of all in OpenCL. It is the reason why we use OpenCL and it is also the most compelling benefit that OpenCL offers. It is, of course, parallelism. But how do we exploit the vast amount of parallelism that GPUs offer? At the simplest level we can do so by exploiting latent areas of parallelism in our host code the simplest of which are loops. In other words – if we can port loops in our host code to the GPU they become parallel and get faster by a factor of the total number of iterations. I demonstrate using a small example.

Host loop

void cpu_3d_loop (int x, int y, int z) {

    for (int i = 0; i < x; i++) {
        for (int j = 0; j < y; j++) {
            for (int k = 0; k < z; k++) {
                printf("CPU %d,%d,%dn", i, j, k);
            }
        }
    }

}

Imagine the loop above in our C++ host code. This is not one loop but in fact three. In other words it has three dimensions. The total number of iterations in this combined loop is x*y*z. If x=4, y=3 and z=2 the total number of iterations would be 4x3x2=24. On the CPU these loops execute serially which is fine for a small number of iterations but for large numbers it becomes a fundamental bottleneck. If this set of loops was ported to the GPU each iteration would run in parallel and the total number of threads in use would be 24 for the previous example.

A small scale example may not seem impressive at first. You could argue that you could just as well run 24 threads on the CPU. But consider this: what happens when you have the above set of loops in your host code performing thousands or even millions of iterations? How are you going to achieve hardware parallelism in this case on the CPU? The answer is you can’t. GPUs each have hundreds of cores and offer a far greater degree of parallelism so loops with a large number of iterations becomes easy work for the GPU which can run thousands or even millions of threads effectively. Below I demonstrate how to port such a loop to OpenCL.

Host binding code

#define __NO_STD_VECTOR
#define __CL_ENABLE_EXCEPTIONS

#include <fstream>
#include <iostream>
#include <iterator>
#include <CL/cl.hpp>
#include <CL/opencl.h>

using namespace cl;

void cpu_3d_loop (int x, int y, int z) {

    for (int i = 0; i < x; i++) {
        for (int j = 0; j < y; j++) {
            for (int k = 0; k < z; k++) {
                printf("CPU %d,%d,%dn", i, j, k);
            }
        }
    }

}

int main () {

    // CPU 3d loop

    int x = 4;
    int y = 3;
    int z = 2;
    cpu_3d_loop(x, y, z);
    std::cout << std::endl;

    // GPU 3d loop

    vector<Platform> platforms;
    vector<Device> devices;
    vector<Kernel> kernels;
    
    try {
    
        // create platform, context and command queue
        Platform::get(&platforms);
        platforms[0].getDevices(CL_DEVICE_TYPE_GPU, &devices);
        Context context(devices);
        CommandQueue queue(context, devices[0]);

        // load opencl source
        std::ifstream cl_file("kernels.cl");
        std::string cl_string(std::istreambuf_iterator<char>(cl_file),
            (std::istreambuf_iterator<char>()));
        Program::Sources source(1, std::make_pair(cl_string.c_str(), 
            cl_string.length() + 1));

        // create program and kernel and set kernel arguments
        Program program(context, source);
        program.build(devices);
        Kernel kernel(program, "ndrange_parallelism");

        // execute kernel and wait for completion
        NDRange global_work_size(x, y, z);
        queue.enqueueNDRangeKernel(kernel, NullRange, global_work_size, NullRange);
        queue.finish();

    } catch (Error e) {
        std::cout << std::endl << e.what() << " : " << e.err() << std::endl;
    }

    return 0;
    
}

The above program runs the cpu loop and then runs the equivalent logic on the gpu. Both cpu and gpu runs produce output to show which iteration they are processing. The key lines of code that demonstrate how to port the loop are below.

NDRange global_work_size(x, y, z);
queue.enqueueNDRangeKernel(kernel, NullRange, global_work_size, NullRange);

Here we set three upper bounds – one for each loop – this is known as the global work size. The kernel can then retrieve values for the currently executing iteration within the kernel itself as shown below. It can then use these indices to do whatever work is inside the loop. In this case we just print the indices for illustration.

Kernel code

The kernel you see below is executed x*y*z times with different values for i, j and k. See? No loops! 🙂

__kernel void ndrange_parallelism () {

	int i = get_global_id(0);
	int j = get_global_id(1);
	int k = get_global_id(2);

	printf("GPU %d,%d,%dn", i, j, k);
	
}

The output of running the above host code is as follows.

CPU 0,0,0
CPU 0,0,1
CPU 0,1,0
CPU 0,1,1
CPU 0,2,0
CPU 0,2,1
CPU 1,0,0
CPU 1,0,1
CPU 1,1,0
CPU 1,1,1
CPU 1,2,0
CPU 1,2,1
CPU 2,0,0
CPU 2,0,1
CPU 2,1,0
CPU 2,1,1
CPU 2,2,0
CPU 2,2,1
CPU 3,0,0
CPU 3,0,1
CPU 3,1,0
CPU 3,1,1
CPU 3,2,0
CPU 3,2,1

GPU 0,0,0
GPU 1,0,0
GPU 2,0,0
GPU 3,0,0
GPU 0,1,0
GPU 1,1,0
GPU 2,1,0
GPU 3,1,0
GPU 0,2,0
GPU 1,2,0
GPU 2,2,0
GPU 3,2,0
GPU 0,0,1
GPU 1,0,1
GPU 2,0,1
GPU 3,0,1
GPU 0,1,1
GPU 1,1,1
GPU 2,1,1
GPU 3,1,1
GPU 0,2,1
GPU 1,2,1
GPU 2,2,1
GPU 3,2,1

NOTE: Although there may appear to be a sequence in the order in which the GPU processes the iterations this is only due to the use of printf(). In reality when not using printf() the order of iterations is completely arbitrary and random. Therefore one must not rely on the order of iterations when porting loops to the GPU. If you need loops to be in a certain order then you can either keep your loops on the host or port only those parts of the loop that do not need to be sequential.

Why use GPU computing?

Although this example is fairly simple it does illustrate the most important value add of GPU computing and OpenCL. Hardware parallelism is the essence of what GPU computing offers and it is the most compelling reason to use it. If you imagine a legacy codebase and all the latent areas of parallelism that are currently running sequentially you can imagine the vast untapped power of GPGPU. Later on in the series we will look at techniques to port existing host code to the GPU. That process can be very difficult but can provide dramatic gains in performance far beyond the limits of CPU computing. Till next time.