Bon. On aurai dit que ca marchait mais... enfait, le VBO ne semble pas modifié par OpenCL, et si je met 32x32 comme mesh, ça segfault.

git-svn-id: file:///var/svn/2013-gpudataviz/trunk@12 371a6b4a-a258-45f8-9dcc-bdd82ce0ac9d

1 files changed, 172 insertions, 18 deletions

diff --git a/tests/poc_c/main.c b/tests/poc_c/main.c
index 4e2be04..a1341c0 100644
--- a/tests/poc_c/main.c
+++ b/tests/poc_c/main.c
@@ -1,5 +1,7 @@
 #include <stdlib.h>
 
+#include <string.h> /* For strlen (to be removed ?) */
+
 #include <gtk/gtk.h>
 #include <gtk/gtkgl.h>
 
@@ -11,7 +13,17 @@
 
 #include <gdk/x11/gdkglx.h> // X11 specific
 
-gboolean gl_area_on_draw (GtkObject* area, GdkEventExpose* event, gpointer data) {
+typedef struct {
+	GLuint gl_vbo;
+	size_t mesh_width, mesh_height;
+} on_expose_data_t;
+
+gboolean gl_area_on_draw (GtkObject* area, GdkEventExpose* event, gpointer void_data) {
+	if (void_data == NULL) {
+		fputs("gl_area_on_draw() : NULL void_data arg\n", stderr);
+		return FALSE;
+	}
+
 	GdkGLDrawable* drawable = gtk_widget_get_gl_drawable (GTK_WIDGET (area));
 	GdkGLContext* context = gtk_widget_get_gl_context (GTK_WIDGET (area));
 
@@ -21,6 +33,8 @@ gboolean gl_area_on_draw (GtkObject* area, GdkEventExpose* event, gpointer data)
 		return FALSE;
 	}
 
+	on_expose_data_t *data = (on_expose_data_t *) void_data;
+
 	GtkAllocation allocation;
 	gtk_widget_get_allocation (GTK_WIDGET (area), &allocation);
 	GLdouble viewport_width = (GLdouble) allocation.width;
@@ -57,13 +71,23 @@ gboolean gl_area_on_draw (GtkObject* area, GdkEventExpose* event, gpointer data)
 
 	// Draw a 1x1 square center on origin and on the x*y plan.
 	// x*y*z, Z negative forward oriented.
+/*
 	glBegin (GL_QUADS);
 	glVertex3d ( 0.5, -0.5, 0.0);
 	glVertex3d ( 0.5,  0.5, 0.0);
 	glVertex3d (-0.5,  0.5, 0.0);
 	glVertex3d (-0.5, -0.5, 0.0);
 	glEnd ();
-
+*/
+	
+	if (data->gl_vbo) {
+		glBindBuffer(GL_ARRAY_BUFFER, data->gl_vbo);
+		glVertexPointer(data->mesh_width * data->mesh_height, GL_FLOAT, 0, (GLvoid *) 0);
+		glEnableClientState(GL_VERTEX_ARRAY);
+		glDrawArrays(GL_POINTS/*GL_QUADS*/, 0, 4);
+		glDisableClientState(GL_COLOR_ARRAY);
+		glBindBuffer(GL_ARRAY_BUFFER, 0);
+	}
 	gdk_gl_drawable_swap_buffers (drawable);
 
 	gdk_gl_drawable_gl_end (drawable);
@@ -71,6 +95,50 @@ gboolean gl_area_on_draw (GtkObject* area, GdkEventExpose* event, gpointer data)
 	return TRUE; // Do not propagate the event.
 }
 
