clray

annotate rt.cl @ 12:85fd61f374d9

find changesets by author, revision, files, or words in the commit message
fixed the bloody intersection bug
author John Tsiombikas <nuclear@member.fsf.org>
date Tue, 03 Aug 2010 13:06:59 +0100
parents a09622aaa043
children 754faf15ba36
rev   line source
nuclear@12 1 /* vim: set ft=opencl:ts=4:sw=4 */
nuclear@12 2
nuclear@2 3 struct RendInfo {
nuclear@2 4 int xsz, ysz;
nuclear@9 5 int num_faces, num_lights;
nuclear@2 6 int max_iter;
nuclear@2 7 };
nuclear@2 8
nuclear@9 9 struct Vertex {
nuclear@2 10 float4 pos;
nuclear@9 11 float4 normal;
nuclear@12 12 float4 tex;
nuclear@12 13 float4 padding;
nuclear@9 14 };
nuclear@9 15
nuclear@9 16 struct Face {
nuclear@9 17 struct Vertex v[3];
nuclear@9 18 float4 normal;
nuclear@9 19 int matid;
nuclear@12 20 int padding[3];
nuclear@9 21 };
nuclear@9 22
nuclear@9 23 struct Material {
nuclear@5 24 float4 kd, ks;
nuclear@9 25 float kr, kt;
nuclear@9 26 float spow;
nuclear@12 27 float padding;
nuclear@2 28 };
nuclear@2 29
nuclear@3 30 struct Light {
nuclear@3 31 float4 pos, color;
nuclear@3 32 };
nuclear@3 33
nuclear@2 34 struct Ray {
nuclear@2 35 float4 origin, dir;
nuclear@2 36 };
nuclear@2 37
nuclear@2 38 struct SurfPoint {
nuclear@2 39 float t;
nuclear@12 40 float4 pos, norm, dbg;
nuclear@9 41 global const struct Face *obj;
nuclear@9 42 global const struct Material *mat;
nuclear@2 43 };
nuclear@2 44
nuclear@2 45 #define EPSILON 1e-6
nuclear@2 46
nuclear@5 47 float4 shade(struct Ray ray, struct SurfPoint sp,
nuclear@5 48 global const struct Light *lights, int num_lights);
nuclear@9 49 bool intersect(struct Ray ray, global const struct Face *face, struct SurfPoint *sp);
nuclear@8 50 float4 reflect(float4 v, float4 n);
nuclear@8 51 float4 transform(float4 v, global const float *xform);
nuclear@12 52 struct Ray transform_ray(global const struct Ray *ray, global const float *xform, global const float *invtrans);
nuclear@12 53 float4 calc_bary(float4 pt, global const struct Face *face, float4 norm);
nuclear@4 54
nuclear@4 55 kernel void render(global float4 *fb,
nuclear@4 56 global const struct RendInfo *rinf,
nuclear@9 57 global const struct Face *faces,
nuclear@9 58 global const struct Material *matlib,
nuclear@4 59 global const struct Light *lights,
nuclear@7 60 global const struct Ray *primrays,
nuclear@12 61 global const float *xform,
nuclear@12 62 global const float *invtrans)
nuclear@2 63 {
nuclear@2 64 int idx = get_global_id(0);
nuclear@2 65
nuclear@12 66 struct Ray ray = transform_ray(primrays + idx, xform, invtrans);
nuclear@8 67
nuclear@4 68 struct SurfPoint sp, sp0;
nuclear@4 69 sp0.t = FLT_MAX;
nuclear@5 70 sp0.obj = 0;
nuclear@4 71
nuclear@9 72 for(int i=0; i<rinf->num_faces; i++) {
nuclear@9 73 if(intersect(ray, faces + i, &sp) && sp.t < sp0.t) {
nuclear@4 74 sp0 = sp;
nuclear@4 75 }
nuclear@2 76 }
nuclear@3 77
nuclear@5 78 if(sp0.obj) {
nuclear@9 79 sp0.mat = matlib + sp0.obj->matid;
nuclear@5 80 fb[idx] = shade(ray, sp0, lights, rinf->num_lights);
nuclear@5 81 } else {
nuclear@5 82 fb[idx] = (float4)(0, 0, 0, 0);
nuclear@5 83 }
nuclear@4 84 }
nuclear@4 85
nuclear@5 86 float4 shade(struct Ray ray, struct SurfPoint sp,
nuclear@5 87 global const struct Light *lights, int num_lights)
nuclear@4 88 {
nuclear@12 89 float4 norm = sp.norm;
nuclear@12 90 bool entering = true;
nuclear@12 91
nuclear@12 92 if(dot(ray.dir, norm) >= 0.0) {
nuclear@12 93 norm = -norm;
nuclear@12 94 entering = false;
nuclear@12 95 }
nuclear@12 96
nuclear@8 97 float4 dcol = (float4)(0, 0, 0, 0);
nuclear@8 98 float4 scol = (float4)(0, 0, 0, 0);
nuclear@5 99
nuclear@5 100 for(int i=0; i<num_lights; i++) {
nuclear@8 101 float4 ldir = normalize(lights[i].pos - sp.pos);
nuclear@8 102 float4 vdir = -normalize(ray.dir);
nuclear@12 103 float4 vref = reflect(vdir, norm);
nuclear@5 104
nuclear@12 105 float diff = fmax(dot(ldir, norm), 0.0f);
nuclear@9 106 float spec = powr(fmax(dot(ldir, vref), 0.0f), sp.mat->spow);
nuclear@5 107
nuclear@9 108 dcol += sp.mat->kd * diff * lights[i].color;
nuclear@12 109 //scol += sp.mat->ks * spec * lights[i].color;
nuclear@5 110 }
nuclear@5 111
nuclear@8 112 return dcol + scol;
nuclear@2 113 }
nuclear@2 114
nuclear@12 115 float dot3(float4 a, float4 b)
nuclear@12 116 {
nuclear@12 117 return a.x * b.x + a.y * b.y + a.z * b.z;
nuclear@12 118 }
nuclear@12 119
nuclear@12 120
nuclear@2 121 bool intersect(struct Ray ray,
nuclear@9 122 global const struct Face *face,
nuclear@2 123 struct SurfPoint *sp)
nuclear@2 124 {
nuclear@12 125 float4 origin = ray.origin;
nuclear@12 126 float4 dir = ray.dir;
nuclear@12 127 float4 norm = face->normal;
nuclear@12 128
nuclear@12 129 float ndotdir = dot3(dir, norm);
nuclear@12 130
nuclear@9 131 if(fabs(ndotdir) <= EPSILON) {
nuclear@9 132 return false;
nuclear@9 133 }
nuclear@2 134
nuclear@9 135 float4 pt = face->v[0].pos;
nuclear@12 136 float4 vec = pt - origin;
nuclear@2 137
nuclear@12 138 float ndotvec = dot3(norm, vec);
nuclear@9 139 float t = ndotvec / ndotdir;
nuclear@2 140
nuclear@2 141 if(t < EPSILON || t > 1.0) {
nuclear@2 142 return false;
nuclear@2 143 }
nuclear@12 144 pt = origin + dir * t;
nuclear@9 145
nuclear@12 146 if(pt.w < 0.0) return false;
nuclear@12 147
nuclear@12 148
nuclear@12 149 float4 bc = calc_bary(pt, face, norm);
nuclear@9 150 float bc_sum = bc.x + bc.y + bc.z;
nuclear@9 151
nuclear@12 152 if(bc_sum < 0.0 || bc_sum > 1.0 + EPSILON) {
nuclear@9 153 return false;
nuclear@12 154 bc *= 1.2;
nuclear@9 155 }
nuclear@2 156
nuclear@2 157 sp->t = t;
nuclear@9 158 sp->pos = pt;
nuclear@12 159 sp->norm = norm;
nuclear@9 160 sp->obj = face;
nuclear@12 161 sp->dbg = bc;
nuclear@2 162 return true;
nuclear@2 163 }
nuclear@5 164
nuclear@8 165 float4 reflect(float4 v, float4 n)
nuclear@5 166 {
nuclear@12 167 float4 res = 2.0f * dot(v, n) * n - v;
nuclear@12 168 return res;
nuclear@5 169 }
nuclear@8 170
nuclear@8 171 float4 transform(float4 v, global const float *xform)
nuclear@8 172 {
nuclear@8 173 float4 res;
nuclear@8 174 res.x = v.x * xform[0] + v.y * xform[4] + v.z * xform[8] + xform[12];
nuclear@8 175 res.y = v.x * xform[1] + v.y * xform[5] + v.z * xform[9] + xform[13];
nuclear@8 176 res.z = v.x * xform[2] + v.y * xform[6] + v.z * xform[10] + xform[14];
nuclear@12 177 res.w = 0.0;
nuclear@8 178 return res;
nuclear@8 179 }
nuclear@8 180
nuclear@12 181 struct Ray transform_ray(global const struct Ray *ray, global const float *xform, global const float *invtrans)
nuclear@8 182 {
nuclear@8 183 struct Ray res;
nuclear@8 184 res.origin = transform(ray->origin, xform);
nuclear@12 185 res.dir = transform(ray->dir, invtrans);
nuclear@8 186 return res;
nuclear@8 187 }
nuclear@9 188
nuclear@12 189 float4 calc_bary(float4 pt, global const struct Face *face, float4 norm)
nuclear@9 190 {
nuclear@12 191 float4 bc = (float4)(0, 0, 0, 0);
nuclear@9 192
nuclear@12 193 // calculate area of the whole triangle
nuclear@12 194 float4 v1 = face->v[1].pos - face->v[0].pos;
nuclear@12 195 float4 v2 = face->v[2].pos - face->v[0].pos;
nuclear@12 196 float4 xv1v2 = cross(v1, v2);
nuclear@12 197
nuclear@12 198 float area = fabs(dot3(xv1v2, norm)) * 0.5;
nuclear@9 199 if(area < EPSILON) {
nuclear@9 200 return bc;
nuclear@9 201 }
nuclear@9 202
nuclear@9 203 float4 pv0 = face->v[0].pos - pt;
nuclear@9 204 float4 pv1 = face->v[1].pos - pt;
nuclear@9 205 float4 pv2 = face->v[2].pos - pt;
nuclear@9 206
nuclear@12 207 // calculate the area of each sub-triangle
nuclear@12 208 float4 x12 = cross(pv1, pv2);
nuclear@12 209 float4 x20 = cross(pv2, pv0);
nuclear@12 210 float4 x01 = cross(pv0, pv1);
nuclear@12 211
nuclear@12 212 float a0 = fabs(dot3(x12, norm)) * 0.5;
nuclear@12 213 float a1 = fabs(dot3(x20, norm)) * 0.5;
nuclear@12 214 float a2 = fabs(dot3(x01, norm)) * 0.5;
nuclear@9 215
nuclear@9 216 bc.x = a0 / area;
nuclear@9 217 bc.y = a1 / area;
nuclear@9 218 bc.z = a2 / area;
nuclear@9 219 return bc;
nuclear@9 220 }