Kernel Specification File

    PySchedCL leverages a specification file (JSON format) for individual OpenCL kernels for carrying out host side management for the computation. The specification file contains parameters required by the framework for dispatching a kernel to available OpenCL devices in the heterogeneous platform. These parameters typically comprise information regarding global work size for the problem, information regarding size and type of buffers and variable arguments for the kernel. A detailed description of each parameter is tabulated below.

Kernel Specification File Parameters:

Attribute Type Description
name String Name of OpenCL Kernel.
src String Path to the source of the OpenCL Kernel file.
workDimension Integer Number of dimensions of the computational kernel.
partition Integer An integer in range [0..10] which describes how to partition the kernel between CPU and GPU.
globalWorkSize String A python list expression denoting the Global Work Size Dimensions of the Kernel.
inputBuffers List A list of dictionaries containing information of read-only buffers. These will be the arguments with global modifier and whose values are not modified in the kernel.
ioBuffers List A list of dictionaries containing information of read-write buffers. These will be the arguments with global modifier and whose values are not modified in the kernel.
outputBuffers List A list of dictionaries containing information of write-only buffers. These will be the arguments with global modifier and whose values are not modified in the kernel.
varArguments List A list of dictionaries containing information regarding variable arguments passed to the kernel.
localArguments List A list of dictionaries containing information regarding local arguments passed to the kernel. These will be the arguments with local modifier.
id String This is the string which should be able to uniquely identify the kernel among the set of kernels. This is an optional field.

An Illustrative Example

Let us consider the standard vector addition kernel from the CUDA SDK below. 
 1
 2
 3
 4
 5
 6
 7
 9
 8
10
11
12
13
14
15
 // OpenCL Kernel Function for element by element vector addition
__kernel void VectorAdd(__global const float* a, __global const float* b, __global float* c, int iNumElements)
{
    // get index into global data array
    int iGID = get_global_id(0);

    // bound check (equivalent to the limit on a 'for' loop for standard/serial C code
    if (iGID >= iNumElements)
    {
        return;
    }

    // add the vector elements
    c[iGID] = a[iGID] + b[iGID];
}

As evident from the above code, the kernel requires two input buffers, one output buffer and one variable denoting the total size of the arrays to be processed. The kernel performs an element-wise addition of numbers in both the input buffers and stores the result in an output buffer. The corresponding specification file is given below. 

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
{

    "globalWorkSize": "[dataset]",
    "inputBuffers": [
        {
            "break": 1,
            "pos": 0,
            "size": "dataset",
            "type": "float"
        },
        {
            "break": 1,
            "pos": 1,
            "size": "dataset",
            "type": "float"
        }
    ],
    "name": "VectorAdd",
    "outputBuffers": [
        {
            "break": 1,
            "pos": 2,
            "size": "dataset",
            "type": "float"
        }
    ],
    "partition": 10,
    "src": "VectorAdd.cl",
    "varArguments": [
        {
            "pos": 3,
            "type": "int",
            "value": "partition_round(dataset, size_percent)"
        }
    ],
    "workDimension": 1
}

The primary reason behind introducing the JSON Specification file is to relieve the end user from designing large complex host side programs and focus more on the algorithm of the application at hand. The specification file shall provide a list of guidance parameters for the framework so that it can carry out the computation of the given kernel.

The guidance parameters can be classified into three primary categories.

  • Source Information:

    These parameters typically represent kernel source file information and associated problem size for the computation. This includes the name of the OpenCL kernel function (name), the name of the kernel source file (src) , the dimension of the kernel (workDimension) and the total number of workitems for each dimension (globalWorkSize)

  • Buffer Information:

    These parameters represent information for every input (inputBuffers) and output (outputBuffers) buffer for the program. It may be observed from the above specification file, each buffer is characterized by a tuple < break, pos, size, type>.
    The break flag indicates whether the buffer is partitionable or not. This is required by the framework to understand whether buffers should be broken while distributing the computation across multiple devices in the target platform.
    The value of pos denotes the positional argument for that buffer in the actual kernel function. This is needed while setting up the kernel arguments for the OpenCL runtime system by the framework.
    The size and type variables denote the size (number of elements) and type of the buffer respectively. Here the size of the buffer and the global work size for the kernel are the same and are represented by a symbolic expression called dataset which may be changed during runtime.

  • Argument Information:

    These parameters represent information for every variable argument (varArguments) for the OpenCL kernel. Each variable argument is characterized by a tuple < pos, type, value>.
    The value of pos like before denotes the positional argument of that variable in the actual kernel function.
    The varArguments field denotes the data type of the variable.
    The value field represents the corresponding value of the variable, which may be a constant or a symbolic expression. The variable kernel argument for the given function represents the total number of elements for each buffer and therefore should have been set to dataset.
    However, we emphasize that this is a partition sensitive parameter and is therefore represented by partition_round(dataset,size_percent). The corresponding function is a pyschedcl internal function used while partitioning applications. The user has the freedom to add custom functions of their own for the same.