hyPACK-2013 GPGPU OpenCL Prog. using AMD-APP

The OpenCL C language code which runs on an OpenCL device is called a program. A program is a collection of functions called kernels, where kernels are units of execution that can be scheduled to run on a device. OpenCL programs are compiled at runtime through a series of API calls and runtime compilation gives the system an opportunity to optimize for a specific device. OpenCL software links only to a common runtime layer (called the ICD), all platform-specific SDK activity is performed at runtime through a dynamaic library interface. The process of creating a kernel is as follows.

• The OpenCL C source code is stored in a character string.

• The source code is turned into a program object, cl_program, by calling cl_CreateProgramWithSource().

• The Program object is then compiled, for one or more OpenCL devices, with clBuildProgram(), If there are compiler errors, they will reported here.

The precise binary representation used is very vendor specific. OpenCL provides a function to return information about the program objects, clGetProgramInfo(). . OpenCL also provides cl_CreateProgramWithBinary(), which takes a list of binaries that matches the device list.

First Step : Obtain a cl_kernel() Object :

To execute kernel on a device, cl_kernel() object can be used which can be extracted from the cl_program(). This can be achieved by exporting from the appropriate object. The name of the kernel that the program exports are used to request it from the compiled program object. The name of the kernelis passed to cl_CreateKernel(), along with program object, and the kernel object will be returned if the program object was valid and the particular is found.

Second Step : Dispatching Kernel through an enqueue function :
The process of setting kernel arguments and information on arguments to be transferred to the device is required. Executing a kernel requires dispatching it though an enqueue function. The fact that the kernel arguments are persistent, we must specify each kernel argument individually using the function clSetKernelArg(). This function takes a kernel object, an index specifying the argument number, the size of the argument, and a pointer to the arguement. When a kernel is executed, this information is used to transfer arguments to the device.

After any required memory objects are transferred to the device, and the kernel arguments are set, the kernel is ready to be executed. Requesting that a device begin executing a kernel is done with a call to clEnqueueNDRangeKernel()

  cl_int
  clEnqueueNDRangeKernel(
                        cl_command_queue command_queue,
                        cl_kernel kernel,
                        cl_unit work_dim,
                        const size_t *global_work_offset,
                        const size_t *global_work_size,
                        const size_t *local_work_offset,
                        cl_unit num_events_in_wait_list,
                        const cl_event *event_wait_list,
                        cl_event *event)

The clEnqueueNDRangeKernel() call is asyhronous and it will return immediately after the corresponding events will have completed before the kernel begins execution.

The OpenCL Platform Layer APIs

• The OpenCL platform layer which implements platform specific features that allow applications to query OpenCL device configuration information and to create OpenCL contexts using one or more devices.

• Querying Platform Info
  
  
The function



cl_int     clGetPlatformInfo
(
cl_platform_info param_name,
size_t param_value_size,
void *param_value,
size_t *param_value_size_ret
)


gets specific information about the OpenCL platform.


• Querying Devices
  
  
The list of devices available can be obtained using he following function



cl_int     clGetDeviceIDs
(
cl_device_type device_type,
cl_uint num_entries,
cl_device_id *devices,
cl_uint *num_devices
)


gets specific information about the list of devices available




• Contexts
  
  The function
  
  

  cl_context     clCreateContext
  (
  cl_context_properties properties,
  cl_uint num_devices,
  const_cl_device_id * devices,
  void (* pfn_notify) , (const char * errinfo,
       const void *prviate_info,
       size_t ch, void *user_data),
  void *user_data ,
  cl_int *errorcode_ret
  )
  

  creates an OpenCL context. An OpenCL context is created with one or more devices. contexts are used by the OpenCL runtime for managing objects such as command-queues, memory, program and kernel objects and for executing kernels on one or more devices specified in the context.
  

