The GenericOCLVolumeFilter class is a facility class that allows easy construction of a volume filter from an existing file containing the OpenCL kernel code for the filter operation. More...

#include <genericoclvolumefilter.h>

Inheritance diagram for CTL::OCL::GenericOCLVolumeFilter:

[legend]

Collaboration diagram for CTL::OCL::GenericOCLVolumeFilter:

[legend]

Public Types
enum	{ Type = 2100 }

Public Types inherited from CTL::AbstractVolumeFilter
enum	{ Type = 2000 }

Public Types inherited from CTL::SerializationInterface
enum	{ Type = -1, UserType = 65536 }

Public Member Functions
int	type () const override

void	filter (VoxelVolume< float > &volume) override
	Filters the input volume by executing the OpenCL kernel of this instance. More...

	GenericOCLVolumeFilter (const std::string &clFileName, const std::vector< float > &arguments={})
	Creates a GenericOCLVolumeFilter instance that is capable of executing the OpenCL kernel (`filter`) of the file clFileName. More...

void	setAdditionalKernelArg (float argument)
	Convenience method for simpler setting of a single additional argument. More...

void	setAdditionalKernelArgs (const std::vector< float > &arguments)
	Sets the additional kernel arguments to arguments. More...

QVariant	parameter () const override
	Returns the parameters of this instance as QVariant. More...

void	setParameter (const QVariant &parameter) override
	Sets the parameters of this instance based on the passed QVariant parameter. More...

Public Member Functions inherited from CTL::AbstractVolumeFilter
void	fromVariant (const QVariant &variant) override

QVariant	toVariant () const override

Public Member Functions inherited from CTL::SerializationInterface
virtual	~SerializationInterface ()=default

Protected Member Functions
	GenericOCLVolumeFilter ()
	Creates a GenericOCLVolumeFilter without any OpenCL kernel associated to it. More...

	GenericOCLVolumeFilter (cl::Kernel *kernel, const std::vector< float > &arguments={})
	Creates a GenericOCLVolumeFilter that executes the OpenCL kernel kernel. More...

virtual cl::NDRange	globalWorksize (const VoxelVolume< float > &volume) const
	Returns the global worksize for the kernel call. More...

virtual cl::NDRange	localWorksize (const VoxelVolume< float > &volume) const
	Returns the local worksize for the kernel call. More...

Protected Member Functions inherited from CTL::AbstractVolumeFilter
	AbstractVolumeFilter (const AbstractVolumeFilter &)=default

	AbstractVolumeFilter (AbstractVolumeFilter &&)=default

AbstractVolumeFilter &	operator= (const AbstractVolumeFilter &)=default

AbstractVolumeFilter &	operator= (AbstractVolumeFilter &&)=default

Protected Member Functions inherited from CTL::SerializationInterface
	SerializationInterface ()=default

	SerializationInterface (const SerializationInterface &)=default

	SerializationInterface (SerializationInterface &&)=default

SerializationInterface &	operator= (const SerializationInterface &)=default

SerializationInterface &	operator= (SerializationInterface &&)=default

Protected Attributes
std::string	_clFileName

std::vector< float >	_additionalArgs

cl::Kernel *	_kernel

cl::CommandQueue	_queue

Static Private Member Functions
static cl::Kernel *	addKernelFromFile (const std::string &clFileName)
	Adds the OpenCL kernel defined in the file clFileName to the OpenCL environment and returns a pointer to the created cl::Kernel object. More...

static cl::CommandQueue	getCommandQueue ()
	Returns the command queue to be used by this instance. More...

Friends
template<class >
struct	SerializationHelper::RegisterWithSerializationHelper

Detailed Description

The GenericOCLVolumeFilter class is a facility class that allows easy construction of a volume filter from an existing file containing the OpenCL kernel code for the filter operation.

This class provides a means of simple construction of a volume filter that executes the filter method implemented in an OpenCL kernel, based on the file containing the corresponding kernel code. To construct an instance, pass the file name of the .cl file containing the OpenCL kernel code of the filter to this class' constructor. The clFileName (incl. path) is relative to the runtime cl_src folder of the CTL (see ClFileLoader::openCLSourceDir())., i.e. files placed directly in that folder can be addressed simply by their file name (see also 'Notes' below for more details). In addition, you might pass a vector containing an arbitrary number of float values. Each of these values will be passed to the kernel as an additional argument.

