icl::utils::CLProgram Class Reference

Main class for OpenCL based accelleration. More...

#include <CLProgram.h>

Public Member Functions
	CLProgram ()
	Default constructor (creates dummy instance) More...
	CLProgram (const CLDeviceContext &device_context, const string &sourceCode)
	CLProgram (const CLDeviceContext &device_context, ifstream &fileStream)
	CLProgram (const string deviceType, const string &sourceCode)
	create CLProgram with given device type (either "gpu" or "cpu") and souce-code More...
	CLProgram (const string deviceType, ifstream &fileStream)
	create CLProgram with given device type (either "gpu" or "cpu") and souce-code file More...
	CLProgram (const string &sourceCode, const CLProgram &parent)
	create CLProgram with given cl-program parent and souce-code More...
	CLProgram (ifstream &fileStream, const CLProgram &parent)
	create CLProgram with given cl-program parent and souce-code file More...
	CLProgram (const CLProgram &other)
	copy constructor (creating shallow copy) More...
CLProgram const &	operator= (CLProgram const &other)
	assignment operator (perorming shallow copy) More...
	~CLProgram ()
	Destructor. More...
bool	isValid ()
	isValid checks whether this is not a dummy but a valid initialized cl-program More...
bool	isValid () const
	isValid checks whether this is not a dummy but a valid initialized cl-program More...
CLBuffer	createBuffer (const string &accessMode, size_t size, const void *src=0)
	creates a buffer object for memory exchange with graphics card memory More...
CLImage2D	createImage2D (const string &accessMode, const size_t width, const size_t height, int depth, const void *src=0)
	creates a image2D object for memory exchange with graphics card memory More...
CLImage2D	createImage2D (const string &accessMode, const size_t width, const size_t height, int depth, int num_channel, const void *src=0)
CLKernel	createKernel (const string &id)
	extract a kernel from the program More...
CLDeviceContext	getDeviceContext ()
void	listSelectedPlatform ()
	lists various properties of the selected platform More...
void	listSelectedDevice ()
	lists various properties of the selected device More...

Static Public Member Functions
static void	listAllPlatformsAndDevices ()
	lists various properties of all platforms and their devices More...
static void	listAllPlatforms ()
	lists various properties of all platforms More...

Private Attributes
Impl *	impl

Detailed Description

Main class for OpenCL based accelleration.

The CLProgram is the based class for ICL's OpenCL support framework. A Program instance can be used to create all other neccessary OpenCL support types. In contrast to OpenCL's C++ framework, ICL's framework is even settled on a higher level providing easier access to the relavant functionality.

Nomenclature

A CLProgram is – as presented above – the main class for OpenCL acellerated implementations. It allows for selected a particular device, i.e. "cpu" or "gpu" devices and it automatically compiles the given OpenCL source code (either passed as string or as an input stream)

A CLBuffer is a memory segment located in the graphics-card's memory. Buffers are associated with a certain CLProgram and they can only be created by that program. Buffers are either read-only, write-only or read-write. This access mode always refers to how the buffer can be accessed by the OpenCL source code, i.e. a "read-only" buffer becomes a "const" data pointer in the corresponding OpenCL kernel interface. And a "write only" buffer cannot be written but not read by the OpenCL source code (?).

TODO: what are the differences ? Why can i use "r" and "w" without producing error messages? (seems to be an OpenCL bug/feature)

A CLImage2D is similar to the CLBuffer but additionally offers interpolation functionality which can be accessed in the kernel by using the sampler_t type. Also a build-in 2D access to the image pixels is supported.

A CLKernel is a callable OpenCL function. CLKernel instances are also created by a program (see Example (Image Convolution)) and each kernel refers to a single function in the OpenCL source code that is declared as "__kernel".
Before a kernel is called, its arguments are given by using the overloaded CLKernel::setArgs method.

Example (Image Convolution)

#include <ICLUtils/CLProgram.h>
#include <ICLQt/Common.h>

