clray

view rt.cl @ 18:4b1604f9798a

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