1.1 --- a/Makefile	Fri Aug 27 18:30:09 2010 +0100
     1.2 +++ b/Makefile	Fri Aug 27 19:00:14 2010 +0100
     1.3 @@ -4,7 +4,7 @@
     1.4  bin = test
     1.5  
     1.6  CXX = g++
     1.7 -CXXFLAGS = -pedantic -Wall -g
     1.8 +CXXFLAGS = -pedantic -Wall -g $(def)
     1.9  LDFLAGS = $(libgl) $(libcl) -lpthread
    1.10  
    1.11  ifeq ($(shell uname -s), Darwin)
    1.12 @@ -13,6 +13,8 @@
    1.13  else
    1.14  	libgl = -lGL -lglut
    1.15  	libcl = -lOpenCL
    1.16 +
    1.17 +	def = -DCLGL_INTEROP
    1.18  endif
    1.19  
    1.20  $(bin): $(obj)

     2.1 --- a/rt.cl	Fri Aug 27 18:30:09 2010 +0100
     2.2 +++ b/rt.cl	Fri Aug 27 19:00:14 2010 +0100
     2.3 @@ -132,7 +132,6 @@
     2.4  	coord.y = idx / img_x;
     2.5  
     2.6  	write_imagef(fb, coord, pixel);
     2.7 -	//fb[idx] = pixel;
     2.8  }
     2.9  
    2.10  float4 shade(struct Ray ray, struct Scene *scn, const struct SurfPoint *sp)
    2.11 @@ -187,13 +186,6 @@
    2.12  
    2.13  		global const struct KDNode *node = scn->kdtree + idx;
    2.14  
    2.15 -		/*if(get_global_id(0) == 0) {
    2.16 -			for(int i=0; i<top+1; i++) {
    2.17 -				printf("   ");
    2.18 -			}
    2.19 -			printf("(%d) idx: %d (%p) num_faces: %d\n", top+1, idx, node, node->num_faces);
    2.20 -		}*/
    2.21 -
    2.22  		if(intersect_aabb(ray, node->aabb)) {
    2.23  			if(node->left == -1) {
    2.24  				// leaf node... check each face in turn and update the nearest intersection as needed
    2.25 @@ -207,9 +199,6 @@
    2.26  				}
    2.27  			} else {
    2.28  				// internal node... recurse to the children
    2.29 -				/*if(get_global_id(0) == 0) {
    2.30 -					printf("pushing %d's children %d and %d\n", idx, node->left, node->right);
    2.31 -				}*/
    2.32  				idxstack[top++] = node->left;
    2.33  				idxstack[top++] = node->right;
    2.34  			}
    2.35 @@ -227,29 +216,6 @@
    2.36  	return true;
    2.37  }
    2.38  
    2.39 -/*bool find_intersection(struct Ray ray, const struct Scene *scn, struct SurfPoint *spres)
    2.40 -{
    2.41 -	struct SurfPoint sp, sp0;
    2.42 -	sp0.t = 1.0;
    2.43 -	sp0.obj = 0;
    2.44 -
    2.45 -	for(int i=0; i<scn->num_faces; i++) {
    2.46 -		if(intersect(ray, scn->faces + i, &sp) && sp.t < sp0.t) {
    2.47 -			sp0 = sp;
    2.48 -		}
    2.49 -	}
    2.50 -
    2.51 -	if(!sp0.obj) {
    2.52 -		return false;
    2.53 -	}
    2.54 -
    2.55 -	if(spres) {
    2.56 -		*spres = sp0;
    2.57 -		spres->mat = scn->matlib[sp0.obj->matid];
    2.58 -	}
    2.59 -	return true;
    2.60 -}*/
    2.61 -
    2.62  bool intersect(struct Ray ray, global const struct Face *face, struct SurfPoint *sp)
    2.63  {
    2.64  	float4 origin = ray.origin;

     3.1 --- a/src/clray.cc	Fri Aug 27 18:30:09 2010 +0100
     3.2 +++ b/src/clray.cc	Fri Aug 27 19:00:14 2010 +0100
     3.3 @@ -135,7 +135,13 @@
     3.4  
     3.5  void cleanup()
     3.6  {
     3.7 +	printf("destroying renderer...\n");
     3.8  	destroy_renderer();
     3.9 +
    3.10 +	printf("shutting down OpenCL...\n");
    3.11 +	destroy_opencl();
    3.12 +
    3.13 +	printf("done cleaning up\n");
    3.14  }
    3.15  
    3.16  static Matrix4x4 mat, inv_mat, inv_trans;

     4.1 --- a/src/ocl.cc	Fri Aug 27 18:30:09 2010 +0100
     4.2 +++ b/src/ocl.cc	Fri Aug 27 19:00:14 2010 +0100
     4.3 @@ -53,6 +53,11 @@
     4.4  		return false;
     4.5  	}
     4.6  
     4.7 +#ifndef CLGL_INTEROP
     4.8 +	cl_context_properties *prop = 0;
     4.9 +
    4.10 +#else
    4.11 +
    4.12  #if defined(__APPLE__)
    4.13  #error "CL/GL context sharing not implemented on MacOSX yet"
    4.14  #elif defined(unix) || defined(__unix__)
    4.15 @@ -72,6 +77,8 @@
    4.16  #error "unknown or unsupported platform"
    4.17  #endif
    4.18  
    4.19 +#endif	/* CLGL_INTEROP */
    4.20 +
    4.21  	if(!(ctx = clCreateContext(prop, 1, &devinf.id, 0, 0, 0))) {
    4.22  		fprintf(stderr, "failed to create opencl context\n");
    4.23  		return false;
    4.24 @@ -84,6 +91,19 @@
    4.25  	return true;
    4.26  }
    4.27  
    4.28 +void destroy_opencl()
    4.29 +{
    4.30 +	if(cmdq) {
    4.31 +		clReleaseCommandQueue(cmdq);
    4.32 +		cmdq = 0;
    4.33 +	}
    4.34 +
    4.35 +	if(ctx) {
    4.36 +		clReleaseContext(ctx);
    4.37 +		ctx = 0;
    4.38 +	}
    4.39 +}
    4.40 +
    4.41  
    4.42  CLMemBuffer *create_mem_buffer(int rdwr, size_t sz, const void *buf)
    4.43  {
    4.44 @@ -238,15 +258,14 @@
    4.45  		clReleaseEvent(wait_event);
    4.46  	}
    4.47  	if(last_event) {
    4.48 +		clWaitForEvents(1, &last_event);
    4.49  		clReleaseEvent(last_event);
    4.50  	}
    4.51  
    4.52  	if(prog) {
    4.53 -
    4.54  		clReleaseProgram(prog);
    4.55  	}
    4.56  	if(kernel) {
    4.57 -
    4.58  		clReleaseKernel(kernel);
    4.59  	}
    4.60  	for(size_t i=0; i<args.size(); i++) {
    4.61 @@ -449,12 +468,6 @@
    4.62  
    4.63  bool CLProgram::run(int dim, ...) const
    4.64  {
    4.65 -	if(!built) {
    4.66 -		if(!((CLProgram*)this)->build()) {
    4.67 -			return false;
    4.68 -		}
    4.69 -	}
    4.70 -
    4.71  	va_list ap;
    4.72  	size_t *global_size = (size_t*)alloca(dim * sizeof *global_size);
    4.73  

     5.1 --- a/src/ocl.h	Fri Aug 27 18:30:09 2010 +0100
     5.2 +++ b/src/ocl.h	Fri Aug 27 19:00:14 2010 +0100
     5.3 @@ -31,6 +31,7 @@
     5.4  
     5.5  
     5.6  bool init_opencl();
     5.7 +void destroy_opencl();
     5.8  
     5.9  CLMemBuffer *create_mem_buffer(int rdwr, size_t sz, const void *buf);
    5.10  CLMemBuffer *create_mem_buffer(int rdwr, unsigned int tex);

     6.1 --- a/src/rt.cc	Fri Aug 27 18:30:09 2010 +0100
     6.2 +++ b/src/rt.cc	Fri Aug 27 19:00:14 2010 +0100
     6.3 @@ -129,6 +129,7 @@
     6.4  
     6.5  	long tm0 = get_msec();
     6.6  
     6.7 +#ifdef CLGL_INTEROP
     6.8  	cl_event ev;
     6.9  	CLMemBuffer *texbuf = prog->get_arg_buffer(KARG_FRAMEBUFFER);
    6.10  
    6.11 @@ -138,20 +139,24 @@
    6.12  
    6.13  	// make sure that we will wait for the acquire to finish before running
    6.14  	prog->set_wait_event(ev);
    6.15 +#endif
    6.16  
    6.17  	if(!prog->run(1, global_size)) {
    6.18  		return false;
    6.19  	}
    6.20  
    6.21 +#ifdef CLGL_INTEROP
    6.22  	if(!release_gl_object(texbuf, &ev)) {
    6.23  		return false;
    6.24  	}
    6.25  	clWaitForEvents(1, &ev);
    6.26 +#endif
    6.27  
    6.28 -	printf("rendered in %ld msec\n", get_msec() - tm0);
    6.29 -
    6.30 -	/*long tm_run = get_msec() - tm0;
    6.31 -
    6.32 +#ifndef CLGL_INTEROP
    6.33 +	/* if we don't compile in CL/GL interoperability support, we need
    6.34 +	 * to copy the output buffer to the OpenGL texture used to displaying
    6.35 +	 * the image.
    6.36 +	 */
    6.37  	CLMemBuffer *mbuf = prog->get_arg_buffer(KARG_FRAMEBUFFER);
    6.38  	void *fb = map_mem_buffer(mbuf, MAP_RD);
    6.39  	if(!fb) {
    6.40 @@ -161,11 +166,9 @@
    6.41  
    6.42  	glTexSubImage2D(GL_TEXTURE_2D, 0, 0, 0, rinf.xsz, rinf.ysz, GL_RGBA, GL_FLOAT, fb);
    6.43  	unmap_mem_buffer(mbuf);
    6.44 +#endif
    6.45  
    6.46 -	long tm_upd = get_msec() - tm0 - tm_run;
    6.47 -
    6.48 -	printf("render %ld msec (%ld run, %ld upd)\n", tm_run + tm_upd, tm_run, tm_upd);
    6.49 -	*/
    6.50 +	printf("rendered in %ld msec\n", get_msec() - tm0);
    6.51  	return true;
    6.52  }
    6.53