# HG changeset patch
# User John Tsiombikas <nuclear@member.fsf.org>
# Date 1278920287 -10800
# Node ID 41d6253492ad6e027cc8ad573ca788bccbcadddd
# Parent  0b0e4d18d53f6318cd681b1d11bbc93dc8369e29
pfffff

diff -r 0b0e4d18d53f -r 41d6253492ad rt.cl
--- /dev/null	Thu Jan 01 00:00:00 1970 +0000
+++ b/rt.cl	Mon Jul 12 10:38:07 2010 +0300
@@ -0,0 +1,77 @@
+struct RendInfo {
+	int xsz, ysz;
+	int num_sph;
+	int max_iter;
+};
+
+struct Sphere {
+	float4 pos;
+	float radius;
+	float4 color;
+};
+
+struct Ray {
+	float4 origin, dir;
+};
+
+struct SurfPoint {
+	float t;
+	float3 pos, norm;
+};
+
+#define EPSILON 1e-6
+
+bool intersect(struct Ray ray, __global const struct Sphere *sph, struct SurfPoint *sp);
+
+__kernel void render(__global float4 *fb,
+		__global const struct RendInfo *rinf,
+		__global const struct Sphere *sphlist,
+		__global const struct Ray *primrays)
+{
+	int idx = get_global_id(0);
+
+	struct Ray ray = primrays[idx];
+	struct SurfPoint sp;
+
+	if(intersect(ray, sphlist, &sp)) {
+		fb[idx] = (float4)(1, 0, 0, 1);
+	} else {
+		fb[idx] = (float4)(0, 0, 0, 1);
+	}
+}
+
+bool intersect(struct Ray ray,
+		__global const struct Sphere *sph,
+		struct SurfPoint *sp)
+{
+	float3 dir = ray.dir.xyz;
+	float3 orig = ray.origin.xyz;
+	float3 spos = sph->pos.xyz;
+
+	float a = dot(dir, dir);
+	float b = 2.0 * dir.x * (orig.x - spos.x) +
+		2.0 * dir.y * (orig.y - spos.y) +
+		2.0 * dir.z * (orig.z - spos.z);
+	float c = dot(spos, spos) + dot(orig, orig) +
+		2.0 * dot(-spos, orig) - sph->radius * sph->radius;
+
+	float d = b * b - 4.0 * a * c;
+	if(d < 0.0) return false;
+
+	float sqrt_d = sqrt(d);
+	float t1 = (-b + sqrt_d) / (2.0 * a);
+	float t2 = (-b - sqrt_d) / (2.0 * a);
+
+	if(t1 < EPSILON) t1 = t2;
+	if(t2 < EPSILON) t2 = t1;
+	float t = t1 < t2 ? t1 : t2;
+
+	if(t < EPSILON || t > 1.0) {
+		return false;
+	}
+
+	sp->t = t;
+	sp->pos = orig + dir * sp->t;
+	sp->norm = (sp->pos - spos) / sph->radius;
+	return true;
+}
diff -r 0b0e4d18d53f -r 41d6253492ad src/clray.cc
--- a/src/clray.cc	Mon Jul 12 07:00:19 2010 +0300
+++ b/src/clray.cc	Mon Jul 12 10:38:07 2010 +0300
@@ -1,7 +1,16 @@
 #include <stdio.h>
+#include <string.h>
+#include <math.h>
+#include <errno.h>
 #include <assert.h>
 #include "ocl.h"
 
+struct RendInfo {
+	int xsz, ysz;
+	int num_sph;
+	int max_iter;
+} __attribute__((packed));
+
 struct Sphere {
 	cl_float4 pos;
 	cl_float radius;
@@ -14,43 +23,90 @@
 } __attribute__((packed));
 
 