struct Conv{
  CLProgram program;          // main class
  CLBuffer input,output,mask; // buffers for image input/output and the 3x3-convolution mask
  CLKernel kernel;            // the OpenCL function
  Size size;                  // ensure correct buffer sizes
  Conv(const float m[9], const Size &s):size(s){
    // source code
    static const char *k = ("__kernel void convolve(constant float *m,                                \n"
                            "                        const __global unsigned char *in,                \n"
                            "                        __global unsigned char *out,                     \n"
                            "                        __local float *localMem){     // unused          \n"
                            "    const int w = get_global_size(0);                                    \n"
                            "    const int h = get_global_size(1);                                    \n"
                            "    const int x = get_global_id(0);                                      \n"
                            "    const int y = get_global_id(1);                                      \n"
                            "    if(x && y && x<w-1 && y<h-1){                                        \n"
                            "      const int idx = x+w*y;                                             \n"
                            "      out[idx] =  m[0]*in[idx-1-w] + m[1]*in[idx-w] + m[2]*in[idx-w+1]   \n"
                            "                + m[3]*in[idx-1]   + m[4]*in[idx]   + m[5]*in[idx+1]     \n"
                            "                + m[6]*in[idx-1+w] + m[7]*in[idx+w] + m[8]*in[idx+w+1];  \n"
                            "    }                                                                    \n"
                            "}                                                                        \n");

    program = CLProgram("gpu",k);                          // create program running on CPU-device
    program.listSelectedDevice();                          // show device seledted
    const int dim = s.getDim();                            // get image dimension
    input = program.createBuffer("r",dim);                 // create input image buffer
    output = program.createBuffer("w",dim);                // create output image buffer
    mask = program.createBuffer("r",9*sizeof(float),m);    // create buffer for the 3x3 conv. mask
    kernel = program.createKernel("convolve");             // create the OpenCL kernel
  }

  void apply(const Img8u &src, Img8u &dst){
    ICLASSERT_THROW(src.getSize() == size && dst.getSize() == size, ICLException("wrong size"));
    input.write(src.begin(0),src.getDim());           // write input image to graphics memory
    CLKernel::LocalMemory lMem(9*sizeof(float));      // create local memory (unused example)
    kernel.setArgs(mask, input, output, lmem);        // set kernel arguments
    kernel.apply(src.getWidth(), src.getHeight(), 0); // apply the kernel (using WxH threads max.
                                                      // i.e. one per pixel)
    output.read(dst.begin(0),dst.getDim());           // read output buffer to destination image
  }
} *conv = 0;


GUI gui;
GenericGrabber grabber;

void init(){
  grabber.init(pa("-i"));

  gui << Image().handle("image") << Show();
  const ImgBase &image = *grabber.grab();
  const float mask[] = { 0.25, 0, -0.25,
                         0.50, 0, -0.50,
                         0.25, 0, -0.25 };
  grabber.useDesired(image.getSize());
  grabber.useDesired(depth8u);
  grabber.useDesired(formatGray);

  conv = new Conv(mask,image.getSize());
}

void run(){
  const Img8u &image = *grabber.grab()->as8u();
  static Img8u res(image.getParams());

  conv->apply(image,res);
  gui["image"] = res;
}

int main(int n, char **args){
  return ICLApp(n,args,"-input|-i(2)",init,run).exec();
}

The same example but this time CLImage2D is used for the memory access instead of CLBuffer.

#include <ICLUtils/CLProgram.h>
#include <ICLQt/Common.h>

