OpenCL/OpenGL Interoperation with Textures

Download source code for this section

1. Introduction

In previous Interop tutorials we discussed how to interoperate OpenGL vertex buffer objects (VBOs) with OpenCL. In this section we’ll demonstrate how to use OpenCL to manipulate OpenGL texture objects.

As in the case of VBOs, the main advantage of using OpenCL to manipulate OpenGL VBOs is that it is not necessary to copy data from and to the GPU insofar as the Texture object can be manipulated directly in the Device memory.

We’ll first cover how to interoperate textures using the GLRender framework and OpenCLTemplate and then follow up with a more in-depth discussion of how to deal with the interoperation using Cloo and the OpenCL API itself.

The video below shows examples of what can be done using the interop: real-time blur and direct pixel manipulation in which pixels are sorted in the image’s lines according to their Red component (pixels with greater Red component get moved to the left).

2. Texture Structure

Since we’re using C# we decided to follow its BGRA color structure to optimize performance and this leads to a somewhat confusing data structure from the OpenCL standpoint which we’ll try to clarify in this section.

2.1 OpenCL C99 Standpoint

From OpenCL C99 standpoint, i.e., when using OpenCL kernels to manipulate image2Ds created from GL textures, the images should be treated as 4-channel float images in RGBA color order. Assuming that we’re using

__kernel void myKernel(__read_only image2d_t bmp)
const sampler_t smp = CLK_NORMALIZED_COORDS_FALSE | //Natural coordinates
CLK_ADDRESS_CLAMP | //Clamp to zeros
CLK_FILTER_NEAREST; //Don't interpolate
This means that:

  1. Pixels need to be read using float4 val = read_imagef(bmp, smp, coords) function;
  2. In the above case, val.x = RED component, val.y = GREEN component and val.z = BLUE component;
  3. In the same manner, pixels need to be written using write_imagef(bmpCopy, coords, val);

As an example, the following kernel would make all pixels RED in a CLGL interop Texture:

__kernel void TurnRED(__write_only image2d_t bmp)
   int x = get_global_id(0);
   int y = get_global_id(1); 
   int2 coords = (int2)(x,y);
   //Attention to RGBA order
   float4 val = (float4)(1.0f, 0.0f, 0.0f, 1.0f);
   write_imagef(bmp, coords, val);

2.2 Host Standpoint

To us, a CLGL shared texture should ideally be manipulated directly in Device memory without copies from/to host memory. Nonetheless, we’ll explain the texture structure from Host standpoint for the sake of completeness.

From the Host code standpoint, a CLGL shared texture (as created in OpenCLTemplate) is a BYTE array with channel order ABGR (alpha, blue, green, red). This means that to write to a shared CLGL Texture from the Host it’s necessary to create a BYTE[4*Width*Height] array in Alpha Blue Green Red order and the same logic applies when reading data. For an in-depth analysis please study the DemoCLGL code included in the source code for this section. It demonstrates how written values are read back to Host memory.

2.3 Notes on Texture Structure

Please bear in mind that this structure was CHOSEN by us and that it is possible to choose other ordering structures as well. Below are the reasons why we chose the presented structure:

  1. First and foremost, because it works;
  2. Using this structure it’s not necessary to swap components of an original bitmap when loading the texture to OpenGL memory, which would cost performance;
  3. A CLGL shared texture won’t usually be read back to the Device as its main purpose is onscreen display;
  4. Manipulating vector components in (float4) BGRA order is just as easy as manipulating them in (float4) RGBA mode.

3. Manipulating the Texture with OpenCL in GLRender Framework

Now that the “boring” part is over, we’ll show how to perform texture manipulation using the GLRender framework. Create a new C# project and add the following code to initialize a CLGL shared context:

GLRender glw;
GLRender.GLVBOModel plane;
CLCalc.Program.Image2D CLBmp;
private void Form1_Load(object sender, EventArgs e)
   //Creates OpenCL/GL shared environment. 1 line :-)
   glw = new GLRender(this, true, -1);
   //Sets background color to light blue
   glw.ClearColor = new float[] { 0.5f, 0.5f, 1.0f };

Now we’ll create a plane in XY and attach an empty C# texture to it. Notice that it would be possible to bind a texture loaded from an image file using C# new Bitmap(string path) constructor. The GLVBOModel creator from equations has a built-in structure to compute Texture Coordinates automatically. If this was not the case it would be necessary to input the OpenGL TexCoords VBO manually.

//Creates plane object to hold textures
plane = GLRender.GLVBOModel.CreateSurface(
    new float[] { -18, 18, 2 },
    new float[] { -10, 10, 2 }
    new string[] { "u", "v", "0.0f" },
    new string[] { "0.0f", "0.0f", "1.0f"});
//Inserts a dummy bitmap into OpenGL memory
plane.SetTexture(new Bitmap(180, 100));
//Adds plane to models being displayed

Let us now retrieve a handle to an OpenCL Image2D from the texture:

//Acquires OpenCL image from OpenGL texture
CLBmp = plane.GetCLTexture2D();

From this point on, it is possible to treat CLBmp as a regular OpenCL Image2D memory object consisting of one float4 vector for each pixel. Just remember to Acquire the Image2D before using it and Release it after manipulation. You may execute the code if you wish and check that a black plane will be drawn in a light blue background.

Let’s create a simple kernel that draws a gradient into the texture.

