Channels ▼
RSS

Tools

Programming with OpenCL 1.2


OpenCL 1.2 is the latest stable OpenCL release with available drivers for both CPU and GPU support. It is an interesting and useful enhancement to OpenCL 1.1 that adds several useful features — some of which are available extensions to earlier OpenCL versions. In this article, I cover the addition of the simplest debug tool ever, which makes it easier to debug OpenCL kernels, and the new "device fission," which enables you to partition a single compute device into many subdevices.

The Simplest Debug Tool: printf

OpenCL 1.2 includes a built-in printf function to the OpenCL C programming language. In previous OpenCL versions, only a few specific vendors provided an equivalent function through an extension, such as the cl_intel_printf extension found in the Intel drivers. The built-in printf function is very similar to the printf defined in the C99 standard, but there are differences that you should review by checking the documentation. For example, OpenCL C printf returns 0 for a successful execution and -1 otherwise, while C99 printf returns the number of printed characters for a successful execution. At the time of writing, it was possible to use printf with OpenCL 1.2 drivers for CPU targets and also for AMD GPUs.

The printf function buffers output until the kernel execution completes, then transfers the output back to the host. Thus, you have to be extremely careful to send text to the standard output when adding calls to printf in your kernels. As you might guess, the call to this function within a massively parallel execution has undesirable side effects, including a big impact on performance and memory usage. You should use printf only for specific debugging purposes on reduced data sets and then remove calls to this function when you want to execute the kernels on the entire data set.

The following lines show the code for a very simple OpenCL kernel that computes the product of a matrix and a vector:

__kernel void matrix_dot_vector(__global const float4 *matrix,
        __global const float4 *vector, __global float *result)
{
    int gid = get_global_id(0);
    result[gid] = dot(matrix[gid], vector[0]);
}

You can add the following line after the assignment to result[gid] to write data about the generated result index and value to the standard output:

printf("result[%d] %f \n", gid, result[i]);

This way, you can easily gather data from the kernel with a function that is familiar.

Working with Device Fission

One of the classic problems that arises when you target CPUs with OpenCL is that the execution of a kernel uses all the available cores and doesn't leave any core free to execute other processes. The cl_ext_device_fission extension provides an interface for subdividing an OpenCL device into multiple subdevices. OpenCL 1.2 has incorporated device fission in the specification, and all the functions that assigned work to devices now accept subdevices.

With the use of device fission, you can control which of the available compute units you want the OpenCL runtime to use in order to execute kernels. For example, if you have an eight-core CPU, you can use device fission to create a subdevice with six cores and leave two cores for the operating system tasks and other processes that require CPU usage. As a developer, device fission is extremely useful because you can continue working with your IDE and other development tools while executing OpenCL kernels that target the CPU.

You can also create sets composed of one or more subdevices with their own command queue. With this feature, you can use device fission to control the queues and dispatch work to each set. The feature is useful when you have algorithms that benefit from distributing work among many sets of subdevices. You have a lot of different possibilities for partitioning tasks and work with advanced task parallelism scenarios.