Notes:

The OpenCL kernel must have the following signature: kernel void filter( read_only image3d_t, global float*, uint, ...)
- The first argument contains the input volume.
- The second argument is a pointer to the buffer into which the output result (filtered volume) must be written. The buffer is a one-dimensional memory block representing the volume in row-major order. You might consider using the facility method provided by OpenCLFunctions::write_bufferf() for simple writes to the buffer; it will be added to the kernel automatically.
- The third argument holds the index of the slice that is to be processed by the kernel instance.
- The '...' stands for an arbitrary number of individual float parameters. Note that when using additional parameters in the kernel, the same amount of parameters needs to be passed to the ctor of this class (or set through setAdditionalKernelArgs()) for the filter to be functional.
The .cl file should be placed somewhere inside the cl_src folder of the CTL: [...]/ctl/modules/src/ocl/cl_src. It will then be copied to the target directory automatically when building the program and the GenericOCLProjectionFilter class is able to find it. Alternatively, the .cl file can also be copied to the runtime OpenCL source code folder (see ClFileLoader::openCLSourceDir()) manually.

Each kernel call must process an entire z-slice of the input volume. The degree of parallelization (i.e. global and local worksizes) are defined by globalWorksize() and localWorksize(). The global worksize defaults to the total number of voxels in x and y direction in the input volume (2D worksize), whereas the local worksize is a NullRange by default. Re-implement the corresponding methods in sub-classes to change this behavior. The kernel will automatically be called for each z-slice when filter() is executed.

Note that this class always uses the first OpenCL device found in OpenCLConfig::instance().devices(). Configure the device list accordingly before instantiating this class if you want to utilize a particular device.

Example:

We want to demonstrate the use of this class by the example of a kernel that simply multiplies the input volume by a constant factor that we define as one additional kernel parameter.

Our kernel implementation looks like this:

// simple_volume_filter.cl
// OpenCL kernel - with one additional argument
kernel void filter( read_only image3d_t inputVol,   /* the input volume */
                    global float* filteredVol,      /* the output buffer */
                    uint z,                         /* the index of the z-slice to be processed */
                    float param)                    /* our additional parameter (scaling factor) */
{
    // get IDs
    const int x = get_global_id(0);
    const int y = get_global_id(1);
    const int4 voxIdx = (int4)(x, y, z, 0);
    // read out the value of the voxel in 'inputVol'
    const float voxVal = read_imagef(inputVol, voxIdx).x; // can choose any component, since input is CL_INTENSITY
    // multiply the input voxel by 'param' and store result in 'filteredVol'
    // NOTE: we use the convenience method 'write_bufferf' here, see documentation of OpenCLFunctions::write_bufferf()
    write_bufferf(filteredVol, voxIdx, param * voxVal, inputVol); // passing 'inputVol' here provides the dimensions of the volume 
}

We first get the global IDs to identify which voxel we are supposed to process. Note that the z-coordinate is fixed within a kernel call and is passed to it as the third argument. We continue by reading out the value of the voxel in the input volume. This value will be multiplied by our scaling factor that gets passed to the kernel as its fourth argument. Finally, we utilize the write_bufferf method (see OpenCLFunctions::write_bufferf()) to easily write the result value to the output buffer.

In the .cpp code, using our kernel in a GenericOCLVolumeFilter looks something like this:

// create a simple cylinder volume filled with value 1.0
auto volume = VoxelVolume<float>::cylinderZ(10.0f, 10.0f, 1.0f, 1.0f);
qInfo() << volume.min() << volume.max();
// output: 0 1
// optionally in a try-catch block (see below)
// create the GenericOCLVolumeFilter with additional parameter 2.0
OCL::GenericOCLVolumeFilter filt("simple_volume_filter.cl", { 2.0f });
    // alternative(s):
    // filt.setAdditionalKernelArg(2.0f);
    // filt.setAdditionalKernelArgs({ 2.0f });
