clray

view rt.cl @ 12:85fd61f374d9

find changesets by author, revision, files, or words in the commit message
fixed the bloody intersection bug
author John Tsiombikas <nuclear@member.fsf.org>
date Tue, 03 Aug 2010 13:06:59 +0100
parents a09622aaa043
children 754faf15ba36
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 };
8
9 struct Vertex {
10 float4 pos;
11 float4 normal;
12 float4 tex;
13 float4 padding;
14 };
15
16 struct Face {
17 struct Vertex v[3];
18 float4 normal;
19 int matid;
20 int padding[3];
21 };
22
23 struct Material {
24 float4 kd, ks;
25 float kr, kt;
26 float spow;
27 float padding;
28 };
29
30 struct Light {
31 float4 pos, color;
32 };
33
34 struct Ray {
35 float4 origin, dir;
36 };
37
38 struct SurfPoint {
39 float t;
40 float4 pos, norm, dbg;
41 global const struct Face *obj;
42 global const struct Material *mat;
43 };
44
45 #define EPSILON 1e-6
46
47 float4 shade(struct Ray ray, struct SurfPoint sp,
48 global const struct Light *lights, int num_lights);
49 bool intersect(struct Ray ray, global const struct Face *face, struct SurfPoint *sp);
50 float4 reflect(float4 v, float4 n);
51 float4 transform(float4 v, global const float *xform);
52 struct Ray transform_ray(global const struct Ray *ray, global const float *xform, global const float *invtrans);
53 float4 calc_bary(float4 pt, global const struct Face *face, float4 norm);
54
55 kernel void render(global float4 *fb,
56 global const struct RendInfo *rinf,
57 global const struct Face *faces,
58 global const struct Material *matlib,
59 global const struct Light *lights,
60 global const struct Ray *primrays,
61 global const float *xform,
62 global const float *invtrans)
63 {
64 int idx = get_global_id(0);
65
66 struct Ray ray = transform_ray(primrays + idx, xform, invtrans);
67
68 struct SurfPoint sp, sp0;
69 sp0.t = FLT_MAX;
70 sp0.obj = 0;
71
72 for(int i=0; i<rinf->num_faces; i++) {
73 if(intersect(ray, faces + i, &sp) && sp.t < sp0.t) {
74 sp0 = sp;
75 }
76 }
77
78 if(sp0.obj) {
79 sp0.mat = matlib + sp0.obj->matid;
80 fb[idx] = shade(ray, sp0, lights, rinf->num_lights);
81 } else {
82 fb[idx] = (float4)(0, 0, 0, 0);
83 }
84 }
85
86 float4 shade(struct Ray ray, struct SurfPoint sp,
87 global const struct Light *lights, int num_lights)
88 {
89 float4 norm = sp.norm;
90 bool entering = true;
91
92 if(dot(ray.dir, norm) >= 0.0) {
93 norm = -norm;
94 entering = false;
95 }
96
97 float4 dcol = (float4)(0, 0, 0, 0);
98 float4 scol = (float4)(0, 0, 0, 0);
99
100 for(int i=0; i<num_lights; i++) {
101 float4 ldir = normalize(lights[i].pos - sp.pos);
102 float4 vdir = -normalize(ray.dir);
103 float4 vref = reflect(vdir, norm);
104
105 float diff = fmax(dot(ldir, norm), 0.0f);
106 float spec = powr(fmax(dot(ldir, vref), 0.0f), sp.mat->spow);
107
108 dcol += sp.mat->kd * diff * lights[i].color;
109 //scol += sp.mat->ks * spec * lights[i].color;
110 }
111
112 return dcol + scol;
113 }
114
115 float dot3(float4 a, float4 b)
116 {
117 return a.x * b.x + a.y * b.y + a.z * b.z;
118 }
119
120
121 bool intersect(struct Ray ray,
122 global const struct Face *face,
123 struct SurfPoint *sp)
124 {
125 float4 origin = ray.origin;
126 float4 dir = ray.dir;
127 float4 norm = face->normal;
128
129 float ndotdir = dot3(dir, norm);
130
131 if(fabs(ndotdir) <= EPSILON) {
132 return false;
133 }
134
135 float4 pt = face->v[0].pos;
136 float4 vec = pt - origin;
137
138 float ndotvec = dot3(norm, vec);
139 float t = ndotvec / ndotdir;
140
141 if(t < EPSILON || t > 1.0) {
142 return false;
143 }
144 pt = origin + dir * t;
145
146 if(pt.w < 0.0) return false;
147
148
149 float4 bc = calc_bary(pt, face, norm);
150 float bc_sum = bc.x + bc.y + bc.z;
151
152 if(bc_sum < 0.0 || bc_sum > 1.0 + EPSILON) {
153 return false;
154 bc *= 1.2;
155 }
156
157 sp->t = t;
158 sp->pos = pt;
159 sp->norm = norm;
160 sp->obj = face;
161 sp->dbg = bc;
162 return true;
163 }
164
165 float4 reflect(float4 v, float4 n)
166 {
167 float4 res = 2.0f * dot(v, n) * n - v;
168 return res;
169 }
170
171 float4 transform(float4 v, global const float *xform)
172 {
173 float4 res;
174 res.x = v.x * xform[0] + v.y * xform[4] + v.z * xform[8] + xform[12];
175 res.y = v.x * xform[1] + v.y * xform[5] + v.z * xform[9] + xform[13];
176 res.z = v.x * xform[2] + v.y * xform[6] + v.z * xform[10] + xform[14];
177 res.w = 0.0;
178 return res;
179 }
180
181 struct Ray transform_ray(global const struct Ray *ray, global const float *xform, global const float *invtrans)
182 {
183 struct Ray res;
184 res.origin = transform(ray->origin, xform);
185 res.dir = transform(ray->dir, invtrans);
186 return res;
187 }
188
189 float4 calc_bary(float4 pt, global const struct Face *face, float4 norm)
190 {
191 float4 bc = (float4)(0, 0, 0, 0);
192
193 // calculate area of the whole triangle
194 float4 v1 = face->v[1].pos - face->v[0].pos;
195 float4 v2 = face->v[2].pos - face->v[0].pos;
196 float4 xv1v2 = cross(v1, v2);
197
198 float area = fabs(dot3(xv1v2, norm)) * 0.5;
199 if(area < EPSILON) {
200 return bc;
201 }
202
203 float4 pv0 = face->v[0].pos - pt;
204 float4 pv1 = face->v[1].pos - pt;
205 float4 pv2 = face->v[2].pos - pt;
206
207 // calculate the area of each sub-triangle
208 float4 x12 = cross(pv1, pv2);
209 float4 x20 = cross(pv2, pv0);
210 float4 x01 = cross(pv0, pv1);
211
212 float a0 = fabs(dot3(x12, norm)) * 0.5;
213 float a1 = fabs(dot3(x20, norm)) * 0.5;
214 float a2 = fabs(dot3(x01, norm)) * 0.5;
215
216 bc.x = a0 / area;
217 bc.y = a1 / area;
218 bc.z = a2 / area;
219 return bc;
220 }