clray

view rt.cl @ 15:754faf15ba36

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