+cl_int execKernel(cl_context cl_ctx, cl_command_queue commandQueue, cl_kernel kernel, size_t mesh_width, size_t mesh_height, size_t group_size, cl_mem cl_vbo, float time) {
+
+	cl_int res;
+  cl_event  eventND[1];
+
+	// Set local and global work group sizes
+	size_t globalWorkSize[2], localWorkSize[2];
+	globalWorkSize[0]=mesh_width;
+	globalWorkSize[1]=mesh_height;
+	localWorkSize[0]=group_size;
+	localWorkSize[1]=1;
+
+  res=clEnqueueAcquireGLObjects(commandQueue, 1, &cl_vbo, 0, 0, NULL);
+
+	res=clSetKernelArg(kernel, 0, sizeof(cl_mem), &cl_vbo); /* float4 *pos */
+	res=clSetKernelArg(kernel, 1, sizeof(cl_uint), (void *)&mesh_width);
+	res=clSetKernelArg(kernel, 2, sizeof(cl_uint), (void *)&mesh_height);
+	res=clSetKernelArg(kernel, 3, sizeof(float), &time);
+
+  // Execute kernel on given device
+  res=clEnqueueNDRangeKernel(commandQueue, kernel, 2, NULL, globalWorkSize, localWorkSize, 0, NULL, eventND);
+	/* cl_int clEnqueueNDRangeKernel (cl_command_queue command_queue,
+  	cl_kernel kernel,
+  	cl_uint work_dim,
+  	const size_t *global_work_offset, 1.0 spec say "must be NULL" (for future usage)
+  	const size_t *global_work_size,
+  	const size_t *local_work_size,
+  	cl_uint num_events_in_wait_list,
+  	const cl_event *event_wait_list,
+  	cl_event *event)
+	*/
+	if ( res != CL_SUCCESS ) {
+		fputs("Failed to clEnqueueNDRangeKernel()\n", stderr);
+		return 1;
+	}
+
+  res=clFlush(commandQueue);
+	res=clWaitForEvents(1,eventND); /*XXX: SimpleGL utilie une attente active, pourquoi ? */
+	res=clReleaseEvent(eventND[0]);
+  res=clEnqueueReleaseGLObjects(commandQueue, 1, &cl_vbo, 0, 0, 0);
+  res=clFinish(commandQueue);
+
+	return CL_SUCCESS;
+}
 
 int main(int argc, char *argv[]) {
 	GtkWidget* main_win;
@@ -84,11 +152,10 @@ int main(int argc, char *argv[]) {
 	Display *glx_display; // X11 specific
 	gboolean res;
 
-	glewInit();
-	if ( !glewIsSupported("GL_ARB_vertex_buffer_object") ) {
-		fputs ("OpenGL extension GL_ARB_vertex_buffer_object is mandatory for this application\n", stderr);
-		return EXIT_FAILURE;
-	}
+	GLuint gl_vbo;
+	on_expose_data_t on_expose_data;
+	on_expose_data.gl_vbo = 0;
+
 
 	gtk_init(&argc, &argv);
 	if (gdk_gl_init_check(&argc, &argv) == FALSE ) {
@@ -114,7 +181,7 @@ int main(int argc, char *argv[]) {
 		fputs ("Failed to set_gl_capability(gl_area)\n", stderr);
 		return EXIT_FAILURE;
 	}
-	gtk_signal_connect(GTK_OBJECT(gl_area), "expose-event", GTK_SIGNAL_FUNC(gl_area_on_draw), NULL);
+	gtk_signal_connect(GTK_OBJECT(gl_area), "expose-event", GTK_SIGNAL_FUNC(gl_area_on_draw), &on_expose_data);
 	gtk_box_pack_start(GTK_BOX(main_box), GTK_WIDGET(gl_area), TRUE, TRUE, 0);
 
 	gtk_widget_show_all(main_win);
@@ -128,6 +195,18 @@ int main(int argc, char *argv[]) {
 	drawable = gtk_widget_get_gl_drawable(GTK_WIDGET(gl_area));
 	res = gdk_gl_drawable_gl_begin(drawable, gl_context);
 
+	GLenum err = glewInit();
+	if (GLEW_OK != err) {
+  	fprintf(stderr, "glewInit() failure : %s\n", glewGetErrorString(err));
+		return EXIT_FAILURE;
+	}
+
+	if ( ! glewIsSupported("GL_ARB_vertex_buffer_object") ) {
+		fputs ("OpenGL extension GL_ARB_vertex_buffer_object is mandatory for this application\n", stderr);
+		return EXIT_FAILURE;
+	}
+
+
 	/* BEGIN : X11 specific */
 
 	/*
@@ -219,34 +298,109 @@ int main(int argc, char *argv[]) {
 		}
 
 		//TODO : implement selection if multiple devices are usable (instead of taking the first one)
+		printf("cl_device==%p\n", (void *)device);
+
+		cl_ctx = clCreateContext(cpsGL,1,&device,0,0,&res);
+		if ( res!=CL_SUCCESS ) {
+			fputs ("Failed to clCreateContext()\n", stderr);
+			continue;
+		}
+
+		commandQueue = clCreateCommandQueue(cl_ctx,device,0,&res);
+		if ( res!=CL_SUCCESS ) {
+			fputs ("Failed to clCreateContext()\n", stderr);
+			continue;
+		}
+
 		break;
 	}
-	printf("cl_device==%p\n", (void *)device);
 
-	cl_ctx = clCreateContext(cpsGL,1,&device,0,0,&res);
-	commandQueue = clCreateCommandQueue(cl_ctx,device,0,&res);
-
-	GLuint gl_vbo;
-	unsigned int meshWidth=2, meshHeight=2;
-	GLsizeiptr gl_vbo_data_size = meshWidth * meshHeight * sizeof(cl_float4);
+	unsigned int mesh_width=2, mesh_height=2, group_size=2;
+	GLsizeiptr gl_vbo_data_size = mesh_width * mesh_height * sizeof(cl_float4);
+/*
 	float gl_vertex_default_pos[16] = { // to be removed
 		 0.5, -0.5, 0.0, 1.0,
 		 0.5,  0.5, 0.0, 1.0,
-		-0.5, -0.5, 0.0, 1.0,
+		-0.5,  0.5, 0.0, 1.0,
 		-0.5, -0.5, 0.0, 1.0
 	};
-
+*/
 	res = gdk_gl_drawable_gl_begin(drawable, gl_context);
 	glGenBuffers(1, &gl_vbo);
 	glBindBuffer(GL_ARRAY_BUFFER, gl_vbo);
 	glBufferData(GL_ARRAY_BUFFER, gl_vbo_data_size, NULL, GL_STREAM_DRAW /*GL_DYNAMIC_DRAW vu dans SimpleGL de ATI*/);
 
-	// to be removed : default values
+	on_expose_data.gl_vbo = gl_vbo;
+	on_expose_data.mesh_width = mesh_width;
+	on_expose_data.mesh_height = mesh_height;
+
+	/* to be removed : default values
 	glBufferSubData(GL_ARRAY_BUFFER,0, gl_vbo_data_size, gl_vertex_default_pos);
+	*/
+
+	cl_mem cl_vbo = clCreateFromGLBuffer(cl_ctx, CL_MEM_WRITE_ONLY, gl_vbo, &res);
+	if ( res!=CL_SUCCESS ) {
+		fputs ("Failed to clCreateFromGLBuffer()\n", stderr);
+		return EXIT_FAILURE;
+	}
 
 	glBindBuffer(GL_ARRAY_BUFFER, 0); // Unbind buffer (no more current buffer)
 	gdk_gl_drawable_gl_end (drawable);
 
+	const char *source="\
+__kernel void zero_z(__global float4 * pos, unsigned int width, unsigned int height, float time) { \
+	unsigned int x = get_global_id(0); \
+	unsigned int y = get_global_id(1); \
+  /* calculate uv coordinates */ \
+  float u = x / (float) width; \
+  float v = y / (float) height; \
+	/* Only normalized quaterinons here */ \
+	float w = 1.0f; \
+	/* Calculate the desirated value */ \
+	float z = 0.0f; \
+	/* Write output vertex */\
+	pos[y*width+x] = (float4)(u - 0.5f, v - 0.5f, z, w); \
+}\
+";
+	size_t sourceLen=strlen(source);
+
+	cl_program program = clCreateProgramWithSource(cl_ctx, 1, &source, &sourceLen,&res);
+	if ( res!=CL_SUCCESS ) {
+		fputs("Failed to clCreateProgramWithSource()\n", stderr);
+		return EXIT_FAILURE;
+	}
+
+	res = clBuildProgram(program, 1, &device, "", NULL, NULL);
+	if ( res!=CL_SUCCESS ) {
+		fputs("Failed to clBuildProgram()\n", stderr);
+
+		if(res == CL_BUILD_PROGRAM_FAILURE) {
+			char *buildLog = NULL;
+			size_t buildLogSize = 0;
+			clGetProgramBuildInfo(program,device,CL_PROGRAM_BUILD_LOG,buildLogSize,buildLog,&buildLogSize);
+			buildLog = (char*)malloc(buildLogSize);
+			memset(buildLog, 0, buildLogSize);
+			clGetProgramBuildInfo(program,device,CL_PROGRAM_BUILD_LOG,buildLogSize,buildLog,NULL);
+
+			fputs("\n                   BUILD LOG\n", stderr);
+			fputs("************************************************\n", stderr);
+			fputs(buildLog, stderr);
+			free(buildLog);
+			fputs("\n************************************************\n", stderr);
+/*
+			fputs("\n                    SOURCE\n", stderr);
+			fputs("************************************************\n", stderr);
+			fputs(source, stderr);
+			fputs("\n************************************************\n", stderr);
+*/
+			return EXIT_FAILURE;
+		}
+	}
+
+	cl_kernel kernel = clCreateKernel(program,"zero_z",&res);
+	execKernel(cl_ctx, commandQueue, kernel, mesh_width, mesh_height, group_size, cl_vbo, /*time*/0.0f);
+	clReleaseKernel(kernel);
+
 	gtk_main();
 	return EXIT_SUCCESS;
 }