+Ray get_primary_ray(int x, int y, int w, int h, float vfov_deg);
+bool write_ppm(const char *fname, float *fb, int xsz, int ysz);
+
 int main()
 {
-	int data[] = {1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16};
-	int res[16];
-	int count = sizeof data / sizeof *data;
+	const int xsz = 800;
+	const int ysz = 600;
 
-	for(int i=0; i<count; i++) {
-		printf("%d ", data[i]);
+	Sphere sphlist[] = {
+		{{0, 0, 10, 1}, 1.0, {1, 0, 0, 1}}
+	};
+	RendInfo rinf = {xsz, ysz, sizeof sphlist / sizeof *sphlist, 6};
+	Ray *prim_rays = new Ray[xsz * ysz];
+	float *fb = new float[xsz * ysz * 4];
+
+	/* calculate primary rays */
+	for(int i=0; i<ysz; i++) {
+		for(int j=0; j<xsz; j++) {
+			prim_rays[i * xsz + j] = get_primary_ray(j, i, xsz, ysz, 45.0);
+		}
 	}
-	putchar('\n');
 
-	CLProgram prog("test");
-	if(!prog.load("test.cl")) {
-		return 1;
-	}
-	if(!prog.set_arg_buffer(0, ARG_WR, sizeof res, res)) {
-		return 1;
-	}
-	if(!prog.set_arg_buffer(1, ARG_RD, sizeof data, data)) {
-		return 1;
-	}
-	if(!prog.set_argi(2, 100)) {
+	/* setup opencl */
+	CLProgram prog("render");
+	if(!prog.load("rt.cl")) {
 		return 1;
 	}
 
-	if(!prog.run(1, 16)) {
+	prog.set_arg_buffer(0, ARG_WR, xsz * ysz * 4 * sizeof(float), fb);
+	prog.set_arg_buffer(1, ARG_RD, sizeof rinf, &rinf);
+	prog.set_arg_buffer(2, ARG_RD, sizeof sphlist, sphlist);
+	prog.set_arg_buffer(3, ARG_RD, xsz * ysz * sizeof *prim_rays, prim_rays);
+
+	if(!prog.run(1, xsz * ysz)) {
 		return 1;
 	}
 
-	CLMemBuffer *mbuf = prog.get_arg_buffer(1);
+	CLMemBuffer *mbuf = prog.get_arg_buffer(0);
 	map_mem_buffer(mbuf, MAP_RD);
+	if(!write_ppm("out.ppm", fb, xsz, ysz)) {
+		return 1;
+	}
+	unmap_mem_buffer(mbuf);
 
-	for(int i=0; i<count; i++) {
-		printf("%d ", res[i]);
-	}
-	putchar('\n');
-	unmap_mem_buffer(mbuf);
+	delete [] fb;
+	delete [] prim_rays;
 
 	return 0;
 }
+
+Ray get_primary_ray(int x, int y, int w, int h, float vfov_deg)
+{
+	float vfov = M_PI * vfov_deg / 180.0;
+	float aspect = (float)w / (float)h;
+
+	float ysz = 2.0;
+	float xsz = aspect * ysz;
+
+	float px = ((float)x / (float)w) * xsz - xsz / 2.0;
+	float py = 1.0 - ((float)y / (float)h) * ysz;
+	float pz = 1.0 / tan(0.5 * vfov);
+
+	pz *= 1000.0;
+
+	Ray ray = {{0, 0, 0, 1}, {px, py, pz, 1}};
+	return ray;
+}
+
+bool write_ppm(const char *fname, float *fb, int xsz, int ysz)
+{
+	FILE *fp;
+
+	if(!(fp = fopen(fname, "wb"))) {
+		fprintf(stderr, "write_ppm: failed to open file %s for writing: %s\n", fname, strerror(errno));
+		return false;
+	}
+	fprintf(fp, "P6\n%d %d\n255\n", xsz, ysz);
+
+	for(int i=0; i<xsz * ysz * 4; i++) {
+		if(i % 4 == 3) continue;
+
+		unsigned char c = (unsigned char)(fb[i] * 255.0);
+		fputc(c, fp);
+	}
+	fclose(fp);
+	return true;
+}
diff -r 0b0e4d18d53f -r 41d6253492ad src/ocl.cc
--- a/src/ocl.cc	Mon Jul 12 07:00:19 2010 +0300
+++ b/src/ocl.cc	Mon Jul 12 10:38:07 2010 +0300
@@ -242,12 +242,16 @@
 
 bool CLProgram::build()
 {
-	char errlog[512];
+	int err;
 
+	if((err = clBuildProgram(prog, 0, 0, 0, 0, 0)) != 0) {
+		size_t sz;
+		clGetProgramBuildInfo(prog, devinf.id, CL_PROGRAM_BUILD_LOG, 0, 0, &sz);
 
-	if(clBuildProgram(prog, 0, 0, 0, 0, 0) != 0) {
-		clGetProgramBuildInfo(prog, devinf.id, CL_PROGRAM_BUILD_LOG, sizeof errlog, errlog, 0);
-		fprintf(stderr, "failed to build program:\n%s\n", errlog);
+		char *errlog = (char*)alloca(sz + 1);
+		clGetProgramBuildInfo(prog, devinf.id, CL_PROGRAM_BUILD_LOG, sz, errlog, 0);
+		fprintf(stderr, "failed to build program: (%d)\n%s\n", err, errlog);
+
 		clReleaseProgram(prog);
 		prog = 0;
 		return false;