// we now filter our volume
filt.filter(volume);
qInfo() << volume.min() << volume.max();
// output: 0 2

Note that for this to work, you need to copy the example kernel file (can be found in ([...]/ctl/doc/examples/simple_volume_filter.cl) into the cl_src folder of the CTL: [...]/ctl/modules/src/ocl/cl_src. The .cl file will then be copied to the target directory automatically when building the program. Alternatively, you can manually copy the .cl file directly to the runtime OpenCL source code folder (which should be located in your build directory; see also ClFileLoader::openCLSourceDir()).

Tip: wrapping the construction and filtering in a try-catch block can help a lot in figuring out issues

try {
    OCL::GenericOCLVolumeFilter filt("simple_volume_filter.cl", { 2.0f });
    filt.filter(volume);
} catch (const std::exception& err) {
    qCritical() << err.what();
}

Constructor & Destructor Documentation

◆ GenericOCLVolumeFilter() [1/3]

CTL::OCL::GenericOCLVolumeFilter::GenericOCLVolumeFilter	(	const std::string &	clFileName,
		const std::vector< float > &	arguments = `{}`
	)

explicit

Creates a GenericOCLVolumeFilter instance that is capable of executing the OpenCL kernel (filter) of the file clFileName.

The clFileName (incl. path) is relative to the runtime cl_src folder of the CTL (see ClFileLoader::openCLSourceDir()). More details can be found under 'Notes' in the detailed class description). For further details on the requirements for the kernel, please also refer to the detailed class description. Note that, if the kernel specified in clFileName has additional arguments, the same number of arguments must be passed to arguments to make the filter work. Alternatively, additional arguments may be passed later on through setAdditionalKernelArgs().

◆ GenericOCLVolumeFilter() [2/3]

CTL::OCL::GenericOCLVolumeFilter::GenericOCLVolumeFilter ( )

protected

Creates a GenericOCLVolumeFilter without any OpenCL kernel associated to it.

For the purpose of deserialization only.

◆ GenericOCLVolumeFilter() [3/3]

CTL::OCL::GenericOCLVolumeFilter::GenericOCLVolumeFilter	(	cl::Kernel *	kernel,
		const std::vector< float > &	arguments = `{}`
	)

explicitprotected

Creates a GenericOCLVolumeFilter that executes the OpenCL kernel kernel.

For internal use only. Do not use this ctor directly—i.e. without an associated .cl file holding the kernel code—as this would prevent serializability of the GenericOCLVolumeFilter instance.

Member Function Documentation

◆ addKernelFromFile()

cl::Kernel * CTL::OCL::GenericOCLVolumeFilter::addKernelFromFile ( const std::string & clFileName )

staticprivate

Adds the OpenCL kernel defined in the file clFileName to the OpenCL environment and returns a pointer to the created cl::Kernel object.

A two stage approach is used to load the specified file clFileName. If the specified file exists under the specified name (incl. path), it will be loaded directly. Otherwise, an attempt to load the file from the runtime cl_src folder of the CTL (see ClFileLoader::openCLSourceDir()) will be made. See also ClFileLoader for more details.

◆ filter()

void CTL::OCL::GenericOCLVolumeFilter::filter ( VoxelVolume< float > & volume )

overridevirtual

Filters the input volume by executing the OpenCL kernel of this instance.

This method will handle all OpenCL host code required to run the kernel of this instance with the input data from volume and the additional kernel arguments specified in the constructor or from the latest call to setAdditionalKernelArgs().

Individual kernel calls are queued for each z-slice of the input volume. The global and local worksizes for the kernel execution are taken from globalWorksize() and localWorksize(), respectively. The full input volume is transfered at once to the OpenCL device and remains there for all calls of the kernel. It is represented as a cl::Image3D whose dimensions correspond to the dimensions of volume. This allows, for example, for data access outside the currently processed z-slice (e.g. to filter in that direction). Similarly, the output buffer has the full size and stays on the device for the entire execution. A single result readout is performed at the very end (i.e. after all z-slices have been processed) and the results are copied back to volume.

