clray

view rt.cl @ 44:e7f79c6ad246

added maximum node face limit by default
author John Tsiombikas <nuclear@member.fsf.org>
date Sat, 28 Aug 2010 21:50:17 +0100
parents 1bcbb53b3505
children 8047637961a2
line source
1 /* vim: set ft=opencl:ts=4:sw=4 */
3 struct RendInfo {
4 float4 ambient;
5 int xsz, ysz;
6 int num_faces, num_lights;
7 int max_iter;
8 int kd_depth;
9 };
11 struct Vertex {
12 float4 pos;
13 float4 normal;
14 float4 tex;
15 float4 padding;
16 };
18 struct Face {
19 struct Vertex v[3];
20 float4 normal;
21 int matid;
22 int padding[3];
23 };
25 struct Material {
26 float4 kd, ks;
27 float kr, kt;
28 float spow;
29 float padding;
30 };
32 struct Light {
33 float4 pos, color;
34 };
36 struct Ray {
37 float4 origin, dir;
38 };
40 struct SurfPoint {
41 float t;
42 float4 pos, norm, dbg;
43 global const struct Face *obj;
44 struct Material mat;
45 };
47 struct Scene {
48 float4 ambient;
49 global const struct Face *faces;
50 int num_faces;
51 global const struct Light *lights;
52 int num_lights;
53 global const struct Material *matlib;
54 //global const struct KDNode *kdtree;
55 };
57 struct AABBox {
58 float4 min, max;
59 };
61 #define MAX_NODE_FACES 32
62 struct KDNode {
63 struct AABBox aabb;
64 int face_idx[MAX_NODE_FACES];
65 int num_faces;
66 int left, right;
67 int padding;
68 };
70 #define RAY_MAG 500.0
71 #define MIN_ENERGY 0.001
72 #define EPSILON 1e-5
74 float4 shade(struct Ray ray, struct Scene *scn, const struct SurfPoint *sp, read_only image2d_t kdimg);
75 bool find_intersection(struct Ray ray, const struct Scene *scn, struct SurfPoint *sp, read_only image2d_t kdimg);
76 bool intersect(struct Ray ray, global const struct Face *face, struct SurfPoint *sp);
77 bool intersect_aabb(struct Ray ray, struct AABBox aabb);
79 float4 reflect(float4 v, float4 n);
80 float4 transform(float4 v, global const float *xform);
81 void transform_ray(struct Ray *ray, global const float *xform, global const float *invtrans);
82 float4 calc_bary(float4 pt, global const struct Face *face, float4 norm);
83 float mean(float4 v);
85 void read_kdnode(int idx, struct KDNode *node, read_only image2d_t kdimg);
88 kernel void render(write_only image2d_t fb,
89 global const struct RendInfo *rinf,
90 global const struct Face *faces,
91 global const struct Material *matlib,
92 global const struct Light *lights,
93 global const struct Ray *primrays,
94 global const float *xform,
95 global const float *invtrans,
96 //global const struct KDNode *kdtree
97 read_only image2d_t kdtree_img)
98 {
99 int idx = get_global_id(0);
101 struct Scene scn;
102 scn.ambient = rinf->ambient;
103 scn.faces = faces;
104 scn.num_faces = rinf->num_faces;
105 scn.lights = lights;
106 scn.num_lights = rinf->num_lights;
107 scn.matlib = matlib;
108 //scn.kdtree_img = kdtree_img;
110 struct Ray ray = primrays[idx];
111 transform_ray(&ray, xform, invtrans);
113 float4 pixel = (float4)(0, 0, 0, 0);
114 float4 energy = (float4)(1.0, 1.0, 1.0, 0.0);
115 int iter = 0;
117 while(iter++ < rinf->max_iter && mean(energy) > MIN_ENERGY) {
118 struct SurfPoint sp;
119 if(find_intersection(ray, &scn, &sp, kdtree_img)) {
120 pixel += shade(ray, &scn, &sp, kdtree_img) * energy;
122 float4 refl_col = sp.mat.ks * sp.mat.kr;
124 ray.origin = sp.pos;
125 ray.dir = reflect(-ray.dir, sp.norm);
127 energy *= refl_col;
128 } else {
129 energy = (float4)(0.0, 0.0, 0.0, 0.0);
130 }
131 }
133 int2 coord;
134 coord.x = idx % rinf->xsz;
135 coord.y = idx / rinf->xsz;
137 write_imagef(fb, coord, pixel);
138 }
140 float4 shade(struct Ray ray, struct Scene *scn, const struct SurfPoint *sp, read_only image2d_t kdimg)
141 {
142 float4 norm = sp->norm;
143 //bool entering = true;
145 if(dot(ray.dir, norm) >= 0.0) {
146 norm = -norm;
147 //entering = false;
148 }
150 float4 dcol = scn->ambient * sp->mat.kd;
151 float4 scol = (float4)(0, 0, 0, 0);
153 for(int i=0; i<scn->num_lights; i++) {
154 float4 ldir = scn->lights[i].pos - sp->pos;
156 struct Ray shadowray;
157 shadowray.origin = sp->pos;
158 shadowray.dir = ldir;
160 if(!find_intersection(shadowray, scn, 0, kdimg)) {
161 ldir = normalize(ldir);
162 float4 vdir = -ray.dir;
163 vdir.x = native_divide(vdir.x, RAY_MAG);
164 vdir.y = native_divide(vdir.y, RAY_MAG);
165 vdir.z = native_divide(vdir.z, RAY_MAG);
166 float4 vref = reflect(vdir, norm);
168 float diff = fmax(dot(ldir, norm), 0.0f);
169 dcol += sp->mat.kd /* scn->lights[i].color*/ * diff;
171 float spec = native_powr(fmax(dot(ldir, vref), 0.0f), sp->mat.spow);
172 scol += sp->mat.ks /* scn->lights[i].color*/ * spec;
173 }
174 }
176 return dcol + scol;
177 }
179 #define STACK_SIZE 64
180 bool find_intersection(struct Ray ray, const struct Scene *scn, struct SurfPoint *spres, read_only image2d_t kdimg)
181 {
182 struct SurfPoint sp0;
183 sp0.t = 1.0;
184 sp0.obj = 0;
186 int idxstack[STACK_SIZE];
187 int top = 0; // points after the topmost element of the stack
188 idxstack[top++] = 0; // root at tree[0]
190 while(top > 0) {
191 int idx = idxstack[--top]; // remove this index from the stack and process it
193 struct KDNode node;
194 read_kdnode(idx, &node, kdimg);
196 if(intersect_aabb(ray, node.aabb)) {
197 if(node.left == -1) {
198 // leaf node... check each face in turn and update the nearest intersection as needed
199 for(int i=0; i<node.num_faces; i++) {
200 struct SurfPoint spt;
201 int fidx = node.face_idx[i];
203 if(intersect(ray, scn->faces + fidx, &spt) && spt.t < sp0.t) {
204 sp0 = spt;
205 }
206 }
207 } else {
208 // internal node... recurse to the children
209 idxstack[top++] = node.left;
210 idxstack[top++] = node.right;
211 }
212 }
213 }
215 if(!sp0.obj) {
216 return false;
217 }
219 if(spres) {
220 *spres = sp0;
221 spres->mat = scn->matlib[sp0.obj->matid];
222 }
223 return true;
224 }
226 bool intersect(struct Ray ray, global const struct Face *face, struct SurfPoint *sp)
227 {
228 float4 origin = ray.origin;
229 float4 dir = ray.dir;
230 float4 norm = face->normal;
232 float ndotdir = dot(dir, norm);
234 if(fabs(ndotdir) <= EPSILON) {
235 return false;
236 }
238 float4 pt = face->v[0].pos;
239 float4 vec = pt - origin;
241 float ndotvec = dot(norm, vec);
242 float t = native_divide(ndotvec, ndotdir);
244 if(t < EPSILON || t > 1.0) {
245 return false;
246 }
247 pt = origin + dir * t;
250 float4 bc = calc_bary(pt, face, norm);
251 float bc_sum = bc.x + bc.y + bc.z;
253 if(bc_sum < 1.0 - EPSILON || bc_sum > 1.0 + EPSILON) {
254 return false;
255 bc *= 1.2;
256 }
258 sp->t = t;
259 sp->pos = pt;
260 sp->norm = normalize(face->v[0].normal * bc.x + face->v[1].normal * bc.y + face->v[2].normal * bc.z);
261 sp->obj = face;
262 sp->dbg = bc;
263 return true;
264 }
266 bool intersect_aabb(struct Ray ray, struct AABBox aabb)
267 {
268 if(ray.origin.x >= aabb.min.x && ray.origin.y >= aabb.min.y && ray.origin.z >= aabb.min.z &&
269 ray.origin.x < aabb.max.x && ray.origin.y < aabb.max.y && ray.origin.z < aabb.max.z) {
270 return true;
271 }
273 float4 bbox[2] = {
274 aabb.min.x, aabb.min.y, aabb.min.z, 0,
275 aabb.max.x, aabb.max.y, aabb.max.z, 0
276 };
278 int xsign = (int)(ray.dir.x < 0.0);
279 float invdirx = native_recip(ray.dir.x);
280 float tmin = (bbox[xsign].x - ray.origin.x) * invdirx;
281 float tmax = (bbox[1 - xsign].x - ray.origin.x) * invdirx;
283 int ysign = (int)(ray.dir.y < 0.0);
284 float invdiry = native_recip(ray.dir.y);
285 float tymin = (bbox[ysign].y - ray.origin.y) * invdiry;
286 float tymax = (bbox[1 - ysign].y - ray.origin.y) * invdiry;
288 if(tmin > tymax || tymin > tmax) {
289 return false;
290 }
292 if(tymin > tmin) tmin = tymin;
293 if(tymax < tmax) tmax = tymax;
295 int zsign = (int)(ray.dir.z < 0.0);
296 float invdirz = native_recip(ray.dir.z);
297 float tzmin = (bbox[zsign].z - ray.origin.z) * invdirz;
298 float tzmax = (bbox[1 - zsign].z - ray.origin.z) * invdirz;
300 if(tmin > tzmax || tzmin > tmax) {
301 return false;
302 }
304 return tmin < 1.0 && tmax > 0.0;
305 }
307 float4 reflect(float4 v, float4 n)
308 {
309 return 2.0f * dot(v, n) * n - v;
310 }
312 float4 transform(float4 v, global const float *xform)
313 {
314 float4 res;
315 res.x = v.x * xform[0] + v.y * xform[4] + v.z * xform[8] + xform[12];
316 res.y = v.x * xform[1] + v.y * xform[5] + v.z * xform[9] + xform[13];
317 res.z = v.x * xform[2] + v.y * xform[6] + v.z * xform[10] + xform[14];
318 res.w = 0.0;
319 return res;
320 }
322 void transform_ray(struct Ray *ray, global const float *xform, global const float *invtrans)
323 {
324 ray->origin = transform(ray->origin, xform);
325 ray->dir = transform(ray->dir, invtrans);
326 }
328 float4 calc_bary(float4 pt, global const struct Face *face, float4 norm)
329 {
330 float4 bc = (float4)(0, 0, 0, 0);
332 // calculate area of the whole triangle
333 float4 v1 = face->v[1].pos - face->v[0].pos;
334 float4 v2 = face->v[2].pos - face->v[0].pos;
335 float4 xv1v2 = cross(v1, v2);
337 float area = fabs(dot(xv1v2, norm)) * 0.5;
338 if(area < EPSILON) {
339 return bc;
340 }
342 float4 pv0 = face->v[0].pos - pt;
343 float4 pv1 = face->v[1].pos - pt;
344 float4 pv2 = face->v[2].pos - pt;
346 // calculate the area of each sub-triangle
347 float4 x12 = cross(pv1, pv2);
348 float4 x20 = cross(pv2, pv0);
349 float4 x01 = cross(pv0, pv1);
351 float a0 = fabs(dot(x12, norm)) * 0.5;
352 float a1 = fabs(dot(x20, norm)) * 0.5;
353 float a2 = fabs(dot(x01, norm)) * 0.5;
355 bc.x = native_divide(a0, area);
356 bc.y = native_divide(a1, area);
357 bc.z = native_divide(a2, area);
358 return bc;
359 }
361 float mean(float4 v)
362 {
363 return native_divide(v.x + v.y + v.z, 3.0);
364 }
367 const sampler_t kdsampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE | CLK_FILTER_NEAREST;
369 // read a KD-tree node from a texture scanline
370 void read_kdnode(int idx, struct KDNode *node, read_only image2d_t kdimg)
371 {
372 int2 tc;
373 tc.x = 0;
374 tc.y = idx;
376 node->aabb.min = read_imagef(kdimg, kdsampler, tc); tc.x++;
377 node->aabb.max = read_imagef(kdimg, kdsampler, tc);
379 tc.x = 2 + MAX_NODE_FACES / 4;
380 float4 pix = read_imagef(kdimg, kdsampler, tc);
381 node->num_faces = (int)pix.x;
382 node->left = (int)pix.y;
383 node->right = (int)pix.z;
385 tc.x = 2;
386 for(int i=0; i<node->num_faces; i+=4) {
387 float4 pix = read_imagef(kdimg, kdsampler, tc); tc.x++;
388 node->face_idx[i] = (int)pix.x;
389 node->face_idx[i + 1] = (int)pix.y;
390 node->face_idx[i + 2] = (int)pix.z;
391 node->face_idx[i + 3] = (int)pix.w;
392 }
393 }