What is Kernel Mapper

Kernel Mapper

The Kernel Mapper generates a python source code which calls Open CL kernel code. To generate a python code, Users should create a kernel definition file, a JSON file. With the python code, users can just create the object and call the function it exposed which are the kernel functions. It's pretty simple and easy to use.

Kernel Definition File

The kernel definition file contains two parts: structure definitions and function definitions. Each of them are an array of structure definition type and function definition type. An example can be found at:

{
  "name": "name of kernel object",
  "types": [...], // array of structure definition type
  "functions": [...] // array of function definition type
}

structure definition type (optional)

As we known, it is possible to use struct in Open CL. This part defines the structs which are used at kernel functions, see next part. At generating phase, Kernel Mapper generates the struct header for your, like:

#ifndef __PRICE_DATA__
#define __PRICE_DATA__

typedef struct {
  long dataIndex;
  float open;
  float high;
  float low;
  float close;
} PriceData;

#endif

All structs definition will be at a single header file, kernel_structs.h. You can just include this header file to have all structs.

At python source code, you may see the code like this:

        self.PriceData = numpy.dtype([('dataIndex', numpy.int), \
                                      ('open', numpy.float32), \
                                      ('high', numpy.float32), \
                                      ('low', numpy.float32), \
                                      ('close', numpy.float32)])

All structs will be defined as member variables in a single class.

As we mentioned before, this part is optional. If you only use primitive type, like byte, short, int, float, long, double, you don't need to write this part.

definition structure

There are two fields in structure definition, name and fields:

1. name (String): the name of the structure. To be compatiable with Open CL and Python, we only accept text with the pattern, [a-zA-Z_][a-zA-Z_0-9]*.
2. fields (Array of TypeDefinition): the fields in this structure. We only accept an object of type definition.

The TypeDefinition contains two fields:

1. name (String): name field is a String and the same as name field of structure definition.
2. type (String): type field can be another structure or primitive types. The supported primitives are:

• byte (8bits)
• ubyte (8bits)
• short (16bits)
• ushort (16bits)
• int (32bits)
• uint (32bits)
• float (32bits)
• string
• long

Please note that sequence of structure definition affects the generated result at the first version. If you want to use another structure as type, please put the referred structure at top of it. Otherwise, we cannot generate a valid header file.

example

{
  "name": "PriceData",
  "fields": [
    { "name": "dataIndex", "type": "long" },
    { "name": "open", "type": "float" },
    { "name": "high", "type": "float" },
    { "name": "low", "type": "float" },
    { "name": "close", "type": "float" }
  ]
}

This JSON object defines the PriceData we just mentioned in previous section. It uses primitive types as its fields.

kernel function definition

This parts defines the functions we want to implement in Open CL. We can use the primitive types or struct types which are defined at previous part as the function arguments. After the generation, Kernel Mapper generates the function declaration for you, like:

#include "kernel_structs.h"
__kernel void test_donothing(__constant int range,
                             __constant int realSize,
                             __global PriceData* inBuffer,
                             __global PriceData* outBuffer) {

}

At python source code, you may see the code like:

def test_donothing(self, range, realSize, inBuffer, outBuffer):
  ...

def prepareOutBuffer(self, size):
  ...

The prepareOutBuffer function creates a numpy array with PriceData as its type and size as its length. Before calling test_donothing, user should call prepareOutBuffer to obtain a buffer for test_donothing.

definition structure

The kernel function definition type contains two fields:

1. name (String): the name of function. Again, we only accept text with the pattern, [a-zA-Z_][a-zA-Z_0-9]*.
2. arguments (array of ArgumentType): the list of arguments in this function.

The ArgumentType contains three fields:

1. name (String): the name of this argument. Again, we only accept text with the pattern, [a-zA-Z_][a-zA-Z_0-9]*.
2. type (String): the type of this argument. It supports primitive types and structure types.
3. memoryType (String): the memory type of this agument. According to Open CL, we may declare a memory object in global memory, local memory, or constant memory. This field should be one of the following value: "global", "local", "constant". If the value is not in these valuses, we will generate it as global memory.
4. argType (String): the argument is in/out format.

TODO: define memory access type (use host memory or copy to device)

example

{
  "name": "test_donothing",
  "arguments": [
    { "name": "range", "type": "int", "memoryType": "constant", "argType": "in" },
    { "name": "realSize", "type": "int", "memoryType": "constant", "argType": "in" },
    { "name": "inBuffer", "type": { "arrayType": "PriceData" }, "memoryType": "global", "argType": "in" },
    { "name": "outBuffer", "type": { "arrayType": "PriceData" }, "memoryType": "global", "argType": "out" }
  ]
}

This JSON object defines the test_donothing we just mentioned in previous section. It uses two constant variables and two PriceData arrays as its arguments.

Example of Kernel Definition File

{
  "name": "AwesomeJob",
  "types": [{
      "name": "PriceData",
      "fields": [
        { "name": "dataIndex", "type": "long" },
        { "name": "open", "type": "float" },
        { "name": "high", "type": "float" },
        { "name": "low", "type": "float" },
        { "name": "close", "type": "float" }
      ]
    }],
  "functions": [{
    "name": "test_donothing",
    "arguments": [
      { "name": "range", "type": "int", "memoryType": "constant", "argType": "in" },
      { "name": "realSize", "type": "int", "memoryType": "constant", "argType": "in" },
      { "name": "inBuffer", "type": { "arrayType": "PriceData" }, "memoryType": "global", "argType": "in" },
      { "name": "outBuffer", "type": { "arrayType": "PriceData" }, "memoryType": "global", "argType": "out" }
    ]
  }, {
    "name": "test_donothing2",
    "arguments": [
      { "name": "inBuffer", "type": { "arrayType": "PriceData" }, "memoryType": "global", "argType": "in" },
      { "name": "outBuffer", "type": { "arrayType": "PriceData" }, "memoryType": "global", "argType": "out" }
    ]
  }]

The generated files are: | -- awesome_job_structs.h -- awesome_job.c -- awesome_job.py

In the file awesome_job_structs.h, Kernel Mapper generates the content like:

#ifndef __PRICE_DATA_H__
#define __PRICE_DATA_H__

typedef struct {
  long dataIndex;
  float open;
  float high;
  float low;
  float close;
} PriceData;

#endif

In the file awesome_job.c, Kernel Mapper generates a skeleton of kernel functions for you and you should implement the kernel function inside of the function, like:

#include "awesome_job_structs.h"
__kernel void test_donothing(__constant int range,
                             __constant int realSize,
                             __global PriceData* inBuffer,
                             __global PriceData* outBuffer) {

  // Please implement your kernel code here.
}
__kernel void test_donothing2(__global PriceData* inBuffer,
                              __global PriceData* outBuffer) {
  // Please implement your kernel code here.
}

The awesome_job.py is the generated python source code. You don't need to write anything but use it. Just create object and call the function. The content looks like:

class AwesomeJob:
    ... #other members

    self.PriceData = numpy.dtype([('dataIndex', numpy.int), \
                                  ('open', numpy.float32), \
                                  ('high', numpy.float32), \
                                  ('low', numpy.float32), \
                                  ('close', numpy.float32)])
    def test_donothing(self, range, realSize, inBuffer, outBuffer, cb=None):
      ...

    def test_donothing2(self, inBuffer, outBuffer, cb=None):
      ...

    def createBuffer(self, type, size=0, data=None):
      ...

    ... #other functions