clray
view rt.cl @ 31:92786fc3317e
yey! *seems* to work now
| author | John Tsiombikas <nuclear@member.fsf.org> |
|---|---|
| date | Sun, 22 Aug 2010 00:50:47 +0100 |
| parents | 04803c702014 |
| children | 7d77ded5f890 |
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 struct KDNode {
62 struct AABBox aabb;
63 int face_idx[32];
64 int num_faces;
65 int padding[3];
66 };
68 #define MIN_ENERGY 0.001
69 #define EPSILON 1e-5
71 float4 shade(struct Ray ray, struct Scene *scn, const struct SurfPoint *sp);
72 bool find_intersection(struct Ray ray, const struct Scene *scn, struct SurfPoint *sp);
73 bool intersect(struct Ray ray, global const struct Face *face, struct SurfPoint *sp);
74 bool intersect_aabb(struct Ray ray, struct AABBox aabb);
76 float4 reflect(float4 v, float4 n);
77 float4 transform(float4 v, global const float *xform);
78 void transform_ray(struct Ray *ray, global const float *xform, global const float *invtrans);
79 float4 calc_bary(float4 pt, global const struct Face *face, float4 norm);
80 float mean(float4 v);
82 kernel void render(global float4 *fb,
83 global const struct RendInfo *rinf,
84 global const struct Face *faces,
85 global const struct Material *matlib,
86 global const struct Light *lights,
87 global const struct Ray *primrays,
88 global const float *xform,
89 global const float *invtrans,
90 global const struct KDNode *kdtree)
91 {
92 int idx = get_global_id(0);
94 struct Scene scn;
95 scn.ambient = rinf->ambient;
96 scn.faces = faces;
97 scn.num_faces = rinf->num_faces;
98 scn.lights = lights;
99 scn.num_lights = rinf->num_lights;
100 scn.matlib = matlib;
101 scn.kdtree = kdtree;
103 struct Ray ray = primrays[idx];
104 transform_ray(&ray, xform, invtrans);
106 float4 pixel = (float4)(0, 0, 0, 0);
107 float4 energy = (float4)(1.0, 1.0, 1.0, 0.0);
108 int iter = 0;
110 while(iter++ < rinf->max_iter && mean(energy) > MIN_ENERGY) {
111 struct SurfPoint sp;
112 if(find_intersection(ray, &scn, &sp)) {
113 pixel += shade(ray, &scn, &sp) * energy;
115 float4 refl_col = sp.mat.ks * sp.mat.kr;
117 ray.origin = sp.pos;
118 ray.dir = reflect(-ray.dir, sp.norm);
120 energy *= sp.mat.ks * sp.mat.kr;
121 } else {
122 iter = INT_MAX - 1; // to break out of the loop
123 }
124 }
126 fb[idx] = pixel;
127 }
129 float4 shade(struct Ray ray, struct Scene *scn, const struct SurfPoint *sp)
130 {
131 float4 norm = sp->norm;
132 bool entering = true;
134 if(dot(ray.dir, norm) >= 0.0) {
135 norm = -norm;
136 entering = false;
137 }
139 float4 dcol = scn->ambient * sp->mat.kd;
140 float4 scol = (float4)(0, 0, 0, 0);
142 for(int i=0; i<scn->num_lights; i++) {
143 float4 ldir = scn->lights[i].pos - sp->pos;
145 struct Ray shadowray;
146 shadowray.origin = sp->pos;
147 shadowray.dir = ldir;
149 if(!find_intersection(shadowray, scn, 0)) {
150 ldir = normalize(ldir);
151 float4 vdir = -normalize(ray.dir);
152 float4 vref = reflect(vdir, norm);
154 float diff = fmax(dot(ldir, norm), 0.0f);
155 dcol += sp->mat.kd * scn->lights[i].color * diff;
157 float spec = powr(fmax(dot(ldir, vref), 0.0f), sp->mat.spow);
158 scol += sp->mat.ks * scn->lights[i].color * spec;
159 }
160 }
162 return dcol + scol;
163 }
165 #define STACK_SIZE 64
166 bool find_intersection(struct Ray ray, const struct Scene *scn, struct SurfPoint *spres)
167 {
168 struct SurfPoint sp0;
169 sp0.t = 1.0;
170 sp0.obj = 0;
172 int idxstack[STACK_SIZE];
173 int top = 0; // points after the topmost element of the stack
174 idxstack[top++] = 1; // root at tree[1] (heap)
176 while(top > 0) {
177 int idx = idxstack[--top]; // remove this index from the stack and process it
179 global const struct KDNode *node = scn->kdtree + idx;
181 /*if(get_global_id(0) == 0) {
182 for(int i=0; i<top+1; i++) {
183 printf(" ");
184 }
185 printf("(%d) idx: %d (%p) num_faces: %d\n", top+1, idx, node, node->num_faces);
186 }*/
188 if(intersect_aabb(ray, node->aabb)) {
189 if(node->num_faces >= 0) {
190 // leaf node... check each face in turn and update the nearest intersection as needed
191 for(int i=0; i<node->num_faces; i++) {
192 struct SurfPoint spt;
193 int fidx = node->face_idx[i];
195 if(intersect(ray, scn->faces + fidx, &spt) && spt.t < sp0.t) {
196 sp0 = spt;
197 }
198 }
199 } else {
200 // internal node... recurse to the children
201 /*if(get_global_id(0) == 0) {
202 printf("pushing %d's children %d and %d\n", idx, idx * 2, idx * 2 + 1);
203 }*/
204 idxstack[top++] = idx * 2;
205 idxstack[top++] = idx * 2 + 1;
206 }
207 }
208 }
210 if(!sp0.obj) {
211 return false;
212 }
214 if(spres) {
215 *spres = sp0;
216 spres->mat = scn->matlib[sp0.obj->matid];
217 }
218 return true;
219 }
221 /*bool find_intersection(struct Ray ray, const struct Scene *scn, struct SurfPoint *spres)
222 {
223 struct SurfPoint sp, sp0;
224 sp0.t = 1.0;
225 sp0.obj = 0;
227 for(int i=0; i<scn->num_faces; i++) {
228 if(intersect(ray, scn->faces + i, &sp) && sp.t < sp0.t) {
229 sp0 = sp;
230 }
231 }
233 if(!sp0.obj) {
234 return false;
235 }
237 if(spres) {
238 *spres = sp0;
239 spres->mat = scn->matlib[sp0.obj->matid];
240 }
241 return true;
242 }*/
244 bool intersect(struct Ray ray, global const struct Face *face, struct SurfPoint *sp)
245 {
246 float4 origin = ray.origin;
247 float4 dir = ray.dir;
248 float4 norm = face->normal;
250 float ndotdir = dot(dir, norm);
252 if(fabs(ndotdir) <= EPSILON) {
253 return false;
254 }
256 float4 pt = face->v[0].pos;
257 float4 vec = pt - origin;
259 float ndotvec = dot(norm, vec);
260 float t = ndotvec / ndotdir;
262 if(t < EPSILON || t > 1.0) {
263 return false;
264 }
265 pt = origin + dir * t;
268 float4 bc = calc_bary(pt, face, norm);
269 float bc_sum = bc.x + bc.y + bc.z;
271 if(bc_sum < 1.0 - EPSILON || bc_sum > 1.0 + EPSILON) {
272 return false;
273 bc *= 1.2;
274 }
276 sp->t = t;
277 sp->pos = pt;
278 sp->norm = normalize(face->v[0].normal * bc.x + face->v[1].normal * bc.y + face->v[2].normal * bc.z);
279 sp->obj = face;
280 sp->dbg = bc;
281 return true;
282 }
284 bool intersect_aabb(struct Ray ray, struct AABBox aabb)
285 {
286 if(ray.origin.x >= aabb.min.x && ray.origin.y >= aabb.min.y && ray.origin.z >= aabb.min.z &&
287 ray.origin.x < aabb.max.x && ray.origin.y < aabb.max.y && ray.origin.z < aabb.max.z) {
288 return true;
289 }
291 float4 bbox[2] = {
292 aabb.min.x, aabb.min.y, aabb.min.z, 0,
293 aabb.max.x, aabb.max.y, aabb.max.z, 0
294 };
296 int xsign = (int)(ray.dir.x < 0.0);
297 float invdirx = 1.0 / ray.dir.x;
298 float tmin = (bbox[xsign].x - ray.origin.x) * invdirx;
299 float tmax = (bbox[1 - xsign].x - ray.origin.x) * invdirx;
301 int ysign = (int)(ray.dir.y < 0.0);
302 float invdiry = 1.0 / ray.dir.y;
303 float tymin = (bbox[ysign].y - ray.origin.y) * invdiry;
304 float tymax = (bbox[1 - ysign].y - ray.origin.y) * invdiry;
306 if(tmin > tymax || tymin > tmax) {
307 return false;
308 }
310 if(tymin > tmin) tmin = tymin;
311 if(tymax < tmax) tmax = tymax;
313 int zsign = (int)(ray.dir.z < 0.0);
314 float invdirz = 1.0 / ray.dir.z;
315 float tzmin = (bbox[zsign].z - ray.origin.z) * invdirz;
316 float tzmax = (bbox[1 - zsign].z - ray.origin.z) * invdirz;
318 if(tmin > tzmax || tzmin > tmax) {
319 return false;
320 }
322 return tmin < 1.0 && tmax > 0.0;
323 }
325 float4 reflect(float4 v, float4 n)
326 {
327 return 2.0f * dot(v, n) * n - v;
328 }
330 float4 transform(float4 v, global const float *xform)
331 {
332 float4 res;
333 res.x = v.x * xform[0] + v.y * xform[4] + v.z * xform[8] + xform[12];
334 res.y = v.x * xform[1] + v.y * xform[5] + v.z * xform[9] + xform[13];
335 res.z = v.x * xform[2] + v.y * xform[6] + v.z * xform[10] + xform[14];
336 res.w = 0.0;
337 return res;
338 }
340 void transform_ray(struct Ray *ray, global const float *xform, global const float *invtrans)
341 {
342 ray->origin = transform(ray->origin, xform);
343 ray->dir = transform(ray->dir, invtrans);
344 }
346 float4 calc_bary(float4 pt, global const struct Face *face, float4 norm)
347 {
348 float4 bc = (float4)(0, 0, 0, 0);
350 // calculate area of the whole triangle
351 float4 v1 = face->v[1].pos - face->v[0].pos;
352 float4 v2 = face->v[2].pos - face->v[0].pos;
353 float4 xv1v2 = cross(v1, v2);
355 float area = fabs(dot(xv1v2, norm)) * 0.5;
356 if(area < EPSILON) {
357 return bc;
358 }
360 float4 pv0 = face->v[0].pos - pt;
361 float4 pv1 = face->v[1].pos - pt;
362 float4 pv2 = face->v[2].pos - pt;
364 // calculate the area of each sub-triangle
365 float4 x12 = cross(pv1, pv2);
366 float4 x20 = cross(pv2, pv0);
367 float4 x01 = cross(pv0, pv1);
369 float a0 = fabs(dot(x12, norm)) * 0.5;
370 float a1 = fabs(dot(x20, norm)) * 0.5;
371 float a2 = fabs(dot(x01, norm)) * 0.5;
373 bc.x = a0 / area;
374 bc.y = a1 / area;
375 bc.z = a2 / area;
376 return bc;
377 }
379 float mean(float4 v)
380 {
381 return native_divide(v.x + v.y + v.z, 3.0);
382 }