Using read-write images in OpenCL 2.0 / Intel Blog

Prior to OpenCL 2.0, it was not possible to perform image read and write operations within the same kernel. You could declare images as CL_MEM_READ_WRITE , but after transferring the image, the kernel had to choose one of two things : either __read_only (read-only access) or __write_only ( write -only access). OpenCL 2.0 introduces the ability to read and write images within a single core. However, there are several features that we will discuss in detail in this post.

input1 = clCreateImage(
oclobjects.context,
CL_MEM_READ_WRITE|CL_MEM_COPY_HOST_PTR,
&format,
&desc,
&input_data1[0],
&err );
SAMPLE_CHECK_ERRORS( err );

Code Snippet 1. You could create an image buffer using CL_MEM_READ_WRITE

__kernel void Alpha( __read_write image2d_t inputImage1, 
__read_only image2d_t 
inputImage2, 
uint width, 
uint height, 
float alpha, 
float beta, 
int gamma )

Snippet 2. OpenCL 2.0 now has the ability to read and write images within a single core

Benefits of Readable and Writable Images

Image convolution is less efficient when using the new image reading and writing functionality, but other image processing algorithms can benefit significantly from this functionality. One example of such processes may be imaging.

In OpenCL 1.2 and earlier, images could only have __read_only and __write_only qualifiers. In OpenCL 2.0, the __read_write qualifier appeared, the output can be copied to the input buffer. This reduces the amount of resources needed. For any image changes, it is necessary to process the image as a buffer and work with this buffer (see cl_khr_image2d_from_buffer ).

The current solution is to process images as buffers and manage buffers. To process two-dimensional images as buffers, a certain amount of resources is required. It also makes it impossible to use the cutting and filtering capabilities available in read_images. Therefore, it is advisable to use images with the read_write qualifier.

Example Overview

In the example, two bitmaps (input1.bmp and input2.bmp) are buffered. Then these images are superimposed on top of each other based on the alpha value (this is the weight factor in the pixel equation). The alpha value is passed as a parameter.


Figure 1. Alpha = 0.84089642

Input images must be 24-bit or 32-bit. The output is a 24-bit image. Input images must be the same size. Images were in ARGB format, this was taken into account when downloading them.


Figure 2. Alpha = 0.32453

The ARGB format is converted to RGBA. Changing the beta value leads to significant changes in the output image.
Using SDK
The SDK demonstrates overlaying images by reading and writing. The following command line options can be used to control the operation of the sample code.

Parameters	Description
-h, --help	Display this text and exit.
-p, --platform <number or string>	Select the platform whose devices are used.
-t, --type all | cpu | gpu | acc | default | <OpenCL constant for device type>	Select the type of device on which the OpenCL kernel runs.
-d, --device <number or string>	Select the device on which all work is performed.
-i, --infile <24- or 32-bit BMP input file>	The name of the first read file in BMP format. The default is input1.bmp .
-j, --infile <24- or 32-bit BMP input file>	The name of the second read file in BMP format. The default is input2.bmp .
-o, --outfile <24- or 32-bit BMP input file>	The name of the output file to write to. The default is output.bmp for OCL1.2 and 20_output.bmp for OCL2.0.
-a, --alpha <floating point value from zero to one>	A nonzero positive value that determines how much two images will be superimposed on one another when combined. The default alpha value is 0.84089642. The default beta value is 0.15950358.

The sample SDK has default values, so the application can run without any input from the user. Users can use their own input BMP files. All files must be 24-bit or 32-bit. The alpha value determines how much the first image will overlap with the second.

calculatedPixel = ((currentPixelImage1 * alpha) + (currentPixeImage2 * beta) + gamma);

The beta value is equal to the difference between unit and alpha value.

float beta = 1 – alpha;

These two values ​​determine the "weight" of images 1 and 2 in the output image.
You can use the gamma value to change the brightness of each pixel. By default, this value is zero. The user can change the brightness of the entire image.

Program launch example

Figure 3. Running the program on an OpenCL 2.0 device

Limitations of Images Readable and Writable

Limiters cannot be used with images that require synchronization between different workgroups. Image convolution requires synchronization of all streams. Convolution with respect to images usually involves mathematical operations on two matrices and the creation of a third matrix as a result. The image convolution example uses Gaussian blur. Other examples use sharpening images, detecting edges, and embossing.

