clray

annotate rt.cl @ 3:88ac4eb2d18a

find changesets by author, revision, files, or words in the commit message
added OpenGL display of the framebuffer
author John Tsiombikas <nuclear@member.fsf.org>
date Tue, 13 Jul 2010 03:38:29 +0300
parents 41d6253492ad
children 3c95d568d3c7
rev   line source
nuclear@2 1 struct RendInfo {
nuclear@2 2 int xsz, ysz;
nuclear@3 3 int num_sph, num_lights;
nuclear@2 4 int max_iter;
nuclear@2 5 };
nuclear@2 6
nuclear@2 7 struct Sphere {
nuclear@2 8 float4 pos;
nuclear@2 9 float radius;
nuclear@2 10 float4 color;
nuclear@2 11 };
nuclear@2 12
nuclear@3 13 struct Light {
nuclear@3 14 float4 pos, color;
nuclear@3 15 };
nuclear@3 16
nuclear@2 17 struct Ray {
nuclear@2 18 float4 origin, dir;
nuclear@2 19 };
nuclear@2 20
nuclear@2 21 struct SurfPoint {
nuclear@2 22 float t;
nuclear@2 23 float3 pos, norm;
nuclear@2 24 };
nuclear@2 25
nuclear@2 26 #define EPSILON 1e-6
nuclear@2 27
nuclear@2 28 bool intersect(struct Ray ray, __global const struct Sphere *sph, struct SurfPoint *sp);
nuclear@2 29
nuclear@2 30 __kernel void render(__global float4 *fb,
nuclear@2 31 __global const struct RendInfo *rinf,
nuclear@2 32 __global const struct Sphere *sphlist,
nuclear@3 33 __global const struct Light *lights,
nuclear@2 34 __global const struct Ray *primrays)
nuclear@2 35 {
nuclear@2 36 int idx = get_global_id(0);
nuclear@2 37
nuclear@2 38 struct Ray ray = primrays[idx];
nuclear@2 39 struct SurfPoint sp;
nuclear@2 40
nuclear@2 41 if(intersect(ray, sphlist, &sp)) {
nuclear@2 42 fb[idx] = (float4)(1, 0, 0, 1);
nuclear@2 43 } else {
nuclear@2 44 fb[idx] = (float4)(0, 0, 0, 1);
nuclear@2 45 }
nuclear@3 46
nuclear@3 47 fb[idx] = primrays[idx].dir * 0.5 + 0.5;
nuclear@2 48 }
nuclear@2 49
nuclear@2 50 bool intersect(struct Ray ray,
nuclear@2 51 __global const struct Sphere *sph,
nuclear@2 52 struct SurfPoint *sp)
nuclear@2 53 {
nuclear@2 54 float3 dir = ray.dir.xyz;
nuclear@2 55 float3 orig = ray.origin.xyz;
nuclear@2 56 float3 spos = sph->pos.xyz;
nuclear@2 57
nuclear@2 58 float a = dot(dir, dir);
nuclear@2 59 float b = 2.0 * dir.x * (orig.x - spos.x) +
nuclear@2 60 2.0 * dir.y * (orig.y - spos.y) +
nuclear@2 61 2.0 * dir.z * (orig.z - spos.z);
nuclear@2 62 float c = dot(spos, spos) + dot(orig, orig) +
nuclear@2 63 2.0 * dot(-spos, orig) - sph->radius * sph->radius;
nuclear@2 64
nuclear@2 65 float d = b * b - 4.0 * a * c;
nuclear@2 66 if(d < 0.0) return false;
nuclear@2 67
nuclear@2 68 float sqrt_d = sqrt(d);
nuclear@2 69 float t1 = (-b + sqrt_d) / (2.0 * a);
nuclear@2 70 float t2 = (-b - sqrt_d) / (2.0 * a);
nuclear@2 71
nuclear@2 72 if(t1 < EPSILON) t1 = t2;
nuclear@2 73 if(t2 < EPSILON) t2 = t1;
nuclear@2 74 float t = t1 < t2 ? t1 : t2;
nuclear@2 75
nuclear@2 76 if(t < EPSILON || t > 1.0) {
nuclear@2 77 return false;
nuclear@2 78 }
nuclear@2 79
nuclear@2 80 sp->t = t;
nuclear@2 81 sp->pos = orig + dir * sp->t;
nuclear@2 82 sp->norm = (sp->pos - spos) / sph->radius;
nuclear@2 83 return true;
nuclear@2 84 }