• Contexts
  
  The function
  
  

  cl_context     clCreateContextFromtype
  (
  const cl_context_properties ^*properties,
  cl_device_type device_type
  void (^*pfn_notify), (const char ^*errinfo, const void ^*prviate_info,
          size_t cb, void ^*user_data ),
  void ^*user_data ,
  cl_int ^*errorcode_ret
  )
  

  creates an OpenCL context from a device type that identifies the specific device(s) to use. properties specifies a list of context property names and their corresponding values. Each property name is immediately followed by the corresponding desired value.
  

The OpenCL Runtime APIs

• The OpenCL provides the API that calls the manager OpenCL objects such as command-queues, memory objects, program objects, kernel objects for __kernel functions in a program and calls that allow you to enqueue commands to a command-queue such as executing a kernel, reading, writing a memory object.

• Command Queues
  
  The function
  
  

  cl_command_queue     clCreateCommandQueue
  (
  cl_context context,
  cl_device_id devices,
  cl_command_queue_properties properties,
  cl_int *errorcode_ret
  )
  

  creates a command-queue on a specific device.
  


• Memory Objects
  
  

  Memory objects are categorized into two types : buffer and image objects. A buffer object stores a one-dimensional collection of elements whereas an image object is used to store a two- or three- dimensional texture, frame-buffer or image.
  

  A buffer Object is created using the following-command function
  
  

  cl_mem     clCreateBuffer
  (
  cl_context context,
  cl_mem_flags flags,
  size_t size,
  void *host_ptr
  cl_int * errorcode_ret
  )
  

  context is a valid OpenCL context used to create the buffer object.
  


• Reading, Writing and Copying Buffer Objects
  
  

  The following functions enqueue commands to read from a buffer object to host memory or write to a buffer object from the host Memory.

  cl_int     clEnqueueReadBuffer
  (
  cl_command_queue command_queue,
  cl_mem buffer,
  cl_bool blocking_read,
  size_t offset,
  size_t ch,
  void *ptr
  cl_uint mem_events_in_wait_list,
  const cl_event *event_wait_list,
  cl_event *event
  )
  

  
  
  
  

  cl_int     clEnqueueWriteBuffer
  (
  cl_command_queue command_queue,
  cl_mem buffer,
  cl_bool blocking_write,
  size_t offset,
  size_t ch,
  const void *ptr,
  cl_uint mem_events_in_wait_list,
  cl_event *event_wait_list,
  const cl_event *event_wait_list,
  cl_event *event
  )
  

  
  
  
  

  cl_int     clEnqueueCopyBuffer
  (
  cl_command_queue command_queue,
  cl_mem buffer,
  cl_mem dst_buffer,
  size_t src_offset,
  size_t dst_offset,
  size_t ch,
  cl_uint mem_events_in_wait_list,
  cl_event *event_wait_list,
  const cl_event *event_wait_list,
  cl_event *event
  )
  

• Retaining and Releasing Memory Objects
  
  

  The function

  cl_int     clRetainMemObject
  (
  cl_mem memobj
  )
  

  increments the memobj reference count. clRetainMemObject returns CL_SUCCESS if the function is executed successfully. It returns CL_INVALID_MEM_OBJECT if memobj is not a valid memory object.

• Creating Image Objects
  
  
  

  An image (1D or 2D) object is created using the following function

  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 *ptr,
  cl_int * errcode_ret,
  )
  

  
  
  

  An 3D image object is created using the following function

  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_row_pitch,
  size_t image_slice_pitch,
  void *ptr,
  cl_int * errcode_ret,
  )
  

• Querying List of Supported Image Formats
  
  
  

  The function

  cl_int     clGetSupportedImageFormats
  (
  cl_context context,
  cl_mem_flags flags,
  cl_mem_object_type image_type,
  cl_unit num_entries, num_entries,
  cl_image_format *image_formats,
  cl_image_format *image_format,
  cl_uint *num_image_formats
  )
  
  

  can be used to get the list of images format supported by an OpenCL implementation when the following information about an image memory object is specified.
  

  Context

  Image type - 2D or 3D image

  Image object allocation information