string srcGradient = @"
__kernel void Grad(__write_only image2d_t bmp)
    int x = get_global_id(0);
    int y = get_global_id(1);
    int w = get_global_size(0)-1;
    int h = get_global_size(1)-1;
    int2 coords = (int2)(x,y);
    float red = (float)x/(float)w;
    float blue = (float)y/(float)h;
    float4 val = (float4)(red, 0.0f, blue, 1.0f);
    write_imagef(bmp, coords, val);
CLCalc.Program.Kernel kernelGrad = new CLCalc.Program.Kernel("Grad");
//Acquire texture before operating with OpenCL
//Executes kernel
kernelGrad.Execute(new CLCalc.Program.MemoryObject[] { CLBmp }, new int[] { CLBmp.Width, CLBmp.Height }); //Releases texture after manipulation

And this is the result shown in the software:

For more details please check CLGLInteropTextures included in the source code for this section.

4. Creating Image2D Objects from GL Textures in OpenCLTemplate

A new Image2D constructor has been added to OpenCLTemplate in order to create shared OpenCL/GL textures. The first step, of course, is to create the texture in OpenGL and that can be done using CLGLInteropFunctions.ApplyTexture. A typical scenario is the following:

CLGLInteropFunctions.ApplyTexture(bmp, ref GLTextureBuffer);
CLTexture2D = new CLCalc.Program.Image2D(GLTextureBuffer);

In the code above, this.GLTextureBuffer is the integer which holds a reference to a texture number in OpenGL memory and bmp is a System.Drawing.Bitmap C# object.

5. Conclusion

We’ve presented how to use GLRender framework to use OpenCL/GL interoperation with textures. Considering that a 3D shared environment is created automatically and that most of the work to bind the texture to OpenGL and retrieve it using OpenCL is automated in the GLRender.GLVBOModel class, we consider that this is a valuable tool for high-performance scientific 3D visualization software.

A Texture buffer structure was chosen to allow high performance while still maintaining easy usability considering that the user only needs the information that OpenCL Image2Ds created from GL buffers should be treated in C99 kernels as float4 images in BGRA color order. The downside of this structure is the complexity of manipulating the Image2Ds in Host memory but we consider this issue as a minor disadvantage seeing as shared Textures should be primarily used for onscreen information display.

Source code is presented to demonstrate all interoperation steps and readers are encouraged to study the Blur and Sort kernels presented in CLGLInteropTextures code provided with the source code for this section.

Download source code for this section

Appendix A: Sharing Textures Using Cloo and the OpenCL API

This section assumes that the user is familiar with texture creation in OpenGL. If this is not the case I would recommed reading some material about the subject, such as NeHe’s textures tutorials at

A.1 Creating the GL Texture

In order to interoperate with OpenCL, the GL texture needs to be completely defined as per the OpenGL specification, as defined in Khronos’ OpenCL specification v1.1, page 316:

The texture object must be a complete texture as per OpenGL rules on texture completeness. The texture format and dimensions defined by OpenGL for the specified miplevel of the texture will be used to create the 2D image object. Only GL texture objects with an internal format that maps to appropriate image channel order and data type specified intables 5.5 and 5.6 may be used to create a 2D image object.

In practice this means that some texture parameters have to be explicitly defined, as shown in OpenTK code below, which is the core of the function CLGLInteropFunctions.ApplyTexture:

if (ind <= 0) ind = GL.GenTexture();
//texture, if there is one
System.Drawing.Bitmap image = new System.Drawing.Bitmap(TextureBitmap);
System.Drawing.Imaging.BitmapData bitmapdata;
System.Drawing.Rectangle rect = new System.Drawing.Rectangle(0, 0, image.Width, image.Height);
bitmapdata = image.LockBits(
//Acquire texture
GL.BindTexture(TextureTarget.Texture2D, ind);
GL.TexParameter(TextureTarget.Texture2D, TextureParameterName.TextureWrapS, (int)All.ClampToEdge);
GL.TexParameter(TextureTarget.Texture2D, TextureParameterName.TextureWrapT, (int)All.ClampToEdge);
GL.TexParameter(TextureTarget.Texture2D, TextureParameterName.TextureMinFilter, (int)All.Nearest);
GL.TexParameter(TextureTarget.Texture2D, TextureParameterName.TextureMagFilter, (int)All.Nearest);
GL.TexImage2D(TextureTarget.Texture2D, 0, PixelInternalFormat.Rgba8, image.Width, image.Height, 0, (OpenTK.Graphics.OpenGL.PixelFormat)(int)All.BgrExt,PixelType.UnsignedByte, bitmapdata.Scan0);

It is very important to keep track of the texture number generated by glGenTextures because OpenCL will need this reference in order to have access to the texture. In the example above this number is stores in ind.

One important aspect to notice is the use of BgrExt type. In words, this means that we’ll use C# BGRA color structure instead of the usual RGBA mode. It would, of course, be possible to manipulate C# pixels before storing them but since this just involves component orders we at CMSoft consider that there’s no reason to lose performance just to fix a less usual color order.

A.2 Creating the Image2D from GL Texture

Once the texture is completely defined in OpenGL terms, acquisition is fairly straightforward using clCreateFromGLTexture2D or its encapsulated version in Cloo:


One still needs to be careful about the order of components of the GL texture. From this point on, aside from having to Acquire/Release the texture (check OpenCL spec), the memory object behaves as any Image2D object.

Leave a Reply

Your email address will not be published. Required fields are marked *