clray

annotate rt.cl @ 39:980bc07be868

find changesets by author, revision, files, or words in the commit message
Implemented OpenGL/OpenCL interop, and removed the texture copy
author John Tsiombikas <nuclear@member.fsf.org>
date Fri, 27 Aug 2010 18:30:09 +0100
parents 7d77ded5f890
children 1bcbb53b3505
rev   line source
nuclear@12 1 /* vim: set ft=opencl:ts=4:sw=4 */
nuclear@12 2
nuclear@2 3 struct RendInfo {
nuclear@22 4 float4 ambient;
nuclear@2 5 int xsz, ysz;
nuclear@9 6 int num_faces, num_lights;
nuclear@2 7 int max_iter;
nuclear@28 8 int kd_depth;
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@28 54 global const struct KDNode *kdtree;
nuclear@28 55 };
nuclear@28 56
nuclear@28 57 struct AABBox {
nuclear@28 58 float4 min, max;
nuclear@28 59 };
nuclear@28 60
nuclear@28 61 struct KDNode {
nuclear@29 62 struct AABBox aabb;
nuclear@28 63 int face_idx[32];
nuclear@28 64 int num_faces;
nuclear@35 65 int left, right;
nuclear@35 66 int padding;
nuclear@16 67 };
nuclear@2 68
nuclear@16 69 #define MIN_ENERGY 0.001
nuclear@21 70 #define EPSILON 1e-5
nuclear@16 71
nuclear@16 72 float4 shade(struct Ray ray, struct Scene *scn, const struct SurfPoint *sp);
nuclear@16 73 bool find_intersection(struct Ray ray, const struct Scene *scn, struct SurfPoint *sp);
nuclear@9 74 bool intersect(struct Ray ray, global const struct Face *face, struct SurfPoint *sp);
nuclear@28 75 bool intersect_aabb(struct Ray ray, struct AABBox aabb);
nuclear@16 76
nuclear@8 77 float4 reflect(float4 v, float4 n);
nuclear@8 78 float4 transform(float4 v, global const float *xform);
nuclear@16 79 void transform_ray(struct Ray *ray, global const float *xform, global const float *invtrans);
nuclear@12 80 float4 calc_bary(float4 pt, global const struct Face *face, float4 norm);
nuclear@19 81 float mean(float4 v);
nuclear@4 82
nuclear@39 83
nuclear@39 84 kernel void render(write_only image2d_t fb,
nuclear@4 85 global const struct RendInfo *rinf,
nuclear@9 86 global const struct Face *faces,
nuclear@9 87 global const struct Material *matlib,
nuclear@4 88 global const struct Light *lights,
nuclear@7 89 global const struct Ray *primrays,
nuclear@12 90 global const float *xform,
nuclear@28 91 global const float *invtrans,
nuclear@28 92 global const struct KDNode *kdtree)
nuclear@2 93 {
nuclear@2 94 int idx = get_global_id(0);
nuclear@2 95
nuclear@16 96 struct Scene scn;
nuclear@16 97 scn.ambient = rinf->ambient;
nuclear@16 98 scn.faces = faces;
nuclear@16 99 scn.num_faces = rinf->num_faces;
nuclear@16 100 scn.lights = lights;
nuclear@16 101 scn.num_lights = rinf->num_lights;
nuclear@16 102 scn.matlib = matlib;
nuclear@30 103 scn.kdtree = kdtree;
nuclear@8 104
nuclear@16 105 struct Ray ray = primrays[idx];
nuclear@16 106 transform_ray(&ray, xform, invtrans);
nuclear@4 107
nuclear@19 108 float4 pixel = (float4)(0, 0, 0, 0);
nuclear@22 109 float4 energy = (float4)(1.0, 1.0, 1.0, 0.0);
nuclear@19 110 int iter = 0;
nuclear@19 111
nuclear@19 112 while(iter++ < rinf->max_iter && mean(energy) > MIN_ENERGY) {
nuclear@19 113 struct SurfPoint sp;
nuclear@19 114 if(find_intersection(ray, &scn, &sp)) {
nuclear@19 115 pixel += shade(ray, &scn, &sp) * energy;
nuclear@19 116
nuclear@19 117 float4 refl_col = sp.mat.ks * sp.mat.kr;
nuclear@19 118
nuclear@19 119 ray.origin = sp.pos;
nuclear@19 120 ray.dir = reflect(-ray.dir, sp.norm);
nuclear@19 121
nuclear@35 122 energy *= refl_col;
nuclear@19 123 } else {
nuclear@35 124 break;
nuclear@19 125 }
nuclear@17 126 }
nuclear@19 127
nuclear@39 128 int img_x = get_image_width(fb);
nuclear@39 129
nuclear@39 130 int2 coord;
nuclear@39 131 coord.x = idx % img_x;
nuclear@39 132 coord.y = idx / img_x;
nuclear@39 133
nuclear@39 134 write_imagef(fb, coord, pixel);
nuclear@39 135 //fb[idx] = pixel;
nuclear@4 136 }
nuclear@4 137
nuclear@16 138 float4 shade(struct Ray ray, struct Scene *scn, const struct SurfPoint *sp)
nuclear@16 139 {
nuclear@16 140 float4 norm = sp->norm;
nuclear@12 141 bool entering = true;
nuclear@12 142
nuclear@12 143 if(dot(ray.dir, norm) >= 0.0) {
nuclear@12 144 norm = -norm;
nuclear@12 145 entering = false;
nuclear@12 146 }
nuclear@12 147
nuclear@19 148 float4 dcol = scn->ambient * sp->mat.kd;
nuclear@8 149 float4 scol = (float4)(0, 0, 0, 0);
nuclear@5 150
nuclear@16 151 for(int i=0; i<scn->num_lights; i++) {
nuclear@16 152 float4 ldir = scn->lights[i].pos - sp->pos;
nuclear@5 153
nuclear@16 154 struct Ray shadowray;
nuclear@16 155 shadowray.origin = sp->pos;
nuclear@16 156 shadowray.dir = ldir;
nuclear@5 157
nuclear@16 158 if(!find_intersection(shadowray, scn, 0)) {
nuclear@16 159 ldir = normalize(ldir);
nuclear@16 160 float4 vdir = -normalize(ray.dir);
nuclear@16 161 float4 vref = reflect(vdir, norm);
nuclear@16 162
nuclear@16 163 float diff = fmax(dot(ldir, norm), 0.0f);
nuclear@22 164 dcol += sp->mat.kd * scn->lights[i].color * diff;
nuclear@16 165
nuclear@20 166 float spec = powr(fmax(dot(ldir, vref), 0.0f), sp->mat.spow);
nuclear@22 167 scol += sp->mat.ks * scn->lights[i].color * spec;
nuclear@16 168 }
nuclear@16 169 }
nuclear@16 170
nuclear@8 171 return dcol + scol;
nuclear@2 172 }
nuclear@2 173
nuclear@30 174 #define STACK_SIZE 64
nuclear@28 175 bool find_intersection(struct Ray ray, const struct Scene *scn, struct SurfPoint *spres)
nuclear@28 176 {
nuclear@29 177 struct SurfPoint sp0;
nuclear@29 178 sp0.t = 1.0;
nuclear@29 179 sp0.obj = 0;
nuclear@29 180
nuclear@29 181 int idxstack[STACK_SIZE];
nuclear@31 182 int top = 0; // points after the topmost element of the stack
nuclear@35 183 idxstack[top++] = 0; // root at tree[0]
nuclear@29 184
nuclear@31 185 while(top > 0) {
nuclear@31 186 int idx = idxstack[--top]; // remove this index from the stack and process it
nuclear@30 187
nuclear@31 188 global const struct KDNode *node = scn->kdtree + idx;
nuclear@29 189
nuclear@31 190 /*if(get_global_id(0) == 0) {
nuclear@31 191 for(int i=0; i<top+1; i++) {
nuclear@31 192 printf(" ");
nuclear@31 193 }
nuclear@31 194 printf("(%d) idx: %d (%p) num_faces: %d\n", top+1, idx, node, node->num_faces);
nuclear@31 195 }*/
nuclear@29 196
nuclear@29 197 if(intersect_aabb(ray, node->aabb)) {
nuclear@35 198 if(node->left == -1) {
nuclear@31 199 // leaf node... check each face in turn and update the nearest intersection as needed
nuclear@29 200 for(int i=0; i<node->num_faces; i++) {
nuclear@31 201 struct SurfPoint spt;
nuclear@29 202 int fidx = node->face_idx[i];
nuclear@29 203
nuclear@31 204 if(intersect(ray, scn->faces + fidx, &spt) && spt.t < sp0.t) {
nuclear@31 205 sp0 = spt;
nuclear@29 206 }
nuclear@29 207 }
nuclear@31 208 } else {
nuclear@31 209 // internal node... recurse to the children
nuclear@31 210 /*if(get_global_id(0) == 0) {
nuclear@35 211 printf("pushing %d's children %d and %d\n", idx, node->left, node->right);
nuclear@31 212 }*/
nuclear@35 213 idxstack[top++] = node->left;
nuclear@35 214 idxstack[top++] = node->right;
nuclear@29 215 }
nuclear@29 216 }
nuclear@29 217 }
nuclear@29 218
nuclear@29 219 if(!sp0.obj) {
nuclear@29 220 return false;
nuclear@29 221 }
nuclear@29 222
nuclear@29 223 if(spres) {
nuclear@29 224 *spres = sp0;
nuclear@29 225 spres->mat = scn->matlib[sp0.obj->matid];
nuclear@29 226 }
nuclear@29 227 return true;
nuclear@28 228 }
nuclear@16 229
nuclear@28 230 /*bool find_intersection(struct Ray ray, const struct Scene *scn, struct SurfPoint *spres)
nuclear@12 231 {
nuclear@16 232 struct SurfPoint sp, sp0;
nuclear@16 233 sp0.t = 1.0;
nuclear@16 234 sp0.obj = 0;
nuclear@16 235
nuclear@16 236 for(int i=0; i<scn->num_faces; i++) {
nuclear@16 237 if(intersect(ray, scn->faces + i, &sp) && sp.t < sp0.t) {
nuclear@16 238 sp0 = sp;
nuclear@16 239 }
nuclear@16 240 }
nuclear@16 241
nuclear@16 242 if(!sp0.obj) {
nuclear@16 243 return false;
nuclear@16 244 }
nuclear@16 245
nuclear@16 246 if(spres) {
nuclear@16 247 *spres = sp0;
nuclear@19 248 spres->mat = scn->matlib[sp0.obj->matid];
nuclear@16 249 }
nuclear@16 250 return true;
nuclear@28 251 }*/
nuclear@12 252
nuclear@16 253 bool intersect(struct Ray ray, global const struct Face *face, struct SurfPoint *sp)
nuclear@2 254 {
nuclear@12 255 float4 origin = ray.origin;
nuclear@12 256 float4 dir = ray.dir;
nuclear@12 257 float4 norm = face->normal;
nuclear@12 258
nuclear@16 259 float ndotdir = dot(dir, norm);
nuclear@12 260
nuclear@9 261 if(fabs(ndotdir) <= EPSILON) {
nuclear@9 262 return false;
nuclear@9 263 }
nuclear@2 264
nuclear@9 265 float4 pt = face->v[0].pos;
nuclear@12 266 float4 vec = pt - origin;
nuclear@2 267
nuclear@16 268 float ndotvec = dot(norm, vec);
nuclear@9 269 float t = ndotvec / ndotdir;
nuclear@2 270
nuclear@2 271 if(t < EPSILON || t > 1.0) {
nuclear@2 272 return false;
nuclear@2 273 }
nuclear@12 274 pt = origin + dir * t;
nuclear@9 275
nuclear@12 276
nuclear@12 277 float4 bc = calc_bary(pt, face, norm);
nuclear@9 278 float bc_sum = bc.x + bc.y + bc.z;
nuclear@9 279
nuclear@20 280 if(bc_sum < 1.0 - EPSILON || bc_sum > 1.0 + EPSILON) {
nuclear@9 281 return false;
nuclear@12 282 bc *= 1.2;
nuclear@9 283 }
nuclear@2 284
nuclear@2 285 sp->t = t;
nuclear@9 286 sp->pos = pt;
nuclear@21 287 sp->norm = normalize(face->v[0].normal * bc.x + face->v[1].normal * bc.y + face->v[2].normal * bc.z);
nuclear@9 288 sp->obj = face;
nuclear@12 289 sp->dbg = bc;
nuclear@2 290 return true;
nuclear@2 291 }
nuclear@5 292
nuclear@28 293 bool intersect_aabb(struct Ray ray, struct AABBox aabb)
nuclear@28 294 {
nuclear@28 295 if(ray.origin.x >= aabb.min.x && ray.origin.y >= aabb.min.y && ray.origin.z >= aabb.min.z &&
nuclear@28 296 ray.origin.x < aabb.max.x && ray.origin.y < aabb.max.y && ray.origin.z < aabb.max.z) {
nuclear@28 297 return true;
nuclear@28 298 }
nuclear@28 299
nuclear@29 300 float4 bbox[2] = {
nuclear@29 301 aabb.min.x, aabb.min.y, aabb.min.z, 0,
nuclear@29 302 aabb.max.x, aabb.max.y, aabb.max.z, 0
nuclear@29 303 };
nuclear@28 304
nuclear@28 305 int xsign = (int)(ray.dir.x < 0.0);
nuclear@28 306 float invdirx = 1.0 / ray.dir.x;
nuclear@28 307 float tmin = (bbox[xsign].x - ray.origin.x) * invdirx;
nuclear@28 308 float tmax = (bbox[1 - xsign].x - ray.origin.x) * invdirx;
nuclear@28 309
nuclear@28 310 int ysign = (int)(ray.dir.y < 0.0);
nuclear@28 311 float invdiry = 1.0 / ray.dir.y;
nuclear@28 312 float tymin = (bbox[ysign].y - ray.origin.y) * invdiry;
nuclear@28 313 float tymax = (bbox[1 - ysign].y - ray.origin.y) * invdiry;
nuclear@28 314
nuclear@28 315 if(tmin > tymax || tymin > tmax) {
nuclear@28 316 return false;
nuclear@28 317 }
nuclear@28 318
nuclear@28 319 if(tymin > tmin) tmin = tymin;
nuclear@28 320 if(tymax < tmax) tmax = tymax;
nuclear@28 321
nuclear@28 322 int zsign = (int)(ray.dir.z < 0.0);
nuclear@28 323 float invdirz = 1.0 / ray.dir.z;
nuclear@28 324 float tzmin = (bbox[zsign].z - ray.origin.z) * invdirz;
nuclear@28 325 float tzmax = (bbox[1 - zsign].z - ray.origin.z) * invdirz;
nuclear@28 326
nuclear@28 327 if(tmin > tzmax || tzmin > tmax) {
nuclear@28 328 return false;
nuclear@28 329 }
nuclear@28 330
nuclear@29 331 return tmin < 1.0 && tmax > 0.0;
nuclear@28 332 }
nuclear@28 333
nuclear@8 334 float4 reflect(float4 v, float4 n)
nuclear@5 335 {
nuclear@23 336 return 2.0f * dot(v, n) * n - v;
nuclear@5 337 }
nuclear@8 338
nuclear@8 339 float4 transform(float4 v, global const float *xform)
nuclear@8 340 {
nuclear@8 341 float4 res;
nuclear@8 342 res.x = v.x * xform[0] + v.y * xform[4] + v.z * xform[8] + xform[12];
nuclear@8 343 res.y = v.x * xform[1] + v.y * xform[5] + v.z * xform[9] + xform[13];
nuclear@8 344 res.z = v.x * xform[2] + v.y * xform[6] + v.z * xform[10] + xform[14];
nuclear@12 345 res.w = 0.0;
nuclear@8 346 return res;
nuclear@8 347 }
nuclear@8 348
nuclear@16 349 void transform_ray(struct Ray *ray, global const float *xform, global const float *invtrans)
nuclear@8 350 {
nuclear@16 351 ray->origin = transform(ray->origin, xform);
nuclear@16 352 ray->dir = transform(ray->dir, invtrans);
nuclear@8 353 }
nuclear@9 354
nuclear@12 355 float4 calc_bary(float4 pt, global const struct Face *face, float4 norm)
nuclear@9 356 {
nuclear@12 357 float4 bc = (float4)(0, 0, 0, 0);
nuclear@9 358
nuclear@12 359 // calculate area of the whole triangle
nuclear@12 360 float4 v1 = face->v[1].pos - face->v[0].pos;
nuclear@12 361 float4 v2 = face->v[2].pos - face->v[0].pos;
nuclear@12 362 float4 xv1v2 = cross(v1, v2);
nuclear@12 363
nuclear@16 364 float area = fabs(dot(xv1v2, norm)) * 0.5;
nuclear@9 365 if(area < EPSILON) {
nuclear@9 366 return bc;
nuclear@9 367 }
nuclear@9 368
nuclear@9 369 float4 pv0 = face->v[0].pos - pt;
nuclear@9 370 float4 pv1 = face->v[1].pos - pt;
nuclear@9 371 float4 pv2 = face->v[2].pos - pt;
nuclear@9 372
nuclear@12 373 // calculate the area of each sub-triangle
nuclear@12 374 float4 x12 = cross(pv1, pv2);
nuclear@12 375 float4 x20 = cross(pv2, pv0);
nuclear@12 376 float4 x01 = cross(pv0, pv1);
nuclear@12 377
nuclear@16 378 float a0 = fabs(dot(x12, norm)) * 0.5;
nuclear@16 379 float a1 = fabs(dot(x20, norm)) * 0.5;
nuclear@16 380 float a2 = fabs(dot(x01, norm)) * 0.5;
nuclear@9 381
nuclear@9 382 bc.x = a0 / area;
nuclear@9 383 bc.y = a1 / area;
nuclear@9 384 bc.z = a2 / area;
nuclear@9 385 return bc;
nuclear@9 386 }
nuclear@19 387
nuclear@19 388 float mean(float4 v)
nuclear@19 389 {
nuclear@19 390 return native_divide(v.x + v.y + v.z, 3.0);
nuclear@19 391 }