clray

view rt.cl @ 23:51f115e337c2

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