Scalar Data Types

cl_bool is also available but unlike the other cl_ types is not guaranteed to be the same size as the bool in kernels.

Vector Data Types

Built-in vector data types are supported by the oneAPI Construction Kit even if the underlying compute device does not support any or all of the vector data types. Vector widths defined by the standard are 2, 3, 4, 8, and 16.

OpenCL devices can report supporting "FULL_PROFILE" or "EMBEDDED_PROFILE" by passing CL_DEVICE_PROFILE as the param_name to clGetDeviceInfo. oneAPI Construction Kit detects which profile a device supports using the following table, as specified in [OpenCL 1.2][#opencl-1.2]. If any of these limits fall below the "FULL_PROFILE" value then the device will report support for "EMBEDDED_PROFILE".

As cl_device_id’s are created from a mux_device_t, which expose a different set of properties, an implementation of [Mux][../modules/mux/spec] can control which OpenCL profile is reported using the following property mappings.

clCreateProgramWithSource

cl_program clCreateProgramWithSource(cl_context context,
                                     cl_uint count,
                                     const char **strings,
                                     const size_t *lengths,
                                     cl_int *errcode_ret);

Creates an OpenCL program object from OpenCL C source, oneAPI Construction Kit does no compilation at this stage.

clCreateProgramWithBinary

cl_program clCreateProgramWithBinary(cl_context context,
                                     cl_uint num_devices,
                                     const cl_device_id *device_list,
                                     const size_t *lengths,
                                     const unsigned char **binaries,
                                     cl_int *binary_status,
                                     cl_int *errcode_ret);

Create an OpenCL program object from a binary. Because oneAPI Construction Kit supports the cl_khr_spir extension, this can also be a SPIR binary. If the binary is a SPIR binary, then the program must be compiled (with either clBuildProgram(), or clCompileProgram() and clLinkProgram()). If the binary is a pre-compiled binary, then it may be compiled again, but this will have no effect.

clCreateProgramWithBuiltInKernels

cl_program clCreateProgramWithBuiltInKernels(cl_context context,
                                             cl_uint num_devices,
                                             const cl_device_id *device_list,
                                             const char *kernel_names,
                                             cl_int *errcode_ret);

Create an OpenCL program object with built-in kernels. This is useful if you have a configurable, but not programmable, accelerator with predefined computations it is optimized for. oneAPI Construction Kit currently doesn’t support any built-in kernels but this will change in the future.

clCompileProgram

cl_int clCompileProgram(cl_program program,
                        cl_uint num_devices,
                        const cl_device_id *device_list,
                        const char *options,
                        cl_uint num_input_headers,
                        const cl_program *input_headers,
                        const char **header_include_names,
                        void (CL_CALLBACK *pfn_notify) (
                          cl_program program,
                          void *user_data
                        ),
                        void *user_data);

Compile the program object. If the program was created from source oneAPI Construction Kit runs the clang compiler frontend at this point.

clLinkProgram

cl_program clLinkProgram(cl_context context,
                         cl_uint num_devices,
                         const cl_device_id *device_list,
                         const char *options,
                         cl_uint num_input_programs,
                         const cl_program *input_programs,
                         void (CL_CALLBACK *pfn_notify) (
                           cl_program program,
                           void *user_data
                         ),
                         void *user_data,
                         cl_int *errcode_ret);

Link one of more program objects. In the spec wording it mentions that this may create an executable, however in oneAPI Construction Kit we delay the creation of an executable until [clEnqueueNDRangeKernel][clEnqueueNDRange] for performance reasons. We do however link separate LLVM Modules together here.

clBuildProgram

cl_int clBuildProgram(cl_program program,
                      cl_uint num_devices,
                      const cl_device_id *device_list,
                      const char *options,
                      void (CL_CALLBACK *pfn_notify) (
                        cl_program program,
                        void *user_data
                      ),
                      void *user_data);

Build, or compile and link, the program object. This is equivalent to a call to clCompileProgram followed by clLinkProgram. The options string parameter can be used to pass compiler flags, including some only available as Codeplay vendor extensions.

clGetProgramInfo

cl_int clGetProgramInfo(cl_program program,
                        cl_program_info param_name,
                        size_t param_value_size,
                        void *param_value,
                        size_t *param_value_size_ret);

Query the program object for information. One use of this is to retrieve compiled binaries which can we used in future to pass to clCreateProgramFromBinary to save compilation time. In this case oneAPI Construction Kit will return LLVM IR for the binary format.

clGetProgramBuildInfo

cl_int clGetProgramBuildInfo(cl_program program,
                             cl_device_id device,
                             cl_program_build_info param_name,
                             size_t param_value_size,
                             void *param_value,
                             size_t *param_value_size_ret);

Query the program for the latest build information. If compilation failed this function can be used to retrieve the error log, which oneAPI Construction Kit will output this as a LLVM diagnostic.

To best debug an OpenCL kernel the build flags -g, -S, and -cl-opt-disable should be set in the options to clBuildProgram(), where -g and -S are part of our cl_codeplay_extra_build_options platform extension.

• -g enables emission of debug info. Without this flag the debugger user will not be able to set a breakpoint on the kernel name.

• -S <path/to/source/file> points the generated debug info to a source file on disk so that the debugger can display kernel source code. If this file does not exist already the OCL runtime will create a file from the string passed to clCreateProgramWithSource(), if this API call occurred.

• -cl-opt-disable disables performance optimizations. Omitting this flag will result in an inability to reliably inspect frame variables and source step. Analogous to debugging C/C++ code that hasn’t been built with -O0.

Note: This section was created before the -S build option was added. Mapping the source in the debugger can be done as a more invasive fall-back but the suggested technique is to utilize -S.

It is problematic for a debugger to show source code for cl programs since they are created from an API call to clCreateProgramWithSource(). Where the source string passed in could be programmatically generated, not just read directly from a .cl file on disk.

As a placeholder the OCL runtime sets the source filename to kernel.opencl, and directory to wherever libOpenCl.so was dynamically loaded into. This path then gets propagated into the debug info and picked up by the debugger. So if a debugger user wants to set a breakpoint on a file & line, then the file will have to be kernel.opencl. This could be an issue if multiple kernels from different cl programs are being run.

More significant for the user experience is that the debugger can’t show the source code line for the current pc address of a stopped process. A workaround to this problem in the case where the .cl file does exist on disk is to rename the file to kernel.opencl. The directories may still not match however, so an example fix would be to copy the source .cl file to /tmp/kernel.opencl, then remap the directory in the debugger to point to /tmp. lldb uses a path substitution, while in gdb you can provide a catch all directory to search for files in. Note lldb doesn’t substitute environmental variables here so instead of using $HOME or $USER set the absolute path.

(lldb) settings set target.source-map /home/foo/Aorta/build /tmp
(gdb) directory /tmp

Example debugging session of running the convolution UnitCL test and debugging its kernel. Note UnitCL needs to be passed the -g, -S, and -cl-opt-disable options via the --unitcl_build_options command line argument.

We start by setting a function breakpoint on the kernel name with ‘b convolution’, which won’t be resolved immediately since libOpenCL.so hasn’t been dynamically loaded yet.

$ lldb -- build/bin/UnitCL \
          --gtest_filter=Execution.Task_07_04_Convolution \
          --unitcl_build_options="-g -cl-opt-disable -S OCL/test/UnitCL/kernels/task_07.04_convolution.cl"
(lldb) b convolution
Breakpoint 1: no locations (pending)
WARNING:  Unable to resolve breakpoint to any actual locations.
(lldb) r

lldb shows all the active threads which are currently stopped which makes the output cluttered, but it’s important to note that the thread IDs displayed have no relation to OpenCL work item IDs. To see all the threads in lldb run ‘thread list’, while to select an individual thread ‘thread select’ can be used. Depending on the OCL host implementation scheduling each process thread could execute several work items, and so may hit a kernel breakpoint several times.

Next in our debug session we step over the line int x = get_global_id(0); with the ‘thread step-over’ command, also aliased to ‘next’. Then by printing the frame variables we can see that the global id of the current work item is 130.

(lldb) thread select 7
(lldb) thread step-over
(lldb) frame var
(float *) src = 0x0000000000d28e80
(float *) dst = 0x00000000007fd800
(int) x = 130
(int) width = 0

We can also set a breakpoint on a file line, where the file must be set as task_07.04_convolution.cl. Additionally we can narrow this breakpoint scope down to a single work-item with the condition ‘x==4’, since x holds the global id. When the breakpoint is hit lldb can list multiple threads as stopped, and the debugger user may have to cycle through the threads to find the specific one where x==4.

(lldb) break set -f task_07.04_convolution.cl -l 20 -c "x == 4"
(lldb) break del 1
(lldb) continue
convolution(src=0x0000000000d28e80, dst=0x00000000007fd800)
   17  	            sum += weight * src[x + i];
   18  	        }
   19  	        sum /= totalWeight;
