OpenCL - OpenGL 2D texture interop

hello @everyone, i’m trying to write a program where i calculate a 2D texture each frame via OpenCLand display it with OpenGL. I already did something similar with vertex buffers, where I manipulated the vertices each frame with OpenCL and rendered them with OpenGL, and it worked fine - however, I’m having serious issues getting the interop textures to work.

I create an OpenGL 2D texture:


glGenTextures(1, &texID[0]);
glBindTexture(GL_TEXTURE_2D, texID[0]);	
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_REPEAT);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_REPEAT);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_LINEAR);
glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_LINEAR_MIPMAP_LINEAR);	
glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA, 256, 256, 0, GL_RGBA, GL_FLOAT, NULL);
glGenerateMipmap(GL_TEXTURE_2D);

and afterwards, an OpenCL 2D image from the texture:


cl::Image2DGL resultTex = cl::Image2DGL(context, CL_MEM_READ_WRITE, GL_TEXTURE_2D, 0, texID[0], &result);
if (result != CL_SUCCESS) {...}

each frame, I acquire the texture, execute my kernel & release the texture again:


cl_int result = CL_SUCCESS;
cl::Event ev;
cl::NDRange globalRange = cl::NDRange(256, 256);
cl::NDRange localRange = cl::NDRange(16, 16);

std::vector<cl::Memory> memObjs;
memObjs.clear();
memObjs.push_back(resultTex);

result = queue.enqueueAcquireGLObjects(&memObjs, NULL, &ev);
ev.wait();
if (result != CL_SUCCESS) {...}

result = queue.enqueueNDRangeKernel(kernelWriteTest, cl::NullRange, globalRange, localRange, NULL, &ev);
ev.wait();
if (result != CL_SUCCESS) {...}

result = queue.enqueueReleaseGLObjects(&memObjs, NULL, &ev);
ev.wait();
if (result != CL_SUCCESS) {...}

glGenerateMipmap(GL_TEXTURE_2D);

the simplified kernel looks like this:


__kernel void testKernelWrite(__write_only image2d_t resultTexture)
{	
	int2 imgCoords = (int2)(get_global_id(0), get_global_id(1));
	
	float4 imgVal = (float4)(1.0f, 1.0f, 0.0f, 1.0f);

	write_imagef(resultTexture, imgCoords, imgVal);
}

basically, the program compiles & runs without problems - however, the color of the texture when rendered with OpenGL never changes, no matter how many times I execute the kernel. i.e. if I pass NULL during texture creation, the texture stays black; if i initialize it with some image, rendering the texture will show the initial image and nothing else.

I also wrote a 2nd kernel, which I too executed each frame after releasing & acquiring the shared texture again:


#pragma OPENCL EXTENSION cl_amd_printf : enable

__kernel void testKernelRead(__read_only image2d_t inputTexture)
{	
	int2 imgCoords = (int2)(get_global_id(0), get_global_id(1));
	printf("%d %d ", imgCoords.x, imgCoords.y);
	
	const sampler_t smp =  CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;

	float4 imgVal = read_imagef(inputTexture, smp, imgCoords); 
	printf("%f %f %f %f
", imgVal.x, imgVal.y, imgVal.z, imgVal.w);
}

weirdly, when executing this kernel, it prints the correct values…

I really hope someone can help me with this problem (:

Did you remember to do a glFinish() before calling clEnqueueAcquireGLObjects() and clFinish() after calling clEnqueueReleaseGLObjects()?

From section 9.8.6.1:

Prior to calling clEnqueueAcquireGLObjects, the application must ensure that any pending GL
operations which access the objects specified in mem_objects have completed. This may be
accomplished portably by issuing and waiting for completion of a glFinish command on all GL
contexts with pending references to these objects. Implementations may offer more efficient
synchronization methods; for example on some platforms calling glFlush may be sufficient, or
synchronization may be implicit within a thread, or there may be vendor-specific extensions that
enable placing a fence in the GL command stream and waiting for completion of that fence in the
CL command queue. Note that no synchronization methods other than glFinish are portable
between OpenGL implementations at this time.

Similarly, after calling clEnqueueReleaseGLObjects, the application is responsible for ensuring
that any pending OpenCL operations which access the objects specified in mem_objects have
completed prior to executing subsequent GL commands which reference these objects. This
may be accomplished portably by calling clWaitForEvents with the event object returned by
clEnqueueReleaseGLObjects, or by calling clFinish. As above, some implementations may
offer more efficient methods.

If your implementation supports the cl_khr_gl_event and GL_ARB_cl_event extensions, then you can use clCreateEventFromGLsyncKHR() and gl CreateSyncFromCLeventARB() respectively to synchronize instead of using glFinish() and clFinish().

well I did call glFinish before acquiring the shared texture, but not clFinish after releasing it, because i thought waiting for the associated event to complete was good enough; however, adding that clFinish call did not change anything. the sync extensions are unfortunately not supported in my implementation.

ok, i feel stupid now.

i simply forgot that the OpenGL texture creation needs to happen after OpenCL context creation - for some reason, i thought it needed to happen before; everything works after changing this.

I have absolutely the same problem. But from the beginning I have texture initialization after CL context creation. Please, look the whole issue here: https://www.khronos.org/message_boards/viewtopic.php?f=28&t=4638