# HG changeset patch
# User John Tsiombikas <nuclear@member.fsf.org>
# Date 1282932014 -3600
# Node ID 1bcbb53b3505ecb6f65e11d47495684c69d0632c
# Parent  980bc07be868dbf36e80aa605a8677f94850e1a4
segfault on exit?

diff -r 980bc07be868 -r 1bcbb53b3505 Makefile
--- a/Makefile	Fri Aug 27 18:30:09 2010 +0100
+++ b/Makefile	Fri Aug 27 19:00:14 2010 +0100
@@ -4,7 +4,7 @@
 bin = test
 
 CXX = g++
-CXXFLAGS = -pedantic -Wall -g
+CXXFLAGS = -pedantic -Wall -g $(def)
 LDFLAGS = $(libgl) $(libcl) -lpthread
 
 ifeq ($(shell uname -s), Darwin)
@@ -13,6 +13,8 @@
 else
 	libgl = -lGL -lglut
 	libcl = -lOpenCL
+
+	def = -DCLGL_INTEROP
 endif
 
 $(bin): $(obj)
diff -r 980bc07be868 -r 1bcbb53b3505 rt.cl
--- a/rt.cl	Fri Aug 27 18:30:09 2010 +0100
+++ b/rt.cl	Fri Aug 27 19:00:14 2010 +0100
@@ -132,7 +132,6 @@
 	coord.y = idx / img_x;
 
 	write_imagef(fb, coord, pixel);
-	//fb[idx] = pixel;
 }
 
 float4 shade(struct Ray ray, struct Scene *scn, const struct SurfPoint *sp)
@@ -187,13 +186,6 @@
 
 		global const struct KDNode *node = scn->kdtree + idx;
 
-		/*if(get_global_id(0) == 0) {
-			for(int i=0; i<top+1; i++) {
-				printf("   ");
-			}
-			printf("(%d) idx: %d (%p) num_faces: %d\n", top+1, idx, node, node->num_faces);
-		}*/
-
 		if(intersect_aabb(ray, node->aabb)) {
 			if(node->left == -1) {
 				// leaf node... check each face in turn and update the nearest intersection as needed
@@ -207,9 +199,6 @@
 				}
 			} else {
 				// internal node... recurse to the children
-				/*if(get_global_id(0) == 0) {
-					printf("pushing %d's children %d and %d\n", idx, node->left, node->right);
-				}*/
 				idxstack[top++] = node->left;
 				idxstack[top++] = node->right;
 			}
@@ -227,29 +216,6 @@
 	return true;
 }
 
-/*bool find_intersection(struct Ray ray, const struct Scene *scn, struct SurfPoint *spres)
-{
-	struct SurfPoint sp, sp0;
-	sp0.t = 1.0;
-	sp0.obj = 0;
-
-	for(int i=0; i<scn->num_faces; i++) {
-		if(intersect(ray, scn->faces + i, &sp) && sp.t < sp0.t) {
-			sp0 = sp;
-		}
-	}
-
-	if(!sp0.obj) {
-		return false;
-	}
-
-	if(spres) {
-		*spres = sp0;
-		spres->mat = scn->matlib[sp0.obj->matid];
-	}
-	return true;
-}*/
-
 bool intersect(struct Ray ray, global const struct Face *face, struct SurfPoint *sp)
 {
 	float4 origin = ray.origin;
diff -r 980bc07be868 -r 1bcbb53b3505 src/clray.cc
--- a/src/clray.cc	Fri Aug 27 18:30:09 2010 +0100
+++ b/src/clray.cc	Fri Aug 27 19:00:14 2010 +0100
@@ -135,7 +135,13 @@
 
 void cleanup()
 {
+	printf("destroying renderer...\n");
 	destroy_renderer();
+
+	printf("shutting down OpenCL...\n");
+	destroy_opencl();
+
+	printf("done cleaning up\n");
 }
 
 static Matrix4x4 mat, inv_mat, inv_trans;
diff -r 980bc07be868 -r 1bcbb53b3505 src/ocl.cc
--- a/src/ocl.cc	Fri Aug 27 18:30:09 2010 +0100
+++ b/src/ocl.cc	Fri Aug 27 19:00:14 2010 +0100
@@ -53,6 +53,11 @@
 		return false;
 	}
 
