clray

annotate rt.cl @ 7:575383f3a239

find changesets by author, revision, files, or words in the commit message
foo
author John Tsiombikas <nuclear@member.fsf.org>
date Fri, 23 Jul 2010 01:22:03 +0100
parents 9f0ddb701882
children deaf85acf6af
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@5 9 float4 kd, ks;
nuclear@2 10 float radius;
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@7 28 struct Matrix4x4 {
nuclear@7 29 float m[16];
nuclear@7 30 };
nuclear@7 31
nuclear@2 32 #define EPSILON 1e-6
nuclear@2 33
nuclear@5 34 float4 shade(struct Ray ray, struct SurfPoint sp,
nuclear@5 35 global const struct Light *lights, int num_lights);
nuclear@4 36 bool intersect(struct Ray ray, global const struct Sphere *sph, struct SurfPoint *sp);
nuclear@5 37 float3 reflect(float3 v, float3 n);
nuclear@2 38
nuclear@4 39
nuclear@4 40 kernel void render(global float4 *fb,
nuclear@4 41 global const struct RendInfo *rinf,
nuclear@4 42 global const struct Sphere *sphlist,
nuclear@4 43 global const struct Light *lights,
nuclear@7 44 global const struct Ray *primrays,
nuclear@7 45 global const struct Matrix4x4 xform)
nuclear@2 46 {
nuclear@2 47 int idx = get_global_id(0);
nuclear@2 48
nuclear@2 49 struct Ray ray = primrays[idx];
nuclear@4 50 struct SurfPoint sp, sp0;
nuclear@2 51
nuclear@4 52 sp0.t = FLT_MAX;
nuclear@5 53 sp0.obj = 0;
nuclear@4 54
nuclear@4 55 for(int i=0; i<rinf->num_sph; i++) {
nuclear@5 56 if(intersect(ray, sphlist + i, &sp) && sp.t < sp0.t) {
nuclear@4 57 sp0 = sp;
nuclear@4 58 }
nuclear@2 59 }
nuclear@3 60
nuclear@5 61 if(sp0.obj) {
nuclear@5 62 fb[idx] = shade(ray, sp0, lights, rinf->num_lights);
nuclear@5 63 } else {
nuclear@5 64 fb[idx] = (float4)(0, 0, 0, 0);
nuclear@5 65 }
nuclear@4 66 }
nuclear@4 67
nuclear@5 68 float4 shade(struct Ray ray, struct SurfPoint sp,
nuclear@5 69 global const struct Light *lights, int num_lights)
nuclear@4 70 {
nuclear@5 71 float3 dcol = (float3)(0, 0, 0);
nuclear@5 72 float3 scol = (float3)(0, 0, 0);
nuclear@5 73
nuclear@5 74 for(int i=0; i<num_lights; i++) {
nuclear@5 75 float3 ldir = normalize(lights[i].pos.xyz - sp.pos);
nuclear@5 76 float3 vdir = -normalize(ray.dir.xyz);
nuclear@5 77 float3 vref = reflect(vdir, sp.norm);
nuclear@5 78
nuclear@5 79 float diff = fmax(dot(ldir, sp.norm), 0.0f);
nuclear@5 80 float spec = powr(fmax(dot(ldir, vref), 0.0f), sp.obj->spow);
nuclear@5 81
nuclear@5 82 dcol += sp.obj->kd.xyz * diff * lights[i].color.xyz;
nuclear@5 83 scol += sp.obj->ks.xyz * spec * lights[i].color.xyz;
nuclear@5 84 }
nuclear@5 85
nuclear@5 86 return (float4)(dcol + scol, 1.0f);
nuclear@2 87 }
nuclear@2 88
nuclear@2 89 bool intersect(struct Ray ray,
nuclear@4 90 global const struct Sphere *sph,
nuclear@2 91 struct SurfPoint *sp)
nuclear@2 92 {
nuclear@2 93 float3 dir = ray.dir.xyz;
nuclear@2 94 float3 orig = ray.origin.xyz;
nuclear@2 95 float3 spos = sph->pos.xyz;
nuclear@2 96
nuclear@2 97 float a = dot(dir, dir);
nuclear@2 98 float b = 2.0 * dir.x * (orig.x - spos.x) +
nuclear@2 99 2.0 * dir.y * (orig.y - spos.y) +
nuclear@2 100 2.0 * dir.z * (orig.z - spos.z);
nuclear@2 101 float c = dot(spos, spos) + dot(orig, orig) +
nuclear@2 102 2.0 * dot(-spos, orig) - sph->radius * sph->radius;
nuclear@2 103
nuclear@2 104 float d = b * b - 4.0 * a * c;
nuclear@2 105 if(d < 0.0) return false;
nuclear@2 106
nuclear@2 107 float sqrt_d = sqrt(d);
nuclear@2 108 float t1 = (-b + sqrt_d) / (2.0 * a);
nuclear@2 109 float t2 = (-b - sqrt_d) / (2.0 * a);
nuclear@2 110
nuclear@2 111 if(t1 < EPSILON) t1 = t2;
nuclear@2 112 if(t2 < EPSILON) t2 = t1;
nuclear@2 113 float t = t1 < t2 ? t1 : t2;
nuclear@2 114
nuclear@2 115 if(t < EPSILON || t > 1.0) {
nuclear@2 116 return false;
nuclear@2 117 }
nuclear@2 118
nuclear@2 119 sp->t = t;
nuclear@2 120 sp->pos = orig + dir * sp->t;
nuclear@2 121 sp->norm = (sp->pos - spos) / sph->radius;
nuclear@5 122 sp->obj = sph;
nuclear@2 123 return true;
nuclear@2 124 }
nuclear@5 125
nuclear@5 126 float3 reflect(float3 v, float3 n)
nuclear@5 127 {
nuclear@5 128 return 2.0f * dot(v, n) * n - v;
nuclear@5 129 }