-> 20  	        dst[x] = sum;
   21  	    }
   22  	    else
   23  	    {

(lldb) print x
(int) $0 = 4
(lldb) print dst[x]
(float) $1 = 0
(lldb) next
(lldb) print dst[x]
(float) $2 = 95.9932632

Here we run through the same debugging scenario as for lldb, but using the equivalent gdb command syntax.

$ gdb --args build/bin/UnitCL \
             --gtest_filter=Execution.Task_07_04_Convolution \
             --unitcl_build_options="-g -cl-opt-disable -S OCL/test/UnitCL/kernels/task_07.04_convolution.cl"
(gdb) b convolution
(gdb) run

Unlike lldb, gdb only displays a single thread when the process is stopped. But you can see them all with command ‘info threads’, and select the individual thread with thread $tid.

(gdb) thread 6
[Switching to thread 6 (Thread 0x7fffeb8f2700 (LWP 4090))]
#0  convolution (src=0xd28e80, dst=0x7fd800) at task_07.04_convolution.cl:7
7	    int x = get_global_id(0);
(gdb) next
8	    int width = get_global_size(0);
(gdb) info locals
x = 130
width = 0
(gdb) info args
src = 0xd28e80
dst = 0x7fd800

We can set a conditional breakpoint on a file line only for work-item 4 using the below gdb command syntax.

(gdb) break task_07.04_convolution.cl:20 if x == 4
(gdb) delete 1
(gdb) continue
Continuing.
[Switching to Thread 0x7fffeb8f2700 (LWP 4106)]

Breakpoint 2, convolution (src=0xd28e80, dst=0x7fd800) at task_07.04_convolution.cl:20
20	        dst[x] = sum;
(gdb) info locals
sum = 0.833333313
totalWeight = 12
x = 4
width = 256
(gdb) print  dst[x]
$1 = 0
(gdb) next
7	    int x = get_global_id(0);
(gdb) print  dst[x]
$2 = 0.833333313

clGetExtensionFunctionAddress

void* clGetExtensionFunctionAddress(const char *func_name);

Query the platform for address of extension function, deprecated in OpenCL 1.2. Replaced by clGetExtensionFunctionAddressForPlatform.

clEnqueueWaitForEvents

cl_int clEnqueueWaitForEvents(cl_command_queue queue,
                              cl_uint num_events,
                              const cl_event *event_list);

Enqueue an event wait list, deprecated in OpenCL 1.2. Replaced by clWaitForEvents.

clEnqueueBarrier

cl_int clEnqueueBarrier(cl_command_queue queue);

Enqueue a barrier on the command queue, deprecated in OpenCL 1.2. Replaced by clEnqueueBarrierWithWaitList

clEnqueueMarker

cl_int clEnqueueMarker(cl_command_queue queue,
                       cl_event *event);

Enqueue a marker on the command queue, deprecated in OpenCL 1.2. Replaced by clEnqueueMarkerWithWaitList

clUnloadCompiler

cl_int clUnloadCompiler();

Unload the compiler, deprecated in OpenCL 1.2. This is a hint from the application and does not guarantee that the compiler will actually be unloaded by the implementation. As this is just a hint oneAPI Construction Kit ignores it.

Replaced by clUnloadPlatformCompiler.

clCreateImage2D

cl_mem clCreateImage2D(cl_context context,
                       cl_mem_flags flags,
                       const cl_image_format *image_format,
                       size_t image_width,
                       size_t image_height,
                       size_t image_row_pitch,
                       void *host_ptr,
                       cl_int *errcode_ret);

Create a 2D image memory object, deprecated in OpenCL 1.2. Replaced by clCreateImage.

clCreateImage3D

cl_mem clCreateImage3D(cl_context context,
                       cl_mem_flags flags,
                       const cl_image_format *image_format,
                       size_t image_width,
                       size_t image_height,
                       size_t image_depth,
                       size_t image_row_pitch,
                       size_t image_slice_pitch,
                       void *host_ptr,
                       cl_int *errcode_ret);

Create a 3D image memory object, deprecated in OpenCL 1.2. Replaced by clCreateImage.