Hello all!
I tried to use OpenCL to draw into an OpenGL render-buffer and display the content of the render-buffer with glBlitFramebuffer.
I’ve checked the return value of every OpenCL function, but everyone returns CL_SUCCESS. I’ve also tried to use the debugger “gdebugger” to find out what’s going on:
First, the window is filled with cyan color using glClear, then the render-buffer is filled green using glClear.
After clEnqueueAcquireGLObjects, “gdebugger” shows me the green filled render-buffer as an image in the OpenCL.
I Expected my OpenCL kernel to overwrite the color of the whole render-buffer with red after calling clEnqueueNDRangeKernel, but nothing happened.
Finally, glBlitFramebuffer copy the still green filled render-buffer into the window.
I can’t find my mistake, please help me!
The Program:
#include <stdio.h>
#include <stdbool.h>
#include <X11/X.h>
#include <X11/Xlib.h>
#include <GL/gl.h>
#include <GL/glx.h>
#include <GL/glext.h>
#include <CL/cl_gl.h>
#ifndef cl_khr_gl_sharing
#error "cl_khr_gl_sharing isn't set"
#endif
#define STR_EVAL(x) #x
#define STR(x) STR_EVAL(x)
#define Error(...) fprintf(stderr,"An error occured in \"" __FILE__ "\" on line " STR(__LINE__) ": " __VA_ARGS__)
struct datas {
Display* dpy;
Window win;
GLXContext glc;
cl_context clc;
cl_device_id deviceId;
cl_command_queue clCommandQueue;
cl_program program;
cl_kernel kernel;
cl_mem cl_color;
GLuint gl_frame;
GLuint gl_color;
};
static bool initOpenCL(struct datas* p){
printf("
--- initializing OpenCL ---
");
cl_uint platformIdCount = 0;
clGetPlatformIDs (0, 0x0, &platformIdCount);
if(!platformIdCount){
Error("No platforms found
");
goto error;
}
printf("Using 1st platform
");
cl_platform_id* platformIds = malloc(sizeof(cl_platform_id)*platformIdCount);
if(!platformIds){
Error("e_initOpenCL: malloc failed
");
goto error;
}
clGetPlatformIDs(platformIdCount, platformIds, 0x0);
printf("platformIdCount: %d
", (int)platformIdCount);
cl_platform_id platformId = *platformIds;
free(platformIds);
cl_context_properties props[] = {
CL_GL_CONTEXT_KHR, (cl_context_properties)p->glc,
CL_GLX_DISPLAY_KHR, (cl_context_properties)p->dpy,
CL_CONTEXT_PLATFORM, (cl_context_properties)platformId,
0
};
cl_device_id deviceId;
clGetGLContextInfoKHR(props, CL_CURRENT_DEVICE_FOR_GL_CONTEXT_KHR, sizeof(deviceId), &deviceId, NULL);
p->deviceId = deviceId;
#define CLDIS_IDN(x) x
#define CLDIS_STR(x) #x
#define CLDIS(x) \
x(CL_DEVICE_NAME), \
x(CL_DEVICE_VENDOR), \
x(CL_DEVICE_VERSION), \
x(CL_DEVICE_TYPE), \
x(CL_DRIVER_VERSION), \
x(CL_DEVICE_PROFILE), \
x(CL_DEVICE_MAX_COMPUTE_UNITS), \
x(CL_DEVICE_MAX_CLOCK_FREQUENCY), \
x(CL_DEVICE_MEM_BASE_ADDR_ALIGN), \
x(CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE)
cl_device_info params[] = {
CLDIS(CLDIS_IDN)
};
const char*const param_names[] = {
CLDIS(CLDIS_STR)
};
#undef CLDIS
#undef CLDIS_IDN
#undef CLDIS_STR
unsigned char param_value[256];
size_t param_value_size_ret;
for(size_t i=0;i<sizeof(param_names)/sizeof(*param_names);i++){
clGetDeviceInfo(deviceId,params[i],sizeof(param_value),param_value,¶m_value_size_ret);
switch(params[i]){
case CL_DEVICE_VENDOR:
case CL_DEVICE_VERSION:
case CL_DRIVER_VERSION:
case CL_DEVICE_PROFILE:
case CL_DEVICE_NAME:
printf(
" %s: %.*s
",
param_names[i],
(int)( param_value_size_ret > sizeof(param_value) ? sizeof(param_value) : param_value_size_ret ),
param_value
);
break;
case CL_DEVICE_MAX_COMPUTE_UNITS: //uint
case CL_DEVICE_MAX_CLOCK_FREQUENCY:
case CL_DEVICE_MEM_BASE_ADDR_ALIGN:
case CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE:
printf(
" %s: %lu
",
param_names[i],
(unsigned long)*(cl_uint*)param_value
);
break;
case CL_DEVICE_TYPE:
printf( " %s: ", param_names[i] );
cl_device_type cldt = *(cl_device_type*)param_value;
#define CLDT(x) \
if((cldt&x)==x){ \
printf(#x); \
}
CLDT(CL_DEVICE_TYPE_CPU)
CLDT(CL_DEVICE_TYPE_GPU)
CLDT(CL_DEVICE_TYPE_ACCELERATOR)
CLDT(CL_DEVICE_TYPE_DEFAULT)
#undef CLDT
printf( "
" );
break;
}
}
p->clc = clCreateContext(props, 1, &deviceId, 0x0, 0x0, 0x0);
p->clCommandQueue = clCreateCommandQueue(p->clc,deviceId,0,0x0);
printf("
--- OpenCL initialized ---
");
return true;
error:
Error("
--- OpenCL initialisation failed ---
");
return false;
}
static inline cl_int buildProgram(struct datas* datas,const char* source,size_t size){
cl_int error = CL_SUCCESS;
datas->program = clCreateProgramWithSource(
datas->clc,
1,
&source,
&size,
&error
);
if(error != CL_SUCCESS){
Error("clCreateProgramWithSource failed: error 0x%x
",error);
return error;
}
error = clBuildProgram(datas->program,1,&datas->deviceId,
"-w -Werror -cl-single-precision-constant -cl-strict-aliasing -cl-fast-relaxed-math"
,0,0);
if(error != CL_SUCCESS){
Error("clCreateProgramWithSource failed: error 0x%x
",error);
char* error_log;
size_t log_size=0;
if(CL_SUCCESS!=clGetProgramBuildInfo(
datas->program,
datas->deviceId,
CL_PROGRAM_BUILD_LOG,
0,
0,
&log_size
)) return error;
error_log = calloc(log_size,1);
if(!error_log)
return error;
if(CL_SUCCESS!=clGetProgramBuildInfo(
datas->program,
datas->deviceId,
CL_PROGRAM_BUILD_LOG,
log_size,
error_log,0
)) return error;
Error("
build log:
%s
",error_log);
free(error_log);
return error;
}
return CL_SUCCESS;
}
bool loadKernel( struct datas* datas, const char* path, const char* name ){
FILE* f = fopen(path,"rb");
if(!f){
Error("Failed to open file \"%s\"!
",path);
return false;
}
fseek(f,0,SEEK_END);
size_t source_code_length = ftell(f);
if(!source_code_length){
printf("Warning: Ignoring empty file \"%s\"
",path);
return false;
}
fseek(f,0,SEEK_SET);
void* source_code = malloc(source_code_length);
if(!source_code){
Error("Failed to alloate memory for file \"%s\"!
",path);
return false;
}
fread(source_code,1,source_code_length,f);
fclose(f);
buildProgram(datas,source_code,source_code_length);
cl_int error_code;
datas->kernel = clCreateKernel( datas->program, name, &error_code );
if( error_code != CL_SUCCESS ){
Error( "Failed to create kernel \"%s\": error 0x%x
", name, error_code );
return false;
}
return true;
}
bool initBuffers(struct datas* datas,size_t width,size_t height){
cl_int error_code;
GLenum errors;
glGenFramebuffers(1, &datas->gl_frame );
glGenRenderbuffers(1, &datas->gl_color );
glBindRenderbuffer( GL_RENDERBUFFER, datas->gl_color );
glRenderbufferStorage( GL_RENDERBUFFER, GL_RGBA8, width, height );
glBindFramebuffer( GL_FRAMEBUFFER, datas->gl_frame );
glFramebufferRenderbuffer( GL_FRAMEBUFFER, GL_COLOR_ATTACHMENT0, GL_RENDERBUFFER, datas->gl_color );
errors = glGetError();
GLenum status = glCheckFramebufferStatus( GL_FRAMEBUFFER );
if( status != GL_FRAMEBUFFER_COMPLETE ){
Error( "glCheckFramebufferStatus error 0x%x
", status );
return false;
}
glBindFramebuffer( GL_FRAMEBUFFER, 0 );
glBindFramebuffer( GL_READ_FRAMEBUFFER, datas->gl_frame );
if(errors!=GL_NO_ERROR){
Error( "glGetError returned 0x%x
", errors );
return false;
}
datas->cl_color = clCreateFromGLRenderbuffer( datas->clc, CL_MEM_READ_WRITE, datas->gl_color, &error_code );
if( error_code != CL_SUCCESS ){
Error( "clCreateFromGLRenderbuffer error 0x%x
", error_code );
return false;
}
return true;
}
bool initWindowAndOpenGL(struct datas* datas,int width,int height){
Window root;
XVisualInfo *vi;
XSetWindowAttributes swa = {0};
datas->dpy = XOpenDisplay(NULL);
if(datas->dpy == NULL) {
Error("cannot connect to X server
");
return false;
}
root = DefaultRootWindow(datas->dpy);
vi = glXChooseVisual(datas->dpy, 0, (GLint[]){
GLX_RGBA,
GLX_RED_SIZE, 1,
GLX_GREEN_SIZE, 1,
GLX_BLUE_SIZE, 1,
GLX_DEPTH_SIZE, 0,
None
});
if(vi == NULL) {
Error("no appropriate visual found
");
return false;
}else{
printf("init: visual %p selected
", (void *)vi->visualid);
}
swa.colormap = XCreateColormap(datas->dpy, root, vi->visual, AllocNone);
swa.event_mask = StructureNotifyMask | KeyPressMask;
datas->win = XCreateWindow(datas->dpy, root, 0, 0, width, height, 0, vi->depth, InputOutput, vi->visual, CWBackPixel | CWColormap | CWEventMask, &swa);
XMapWindow(datas->dpy, datas->win);
datas->glc = glXCreateContext(datas->dpy, vi, NULL, GL_TRUE);
glXMakeCurrent(datas->dpy, datas->win, datas->glc);
XFree(vi);
return true;
}
int main(){
size_t fb_size[] = {
800,600,0
};
struct datas datas = {0};
initWindowAndOpenGL(&datas,fb_size[0],fb_size[1]);
initOpenCL(&datas);
initBuffers(&datas, fb_size[0], fb_size[1]);
GLenum errors = glGetError();
if(errors!=GL_NO_ERROR){
Error( "glGetError returned 0x%x
", errors );
return false;
}
glClearColor(0,1,1,1);
glClear(GL_COLOR_BUFFER_BIT);
glClearColor(0,1,0,1);
glViewport(0, 0, fb_size[0], fb_size[1]);
loadKernel(&datas,"program.cl","myKernel");
while(true){
XSync(datas.dpy,true);
glBindFramebuffer(GL_DRAW_FRAMEBUFFER, datas.gl_frame);
glClear(GL_COLOR_BUFFER_BIT);
glBindFramebuffer(GL_DRAW_FRAMEBUFFER, 0 );
{
cl_int error=0;
glFinish();
error = clEnqueueAcquireGLObjects(datas.clCommandQueue,1,&datas.cl_color,0,0,0);
if(error != CL_SUCCESS){
Error("clEnqueueAcquireGLObjects returned error: 0x%x
",error);
break;
}
{
error = clSetKernelArg(datas.kernel, 0, sizeof(cl_mem), &datas.cl_color);
if(error != CL_SUCCESS){
Error("clSetKernelArg returned error: 0x%x
",error);
break;
}
error = clEnqueueNDRangeKernel( datas.clCommandQueue, datas.kernel, 1, 0, fb_size, 0, 0, 0, 0 );
if(error != CL_SUCCESS){
Error("clEnqueueNDRangeKernel returned error: 0x%x
",error);
break;
}
}
error = clEnqueueReleaseGLObjects(datas.clCommandQueue,1,&datas.cl_color,0,0,0);
if(error != CL_SUCCESS){
Error("clEnqueueReleaseGLObjects returned error: 0x%x
",error);
break;
}
clFinish(datas.clCommandQueue);
}
glBlitFramebuffer(0,0,fb_size[0], fb_size[1],0,0,fb_size[0], fb_size[1],GL_COLOR_BUFFER_BIT,GL_NEAREST);
}
return true;
}
The OpenCL kernel:
kernel void myKernel(global write_only image2d_t fb){
const int2 xy = {
get_global_id(0),
get_global_id(1)
};
const uint4 color = {
255,
0,
0, 127
};
write_imageui(fb, xy, color);
}