clray

annotate rt.cl @ 4:3c95d568d3c7

wow a ball
author John Tsiombikas <nuclear@member.fsf.org>
date Thu, 15 Jul 2010 07:37:05 +0300
parents 88ac4eb2d18a
children 9f0ddb701882
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@4 10 float4 kd, ks;
nuclear@4 11 float spow, kr, kt;
nuclear@2 12 };
nuclear@2 13
nuclear@3 14 struct Light {
nuclear@3 15 float4 pos, color;
nuclear@3 16 };
nuclear@3 17
nuclear@2 18 struct Ray {
nuclear@2 19 float4 origin, dir;
nuclear@2 20 };
nuclear@2 21
nuclear@2 22 struct SurfPoint {
nuclear@2 23 float t;
nuclear@2 24 float3 pos, norm;
nuclear@4 25 global const struct Sphere *obj;
nuclear@2 26 };
nuclear@2 27
nuclear@2 28 #define EPSILON 1e-6
nuclear@2 29
nuclear@4 30 float4 shade(struct Ray ray, struct SurfPoint sp);
nuclear@4 31 bool intersect(struct Ray ray, global const struct Sphere *sph, struct SurfPoint *sp);
nuclear@2 32
nuclear@4 33
nuclear@4 34 kernel void render(global float4 *fb,
nuclear@4 35 global const struct RendInfo *rinf,
nuclear@4 36 global const struct Sphere *sphlist,
nuclear@4 37 global const struct Light *lights,
nuclear@4 38 global const struct Ray *primrays)
nuclear@2 39 {
nuclear@2 40 int idx = get_global_id(0);
nuclear@2 41
nuclear@2 42 struct Ray ray = primrays[idx];
nuclear@4 43 struct SurfPoint sp, sp0;
nuclear@2 44
nuclear@4 45 sp0.t = FLT_MAX;
nuclear@4 46
nuclear@4 47 for(int i=0; i<rinf->num_sph; i++) {
nuclear@4 48 if(intersect(ray, sphlist, &sp) && sp.t < sp0.t) {
nuclear@4 49 sp0 = sp;
nuclear@4 50 }
nuclear@2 51 }
nuclear@3 52
nuclear@4 53 fb[idx] = shade(ray, sp0);
nuclear@4 54 }
nuclear@4 55
nuclear@4 56 float4 shade(struct Ray ray, struct SurfPoint sp)
nuclear@4 57 {
nuclear@4 58 return sp.obj->kd;
nuclear@2 59 }
nuclear@2 60
nuclear@2 61 bool intersect(struct Ray ray,
nuclear@4 62 global const struct Sphere *sph,
nuclear@2 63 struct SurfPoint *sp)
nuclear@2 64 {
nuclear@2 65 float3 dir = ray.dir.xyz;
nuclear@2 66 float3 orig = ray.origin.xyz;
nuclear@2 67 float3 spos = sph->pos.xyz;
nuclear@2 68
nuclear@2 69 float a = dot(dir, dir);
nuclear@2 70 float b = 2.0 * dir.x * (orig.x - spos.x) +
nuclear@2 71 2.0 * dir.y * (orig.y - spos.y) +
nuclear@2 72 2.0 * dir.z * (orig.z - spos.z);
nuclear@2 73 float c = dot(spos, spos) + dot(orig, orig) +
nuclear@2 74 2.0 * dot(-spos, orig) - sph->radius * sph->radius;
nuclear@2 75
nuclear@2 76 float d = b * b - 4.0 * a * c;
nuclear@2 77 if(d < 0.0) return false;
nuclear@2 78
nuclear@2 79 float sqrt_d = sqrt(d);
nuclear@2 80 float t1 = (-b + sqrt_d) / (2.0 * a);
nuclear@2 81 float t2 = (-b - sqrt_d) / (2.0 * a);
nuclear@2 82
nuclear@2 83 if(t1 < EPSILON) t1 = t2;
nuclear@2 84 if(t2 < EPSILON) t2 = t1;
nuclear@2 85 float t = t1 < t2 ? t1 : t2;
nuclear@2 86
nuclear@2 87 if(t < EPSILON || t > 1.0) {
nuclear@2 88 return false;
nuclear@2 89 }
nuclear@2 90
nuclear@2 91 sp->t = t;
nuclear@2 92 sp->pos = orig + dir * sp->t;
nuclear@2 93 sp->norm = (sp->pos - spos) / sph->radius;
nuclear@2 94 return true;
nuclear@2 95 }