clray

annotate rt.cl @ 19:8baea9b66b50

find changesets by author, revision, files, or words in the commit message
added reflection
author John Tsiombikas <nuclear@member.fsf.org>
date Mon, 09 Aug 2010 06:45:57 +0100
parents 4b1604f9798a
children 63a6b46f58a0
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@16 7 float4 ambient;
John@15 8 int dbg;
nuclear@2 9 };
nuclear@2 10
nuclear@9 11 struct Vertex {
nuclear@2 12 float4 pos;
nuclear@9 13 float4 normal;
nuclear@12 14 float4 tex;
nuclear@12 15 float4 padding;
nuclear@9 16 };
nuclear@9 17
nuclear@9 18 struct Face {
nuclear@9 19 struct Vertex v[3];
nuclear@9 20 float4 normal;
nuclear@9 21 int matid;
nuclear@12 22 int padding[3];
nuclear@9 23 };
nuclear@9 24
nuclear@9 25 struct Material {
nuclear@5 26 float4 kd, ks;
nuclear@9 27 float kr, kt;
nuclear@9 28 float spow;
nuclear@12 29 float padding;
nuclear@2 30 };
nuclear@2 31
nuclear@3 32 struct Light {
nuclear@3 33 float4 pos, color;
nuclear@3 34 };
nuclear@3 35
nuclear@2 36 struct Ray {
nuclear@2 37 float4 origin, dir;
nuclear@2 38 };
nuclear@2 39
nuclear@2 40 struct SurfPoint {
nuclear@2 41 float t;
nuclear@12 42 float4 pos, norm, dbg;
nuclear@9 43 global const struct Face *obj;
nuclear@19 44 struct Material mat;
nuclear@2 45 };
nuclear@2 46
nuclear@16 47 struct Scene {
nuclear@16 48 float4 ambient;
nuclear@16 49 global const struct Face *faces;
nuclear@16 50 int num_faces;
nuclear@16 51 global const struct Light *lights;
nuclear@16 52 int num_lights;
nuclear@16 53 global const struct Material *matlib;
nuclear@16 54 };
nuclear@2 55
nuclear@16 56 #define MIN_ENERGY 0.001
nuclear@16 57 #define EPSILON 1e-6
nuclear@16 58
nuclear@17 59 //float4 trace(struct Ray ray, struct Scene *scn);
nuclear@16 60 float4 shade(struct Ray ray, struct Scene *scn, const struct SurfPoint *sp);
nuclear@16 61 bool find_intersection(struct Ray ray, const struct Scene *scn, struct SurfPoint *sp);
nuclear@9 62 bool intersect(struct Ray ray, global const struct Face *face, struct SurfPoint *sp);
nuclear@16 63
nuclear@8 64 float4 reflect(float4 v, float4 n);
nuclear@8 65 float4 transform(float4 v, global const float *xform);
nuclear@16 66 void transform_ray(struct Ray *ray, global const float *xform, global const float *invtrans);
nuclear@12 67 float4 calc_bary(float4 pt, global const struct Face *face, float4 norm);
nuclear@19 68 float mean(float4 v);
nuclear@4 69
nuclear@4 70 kernel void render(global float4 *fb,
nuclear@4 71 global const struct RendInfo *rinf,
nuclear@9 72 global const struct Face *faces,
nuclear@9 73 global const struct Material *matlib,
nuclear@4 74 global const struct Light *lights,
nuclear@7 75 global const struct Ray *primrays,
nuclear@12 76 global const float *xform,
John@15 77 global const float *invtrans,
John@15 78 global struct Face *outfaces)
nuclear@2 79 {
nuclear@2 80 int idx = get_global_id(0);
nuclear@2 81
nuclear@18 82 if(!idx) {
nuclear@18 83 for(int i=0; i<rinf->num_faces; i++) {
nuclear@18 84 outfaces[i] = faces[i];
nuclear@18 85 }
nuclear@18 86 }
nuclear@18 87
nuclear@16 88 struct Scene scn;
nuclear@16 89 scn.ambient = rinf->ambient;
nuclear@16 90 scn.faces = faces;
nuclear@16 91 scn.num_faces = rinf->num_faces;
nuclear@16 92 scn.lights = lights;
nuclear@16 93 scn.num_lights = rinf->num_lights;
nuclear@16 94 scn.matlib = matlib;
nuclear@8 95
nuclear@16 96 struct Ray ray = primrays[idx];
nuclear@16 97 transform_ray(&ray, xform, invtrans);
nuclear@4 98
nuclear@17 99 //fb[idx] = trace(ray, &scn);
nuclear@17 100
nuclear@19 101 float4 pixel = (float4)(0, 0, 0, 0);
nuclear@19 102 float4 energy = (float4)(1.0, 1.0, 1.0, 1.0);
nuclear@19 103 int iter = 0;
nuclear@19 104
nuclear@19 105 while(iter++ < rinf->max_iter && mean(energy) > MIN_ENERGY) {
nuclear@19 106 struct SurfPoint sp;
nuclear@19 107 if(find_intersection(ray, &scn, &sp)) {
nuclear@19 108 pixel += shade(ray, &scn, &sp) * energy;
nuclear@19 109
nuclear@19 110 float4 refl_col = sp.mat.ks * sp.mat.kr;
nuclear@19 111
nuclear@19 112 ray.origin = sp.pos;
nuclear@19 113 ray.dir = reflect(-ray.dir, sp.norm);
nuclear@19 114
nuclear@19 115 energy *= sp.mat.ks * sp.mat.kr;
nuclear@19 116 } else {
nuclear@19 117 iter = INT_MAX - 1; // to break out of the loop
nuclear@19 118 }
nuclear@17 119 }
nuclear@19 120
nuclear@19 121 fb[idx] = pixel;
nuclear@4 122 }
nuclear@4 123
nuclear@17 124 /*float4 trace(struct Ray ray, struct Scene *scn)
nuclear@4 125 {
nuclear@16 126 float4 color;
nuclear@16 127 struct SurfPoint sp;
nuclear@16 128
nuclear@16 129 if(find_intersection(ray, scn, &sp)) {
nuclear@16 130 color = shade(ray, scn, &sp);
nuclear@16 131 } else {
nuclear@16 132 color = (float4)(0, 0, 0, 0);
nuclear@16 133 }
nuclear@16 134 return color;
nuclear@17 135 }*/
nuclear@16 136
nuclear@16 137 float4 shade(struct Ray ray, struct Scene *scn, const struct SurfPoint *sp)
nuclear@16 138 {
nuclear@16 139 float4 norm = sp->norm;
nuclear@12 140 bool entering = true;
nuclear@12 141
nuclear@12 142 if(dot(ray.dir, norm) >= 0.0) {
nuclear@12 143 norm = -norm;
nuclear@12 144 entering = false;
nuclear@12 145 }
nuclear@12 146
nuclear@19 147 float4 dcol = scn->ambient * sp->mat.kd;
nuclear@8 148 float4 scol = (float4)(0, 0, 0, 0);
nuclear@5 149
nuclear@16 150 for(int i=0; i<scn->num_lights; i++) {
nuclear@16 151 float4 ldir = scn->lights[i].pos - sp->pos;
nuclear@5 152
nuclear@16 153 struct Ray shadowray;
nuclear@16 154 shadowray.origin = sp->pos;
nuclear@16 155 shadowray.dir = ldir;
nuclear@5 156
nuclear@16 157 if(!find_intersection(shadowray, scn, 0)) {
nuclear@16 158 ldir = normalize(ldir);
nuclear@16 159 float4 vdir = -normalize(ray.dir);
nuclear@16 160 float4 vref = reflect(vdir, norm);
nuclear@16 161
nuclear@16 162 float diff = fmax(dot(ldir, norm), 0.0f);
nuclear@19 163 dcol += sp->mat.kd * diff * scn->lights[i].color;
nuclear@16 164
nuclear@16 165 //float spec = powr(fmax(dot(ldir, vref), 0.0f), mat.spow);
nuclear@19 166 //scol += sp->mat.ks * spec * scn->lights[i].color;
nuclear@16 167 }
nuclear@16 168 }
nuclear@16 169
nuclear@8 170 return dcol + scol;
nuclear@2 171 }
nuclear@2 172
nuclear@16 173
nuclear@16 174 bool find_intersection(struct Ray ray, const struct Scene *scn, struct SurfPoint *spres)
nuclear@12 175 {
nuclear@16 176 struct SurfPoint sp, sp0;
nuclear@16 177 sp0.t = 1.0;
nuclear@16 178 sp0.obj = 0;
nuclear@16 179
nuclear@16 180 for(int i=0; i<scn->num_faces; i++) {
nuclear@16 181 if(intersect(ray, scn->faces + i, &sp) && sp.t < sp0.t) {
nuclear@16 182 sp0 = sp;
nuclear@16 183 }
nuclear@16 184 }
nuclear@16 185
nuclear@16 186 if(!sp0.obj) {
nuclear@16 187 return false;
nuclear@16 188 }
nuclear@16 189
nuclear@16 190 if(spres) {
nuclear@16 191 *spres = sp0;
nuclear@19 192 spres->mat = scn->matlib[sp0.obj->matid];
nuclear@16 193 }
nuclear@16 194 return true;
nuclear@12 195 }
nuclear@12 196
nuclear@16 197 bool intersect(struct Ray ray, global const struct Face *face, struct SurfPoint *sp)
nuclear@2 198 {
nuclear@12 199 float4 origin = ray.origin;
nuclear@12 200 float4 dir = ray.dir;
nuclear@12 201 float4 norm = face->normal;
nuclear@12 202
nuclear@16 203 float ndotdir = dot(dir, norm);
nuclear@12 204
nuclear@9 205 if(fabs(ndotdir) <= EPSILON) {
nuclear@9 206 return false;
nuclear@9 207 }
nuclear@2 208
nuclear@9 209 float4 pt = face->v[0].pos;
nuclear@12 210 float4 vec = pt - origin;
nuclear@2 211
nuclear@16 212 float ndotvec = dot(norm, vec);
nuclear@9 213 float t = ndotvec / ndotdir;
nuclear@2 214
nuclear@2 215 if(t < EPSILON || t > 1.0) {
nuclear@2 216 return false;
nuclear@2 217 }
nuclear@12 218 pt = origin + dir * t;
nuclear@9 219
nuclear@12 220
nuclear@12 221 float4 bc = calc_bary(pt, face, norm);
nuclear@9 222 float bc_sum = bc.x + bc.y + bc.z;
nuclear@9 223
nuclear@12 224 if(bc_sum < 0.0 || bc_sum > 1.0 + EPSILON) {
nuclear@9 225 return false;
nuclear@12 226 bc *= 1.2;
nuclear@9 227 }
nuclear@2 228
nuclear@2 229 sp->t = t;
nuclear@9 230 sp->pos = pt;
nuclear@12 231 sp->norm = norm;
nuclear@9 232 sp->obj = face;
nuclear@12 233 sp->dbg = bc;
nuclear@2 234 return true;
nuclear@2 235 }
nuclear@5 236
nuclear@8 237 float4 reflect(float4 v, float4 n)
nuclear@5 238 {
nuclear@12 239 float4 res = 2.0f * dot(v, n) * n - v;
nuclear@12 240 return res;
nuclear@5 241 }
nuclear@8 242
nuclear@8 243 float4 transform(float4 v, global const float *xform)
nuclear@8 244 {
nuclear@8 245 float4 res;
nuclear@8 246 res.x = v.x * xform[0] + v.y * xform[4] + v.z * xform[8] + xform[12];
nuclear@8 247 res.y = v.x * xform[1] + v.y * xform[5] + v.z * xform[9] + xform[13];
nuclear@8 248 res.z = v.x * xform[2] + v.y * xform[6] + v.z * xform[10] + xform[14];
nuclear@12 249 res.w = 0.0;
nuclear@8 250 return res;
nuclear@8 251 }
nuclear@8 252
nuclear@16 253 void transform_ray(struct Ray *ray, global const float *xform, global const float *invtrans)
nuclear@8 254 {
nuclear@16 255 ray->origin = transform(ray->origin, xform);
nuclear@16 256 ray->dir = transform(ray->dir, invtrans);
nuclear@8 257 }
nuclear@9 258
nuclear@12 259 float4 calc_bary(float4 pt, global const struct Face *face, float4 norm)
nuclear@9 260 {
nuclear@12 261 float4 bc = (float4)(0, 0, 0, 0);
nuclear@9 262
nuclear@12 263 // calculate area of the whole triangle
nuclear@12 264 float4 v1 = face->v[1].pos - face->v[0].pos;
nuclear@12 265 float4 v2 = face->v[2].pos - face->v[0].pos;
nuclear@12 266 float4 xv1v2 = cross(v1, v2);
nuclear@12 267
nuclear@16 268 float area = fabs(dot(xv1v2, norm)) * 0.5;
nuclear@9 269 if(area < EPSILON) {
nuclear@9 270 return bc;
nuclear@9 271 }
nuclear@9 272
nuclear@9 273 float4 pv0 = face->v[0].pos - pt;
nuclear@9 274 float4 pv1 = face->v[1].pos - pt;
nuclear@9 275 float4 pv2 = face->v[2].pos - pt;
nuclear@9 276
nuclear@12 277 // calculate the area of each sub-triangle
nuclear@12 278 float4 x12 = cross(pv1, pv2);
nuclear@12 279 float4 x20 = cross(pv2, pv0);
nuclear@12 280 float4 x01 = cross(pv0, pv1);
nuclear@12 281
nuclear@16 282 float a0 = fabs(dot(x12, norm)) * 0.5;
nuclear@16 283 float a1 = fabs(dot(x20, norm)) * 0.5;
nuclear@16 284 float a2 = fabs(dot(x01, norm)) * 0.5;
nuclear@9 285
nuclear@9 286 bc.x = a0 / area;
nuclear@9 287 bc.y = a1 / area;
nuclear@9 288 bc.z = a2 / area;
nuclear@9 289 return bc;
nuclear@9 290 }
nuclear@19 291
nuclear@19 292 float mean(float4 v)
nuclear@19 293 {
nuclear@19 294 return native_divide(v.x + v.y + v.z, 3.0);
nuclear@19 295 }