• Reading, Writing and Copying Image Objects
  
  
  

  The following functions enqueue commands to read from 2D or 3D image object to host memory or write to a 2D or 3D image object from host memory.

  cl_int     clEnqueueReadImage
  (
  cl_command_queue command_queue,
  cl_mem image,
  cl_bool blocking_read,
  const size_t t_origin[3],
  const size_t t_region[3] ,
  size_t row_pitch
  size_t slice_pitch,
  void *ptr,
  cl_unit num_events_in_wait_list,
  const cl event *event_wait_list,
  cl_event *event
  cl_event *event
  )
  
  

  
  
  

  cl_int     clEnqueueWriteImage
  (
  cl_command_queue command_queue,
  cl_mem image,
  cl_bool blocking_read,
  const size_t t_origin[3],
  const size_t t_region[3] ,
  size_t input_row_pitch
  size_t input_slice_pitch,
  const void *ptr,
  cl_unit num_events_in_wait_list,
  const cl event *event_wait_list,
  cl_event *event
  )
  
  

  

  The function

  cl_int     clEnqueueCopyImage
  (
  cl_command_queue command_queue,
  cl_mem src_image,
  cl_mem dst_image,
  const size_t src_origin[3],
  const size_t dst_region[3],
  const size_t region[3],
  cl_unit num_events_in_wait_list,
  const cl_event *event_wait_list,
  cl_event *event
  )

  enqueues a command to copy image obehcts. src_image and dst_image can be 2D or 3D image objects allowing us to perform the following actions.

• Copy a 2D image object to a 2D image object
• Copy a 2D image object to a 2D slice of a 3D image Object
• Copy a 2D slice of a 3D image object to a 2D image object
• Copy a 3D image object to a 3D image object.

• Copying between Image and Buffer Objects
  
  

  The function

  cl_int     clEnqueueCopyImageToBuffer
  (
  cl_command_queue command_queue,
  cl_mem src_image,
  cl_mem dst_image,
  const size_t src_origin[3],
  const size_t region[3],
  size_t dst_offset,
  cl_unit num_events_in_wait_list,
  const cl_event *event_wait_list,
  cl_event *event
  )
  
  

  enqueues a command to copy an image object to a buffer object.

  

  cl_int     clEnqueueCopyBufferToImage
  (
  cl_command_queue command_queue,
  cl_mem src_buffer,
  cl_mem dst_image,
  size_t src_offset,
  const size_t dst_origin[3],
  const size_t region[3],
  cl_unit num_events_in_wait_list,
  const cl_event *event_wait_list,
  cl_event *event
  )
  
  

  enqueues a command to copy a buffer object to an image object.

  

• Mapping and Unmapping Memory Objects
  
  
  

  The function

  void *     clEnqueueMapBuffer
  (
  cl_command_queue command_queue,
  cl_mem image,
  cl_bool blocking_read,
  cl_map_flags map_flags,
  size_t offset,
  size_t cb,
  cl_unit num_events_in_wait_list,
  const cl_event *event_wait_list,
  cl_event *event
  cl_int *errcode_ret
  )
  
  

  enqueues a command to map a region of the buffer object given by buffer inro the host address space and returns a pointer to this mapped region.

  
  

  The function

  void *     clEnqueueMapImage
  (
  cl_command_queue command_queue,
  cl_mem image,
  cl_bool blocking_read,
  cl_map_flags map_flags,
  const size_t origin[3],
  const size_t region[3],
  size_t *image_row_pitch,
  size_t *image_slice_pitch,
  cl_unit num_events_in_wait_list,
  const cl_event *event_wait_list,
  cl_event *event
  cl_int *errcode_ret
  )
  
  

  enqueues a command to map a region into the image object given by image into the host address space and returns a pointer to this mapped region.

  

• Memory Objects Queries
  
  
  

  To get information that is common to all memory objects (buffer and image objects), use the following function

  cl_int     clGetMemObjectInfo
  (
  cl_mem memobj,
  cl_mem_info param_name,
  size_t param_value_size,
  void *param_value,
  size_t *param_value_size_ret
  )
  
  

  memobj specifies the memory object being queried.

  

  cl_int     clGetgetImageInfo
  (
  cl_mem memobj,
  cl_image_info param_name,
  size_t param_value_size,
  void *param_value,
  size_t *param_value_size_ret
  )
  
  

  To get Information specific to an image object created with clCreateImage{2D|3D}.

  