struct Conv{
  CLProgram program;           // main class
  CLImage2D input,output,mask; // images for input/output and the 3x3-convolution mask
  CLKernel kernel;             // the OpenCL function
  Size size;                   // ensure correct buffer sizes
  Conv(const float m[9], const Size &s):size(s){
    // source code
    static const char *k = (
                            "__kernel void convolve(__read_only image2d_t m,                            \n"
                            "                        __read_only image2d_t  in,                         \n"
                            "                         __write_only image2d_t out){                      \n"
                            "    const int x = get_global_id(0);                                        \n"
                            "    const int y = get_global_id(1);                                        \n"
                            "    const int w = get_global_size(0);                                      \n"
                            "    const int h = get_global_size(1);                                      \n"
                            "    if(x && y && x<w-1 && y<h-1){                                          \n"
                            "      const sampler_t sampler= CLK_NORMALIZED_COORDS_FALSE |               \n"
                            "                               CLK_ADDRESS_CLAMP |                         \n"
                            "                               CLK_FILTER_NEAREST;                         \n"
                            "      uint4 outPixel = 0;                                                  \n"
                            "      for (int mx = 0; mx < 3; mx++) {                                     \n"
                            "        for (int my = 0; my < 3; my++) {                                   \n"
                            "          uint4 inPixel = read_imageui(in, sampler, (int2)(x-mx-1,y-my-1));\n"
                            "          float4 mValue = read_imagef(m, sampler, (int2)(mx,my));          \n"
                            "          outPixel.s0 += mValue.s0 * inPixel.s0;                           \n"
                            "        }                                                                  \n"
                            "      }                                                                    \n"
                            "      write_imageui(out, (int2)(x,y), outPixel);                           \n"
                            "  } \n"
                            "}                                                                        \n");

    program = CLProgram("gpu",k);                              // create program running on CPU-device
    program.listSelectedDevice();                              // show device seledted
    input = program.createImage2D("r", s.width, s.height, 0);  // create input image
    output = program.createImage2D("w", s.width, s.height, 0); // create output image
    mask = program.createImage2D("r",3, 3, 3, m);              // create image for the 3x3 conv. mask
    kernel = program.createKernel("convolve");                 // create the OpenCL kernel
  }

  void apply(const Img8u &src, Img8u &dst){
    ICLASSERT_THROW(src.getSize() == size && dst.getSize() == size, ICLException("wrong size"));
    input.write(src.begin(0));                                 // write input image to graphics memory
    kernel.setArgs(mask, input, output);                       // set kernel arguments
    kernel.apply(src.getWidth(), src.getHeight(), 0);          // apply the kernel (using WxH threads max.
                                                               // i.e. one per pixel)
    output.read(dst.begin(0));                                 // read output image to destination image
  }
} *conv = 0;


GUI gui;
GenericGrabber grabber;

void init(){
  grabber.init(pa("-i"));

  gui << Image().handle("image") << Show();
  const ImgBase &image = *grabber.grab();
  const float mask[]  = { 0.25, 0, -0.25,
                          0.5,  0, -0.5,
                          0.25, 0, -0.25};
  grabber.useDesired(image.getSize());
  grabber.useDesired(depth8u);
  grabber.useDesired(formatGray);

  conv = new Conv(mask,image.getSize());
}

void run(){
  const Img8u &image = *grabber.grab()->as8u();
  static Img8u res(image.getParams());

  conv->apply(image,res);
  gui["image"] = res;
}

int main(int n, char **args){
  return ICLApp(n,args,"-input|-i(2)",init,run).exec();
}

Constructor & Destructor Documentation

CLProgram() [1/8]

icl::utils::CLProgram::CLProgram

(

)

Default constructor (creates dummy instance)

CLProgram() [2/8]

icl::utils::CLProgram::CLProgram	(	const CLDeviceContext &	device_context,
		const string &	sourceCode
	)

CLProgram() [3/8]

icl::utils::CLProgram::CLProgram	(	const CLDeviceContext &	device_context,
		ifstream &	fileStream
	)

CLProgram() [4/8]

icl::utils::CLProgram::CLProgram	(	const string	deviceType,
		const string &	sourceCode
	)

create CLProgram with given device type (either "gpu" or "cpu") and souce-code

CLProgram() [5/8]

icl::utils::CLProgram::CLProgram	(	const string	deviceType,
		ifstream &	fileStream
	)

create CLProgram with given device type (either "gpu" or "cpu") and souce-code file

