summaryrefslogtreecommitdiff
path: root/tests/poc_c
diff options
context:
space:
mode:
authorLudovic Pouzenc <ludovic@pouzenc.fr>2013-02-23 22:22:41 +0000
committerLudovic Pouzenc <ludovic@pouzenc.fr>2013-02-23 22:22:41 +0000
commitb47a947b3b5587ae496d93fecb8ff3345225c905 (patch)
tree5d467dc2ab20faf513dfc7d3d3a7c8c5c850747e /tests/poc_c
parentaa769549e1268c67f3dbd62191a4071aead0f877 (diff)
download2013-gpudataviz-b47a947b3b5587ae496d93fecb8ff3345225c905.tar.gz
2013-gpudataviz-b47a947b3b5587ae496d93fecb8ff3345225c905.tar.bz2
2013-gpudataviz-b47a947b3b5587ae496d93fecb8ff3345225c905.zip
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
Diffstat (limited to 'tests/poc_c')
-rw-r--r--tests/poc_c/main.c190
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;
}