• Memory Objects Queries
  
  
  

  A sampler object describes how to sample an image when the image is read in the kernel. The OpenCL functions are used to create sampler objects.

  cl_sampler     clCreateSampler
  (
  cl_context context,
  cl_bool normalized_coords,
  cl_addressing_mode addressing_mode,
  cl_filer_mode filer_mode,
  cl_int *errorcode_ret
  )
  
  

  The above function creates a sample object.

  

  cl_int     clGetSamplerInfo
  (
  cl_sampler sampler,
  cl_sampler_info param_name,
  size_t param_value_size,
  void *param_value,
  size_t *param_value_size_ret
  )
  
  

  To get Information specific to an image object created with clCreateImage{2D|3D}.

  

• Program Objects
  
  
  

  An OpenCL program consists of a set of kernels that are identified as function declared with the __kernel qualifier in the program source. OpenCL programs may also contain auxiliary functions and constant data that can be used by __kernel functions.

  Creating Program Objects
  
  

  The function

  cl_program     clCreateProgramWithSource
  (
  cl_contex tcontext
  const char **strings,
  const size_t *lengths,
  cl_int *errorcode_ret
  )
  
  

  The above function creates a program object for a context, and loads the source code specified by the text strings in the string array into the program object. The devices associated with the program object are the devices associated with the context

  

  cl_program     clCreateProgramWithBinary
  (
  cl_context context,
  cl_uint num_devices,
  const cl_device_id *device_list,
  const size_t *lengths,
  const void **binaries,
  const void *errorcode_ret
  cl_int *binary_status,
  cl_int *errcode_ret
  )
  
  

  The above function creates a program object for a context, and loads the binary bits specified by binary into the program object.

  
  Building Program Executables
  
  

  The function

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

  builds (compiles & links) a program executable from the program source or binary for all the devices or a specific device(s) in the OpenCL context associated with program.

  
  Building Options
  
  

  The build options are categorized as pre-processor options, options for math intrinsics, options that control optimization and miscellaneous options.

  

• Kernel Objects
  
  
  

  A kernel is a function declared in a program. A kernel is identified by the __kernel qualifier applied to any function in a program. A kernel object encapsulates the specific __kernel function declared in a program and the argument values to be used when executing this __kernel function.

  Creating Kernel Objects
  
  

  To create a kernel object, use the function

  cl_kernel     clCreateKernel
  (
  cl_program program
  const char **kernel_name,
  const size_t *lengths,
  cl_int *errorcode_ret
  )
  
  





To create kernel objects for all kernel functions in the program, use the function

cl_int     clCreateKernelsInProgram
(
cl_program program
cl_uint num_kernels,
cl_kernel *kernels,
cl_uint *num_kernels_ret
)



• Kernel Objects
  
  

  To create kernel objects for all kernel functions in the program, use the function

  Setting Kernel Arguments
  
  

  To execute a kernel, the kernel arguments must be set. The function

  cl_int     clSetKernelArg
  (
  cl_kernel kernel,
  cl_uint arg_index,
  size_t arg_size,
  const void *arg_values
  )
  
  

  
  

  is used to set the argument value for a specific argument of a kernel.