CLProgram() [6/8]

icl::utils::CLProgram::CLProgram	(	const string &	sourceCode,
		const CLProgram &	parent
	)

create CLProgram with given cl-program parent and souce-code

CLProgram() [7/8]

icl::utils::CLProgram::CLProgram	(	ifstream &	fileStream,
		const CLProgram &	parent
	)

create CLProgram with given cl-program parent and souce-code file

CLProgram() [8/8]

icl::utils::CLProgram::CLProgram

(

const CLProgram &

other

)

copy constructor (creating shallow copy)

~CLProgram()

icl::utils::CLProgram::~CLProgram

(

)

Destructor.

Member Function Documentation

createBuffer()

CLBuffer icl::utils::CLProgram::createBuffer	(	const string &	accessMode,
		size_t	size,
		const void *	src = 0
	)

creates a buffer object for memory exchange with graphics card memory

acessMode can either be "r", "w" or "rw", which refers to the readibility of the data by the OpenCL source code (actually this seems to be not relevant since all buffers can be read and written).
Each buffer has a fixed size (given in bytes). Optionally an initial source pointer can be passed that is then automatically uploaded to the buffer exisiting in the graphics memory.

createImage2D() [1/2]

CLImage2D icl::utils::CLProgram::createImage2D	(	const string &	accessMode,
		const size_t	width,
		const size_t	height,
		int	depth,
		const void *	src = 0
	)

creates a image2D object for memory exchange with graphics card memory

acessMode can either be "r", "w" or "rw", which refers to the readibility of the data by the OpenCL source code (actually this seems to be not relevant since all images can be read and written).
Optionally an initial source pointer can be passed that is then automatically uploaded to the image exisiting in the graphics memory.

various image depths can be used
depth8u = 0, < 8Bit unsigned integer values range {0,1,...255}
depth16s = 1, < 16Bit signed integer values
depth32s = 2, < 32Bit signed integer values
depth32f = 3, < 32Bit floating point values
depth64f = 4, < 64Bit floating point values

createImage2D() [2/2]

CLImage2D icl::utils::CLProgram::createImage2D	(	const string &	accessMode,
		const size_t	width,
		const size_t	height,
		int	depth,
		int	num_channel,
		const void *	src = 0
	)

createKernel()

CLKernel icl::utils::CLProgram::createKernel

(

const string &

id

)

extract a kernel from the program

Kernels in the CLProgram's source code have to be qualified with the __kernel qualifier. The kernel (aka function) name in the OpenCL source code is used as id.

getDeviceContext()

CLDeviceContext icl::utils::CLProgram::getDeviceContext

(

)

isValid() [1/2]

bool icl::utils::CLProgram::isValid

(

)

isValid checks whether this is not a dummy but a valid initialized cl-program

Returns

true if the instance is a initialized with valid cl-code and device

isValid() [2/2]

bool icl::utils::CLProgram::isValid

(

)

const

isValid checks whether this is not a dummy but a valid initialized cl-program

Returns

true if the instance is a initialized with valid cl-code and device

listAllPlatforms()

static void icl::utils::CLProgram::listAllPlatforms ( )

static

lists various properties of all platforms

listAllPlatformsAndDevices()

static void icl::utils::CLProgram::listAllPlatformsAndDevices ( )

static

lists various properties of all platforms and their devices

listSelectedDevice()

void icl::utils::CLProgram::listSelectedDevice

(

)

lists various properties of the selected device

listSelectedPlatform()

void icl::utils::CLProgram::listSelectedPlatform

(

)

lists various properties of the selected platform

operator=()

CLProgram const& icl::utils::CLProgram::operator=

(

CLProgram const &

other

)

assignment operator (perorming shallow copy)

Member Data Documentation

impl

Impl* icl::utils::CLProgram::impl

private

---

The documentation for this class was generated from the following file:

• /Users/alneuman/vm/icl/ICLUtils/src/ICLUtils/CLProgram.h