1.1 --- a/Makefile	Fri Jul 23 01:22:03 2010 +0100
     1.2 +++ b/Makefile	Fri Jul 23 19:48:43 2010 +0100
     1.3 @@ -4,7 +4,15 @@
     1.4  
     1.5  CXX = g++
     1.6  CXXFLAGS = -pedantic -Wall -g
     1.7 -LDFLAGS = -framework OpenGL -framework GLUT -framework OpenCL
     1.8 +LDFLAGS = $(libgl) $(libcl)
     1.9 +
    1.10 +ifeq ($(shell uname -s), Darwin)
    1.11 +	libgl = -framework OpenGL -framework GLUT
    1.12 +	libcl = -framework OpenCL
    1.13 +else
    1.14 +	libgl = -lGL -lglut
    1.15 +	libcl = -lOpenCL
    1.16 +endif
    1.17  
    1.18  $(bin): $(obj)
    1.19  	$(CXX) -o $@ $(obj) $(LDFLAGS)

     2.1 --- a/rt.cl	Fri Jul 23 01:22:03 2010 +0100
     2.2 +++ b/rt.cl	Fri Jul 23 19:48:43 2010 +0100
     2.3 @@ -21,20 +21,18 @@
     2.4  
     2.5  struct SurfPoint {
     2.6  	float t;
     2.7 -	float3 pos, norm;
     2.8 +	float4 pos, norm;
     2.9  	global const struct Sphere *obj;
    2.10  };
    2.11  
    2.12 -struct Matrix4x4 {
    2.13 -	float m[16];
    2.14 -};
    2.15 -
    2.16  #define EPSILON 1e-6
    2.17  
    2.18  float4 shade(struct Ray ray, struct SurfPoint sp,
    2.19  		global const struct Light *lights, int num_lights);
    2.20  bool intersect(struct Ray ray, global const struct Sphere *sph, struct SurfPoint *sp);
    2.21 -float3 reflect(float3 v, float3 n);
    2.22 +float4 reflect(float4 v, float4 n);
    2.23 +float4 transform(float4 v, global const float *xform);
    2.24 +struct Ray transform_ray(global const struct Ray *ray, global const float *xform);
    2.25  
    2.26  
    2.27  kernel void render(global float4 *fb,
    2.28 @@ -42,13 +40,13 @@
    2.29  		global const struct Sphere *sphlist,
    2.30  		global const struct Light *lights,
    2.31  		global const struct Ray *primrays,
    2.32 -		global const struct Matrix4x4 xform)
    2.33 +		global const float *xform)
    2.34  {
    2.35  	int idx = get_global_id(0);
    2.36  
    2.37 -	struct Ray ray = primrays[idx];
    2.38 +	struct Ray ray = transform_ray(primrays + idx, xform);
    2.39 +
    2.40  	struct SurfPoint sp, sp0;
    2.41 -
    2.42  	sp0.t = FLT_MAX;
    2.43  	sp0.obj = 0;
    2.44  
    2.45 @@ -68,31 +66,31 @@
    2.46  float4 shade(struct Ray ray, struct SurfPoint sp,
    2.47  		global const struct Light *lights, int num_lights)
    2.48  {
    2.49 -	float3 dcol = (float3)(0, 0, 0);
    2.50 -	float3 scol = (float3)(0, 0, 0);
    2.51 +	float4 dcol = (float4)(0, 0, 0, 0);
    2.52 +	float4 scol = (float4)(0, 0, 0, 0);
    2.53  
    2.54  	for(int i=0; i<num_lights; i++) {
    2.55 -		float3 ldir = normalize(lights[i].pos.xyz - sp.pos);
    2.56 -		float3 vdir = -normalize(ray.dir.xyz);
    2.57 -		float3 vref = reflect(vdir, sp.norm);
    2.58 +		float4 ldir = normalize(lights[i].pos - sp.pos);
    2.59 +		float4 vdir = -normalize(ray.dir);
    2.60 +		float4 vref = reflect(vdir, sp.norm);
    2.61  
    2.62  		float diff = fmax(dot(ldir, sp.norm), 0.0f);
    2.63  		float spec = powr(fmax(dot(ldir, vref), 0.0f), sp.obj->spow);
    2.64  
    2.65 -		dcol += sp.obj->kd.xyz * diff * lights[i].color.xyz;
    2.66 -		scol += sp.obj->ks.xyz * spec * lights[i].color.xyz;
    2.67 +		dcol += sp.obj->kd * diff * lights[i].color;
    2.68 +		scol += sp.obj->ks * spec * lights[i].color;
    2.69  	}
    2.70  
    2.71 -	return (float4)(dcol + scol, 1.0f);
    2.72 +	return dcol + scol;
    2.73  }
    2.74  
    2.75  bool intersect(struct Ray ray,
    2.76  		global const struct Sphere *sph,
    2.77  		struct SurfPoint *sp)
    2.78  {
    2.79 -	float3 dir = ray.dir.xyz;
    2.80 -	float3 orig = ray.origin.xyz;
    2.81 -	float3 spos = sph->pos.xyz;
    2.82 +	float4 dir = ray.dir;
    2.83 +	float4 orig = ray.origin;
    2.84 +	float4 spos = sph->pos;
    2.85  
    2.86  	float a = dot(dir, dir);
    2.87  	float b = 2.0 * dir.x * (orig.x - spos.x) +
    2.88 @@ -123,7 +121,33 @@
    2.89  	return true;
    2.90  }
    2.91  
    2.92 -float3 reflect(float3 v, float3 n)
    2.93 +float4 reflect(float4 v, float4 n)
    2.94  {
    2.95  	return 2.0f * dot(v, n) * n - v;
    2.96  }
    2.97 +
    2.98 +float4 transform(float4 v, global const float *xform)
    2.99 +{
   2.100 +	float4 res;
   2.101 +	res.x = v.x * xform[0] + v.y * xform[4] + v.z * xform[8] + xform[12];
   2.102 +	res.y = v.x * xform[1] + v.y * xform[5] + v.z * xform[9] + xform[13];
   2.103 +	res.z = v.x * xform[2] + v.y * xform[6] + v.z * xform[10] + xform[14];
   2.104 +	res.w = 1.0;
   2.105 +	return res;
   2.106 +}
   2.107 +
   2.108 +struct Ray transform_ray(global const struct Ray *ray, global const float *xform)
   2.109 +{
   2.110 +	struct Ray res;
   2.111 +	float rot[16];
   2.112 +
   2.113 +	for(int i=0; i<16; i++) {
   2.114 +		rot[i] = xform[i];
   2.115 +	}
   2.116 +	rot[3] = rot[7] = rot[11] = rot[12] = rot[13] = rot[14] = 0.0f;
   2.117 +	rot[15] = 1.0f;
   2.118 +
   2.119 +	res.origin = transform(ray->origin, xform);
   2.120 +	res.dir = transform(ray->dir, xform);
   2.121 +	return res;
   2.122 +}

     3.1 --- a/src/clray.cc	Fri Jul 23 01:22:03 2010 +0100
     3.2 +++ b/src/clray.cc	Fri Jul 23 19:48:43 2010 +0100
     3.3 @@ -21,6 +21,9 @@
     3.4  static int xsz, ysz;
     3.5  static bool need_update = true;
     3.6  
     3.7 +static float cam_theta, cam_phi = 25.0;
     3.8 +static float cam_dist = 10.0;
     3.9 +
    3.10  int main(int argc, char **argv)
    3.11  {
    3.12  	glutInitWindowSize(800, 600);
    3.13 @@ -63,7 +66,21 @@
    3.14  
    3.15  void disp()
    3.16  {
    3.17 +	glMatrixMode(GL_MODELVIEW);
    3.18 +	glLoadIdentity();
    3.19 +
    3.20  	if(need_update) {
    3.21 +		float mat[16];
    3.22 +
    3.23 +		glPushMatrix();
    3.24 +		glRotatef(cam_theta, 0, 1, 0);
    3.25 +		glRotatef(cam_phi, 1, 0, 0);
    3.26 +		glTranslatef(0, 0, -cam_dist);
    3.27 +
    3.28 +		glGetFloatv(GL_MODELVIEW_MATRIX, mat);
    3.29 +		set_xform(mat);
    3.30 +		glPopMatrix();
    3.31 +
    3.32  		render();
    3.33  		glTexSubImage2D(GL_TEXTURE_2D, 0, 0, 0, xsz, ysz, GL_RGBA, GL_FLOAT, fb);
    3.34  		need_update = false;
    3.35 @@ -115,12 +132,47 @@
    3.36  	}
    3.37  }
    3.38  
    3.39 +static bool bnstate[32];
    3.40 +static int prev_x, prev_y;
    3.41 +
    3.42  void mouse(int bn, int state, int x, int y)
    3.43  {
    3.44 +	if(state == GLUT_DOWN) {
    3.45 +		prev_x = x;
    3.46 +		prev_y = y;
    3.47 +		bnstate[bn] = true;
    3.48 +	} else {
    3.49 +		bnstate[bn] = false;
    3.50 +	}
    3.51  }
    3.52  
    3.53 +#define ROT_SCALE	0.5
    3.54 +#define PAN_SCALE	0.1
    3.55 +
    3.56  void motion(int x, int y)
    3.57  {
    3.58 +	int dx = x - prev_x;
    3.59 +	int dy = y - prev_y;
    3.60 +	prev_x = x;
    3.61 +	prev_y = y;
    3.62 +
    3.63 +	if(bnstate[0]) {
    3.64 +		cam_theta += dx * ROT_SCALE;
    3.65 +		cam_phi += dy * ROT_SCALE;
    3.66 +
    3.67 +		if(cam_phi < -89) cam_phi = 89;
    3.68 +		if(cam_phi > 89) cam_phi = 89;
    3.69 +
    3.70 +		need_update = true;
    3.71 +		glutPostRedisplay();
    3.72 +	}
    3.73 +	if(bnstate[2]) {
    3.74 +		cam_dist += dy * PAN_SCALE;
    3.75 +		if(cam_dist < 0) cam_dist = 0;
    3.76 +
    3.77 +		need_update = true;
    3.78 +		glutPostRedisplay();
    3.79 +	}
    3.80  }
    3.81  
    3.82  bool write_ppm(const char *fname, float *fb, int xsz, int ysz)

     4.1 --- a/src/mesh.cc	Fri Jul 23 01:22:03 2010 +0100
     4.2 +++ b/src/mesh.cc	Fri Jul 23 19:48:43 2010 +0100
     4.3 @@ -2,6 +2,7 @@
     4.4  #include <stdlib.h>
     4.5  #include <string.h>
     4.6  #include <errno.h>
     4.7 +#include <limits.h>
     4.8  #include <string>
     4.9  #include <vector>
    4.10  #include <map>
    4.11 @@ -9,7 +10,7 @@
    4.12  
    4.13  using namespace std;
    4.14  
    4.15 -#define COMMANDS 	\
    4.16 +#define COMMANDS	\
    4.17  	CMD(V),			\
    4.18  	CMD(VN),		\
    4.19  	CMD(VT),		\
    4.20 @@ -136,7 +137,7 @@
    4.21  	char cur_name[16];
    4.22  
    4.23  	obj_file obj;
    4.24 -	
    4.25 +
    4.26  	sprintf(cur_name, "default%02d.obj", seq++);
    4.27  	obj.cur_obj = cur_name;
    4.28  
    4.29 @@ -144,7 +145,7 @@
    4.30  	for(;;) {
    4.31  		Vector3 vec;
    4.32  		obj_face face;
    4.33 -		
    4.34 +
    4.35  		char line[BUF_SZ];
    4.36  		fgets(line, sizeof line, fp);
    4.37  		if(feof(fp)) {
    4.38 @@ -246,7 +247,7 @@
    4.39  					mat.kt = 1.0 - vmtl[i].alpha;
    4.40  					mat.kr = 0.0;	// TODO
    4.41  					mat.spow = vmtl[i].shininess;
    4.42 -					
    4.43 +
    4.44  					matnames[vmtl[i].name] = i;
    4.45  				}
    4.46  			}
    4.47 @@ -356,7 +357,7 @@
    4.48  	if(added_tc) {
    4.49  		obj->vt.pop_back();
    4.50  	}
    4.51 -	
    4.52 +
    4.53  	return mesh;
    4.54  }
    4.55  
    4.56 @@ -445,7 +446,7 @@
    4.57  static int get_cmd(char *str)
    4.58  {
    4.59  	char *s = str;
    4.60 -	while((*s = toupper(*s))) s++; 
    4.61 +	while((*s = toupper(*s))) s++;
    4.62  
    4.63  	for(int i=0; cmd_names[i]; i++) {
    4.64  		if(strcmp(str, cmd_names[i]) == 0) {

     5.1 --- a/src/ocl.cc	Fri Jul 23 01:22:03 2010 +0100
     5.2 +++ b/src/ocl.cc	Fri Jul 23 19:48:43 2010 +0100
     5.3 @@ -1,10 +1,14 @@
     5.4 +#define OCL_CC_
     5.5 +
     5.6  #include <stdio.h>
     5.7  #include <stdlib.h>
     5.8  #include <string.h>
     5.9 +#include <stdarg.h>
    5.10  #include <errno.h>
    5.11  #include <alloca.h>
    5.12  #include <sys/stat.h>
    5.13  #include "ocl.h"
    5.14 +#include "ocl_errstr.h"
    5.15  
    5.16  
    5.17  class InitCL {
    5.18 @@ -31,6 +35,7 @@
    5.19  static int devcmp(struct device_info *a, struct device_info *b);
    5.20  static const char *devtypestr(cl_device_type type);
    5.21  static void print_memsize(FILE *out, unsigned long memsz);
    5.22 +static const char *clstrerror(int err);
    5.23  
    5.24  
    5.25  static InitCL initcl;
    5.26 @@ -72,7 +77,7 @@
    5.27  
    5.28  
    5.29  	if(!(mem = clCreateBuffer(ctx, rdwr | CL_MEM_USE_HOST_PTR, sz, buf, &err))) {
    5.30 -		fprintf(stderr, "failed to create memory buffer (%d)\n", err);
    5.31 +		fprintf(stderr, "failed to create memory buffer: %s\n", clstrerror(err));
    5.32  		return 0;
    5.33  	}
    5.34  
    5.35 @@ -98,7 +103,7 @@
    5.36  	int err;
    5.37  	mbuf->ptr = clEnqueueMapBuffer(cmdq, mbuf->mem, 1, rdwr, 0, mbuf->size, 0, 0, 0, &err);
    5.38  	if(!mbuf->ptr) {
    5.39 -		fprintf(stderr, "failed to map buffer (%d)\n", err);
    5.40 +		fprintf(stderr, "failed to map buffer: %s\n", clstrerror(err));
    5.41  		return 0;
    5.42  	}
    5.43  	return mbuf->ptr;
    5.44 @@ -116,7 +121,7 @@
    5.45  
    5.46  	int err;
    5.47  	if((err = clEnqueueWriteBuffer(cmdq, mbuf->mem, 1, 0, sz, src, 0, 0, 0)) != 0) {
    5.48 -		fprintf(stderr, "failed to write buffer (%d)\n", err);
    5.49 +		fprintf(stderr, "failed to write buffer: %s\n", clstrerror(err));
    5.50  		return false;
    5.51  	}
    5.52  	return true;
    5.53 @@ -128,7 +133,7 @@
    5.54  
    5.55  	int err;
    5.56  	if((err = clEnqueueReadBuffer(cmdq, mbuf->mem, 1, 0, sz, dest, 0, 0, 0)) != 0) {
    5.57 -		fprintf(stderr, "failed to read buffer (%d)\n", err);
    5.58 +		fprintf(stderr, "failed to read buffer: %s\n", clstrerror(err));
    5.59  		return false;
    5.60  	}
    5.61  	return true;
    5.62 @@ -250,7 +255,7 @@
    5.63  
    5.64  		char *errlog = (char*)alloca(sz + 1);
    5.65  		clGetProgramBuildInfo(prog, devinf.id, CL_PROGRAM_BUILD_LOG, sz, errlog, 0);
    5.66 -		fprintf(stderr, "failed to build program: (%d)\n%s\n", err, errlog);
    5.67 +		fprintf(stderr, "failed to build program: %s\n%s\n", clstrerror(err), errlog);
    5.68  
    5.69  		clReleaseProgram(prog);
    5.70  		prog = 0;
    5.71 @@ -275,14 +280,14 @@
    5.72  		switch(args[i].type) {
    5.73  		case ARGTYPE_INT:
    5.74  			if((err = clSetKernelArg(kernel, i, sizeof(int), &args[i].v.ival)) != 0) {
    5.75 -				fprintf(stderr, "failed to bind kernel argument: %d (%d)\n", (int)i, err);
    5.76 +				fprintf(stderr, "failed to bind kernel argument %d: %s\n", (int)i, clstrerror(err));
    5.77  				goto fail;
    5.78  			}
    5.79  			break;
    5.80  
    5.81  		case ARGTYPE_FLOAT:
    5.82  			if((err = clSetKernelArg(kernel, i, sizeof(float), &args[i].v.fval)) != 0) {
    5.83 -				fprintf(stderr, "failed to bind kernel argument: %d (%d)\n", (int)i, err);
    5.84 +				fprintf(stderr, "failed to bind kernel argument %d: %s\n", (int)i, clstrerror(err));
    5.85  				goto fail;
    5.86  			}
    5.87  			break;
    5.88 @@ -292,7 +297,7 @@
    5.89  				CLMemBuffer *mbuf = args[i].v.mbuf;
    5.90  
    5.91  				if((err = clSetKernelArg(kernel, i, sizeof mbuf->mem, &mbuf->mem)) != 0) {
    5.92 -					fprintf(stderr, "failed to bind kernel argument: %d (%d)\n", (int)i, err);
    5.93 +					fprintf(stderr, "failed to bind kernel argument %d: %s\n", (int)i, clstrerror(err));
    5.94  					goto fail;
    5.95  				}
    5.96  			}
    5.97 @@ -338,7 +343,7 @@
    5.98  
    5.99  	int err;
   5.100  	if((err = clEnqueueNDRangeKernel(cmdq, kernel, dim, 0, global_size, 0, 0, 0, 0)) != 0) {
   5.101 -		fprintf(stderr, "error executing kernel (%d)\n", err);
   5.102 +		fprintf(stderr, "error executing kernel: %s\n", clstrerror(err));
   5.103  		return false;
   5.104  	}
   5.105  	return true;
   5.106 @@ -346,13 +351,36 @@
   5.107  
   5.108  static int select_device(struct device_info *dev_inf, int (*devcmp)(struct device_info*, struct device_info*))
   5.109  {
   5.110 -	unsigned int i, j, num_dev, sel;
   5.111 +	unsigned int i, j, num_dev, num_plat, sel, ret;
   5.112  	cl_device_id dev[32];
   5.113 +	cl_platform_id plat[32];
   5.114  
   5.115  	dev_inf->work_item_sizes = 0;
   5.116  
   5.117 +	if((ret = clGetPlatformIDs(32, plat, &num_plat)) != 0) {
   5.118 +		fprintf(stderr, "clGetPlatformIDs failed: %s\n", clstrerror(ret));
   5.119 +		return -1;
   5.120 +	}
   5.121 +	if(!num_plat) {
   5.122 +		fprintf(stderr, "OpenCL not available!\n");
   5.123 +		return -1;
   5.124 +	}
   5.125  
   5.126 -	clGetDeviceIDs(0, CL_DEVICE_TYPE_ALL, 32, dev, &num_dev);
   5.127 +	for(i=0; i<num_plat; i++) {
   5.128 +		char buf[512];
   5.129 +
   5.130 +		clGetPlatformInfo(plat[i], CL_PLATFORM_NAME, sizeof buf, buf, 0);
   5.131 +		printf("[%d]: %s", i, buf);
   5.132 +		clGetPlatformInfo(plat[i], CL_PLATFORM_VENDOR, sizeof buf, buf, 0);
   5.133 +		printf(", %s", buf);
   5.134 +		clGetPlatformInfo(plat[i], CL_PLATFORM_VERSION, sizeof buf, buf, 0);
   5.135 +		printf(" (%s)\n", buf);
   5.136 +	}
   5.137 +
   5.138 +	if((ret = clGetDeviceIDs(plat[0], CL_DEVICE_TYPE_ALL, 32, dev, &num_dev)) != 0) {
   5.139 +		fprintf(stderr, "clGetDeviceIDs failed: %s\n", clstrerror(ret));
   5.140 +		return -1;
   5.141 +	}
   5.142  	printf("found %d cl devices.\n", num_dev);
   5.143  
   5.144  	for(i=0; i<num_dev; i++) {
   5.145 @@ -457,3 +485,14 @@
   5.146  		memsz /= 1024;
   5.147  	}
   5.148  }
   5.149 +
   5.150 +static const char *clstrerror(int err)
   5.151 +{
   5.152 +	if(err > 0) {
   5.153 +		return "<invalid error code>";
   5.154 +	}
   5.155 +	if(err <= -(int)(sizeof ocl_errstr / sizeof *ocl_errstr)) {
   5.156 +		return "<unknown error>";
   5.157 +	}
   5.158 +	return ocl_errstr[-err];
   5.159 +}

     6.1 --- a/src/ocl.h	Fri Jul 23 01:22:03 2010 +0100
     6.2 +++ b/src/ocl.h	Fri Jul 23 19:48:43 2010 +0100
     6.3 @@ -4,7 +4,7 @@
     6.4  #include <vector>
     6.5  #include <string>
     6.6  #ifndef __APPLE__
     6.7 -#include <CL/opencl.h>
     6.8 +#include <CL/cl.h>
     6.9  #else
    6.10  #include <OpenCL/opencl.h>
    6.11  #endif

     7.1 --- /dev/null	Thu Jan 01 00:00:00 1970 +0000
     7.2 +++ b/src/ocl_errstr.h	Fri Jul 23 19:48:43 2010 +0100
     7.3 @@ -0,0 +1,64 @@
     7.4 +#ifndef OCL_CC_
     7.5 +#error "don't include ocl_errstr.h"
     7.6 +#endif
     7.7 +
     7.8 +static const char *ocl_errstr[] = {
     7.9 +	"CL_SUCCESS",
    7.10 +	"CL_DEVICE_NOT_FOUND",
    7.11 +	"CL_DEVICE_NOT_AVAILABLE",
    7.12 +	"CL_COMPILER_NOT_AVAILABLE",
    7.13 +	"CL_MEM_OBJECT_ALLOCATION_FAILURE",
    7.14 +	"CL_OUT_OF_RESOURCES",
    7.15 +	"CL_OUT_OF_HOST_MEMORY",
    7.16 +	"CL_PROFILING_INFO_NOT_AVAILABLE",
    7.17 +	"CL_MEM_COPY_OVERLAP",
    7.18 +	"CL_IMAGE_FORMAT_MISMATCH",
    7.19 +	"CL_IMAGE_FORMAT_NOT_SUPPORTED",
    7.20 +	"CL_BUILD_PROGRAM_FAILURE",
    7.21 +	"CL_MAP_FAILURE",						/* 12 */
    7.22 +
    7.23 +	"<unknown error>", "<unknown error>",	/* 13, 14 */
    7.24 +	"<unknown error>", "<unknown error>",	/* 15, 16 */
    7.25 +	"<unknown error>", "<unknown error>",	/* 17, 18 */
    7.26 +	"<unknown error>", "<unknown error>",	/* 19, 20 */
    7.27 +	"<unknown error>", "<unknown error>",	/* 21, 22 */
    7.28 +	"<unknown error>", "<unknown error>",	/* 23, 24 */
    7.29 +	"<unknown error>", "<unknown error>",	/* 25, 26 */
    7.30 +	"<unknown error>", "<unknown error>",	/* 27, 28 */
    7.31 +	"<unknown error>",						/* 29 */
    7.32 +
    7.33 +	"CL_INVALID_VALUE",						/* 30 */
    7.34 +	"CL_INVALID_DEVICE_TYPE",
    7.35 +	"CL_INVALID_PLATFORM",
    7.36 +	"CL_INVALID_DEVICE",
    7.37 +	"CL_INVALID_CONTEXT",
    7.38 +	"CL_INVALID_QUEUE_PROPERTIES",
    7.39 +	"CL_INVALID_COMMAND_QUEUE",
    7.40 +	"CL_INVALID_HOST_PTR",
    7.41 +	"CL_INVALID_MEM_OBJECT",
    7.42 +	"CL_INVALID_IMAGE_FORMAT_DESCRIPTOR",
    7.43 +	"CL_INVALID_IMAGE_SIZE",
    7.44 +	"CL_INVALID_SAMPLER",
    7.45 +	"CL_INVALID_BINARY",
    7.46 +	"CL_INVALID_BUILD_OPTIONS",
    7.47 +	"CL_INVALID_PROGRAM",
    7.48 +	"CL_INVALID_PROGRAM_EXECUTABLE",
    7.49 +	"CL_INVALID_KERNEL_NAME",
    7.50 +	"CL_INVALID_KERNEL_DEFINITION",
    7.51 +	"CL_INVALID_KERNEL",
    7.52 +	"CL_INVALID_ARG_INDEX",
    7.53 +	"CL_INVALID_ARG_VALUE",
    7.54 +	"CL_INVALID_ARG_SIZE",
    7.55 +	"CL_INVALID_KERNEL_ARGS",
    7.56 +	"CL_INVALID_WORK_DIMENSION",
    7.57 +	"CL_INVALID_WORK_GROUP_SIZE",
    7.58 +	"CL_INVALID_WORK_ITEM_SIZE",
    7.59 +	"CL_INVALID_GLOBAL_OFFSET",
    7.60 +	"CL_INVALID_EVENT_WAIT_LIST",
    7.61 +	"CL_INVALID_EVENT",
    7.62 +	"CL_INVALID_OPERATION",
    7.63 +	"CL_INVALID_GL_OBJECT",
    7.64 +	"CL_INVALID_BUFFER_SIZE",
    7.65 +	"CL_INVALID_MIP_LEVEL",
    7.66 +	"CL_INVALID_GLOBAL_WORK_SIZE"
    7.67 +};

     8.1 --- a/src/rt.cc	Fri Jul 23 01:22:03 2010 +0100
     8.2 +++ b/src/rt.cc	Fri Jul 23 19:48:43 2010 +0100
     8.3 @@ -1,4 +1,5 @@
     8.4  #include <stdio.h>
     8.5 +#include <string.h>
     8.6  #include <math.h>
     8.7  #include <assert.h>
     8.8  #include "ocl.h"
     8.9 @@ -10,25 +11,21 @@
    8.10  } __attribute__((packed));
    8.11  
    8.12  struct Sphere {
    8.13 -	cl_float4 pos;
    8.14 -	cl_float4 kd, ks;
    8.15 -	cl_float radius;
    8.16 -	cl_float spow;
    8.17 -	cl_float kr, kt;
    8.18 +	float pos[4];
    8.19 +	float kd[4], ks[4];
    8.20 +	float radius;
    8.21 +	float spow;
    8.22 +	float kr, kt;
    8.23  } __attribute__((packed));
    8.24  
    8.25  struct Ray {
    8.26 -	cl_float4 origin, dir;
    8.27 +	float origin[4], dir[4];
    8.28  } __attribute__((packed));
    8.29  
    8.30  struct Light {
    8.31 -	cl_float4 pos, color;
    8.32 +	float pos[4], color[4];
    8.33  } __attribute__((packed));
    8.34  
    8.35 -struct Matrix4x4 {
    8.36 -	cl_float m[16];
    8.37 -};
    8.38 -
    8.39  static Ray get_primary_ray(int x, int y, int w, int h, float vfov_deg);
    8.40  
    8.41  static Ray *prim_rays;
    8.42 @@ -36,16 +33,16 @@
    8.43  static int global_size;
    8.44  
    8.45  static Sphere sphlist[] = {
    8.46 -	{{0, 0, 8, 1}, {0.7, 0.2, 0.15, 1}, {1, 1, 1, 1}, 1.0, 60, 0, 0},
    8.47 -	{{-0.2, 0.4, 5, 1}, {0.2, 0.9, 0.3, 1}, {1, 1, 1, 1}, 0.25, 40, 0, 0}
    8.48 +	{{0, 0, 0, 1}, {0.7, 0.2, 0.15, 1}, {1, 1, 1, 1}, 1.0, 60, 0, 0},
    8.49 +	{{-0.2, 0.4, -3, 1}, {0.2, 0.9, 0.3, 1}, {1, 1, 1, 1}, 0.25, 40, 0, 0}
    8.50  };
    8.51  
    8.52  static Light lightlist[] = {
    8.53  	{{-10, 10, -20, 1}, {1, 1, 1, 1}}
    8.54  };
    8.55  
    8.56 -static Matrix4x4 xform = {
    8.57 -	{1, 0, 0, 0, 0, 1, 0, 0, 0, 0, 1, 0, 0, 0, 0, 1}
    8.58 +static float xform[16] = {
    8.59 +	1, 0, 0, 0, 0, 1, 0, 0, 0, 0, 1, 0, 0, 0, 0, 1
    8.60  };
    8.61  
    8.62  static RendInfo rinf;
    8.63 @@ -72,7 +69,7 @@
    8.64  	/* setup opencl */
    8.65  	prog = new CLProgram("render");
    8.66  	if(!prog->load("rt.cl")) {
    8.67 -		return 1;
    8.68 +		return false;
    8.69  	}
    8.70  
    8.71  	/* setup argument buffers */
    8.72 @@ -108,6 +105,16 @@
    8.73  	return true;
    8.74  }
    8.75  
    8.76 +void set_xform(float *matrix)
    8.77 +{
    8.78 +	CLMemBuffer *mbuf = prog->get_arg_buffer(5);
    8.79 +	assert(mbuf);
    8.80 +
    8.81 +	assert(map_mem_buffer(mbuf, MAP_WR) == xform);
    8.82 +	memcpy(xform, matrix, sizeof xform);
    8.83 +	unmap_mem_buffer(mbuf);
    8.84 +}
    8.85 +
    8.86  static Ray get_primary_ray(int x, int y, int w, int h, float vfov_deg)
    8.87  {
    8.88  	float vfov = M_PI * vfov_deg / 180.0;

     9.1 --- a/src/rt.h	Fri Jul 23 01:22:03 2010 +0100
     9.2 +++ b/src/rt.h	Fri Jul 23 19:48:43 2010 +0100
     9.3 @@ -4,5 +4,6 @@
     9.4  bool init_renderer(int xsz, int ysz, float *fb);
     9.5  void destroy_renderer();
     9.6  bool render();
     9.7 +void set_xform(float *matrix);
     9.8  
     9.9  #endif	/* RT_H_ */