• Execution Kernel
  
  
  

  The function

  cl_int     clEnqueueNDRangeKernel
  (
  cl_command_queue command_queue ,
  cl_kernel kernel,
  cl_uint work_dim ,
  const size_t *global_work_offset,
  const size_t *global_work_size,
  const size_t *local_work_size,
  cl_uint num_events_in_wait_list,
  const cl_event *event_wait_list,
  cl_event *event
  )
  
  

  
  

  enqueues a command to execute a kernel on a device.

  
  
  
  
  

  The function

  cl_int     clEnqueueTask
  (
  cl_command_queue command_queue ,
  cl_kernel kernel,
  cl_uint work_dim,
  cl_uint num_events_in_wait_list,
  const cl_event *event_wait_list,
  cl_event *event
  )
  
  

  
  

  enqueues a command to execute a kernel on a device.

  
  

  The function

  cl_int     clEnqueueNativeKernel
  (
  cl_command_queue command_queue ,
  void (*user_func) (void(*)
  void (*args), size_t cb_args,
  cl_uint num_mem_objects,
  size_t arg_size,
  const cl_mem **mem_list
  const void ***args_mem_loc
  cl_uint num_events_in_wait_list ,
  const cl_event *event_wait_list,
  cl_event *event
  )
  

  
  

  enqueues a command to execute a native C/C++ function not compiled using the OpenCL Compiler.

Event objects can be used to refer to a kernel execution command ( clEnqueueNDRangeKernel, clEnqueueTask, clEnqueueNativeKernel)
or read, write, map and copy commands on memory objects (
( clEnqueue{Read|Write|Map|} {Buffer|Image}, clEnqueueCopy{Buffer|Image}, clEnqueueCopyBufferToImage, or ( clEnqueueCopyImageToBuffer).

An event object can be used to track the execution of a command.

The function

cl_int     clWaitForEvents
(
cl_uint num_events,
const cl_event *event_list,
)





waits on the host thread for commands identified by event objects in event_list to complete.

The function

cl_int     clGetEventsInfo
(
const cl_event event,
cl_command_queue command_queue ,
cl_event_info param_name,
size_t param_value_size,
void **param_value
size_t *param_value_size_ret,
)





returns information about the event object.

• Out-of-Order Execution of Kernels and Memory Object Commands
  
  
  

  The OpenCL functions that are submitted to a command-queue are queued in the order the calls are made but can be configured to execute in-order or out-of-order. The properties arguments in clCreateCommandQueue, can be used to specify the execution order.

• Profiling Operations on Memory Objects and Kernels
  
  

  Supports the profiling of OpenCL functions that are enqueued as commands to command-queue. The enqueued commands are identified by unique event objects. Event objects can be used to capture profiling information that measure execution times of a command.

• The function

  cl_int     clFlush
  (
  cl_command_queue command_queue,
  )
  

  
  

  issues all previously queued OpenCL commands in command_queue to all devices associated with command_queue

Given below, illustrate the basic programming steps required for a minimum amount of code. Many test programs might require similar steps and these steps do not include error checks.

• First querying the runtime to determine which platforms are present. There can be any number of different OpenCL implementations are present.

• Create a context (The OpenCL Context has associated with it a number of compute devices such as CPU or GPU devices)

  Within a context, OpenCL guarantees a relaxed consistency between these devices. This means that memory objects, such as buffers or images, are allocated per context, but changes made by one device are only guaranteed to be visible by another device at well-defined synchronization points.

• OpenCL provides events, with the ability to synchronization on a given event to enforce the correct order of execution,

• Many operations are performed with respect to a given context: there are many operations specific to a device. For example, program compilation and kernel execution are done on a peer-device basis.

• Performing work with a device, such as executing kernels or moving data to end from the device's local memory is done using a corresponding a command queue.

• A command queue is associated with a single device and a given context. : all work for a specific device is done through this interface. Note that while a single command queue can be associated with only a single device. For example, it is possible to have one command queue for executing kernels and a command kernel for managing data transfers between the host and the device.

Most OpenCL program follows the same pattern. Given a specific platform, select a device or devices to create a context, allocate memory, create device-specific command queues to create a context, allocate memory, create device-specific command queues, and perform data transfers & computations.

Generally, the platform is the gateway to accessing specific devices, given these devices and a corresponding context the application is independent of the platform. Given a context, the application can:

• Create a command queues

• Create programs to run on one or more associated devices

• Create kernels within those programs

• Allocate memory buffers or image, either on the host or on the device(s) (memory can be copied between the host and device)

• Write data to the device

• Submit the kernel (with appropriate arguments) to the command queue for execution.

• Read data back to the host form the device.

The relationship between context(s), device(s), buffer(s), program(s), kernel(s), and command queue(S) is best seen by looking at simple code.

• The OpenCL Framework allows applications to use host and one or more OpenCL devices as a single heterogeneous parallel computer system. The framework contains the components OpenCL Platform layer, OpenCL Runtime, and OpenCL Compiler.

The OpenCL Platform Layer

• The OpenCL platform layer which implements platform specific features that allow applications to query OpenCL device configuration information and to create OpenCL contexts using one or more devices.

Querying Platform Info

Querying Devices

Contexts

The OpenCL Runtime

Command Queues

Memory Objects


- Creating Buffer Objects,
- Reading, Writing and copying Buffer Objects
- Retaining and Releasing Memory Objects
- Creating Image Objects,
- Querying List of Supported Image formats,
- Reading, Writing and Copying Image objects
- Copying between Image and Buffer Objects
- Mapping and Unmapping Memory Objects
- Memory Object Queries

Sampler Objects


Program Objects


- Creating Program Objects
- Building Program Executables


Build Options


- Options (Preprocessor, Math Intrinsic, Optimization)
- Uploading the OpenCL compiler
- Program Object Queries


Kernel Objects


- Creating Kernel Objects
- Setting Kernel Arguments
- Kernel Object Queries


Executing Kernels


Event Objects


Profiling Operations on Memory Objects and Kernels


Flush and Finish

The OpenCL Compilation

• The OpenCL compiler provides binaries either CPU binaries or GPU binaries, depending on that target device. Refer OpenCL Compilation and Creating an OpenCL Progam Object

The OpenCL Compilation :

The Compiler tool-chain provides a common framework for both CPUs & GPUs, sharing the front-end and some high-level compiler transformations. The back-ends are optimized for the device type (CPU or GPU). Most of the application remains same, but OpenCL APIs are included at various parts of the code. The kernels are compiled by the OpenCL compiler to either CPU binaries or GPU binaries, depending on that target device.

• CPU Processing :For CPU processing, the OpenCL runtime uses the LLVM AS ( Low-level virtual Machine ) to generate x86 binaries. The OpenCL runtime automatically determines the number of processing elements or cores, present in the CPU and distributes the OpenCL kernel between them.

• GPU Processing : For GPU processing, the OpenCL runtime layer generates GPU specific AMD -ATI binaries with CAL or CUDA enabled NVIDIA architecture GPU binaries.

Compilation Program

An OpenCL application consists of a host program (C/C++) and an optional kernel program (.cl). To compile an OpenCL application, the host program must be compiled and this can be done using the off-the-shelf compiler such as g++ or MSV++. The application kernels are compiled into device-specific binaries using the OpenCL compiler. The compiler uses a standard C front-end as well as the LLVM framework, with extensions for OpenCL.

To compile OpenCL applications on Windows requires that Visual Studio 2008 Professional Edition and the Intel C compiler and all C++ files must be added with appropriate settings. To compile OpenCL applications on Linux requires that the gcc or the Intel C compiler is installed and all C++ files must be compiled with appropriate settings on 32-bit /64-bit systems.

The OpenCL Library and runtime environment depends upon the target GPU (i.e CUDA enabled NVIDIA or AMD ATI - Stream DSK).

Running Program

An OpenCL application is compiled on the target system, the runtime system assigns the work in the command queues to the underling devices. Commands are placed into the queue using the clEnqueue commands shown below. The commands can be broadly classified into three categories.

• kernel commands (for example, clEnqueueNDRangeKernel(), etc.).

• Memory commands (for example, clEnqueueNDReadBuffer(), etc.), and

• Memory commands (for example, clEnqueueWaitForEvents(),etc.

An OpenCL application can create multiple command queues and please refer OpenCL specification or OpenCL Programming Guide for the CUDA Architecture or AMD ATI Stream computing OpenCL Programming Guide.

To execute a OpenCL Program,change the directory to “/bin” directory & give the name of the executable at command prompt.

Read the instructions displayed on screen to execute the programs.

NOTE : For Compilation & execution of the OpenCL program : OpenCL Driver should be installed on the system properly. OpenCL Library file path should be included in environment variable “LD_LIBRARY_PATH”.

The Simple Vector Vector addition example demonstrates the simplest aproach for an OpenCL implementation in which approriate platform is chosen. The The main() function is reproduced along with source code kernel. The main() function of Vector Vector Addition example either implement or calls the functions that perform the following operations.

• Create an OpenCL context on the first available device
  

• Create a command \96queue on the first available device
  

• Load a kernel file ( vect-vect-addition.cl ) and build it into a program object
  

• Create a kernel object for the kernel function vect-vect-addition()
  

• Query the kernel for execution
  

• Read the results of the kernel back into the result buffer
  

 
 VectVectAddition.cl :
    __kernel void VectVectAddition_kernel( __global const float *a,
                                          __global const float *b,
                                          __global float *result)

 {
          int gid = get_global_id(0);
          result[gid] = a[gid] + b[gid];
 }

 //
 // Cleanup any created OpenCL resources
 //
  void Cleanup( cl_context context,
                 cl_command_queue commandQueue,
                             cl_program program, cl_kernel kernel,
                             cl_mem memObjects[3] )
  {

  for(int i = 0; i < ARRAY_SIZE; i++) {
        if (memObjects[i] != 0)
               clReleaseMemObject(memObjects[i]);
  }

        if (commandQueue != 0)
               clReleaseCommandQueue(commandQueue);

        if (kenrel != 0)
               clReleaseKernel(kernel);

        if (program != 0)
               clReleaseProgram(program);

        if (context != 0)
               clReleaseContext(context);

  cl_context context = 0;
  cl_command_queue commandQueue = 0;
  cl_program program = 0;
  cl_device_id device = 0;
  cl_kernel kernel = 0;
  cl_mem memObjects[3] = ( 0, 0, 0 );
  cl_int errNum;

  // Create an OpenCL context on first available platform
  context = CreateContext();
  if(context == NULL)
  {
        cerr << “Failed to create OpenCL context.” << endl;
        return 1;
  }

  // Create a command-queue on the first device available
  // on the created context
  commandQueue = CreateCommandQueue(context, sdevice);
  if(commandQueue == NULL)
  {
        Cleanup(context, commandQueue, program, kernel,
                    memObjects);
        return 1;
  }

  // Create OpenCL program from VectVectAddition.cl kernel source
  program =   CreateProgram(context, device, “VectVectAddition.cl”);
  if (program == NULL)
  {
        Cleanup(context, commandQueue, program, kernel,
                    memObjects);
        return 1;
  }

  // Create OpenCL kernel
  kernel = clCreateKernel(program, "VectVectAddition_kernel", NULL);
  if(kernel == NULL)
  {
        cerr << “Failed to create OpenCL context.” << endl;
        return 1;
  }

  // Create memory objects that will be used as arguments to
  // kernel. First create host memory arrays that will be
  // used to store the arguments to the kernel
  float result[ARRAY_SIZE];
  float a[ARRAY_SIZE];
  float b[ARRAY_SIZE];
  //Fill Arrays 'a' and 'b' on the host
  for(int i = 0; i < ARRAY_SIZE; i++) {
        a[i] = i;
        b[i] = i+2;
  }


  // Copy the arrays into memory objects that will be passed
  // to this kernel
  if (!CreateMemObjects(context, memObjects, a, b) )
  {
        Cleanup(context, commandQueue, program, kernel,
                    memObjects);
        return 1;
  }

  // Set the kernel arguments (result, a, b)
  errNum = clSetKernelArg(kernel, 0,
                                      sizeof ( cl_mem ), &memObjects[0]);
  errNum = clSetKernelArg(kernel, 1,
                                      sizeof ( cl_mem ), &memObjects[1]);
  errNum = clSetKernelArg(kernel, 2,
                                      sizeof ( cl_mem ), &memObjects[2]);
  if(errNum != CL_SUCCESS)
  {
        cerr << “Error setting kernel arguments.” << endl;
        Cleanup(context, commandQueue, program, kernel,
                    memObjects);
        return 1;
  }

  size_t globalWorkSize[1] = { ARRAY_SIZE };
  size_t localWorkSize[1] = { 1 };

  // Queue the kernel up for execution across the array
  errNum = clEnqueueNDRangeKernel(commandQueue, kernel, 1, NULL,
                                      globalWorkSize, localWorkSize,
                                      0, NULL, NULL);
  if(errNum != CL_SUCCESS)
  {
        cerr << “Error queuing kernel for execution,” << endl;
        Cleanup(context, commandQueue, program, kernel,
                    memObjects);
        return 1;
  }

  // Read the output buffer back to the Host
  errNum = clEnqueueReadBuffer(commandQueue, memObjects[2],
                                      CL_TRUE, 0, ARRAY_SIZE * sizeof(float)
                                      result, NULL, NULL);
  if(errNum != CL_SUCCESS)
  {
        cerr << “Error queuing kernel result buffer,” << endl;
        Cleanup(context, commandQueue, program, kernel,
                    memObjects);
        return 1;
  }

  // Output the result buffer
  for(int i = 0; i < ARRAY_SIZE; i++) {
        cerr << “Error queuing kernel result[i],” << endl;
  }

  cout << endl;
  cout << "Executed program successfully." << endl;
        Cleanup(context, commandQueue, program, kernel,
                    memObjects);
        return 0;
  }

• Objective

  Write OpenCL program to calculate multiplication of scalar value and vector
  

• Description

  We create a one-dimensional globalWorkSize array that is overlaid on vector. The input vector using Single Precision/Double Precision input data are generated on host-CPU and transfer the vector to device-GPU for scalar-vector multiplication. In global memory,a simple kernel based on the 1- dimension indexspace of work-groups is generated in which work-item is given a unique ID within its work-group. The final resultant vector is generated on device and transferred to Host. This code demonstrates the development of OpencL Kernel for simple computations.

Each work-item performs multiplication of element of a vector with scalar using work item ID.

The choice of work-items in the code are given as
size_t globalWorkSize[3] = [128, 1,1];
Using local memory and choice of multiple work-items in the same workgroup may increase the performance. It is possible to achieve this with an explicit assignment from a global memory pointer as a local memory pointer. The below table indicates the important steps to be performed for exeuciton on device.

Example : Assume that 1024 elements are taken in each vector. The size can be specified as a one- , two- , or three- dimensional vector. The host code to specify an ND Range for 1024 elements is follows :

size_t indexSpaceSize[3] = [1024, 1,1];

Most importantly, dividing the work-items of an NDRange into smaller, equal sized workgroups as shown in Figure.

• An Index space with N dimensions requires workgroups to be specified using N dimensions, thus, a threee -dimensional requires three-dimensional workgroups.

Example :
Perform Barrier Operations & synchronization
work-tems within a workgroup can peform barrier operations to synchronize and they have access to a shared memory address space Because workgroups sizes are fixed, this communication does not have have a need to scale and hence does not affect scalability of a large concurrent dispatch.

For example 1.3, i.e., Vector Vector Addition, the workgroup can be specified as

size_t workGroupSize[3] = [64, 1, 1];

If the total number of work-items per array is 1024, this results in creating 16 work-groups (1024 work-items / 64 per workgroups.

Most importantly, OpenCL requires that the index space sizes are evenely divisible by the work-group sizes in each dimension.

For hardware efficiency, the workgroup size is usually fixed to a favourable size, and we round up the index space size in each dimension to satisfy this divisibility requirement.

• In the kernel code, user can specify that exta work-items in each dimension simply return immediately without outputting any data.

• Many highly data paralle computations in which access of memory for arrays that peforms computation (example Vector-Vector Addition), the OpenCL allows the local workgroup size to be ignoed by the programmer and generated automatically by the implementation ; in this case; the dveloper will pass NULL instead.