clray

annotate rt.cl @ 2:41d6253492ad

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