As an example, consider the Gaussian blur. A Gaussian filter is a low-pass filter that removes high-frequency values. As a result, the level of detail in the image is reduced and a blur effect is obtained. Applying Gaussian blur is the same as transforming an image using the Gaussian distribution function (often called a mask). To demonstrate the functionality of reading and writing images, we had to apply horizontal and vertical blur.

In OpenCL 1.2, this would have to be done in two passes. One core would be used only for horizontal blur, and the other for vertical blur. The result of one blur will be used as input for the next (depending on which blur was the first).

__kernel void GaussianBlurHorizontalPass( __read_only image2d_t inputImage, __write_only image2d_t outputImage, __constant float* mask, int maskSize)
{
    int2 currentPosition = (int2)(get_global_id(0), get_global_id(1));
    float4 currentPixel = (float4)(0,0,0,0);
    float4 calculatedPixel = (float4)(0,0,0,0);
    for(int maskIndex = -maskSize; maskIndex < maskSize+1; ++maskIndex)
    {
        currentPixel = read_imagef(inputImage, imageSampler, currentPosition + (int2)(maskIndex, 0));
        calculatedPixel += currentPixel * mask[maskSize + maskIndex];
    }
    write_imagef(outputImage, currentPosition, calculatedPixel);
}
__kernel void GaussianBlurVerticalPass( __read_only image2d_t inputImage, __write_only image2d_t outputImage, __constant float* mask, int maskSize)
{
    int2 currentPosition = (int2)(get_global_id(0), get_global_id(1));
    float4 currentPixel = (float4)(0,0,0,0);
    float4 calculatedPixel = (float4)(0,0,0,0);  
    for(int maskIndex = -maskSize; maskIndex < maskSize+1; ++maskIndex)
    {
        currentPixel = read_imagef(inputImage, imageSampler, currentPosition + (int2)(0, maskIndex));
        calculatedPixel += currentPixel * mask[maskSize + maskIndex];
    }
    write_imagef(outputImage, currentPosition, calculatedPixel);
}

Code Snippet 3. Gaussian Blur Core in OpenCL 1.2

In OpenCL 2.0, these two kernels can be combined into one. Use the limiter to force the blur to end horizontally or vertically before starting the next blur.

__kernel void GaussianBlurDualPass( __read_only image2d_t inputImage, __read_write image2d_t tempRW, __write_only image2d_t outputImage, __constant float* mask, int maskSize)
{
    int2 currentPosition = (int2)(get_global_id(0), get_global_id(1));
    float4 currentPixel = (float4)(0,0,0,0);   
    float4 calculatedPixel = (float4)(0,0,0,0)
    currentPixel = read_imagef(inputImage, currentPosition);
    for(int maskIndex = -maskSize; maskIndex < maskSize+1; ++maskIndex)
    {
        currentPixel = read_imagef(inputImage, currentPosition + (int2)(maskIndex, 0));      
        calculatedPixel += currentPixel * mask[maskSize + maskIndex];
    }
    write_imagef(tempRW, currentPosition, calculatedPixel);
    barrier(CLK_GLOBAL_MEM_FENCE);
    for(int maskIndex = -maskSize; maskIndex < maskSize+1; ++maskIndex)
    {
        currentPixel = read_imagef(tempRW, currentPosition + (int2)(0, maskIndex));
        calculatedPixel += currentPixel * mask[maskSize + maskIndex];
    }
    write_imagef(outputImage, currentPosition, calculatedPixel);
}

Code Snippet 4. Gaussian Blur Kernel in OpenCL 2.0

It turns out that delimiters are inefficient. Using delimiters does not guarantee that the horizontal blur will be performed before the vertical blur starts (if the horizontal blur was the first). As a result, with several launches, different results were obtained. Limiters can be used to synchronize threads in a group. The cause of the problem is that edge pixels are read from several workgroups, but there is no way to synchronize between multiple workgroups. The initial assumption about the possibility of implementing a single Gaussian blur by reading and writing images turned out to be incorrect, since in OpenCL it is not possible to synchronize data dependencies between work groups.