1.1 --- /dev/null	Thu Jan 01 00:00:00 1970 +0000
     1.2 +++ b/rt.cl	Mon Jul 12 10:38:07 2010 +0300
     1.3 @@ -0,0 +1,77 @@
     1.4 +struct RendInfo {
     1.5 +	int xsz, ysz;
     1.6 +	int num_sph;
     1.7 +	int max_iter;
     1.8 +};
     1.9 +
    1.10 +struct Sphere {
    1.11 +	float4 pos;
    1.12 +	float radius;
    1.13 +	float4 color;
    1.14 +};
    1.15 +
    1.16 +struct Ray {
    1.17 +	float4 origin, dir;
    1.18 +};
    1.19 +
    1.20 +struct SurfPoint {
    1.21 +	float t;
    1.22 +	float3 pos, norm;
    1.23 +};
    1.24 +
    1.25 +#define EPSILON 1e-6
    1.26 +
    1.27 +bool intersect(struct Ray ray, __global const struct Sphere *sph, struct SurfPoint *sp);
    1.28 +
    1.29 +__kernel void render(__global float4 *fb,
    1.30 +		__global const struct RendInfo *rinf,
    1.31 +		__global const struct Sphere *sphlist,
    1.32 +		__global const struct Ray *primrays)
    1.33 +{
    1.34 +	int idx = get_global_id(0);
    1.35 +
    1.36 +	struct Ray ray = primrays[idx];
    1.37 +	struct SurfPoint sp;
    1.38 +
    1.39 +	if(intersect(ray, sphlist, &sp)) {
    1.40 +		fb[idx] = (float4)(1, 0, 0, 1);
    1.41 +	} else {
    1.42 +		fb[idx] = (float4)(0, 0, 0, 1);
    1.43 +	}
    1.44 +}
    1.45 +
    1.46 +bool intersect(struct Ray ray,
    1.47 +		__global const struct Sphere *sph,
    1.48 +		struct SurfPoint *sp)
    1.49 +{
    1.50 +	float3 dir = ray.dir.xyz;
    1.51 +	float3 orig = ray.origin.xyz;
    1.52 +	float3 spos = sph->pos.xyz;
    1.53 +
    1.54 +	float a = dot(dir, dir);
    1.55 +	float b = 2.0 * dir.x * (orig.x - spos.x) +
    1.56 +		2.0 * dir.y * (orig.y - spos.y) +
    1.57 +		2.0 * dir.z * (orig.z - spos.z);
    1.58 +	float c = dot(spos, spos) + dot(orig, orig) +
    1.59 +		2.0 * dot(-spos, orig) - sph->radius * sph->radius;
    1.60 +
    1.61 +	float d = b * b - 4.0 * a * c;
    1.62 +	if(d < 0.0) return false;
    1.63 +
    1.64 +	float sqrt_d = sqrt(d);
    1.65 +	float t1 = (-b + sqrt_d) / (2.0 * a);
    1.66 +	float t2 = (-b - sqrt_d) / (2.0 * a);
    1.67 +
    1.68 +	if(t1 < EPSILON) t1 = t2;
    1.69 +	if(t2 < EPSILON) t2 = t1;
    1.70 +	float t = t1 < t2 ? t1 : t2;
    1.71 +
    1.72 +	if(t < EPSILON || t > 1.0) {
    1.73 +		return false;
    1.74 +	}
    1.75 +
    1.76 +	sp->t = t;
    1.77 +	sp->pos = orig + dir * sp->t;
    1.78 +	sp->norm = (sp->pos - spos) / sph->radius;
    1.79 +	return true;
    1.80 +}

     2.1 --- a/src/clray.cc	Mon Jul 12 07:00:19 2010 +0300
     2.2 +++ b/src/clray.cc	Mon Jul 12 10:38:07 2010 +0300
     2.3 @@ -1,7 +1,16 @@
     2.4  #include <stdio.h>
     2.5 +#include <string.h>
     2.6 +#include <math.h>
     2.7 +#include <errno.h>
     2.8  #include <assert.h>
     2.9  #include "ocl.h"
    2.10  
    2.11 +struct RendInfo {
    2.12 +	int xsz, ysz;
    2.13 +	int num_sph;
    2.14 +	int max_iter;
    2.15 +} __attribute__((packed));
    2.16 +
    2.17  struct Sphere {
    2.18  	cl_float4 pos;
    2.19  	cl_float radius;
    2.20 @@ -14,43 +23,90 @@
    2.21  } __attribute__((packed));
    2.22  
    2.23  
    2.24 +Ray get_primary_ray(int x, int y, int w, int h, float vfov_deg);
    2.25 +bool write_ppm(const char *fname, float *fb, int xsz, int ysz);
    2.26 +
    2.27  int main()
    2.28  {
    2.29 -	int data[] = {1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16};
    2.30 -	int res[16];
    2.31 -	int count = sizeof data / sizeof *data;
    2.32 +	const int xsz = 800;
    2.33 +	const int ysz = 600;
    2.34  
    2.35 -	for(int i=0; i<count; i++) {
    2.36 -		printf("%d ", data[i]);
    2.37 +	Sphere sphlist[] = {
    2.38 +		{{0, 0, 10, 1}, 1.0, {1, 0, 0, 1}}
    2.39 +	};
    2.40 +	RendInfo rinf = {xsz, ysz, sizeof sphlist / sizeof *sphlist, 6};
    2.41 +	Ray *prim_rays = new Ray[xsz * ysz];
    2.42 +	float *fb = new float[xsz * ysz * 4];
    2.43 +
    2.44 +	/* calculate primary rays */
    2.45 +	for(int i=0; i<ysz; i++) {
    2.46 +		for(int j=0; j<xsz; j++) {
    2.47 +			prim_rays[i * xsz + j] = get_primary_ray(j, i, xsz, ysz, 45.0);
    2.48 +		}
    2.49  	}
    2.50 -	putchar('\n');
    2.51  
    2.52 -	CLProgram prog("test");
    2.53 -	if(!prog.load("test.cl")) {
    2.54 -		return 1;
    2.55 -	}
    2.56 -	if(!prog.set_arg_buffer(0, ARG_WR, sizeof res, res)) {
    2.57 -		return 1;
    2.58 -	}
    2.59 -	if(!prog.set_arg_buffer(1, ARG_RD, sizeof data, data)) {
    2.60 -		return 1;
    2.61 -	}
    2.62 -	if(!prog.set_argi(2, 100)) {
    2.63 +	/* setup opencl */
    2.64 +	CLProgram prog("render");
    2.65 +	if(!prog.load("rt.cl")) {
    2.66  		return 1;
    2.67  	}
    2.68  
    2.69 -	if(!prog.run(1, 16)) {
    2.70 +	prog.set_arg_buffer(0, ARG_WR, xsz * ysz * 4 * sizeof(float), fb);
    2.71 +	prog.set_arg_buffer(1, ARG_RD, sizeof rinf, &rinf);
    2.72 +	prog.set_arg_buffer(2, ARG_RD, sizeof sphlist, sphlist);
    2.73 +	prog.set_arg_buffer(3, ARG_RD, xsz * ysz * sizeof *prim_rays, prim_rays);
    2.74 +
    2.75 +	if(!prog.run(1, xsz * ysz)) {
    2.76  		return 1;
    2.77  	}
    2.78  
    2.79 -	CLMemBuffer *mbuf = prog.get_arg_buffer(1);
    2.80 +	CLMemBuffer *mbuf = prog.get_arg_buffer(0);
    2.81  	map_mem_buffer(mbuf, MAP_RD);
    2.82 +	if(!write_ppm("out.ppm", fb, xsz, ysz)) {
    2.83 +		return 1;
    2.84 +	}
    2.85 +	unmap_mem_buffer(mbuf);
    2.86  
    2.87 -	for(int i=0; i<count; i++) {
    2.88 -		printf("%d ", res[i]);
    2.89 -	}
    2.90 -	putchar('\n');
    2.91 -	unmap_mem_buffer(mbuf);
    2.92 +	delete [] fb;
    2.93 +	delete [] prim_rays;
    2.94  
    2.95  	return 0;
    2.96  }
    2.97 +
    2.98 +Ray get_primary_ray(int x, int y, int w, int h, float vfov_deg)
    2.99 +{
   2.100 +	float vfov = M_PI * vfov_deg / 180.0;
   2.101 +	float aspect = (float)w / (float)h;
   2.102 +
   2.103 +	float ysz = 2.0;
   2.104 +	float xsz = aspect * ysz;
   2.105 +
   2.106 +	float px = ((float)x / (float)w) * xsz - xsz / 2.0;
   2.107 +	float py = 1.0 - ((float)y / (float)h) * ysz;
   2.108 +	float pz = 1.0 / tan(0.5 * vfov);
   2.109 +
   2.110 +	pz *= 1000.0;
   2.111 +
   2.112 +	Ray ray = {{0, 0, 0, 1}, {px, py, pz, 1}};
   2.113 +	return ray;
   2.114 +}
   2.115 +
   2.116 +bool write_ppm(const char *fname, float *fb, int xsz, int ysz)
   2.117 +{
   2.118 +	FILE *fp;
   2.119 +
   2.120 +	if(!(fp = fopen(fname, "wb"))) {
   2.121 +		fprintf(stderr, "write_ppm: failed to open file %s for writing: %s\n", fname, strerror(errno));
   2.122 +		return false;
   2.123 +	}
   2.124 +	fprintf(fp, "P6\n%d %d\n255\n", xsz, ysz);
   2.125 +
   2.126 +	for(int i=0; i<xsz * ysz * 4; i++) {
   2.127 +		if(i % 4 == 3) continue;
   2.128 +
   2.129 +		unsigned char c = (unsigned char)(fb[i] * 255.0);
   2.130 +		fputc(c, fp);
   2.131 +	}
   2.132 +	fclose(fp);
   2.133 +	return true;
   2.134 +}

     3.1 --- a/src/ocl.cc	Mon Jul 12 07:00:19 2010 +0300
     3.2 +++ b/src/ocl.cc	Mon Jul 12 10:38:07 2010 +0300
     3.3 @@ -242,12 +242,16 @@
     3.4  
     3.5  bool CLProgram::build()
     3.6  {
     3.7 -	char errlog[512];
     3.8 +	int err;
     3.9  
    3.10 +	if((err = clBuildProgram(prog, 0, 0, 0, 0, 0)) != 0) {
    3.11 +		size_t sz;
    3.12 +		clGetProgramBuildInfo(prog, devinf.id, CL_PROGRAM_BUILD_LOG, 0, 0, &sz);
    3.13  
    3.14 -	if(clBuildProgram(prog, 0, 0, 0, 0, 0) != 0) {
    3.15 -		clGetProgramBuildInfo(prog, devinf.id, CL_PROGRAM_BUILD_LOG, sizeof errlog, errlog, 0);
    3.16 -		fprintf(stderr, "failed to build program:\n%s\n", errlog);
    3.17 +		char *errlog = (char*)alloca(sz + 1);
    3.18 +		clGetProgramBuildInfo(prog, devinf.id, CL_PROGRAM_BUILD_LOG, sz, errlog, 0);
    3.19 +		fprintf(stderr, "failed to build program: (%d)\n%s\n", err, errlog);
    3.20 +
    3.21  		clReleaseProgram(prog);
    3.22  		prog = 0;
    3.23  		return false;