PDA

View Full Version : Wrong OpenGL Texture Format Type for CL/GL-Interop?



rdoubleui
01-31-2012, 09:04 AM
I'm trying the OpenCL-OpenGL interop for textures on my Geforce 330M with CUDA Toolkit 4.0.

I want to capture a frame, use that data as an input image (Image2D) to a OpenCL Kernel. The Kernel should manipulate the data and write it to an Image2DGL, which is an image object with an attached OpenGL texture. Basically it looks like that:


_______________ RGB _______________
| | uint8* | | CL_RGBA / CL_UNORM_INT8
| Grabber | ------------> | Image2D | -------------------------.
| avcodec | | [input] | |
|_______________| |_______________| |
|
V
_______________ _______________ _______________
| | | | | |
| Texture | ------------> | Image2DGL | <-----------------> | Kernel |
|_______________| | [output] | |_______________|
|_______________|
Internal
Format: GL_RGBA
Format: GL_RGBA
Type: ?

I'm initializing the texture like that:



GLuint tex = 0;

void initTexture( int width, int height )
{
glGenTextures(1, &tex);
glBindTexture(GL_TEXTURE_RECTANGLE, tex);
// now here is where I need assistance: The type parameter of the Texture (GL_UNSIGNED_BYTE)
glTexImage2D(GL_TEXTURE_RECTANGLE, 0, GL_RGBA, width, height, 0, GL_UNSIGNED_BYTE, GL_FLOAT, NULL );
}


Then I create the shared image (Image2DGL):



texMems.push_back(Image2DGL(clw->context, CL_MEM_READ_WRITE, GL_TEXTURE_RECTANGLE, 0, tex, &err));


Then I create the source image (input image):



ImageFormat format;
format.image_channel_data_type = CL_UNORM_INT8;
format.image_channel_order = CL_RGBA;
srcImgBuffer = Image2D(clw->context, CL_MEM_READ_WRITE, format, width, height, 0, NULL, &err);


In every render loop I'm writing the data into the srcImgBuffer:



// write the frame to the image buffer
clw->queue.enqueueWriteImage(srcImgBuffer, CL_TRUE, origin, region, 0, 0, (void*)data, NULL, NULL);


Also I'm setting the arguments for the kernel:



tex_kernel.setArg(0, texMems[0]);
tex_kernel.setArg(1, srcImgBuffer);
tex_kernel.setArg(2, width);
tex_kernel.setArg(3, height);


Before and after I do acquire and release the GL objects. The testing kernel looks like that:



__kernel void init_texture_kernel(__write_only image2d_t out, __read_only image2d_t in, int w, int h)
{
const sampler_t smp = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;

int2 coords = { get_global_id(0), get_global_id(1) };
float4 pixel = read_imagef(in, smp, coords);
float4 test = { (float)coords.x/(float)w , 0, 0, 1};
write_imagef( out, coords, pixel );
}


The image_channel_data_type can be read as float in the kernel and is interpreted as normalized value. The output image looks not right, I've got a sliced picture (linewise), obviously because of the wrong data interpretation. As I mentioned I assume the error is within the initialization of the texture's type. I also tried GL_FLOAT (as I'm writing as float to the image in the kernel).

http://i.stack.imgur.com/QeG7t.png

The left one is a PPM out of the decoder, the right one is what I'm getting back on my output texture.

If I bind the captured frames to the texture directly, the video plays ok. So it must be related to the CL-GL interface somehow.

The Image2DGL init did not throw any errors.

If someone actually read up to here: Do you have suggestions regarding the texture's type to resolve the issue? Or can anybody point to a more-in-depth explanation of the interop?

rdoubleui
02-01-2012, 09:24 AM
Nevermind, I've found the solution.

The trouble was with the grabbed frames (as probably everybody except me can guess from the screenshot). I grabbed RGB and in the kernel read the pixels using float4s. Of course that messes up the picture.

Plus I probably posted in the wrong forum anyway, as this is about interoperability of OpenCL code on different devices. Sorry.