+#ifndef CLGL_INTEROP
+	cl_context_properties *prop = 0;
+
+#else
+
 #if defined(__APPLE__)
 #error "CL/GL context sharing not implemented on MacOSX yet"
 #elif defined(unix) || defined(__unix__)
@@ -72,6 +77,8 @@
 #error "unknown or unsupported platform"
 #endif
 
+#endif	/* CLGL_INTEROP */
+
 	if(!(ctx = clCreateContext(prop, 1, &devinf.id, 0, 0, 0))) {
 		fprintf(stderr, "failed to create opencl context\n");
 		return false;
@@ -84,6 +91,19 @@
 	return true;
 }
 
+void destroy_opencl()
+{
+	if(cmdq) {
+		clReleaseCommandQueue(cmdq);
+		cmdq = 0;
+	}
+
+	if(ctx) {
+		clReleaseContext(ctx);
+		ctx = 0;
+	}
+}
+
 
 CLMemBuffer *create_mem_buffer(int rdwr, size_t sz, const void *buf)
 {
@@ -238,15 +258,14 @@
 		clReleaseEvent(wait_event);
 	}
 	if(last_event) {
+		clWaitForEvents(1, &last_event);
 		clReleaseEvent(last_event);
 	}
 
 	if(prog) {
-
 		clReleaseProgram(prog);
 	}
 	if(kernel) {
-
 		clReleaseKernel(kernel);
 	}
 	for(size_t i=0; i<args.size(); i++) {
@@ -449,12 +468,6 @@
 
 bool CLProgram::run(int dim, ...) const
 {
-	if(!built) {
-		if(!((CLProgram*)this)->build()) {
-			return false;
-		}
-	}
-
 	va_list ap;
 	size_t *global_size = (size_t*)alloca(dim * sizeof *global_size);
 
diff -r 980bc07be868 -r 1bcbb53b3505 src/ocl.h
--- a/src/ocl.h	Fri Aug 27 18:30:09 2010 +0100
+++ b/src/ocl.h	Fri Aug 27 19:00:14 2010 +0100
@@ -31,6 +31,7 @@
 
 
 bool init_opencl();
+void destroy_opencl();
 
 CLMemBuffer *create_mem_buffer(int rdwr, size_t sz, const void *buf);
 CLMemBuffer *create_mem_buffer(int rdwr, unsigned int tex);
diff -r 980bc07be868 -r 1bcbb53b3505 src/rt.cc
--- a/src/rt.cc	Fri Aug 27 18:30:09 2010 +0100
+++ b/src/rt.cc	Fri Aug 27 19:00:14 2010 +0100
@@ -129,6 +129,7 @@
 
 	long tm0 = get_msec();
 
+#ifdef CLGL_INTEROP
 	cl_event ev;
 	CLMemBuffer *texbuf = prog->get_arg_buffer(KARG_FRAMEBUFFER);
 
@@ -138,20 +139,24 @@
 
 	// make sure that we will wait for the acquire to finish before running
 	prog->set_wait_event(ev);
+#endif
 
 	if(!prog->run(1, global_size)) {
 		return false;
 	}
 
+#ifdef CLGL_INTEROP
 	if(!release_gl_object(texbuf, &ev)) {
 		return false;
 	}
 	clWaitForEvents(1, &ev);
+#endif
 
-	printf("rendered in %ld msec\n", get_msec() - tm0);
-
-	/*long tm_run = get_msec() - tm0;
-
+#ifndef CLGL_INTEROP
+	/* if we don't compile in CL/GL interoperability support, we need
+	 * to copy the output buffer to the OpenGL texture used to displaying
+	 * the image.
+	 */
 	CLMemBuffer *mbuf = prog->get_arg_buffer(KARG_FRAMEBUFFER);
 	void *fb = map_mem_buffer(mbuf, MAP_RD);
 	if(!fb) {
@@ -161,11 +166,9 @@
 
 	glTexSubImage2D(GL_TEXTURE_2D, 0, 0, 0, rinf.xsz, rinf.ysz, GL_RGBA, GL_FLOAT, fb);
 	unmap_mem_buffer(mbuf);
+#endif
 
-	long tm_upd = get_msec() - tm0 - tm_run;
-
-	printf("render %ld msec (%ld run, %ld upd)\n", tm_run + tm_upd, tm_run, tm_upd);
-	*/
+	printf("rendered in %ld msec\n", get_msec() - tm0);
 	return true;
 }