Obviously, the use of device fission requires that you have a good understanding of the details of the underlying hardware. When you create subdevices and dispatch work to them, you must consider many things that might affect performance (if you're writing OpenCL kernels, it means you want to achieve the best performance). If you don't take into account the shared resources for different subdevices, such as a shared cache memory or Non-Uniform Memory Access (NUMA) nodes, you will lose performance. Luckily, device fission in OpenCL 1.2 provides many predefined partitioning types and options that make it easy to specify the ways in which you want to split a device, and many of them allow you to consider the affinity of the compute units to share levels of cache hierarchy or a NUMA node. These predefined partitioning types and options allow you to make an efficient use of shared hardware resources when generating subdevices.

Device fission allows you to partition subdevices. Thus, once you create subdevices, they can be furthered partitioned by creating new subdevices. The relationships of the different subdevices form a tree in which the subdevices have a parent device or subdevice. It is possible to use different partition types and options each time you request OpenCL to split the device or subdevices. For example, you can partition the CPU devices by affinity, then partition one of those subdevices equally into eight subdevices. Obviously, the root device doesn't have a parent.

When you work with device fission for CPUs, you need to take into account that each compute unit is equivalent to a logical core or a hardware thread. So, when you work with Intel CPUs with Hyper-Threading technology enabled, two logical cores or hardware threads share one physical core.

The clGetDeviceInfo function has new associated device property IDs to retrieve information that allows you to plan the partitioning scheme to create subdevices. Before creating subdevices, you can retrieve the following information using the specified device property IDs:

  • The maximum number of subdevices that you can create for a device: CL_DEVICE_PARTITION_MAX_SUB_DEVICES. For example, if a CPU has eight logical cores, the value will be 8.
  • The partition types that the device supports: CL_DEVICE_PARTITION_PROPERTIES. I'll dive deep on partition types later.
  • The affinity domains for partitioning the device that the device supports: CL_DEVICE_PARTITION_AFFINITY_DOMAIN. When you specify a partitioning by affinity domain, you can use any of the affinity domains included in the returned list of supported values.

You can determine whether both the OpenCL implementation and the device support device fission by checking the maximum number of subdevices that you can create for a device (CL_DEVICE_PARTITION_MAX_SUB_DEVICES). Once you are sure that you can create subdevices, you can check the supported partition types and the affinity domains in cases where you want to work with partitioning by affinity domain. With all this information, you can write code that can take into account different hardware architectures and use different partitioning schemes based on the information gathered from the different clGetDeviceInfo calls.

After you make the call clGetDeviceIDs and the necessary calls to clGetDeviceInfo for the selected device, you can start creating the subdevices with calls to the new clCreateSubDevices function. Note that you must create the subdevices before you create the OpenCL context. OpenCL 1.2 devices and subdevices have retain (clRetainDevice) and release (clReleaseDevice) functions that allow you to increment and decrement the reference count as is done on other OpenCL objects.

The following lines show the C declaration of the clCreateSubDevices function:

cl_int clCreateSubDevices(cl_device_id in_device,
    const cl_device_partition_property *properties,
    cl_uint num_devices,
    cl_device_id out_devices,
    cl_uint *num_devices_ret);

The function requires the following arguments:

  • in_device: Indicates the ID of the device (cl_device_id) that you want to split into subdevices.
  • properties: Provides a property list that starts with the desired partition type and then provides additional values required by the selected partition type. The last value of the property list that indicates the end of the properties list must be CL_DEVICE_PARTITION_BY_COUNTS_LIST_END (0). As I explained before, in order to have a successful subdevices creation, you need to make sure that the partition type specified in this property list is supported by the device by making calls to clGetDeviceInfo.
  • num_devices: Specifies the size of the out_devices array.
  • out_devices: Provides a buffer for the generated subdevices with a number of elements specified by num_devices.
  • num_devices_ret: Returns the number of subdevices that the device may be partitioned into considering the partition type and the other values specified in the property list (properties).


Related Reading


More Insights






Currently we allow the following HTML tags in comments:

Single tags

These tags can be used alone and don't need an ending tag.

<br> Defines a single line break

<hr> Defines a horizontal line

Matching tags

These require an ending tag - e.g. <i>italic text</i>

<a> Defines an anchor

<b> Defines bold text

<big> Defines big text

<blockquote> Defines a long quotation

<caption> Defines a table caption

<cite> Defines a citation

<code> Defines computer code text

<em> Defines emphasized text

<fieldset> Defines a border around elements in a form

<h1> This is heading 1

<h2> This is heading 2

<h3> This is heading 3

<h4> This is heading 4

<h5> This is heading 5

<h6> This is heading 6

<i> Defines italic text

<p> Defines a paragraph

<pre> Defines preformatted text

<q> Defines a short quotation

<samp> Defines sample computer code text

<small> Defines small text

<span> Defines a section in a document

<s> Defines strikethrough text

<strike> Defines strikethrough text

<strong> Defines strong text

<sub> Defines subscripted text

<sup> Defines superscripted text

<u> Defines underlined text

Dr. Dobb's encourages readers to engage in spirited, healthy debate, including taking us to task. However, Dr. Dobb's moderates all comments posted to our site, and reserves the right to modify or remove any content that it determines to be derogatory, offensive, inflammatory, vulgar, irrelevant/off-topic, racist or obvious marketing or spam. Dr. Dobb's further reserves the right to disable the profile of any commenter participating in said activities.

 
Disqus Tips To upload an avatar photo, first complete your Disqus profile. | View the list of supported HTML tags you can use to style comments. | Please read our commenting policy.
 

Video