This implementation requires OpenCL image support, but no image write support.

Implements CTL::AbstractVolumeFilter.

◆ getCommandQueue()

cl::CommandQueue CTL::OCL::GenericOCLVolumeFilter::getCommandQueue ( )

staticprivate

Returns the command queue to be used by this instance.

Note that this will always use the first OpenCL device found in OpenCLConfig::instance().devices(). Configure the device list accordingly before instantiating this class if you want to utilize a particular device.

◆ globalWorksize()

cl::NDRange CTL::OCL::GenericOCLVolumeFilter::globalWorksize ( const VoxelVolume< float > & volume ) const

protectedvirtual

Returns the global worksize for the kernel call.

By default, this returns a two-dimensional cl::NDRange, where the first and second dimension are the number of voxels in x- and y-direction in volume, respectively.

Re-implement in sub-classes to change this behavior.

◆ localWorksize()

cl::NDRange CTL::OCL::GenericOCLVolumeFilter::localWorksize ( const VoxelVolume< float > & volume ) const

protectedvirtual

Returns the local worksize for the kernel call.

By default, this returns cl::NullRange. Re-implement in sub-classes to change this behavior.

◆ parameter()

QVariant CTL::OCL::GenericOCLVolumeFilter::parameter ( ) const

overridevirtual

Returns the parameters of this instance as QVariant.

This returns a QVariantMap with two key-value-pairs:

("Kernel file name", file name of the .cl file containing the kernel),
("Additional arguments", list of the used additional kernel parameters).

Reimplemented from CTL::AbstractVolumeFilter.

◆ setAdditionalKernelArg()

void CTL::OCL::GenericOCLVolumeFilter::setAdditionalKernelArg ( float argument )

Convenience method for simpler setting of a single additional argument.

This is a convenience method to be used only for kernels that have a single additional argument.

Same as: setAdditionalKernelArgs( { argument } );

◆ setAdditionalKernelArgs()

void CTL::OCL::GenericOCLVolumeFilter::setAdditionalKernelArgs ( const std::vector< float > & arguments )

Sets the additional kernel arguments to arguments.

Each element of arguments is passed as an individual argument to the OpenCL kernel of this instance. The index of these arguments starts at 3 (three fixed arguments + additional args from arguments).

Note that, if the kernel used by this instance has additional arguments, the same number of arguments must be passed to arguments to make the filter work.

◆ setParameter()

void CTL::OCL::GenericOCLVolumeFilter::setParameter ( const QVariant & parameter )

overridevirtual

Sets the parameters of this instance based on the passed QVariant parameter.

Parameters must be passed as a QVariantMap with one or both of the follwing key-value-pairs:

("Kernel file name", file name of the .cl file containing the kernel),
("Additional arguments", list of the used additional kernel parameters).

Note that it is not recommended to change settings of an GenericOCLVolumeFilter this way. This methods mainly serves its purpose in deserialization.

Reimplemented from CTL::AbstractVolumeFilter.

◆ type()

int CTL::OCL::GenericOCLVolumeFilter::type ( ) const

inlineoverridevirtual

Returns the type-id of the serializable object. Used in deserialization to determine the proper object type.

Add derived classes to the enumeration using the CTL_TYPE_ID macro.

Reimplemented from CTL::AbstractVolumeFilter.

Reimplemented in CTL::OCL::TVRegularizer, and CTL::OCL::HuberRegularizer.

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

modules/src/processing/genericoclvolumefilter.h
modules/src/processing/genericoclvolumefilter.cpp

Public Types

Public Member Functions

Protected Member Functions

Protected Attributes

Static Private Member Functions

Friends

Detailed Description

Constructor & Destructor Documentation

◆ GenericOCLVolumeFilter() [1/3]

◆ GenericOCLVolumeFilter() [2/3]

◆ GenericOCLVolumeFilter() [3/3]

Member Function Documentation

◆ addKernelFromFile()

◆ filter()

◆ getCommandQueue()

◆ globalWorksize()

◆ localWorksize()

◆ parameter()

◆ setAdditionalKernelArg()

◆ setAdditionalKernelArgs()

◆ setParameter()

◆ type()