clray: e7f79c6ad246 rt.cl

clray

view rt.cl @ 44:e7f79c6ad246

added maximum node face limit by default

author	John Tsiombikas <nuclear@member.fsf.org>
date	Sat, 28 Aug 2010 21:50:17 +0100
parents	1bcbb53b3505
children	8047637961a2

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 #define MAX_NODE_FACES 32

62 struct KDNode {

63 struct AABBox aabb;

64 int face_idx[MAX_NODE_FACES];

65 int num_faces;

66 int left, right;

67 int padding;

68 };

70 #define RAY_MAG 500.0

71 #define MIN_ENERGY 0.001

72 #define EPSILON 1e-5

74 float4 shade(struct Ray ray, struct Scene *scn, const struct SurfPoint *sp, read_only image2d_t kdimg);

75 bool find_intersection(struct Ray ray, const struct Scene *scn, struct SurfPoint *sp, read_only image2d_t kdimg);

76 bool intersect(struct Ray ray, global const struct Face *face, struct SurfPoint *sp);

77 bool intersect_aabb(struct Ray ray, struct AABBox aabb);

79 float4 reflect(float4 v, float4 n);

80 float4 transform(float4 v, global const float *xform);

81 void transform_ray(struct Ray *ray, global const float *xform, global const float *invtrans);

82 float4 calc_bary(float4 pt, global const struct Face *face, float4 norm);

83 float mean(float4 v);

85 void read_kdnode(int idx, struct KDNode *node, read_only image2d_t kdimg);

88 kernel void render(write_only image2d_t fb,

89 global const struct RendInfo *rinf,

90 global const struct Face *faces,

91 global const struct Material *matlib,

92 global const struct Light *lights,

93 global const struct Ray *primrays,

94 global const float *xform,

95 global const float *invtrans,

96 //global const struct KDNode *kdtree

97 read_only image2d_t kdtree_img)

98 {

99 int idx = get_global_id(0);

100

101 struct Scene scn;

102 scn.ambient = rinf->ambient;

103 scn.faces = faces;

104 scn.num_faces = rinf->num_faces;

105 scn.lights = lights;

106 scn.num_lights = rinf->num_lights;

107 scn.matlib = matlib;

108 //scn.kdtree_img = kdtree_img;

109

110 struct Ray ray = primrays[idx];

111 transform_ray(&ray, xform, invtrans);

112

113 float4 pixel = (float4)(0, 0, 0, 0);

114 float4 energy = (float4)(1.0, 1.0, 1.0, 0.0);

115 int iter = 0;

116

117 while(iter++ < rinf->max_iter && mean(energy) > MIN_ENERGY) {

118 struct SurfPoint sp;

119 if(find_intersection(ray, &scn, &sp, kdtree_img)) {

120 pixel += shade(ray, &scn, &sp, kdtree_img) * energy;

121

122 float4 refl_col = sp.mat.ks * sp.mat.kr;

123

124 ray.origin = sp.pos;

125 ray.dir = reflect(-ray.dir, sp.norm);

126

127 energy *= refl_col;

128 } else {

129 energy = (float4)(0.0, 0.0, 0.0, 0.0);

130 }

131 }

132

133 int2 coord;

134 coord.x = idx % rinf->xsz;

135 coord.y = idx / rinf->xsz;

136

137 write_imagef(fb, coord, pixel);

138 }

139

140 float4 shade(struct Ray ray, struct Scene *scn, const struct SurfPoint *sp, read_only image2d_t kdimg)

141 {

142 float4 norm = sp->norm;

143 //bool entering = true;

144

145 if(dot(ray.dir, norm) >= 0.0) {

146 norm = -norm;

147 //entering = false;

148 }

149

150 float4 dcol = scn->ambient * sp->mat.kd;

151 float4 scol = (float4)(0, 0, 0, 0);

152

153 for(int i=0; i<scn->num_lights; i++) {

154 float4 ldir = scn->lights[i].pos - sp->pos;

155

156 struct Ray shadowray;

157 shadowray.origin = sp->pos;

158 shadowray.dir = ldir;

159

160 if(!find_intersection(shadowray, scn, 0, kdimg)) {

161 ldir = normalize(ldir);

162 float4 vdir = -ray.dir;

163 vdir.x = native_divide(vdir.x, RAY_MAG);

164 vdir.y = native_divide(vdir.y, RAY_MAG);

165 vdir.z = native_divide(vdir.z, RAY_MAG);

166 float4 vref = reflect(vdir, norm);

167

168 float diff = fmax(dot(ldir, norm), 0.0f);

169 dcol += sp->mat.kd /* scn->lights[i].color*/ * diff;

170

171 float spec = native_powr(fmax(dot(ldir, vref), 0.0f), sp->mat.spow);

172 scol += sp->mat.ks /* scn->lights[i].color*/ * spec;

173 }

174 }

175

176 return dcol + scol;

177 }

178

179 #define STACK_SIZE 64

180 bool find_intersection(struct Ray ray, const struct Scene *scn, struct SurfPoint *spres, read_only image2d_t kdimg)

181 {

182 struct SurfPoint sp0;

183 sp0.t = 1.0;

184 sp0.obj = 0;

185

186 int idxstack[STACK_SIZE];

187 int top = 0; // points after the topmost element of the stack

188 idxstack[top++] = 0; // root at tree[0]

189

190 while(top > 0) {

191 int idx = idxstack[--top]; // remove this index from the stack and process it

192

193 struct KDNode node;

194 read_kdnode(idx, &node, kdimg);

195

196 if(intersect_aabb(ray, node.aabb)) {

197 if(node.left == -1) {

198 // leaf node... check each face in turn and update the nearest intersection as needed

199 for(int i=0; i<node.num_faces; i++) {

200 struct SurfPoint spt;

201 int fidx = node.face_idx[i];

202

203 if(intersect(ray, scn->faces + fidx, &spt) && spt.t < sp0.t) {

204 sp0 = spt;

205 }

206 }

207 } else {

208 // internal node... recurse to the children

209 idxstack[top++] = node.left;

210 idxstack[top++] = node.right;

211 }

212 }

213 }

214

215 if(!sp0.obj) {

216 return false;

217 }

218

219 if(spres) {

220 *spres = sp0;

221 spres->mat = scn->matlib[sp0.obj->matid];

222 }

223 return true;

224 }

225

226 bool intersect(struct Ray ray, global const struct Face *face, struct SurfPoint *sp)

227 {

228 float4 origin = ray.origin;

229 float4 dir = ray.dir;

230 float4 norm = face->normal;

231

232 float ndotdir = dot(dir, norm);

233

234 if(fabs(ndotdir) <= EPSILON) {

235 return false;

236 }

237

238 float4 pt = face->v[0].pos;

239 float4 vec = pt - origin;

240

241 float ndotvec = dot(norm, vec);

242 float t = native_divide(ndotvec, ndotdir);

243

244 if(t < EPSILON || t > 1.0) {

245 return false;

246 }

247 pt = origin + dir * t;

248

249

250 float4 bc = calc_bary(pt, face, norm);

251 float bc_sum = bc.x + bc.y + bc.z;

252

253 if(bc_sum < 1.0 - EPSILON || bc_sum > 1.0 + EPSILON) {

254 return false;

255 bc *= 1.2;

256 }

257

258 sp->t = t;

259 sp->pos = pt;

260 sp->norm = normalize(face->v[0].normal * bc.x + face->v[1].normal * bc.y + face->v[2].normal * bc.z);

261 sp->obj = face;

262 sp->dbg = bc;

263 return true;

264 }

265

266 bool intersect_aabb(struct Ray ray, struct AABBox aabb)

267 {

268 if(ray.origin.x >= aabb.min.x && ray.origin.y >= aabb.min.y && ray.origin.z >= aabb.min.z &&

269 ray.origin.x < aabb.max.x && ray.origin.y < aabb.max.y && ray.origin.z < aabb.max.z) {

270 return true;

271 }

272

273 float4 bbox[2] = {

274 aabb.min.x, aabb.min.y, aabb.min.z, 0,

275 aabb.max.x, aabb.max.y, aabb.max.z, 0

276 };

277

278 int xsign = (int)(ray.dir.x < 0.0);

279 float invdirx = native_recip(ray.dir.x);

280 float tmin = (bbox[xsign].x - ray.origin.x) * invdirx;

281 float tmax = (bbox[1 - xsign].x - ray.origin.x) * invdirx;

282

283 int ysign = (int)(ray.dir.y < 0.0);

284 float invdiry = native_recip(ray.dir.y);

285 float tymin = (bbox[ysign].y - ray.origin.y) * invdiry;

286 float tymax = (bbox[1 - ysign].y - ray.origin.y) * invdiry;

287

288 if(tmin > tymax || tymin > tmax) {

289 return false;

290 }

291

292 if(tymin > tmin) tmin = tymin;

293 if(tymax < tmax) tmax = tymax;

294

295 int zsign = (int)(ray.dir.z < 0.0);

296 float invdirz = native_recip(ray.dir.z);

297 float tzmin = (bbox[zsign].z - ray.origin.z) * invdirz;

298 float tzmax = (bbox[1 - zsign].z - ray.origin.z) * invdirz;

299

300 if(tmin > tzmax || tzmin > tmax) {

301 return false;

302 }

303

304 return tmin < 1.0 && tmax > 0.0;

305 }

306

307 float4 reflect(float4 v, float4 n)

308 {

309 return 2.0f * dot(v, n) * n - v;

310 }

311

312 float4 transform(float4 v, global const float *xform)

313 {

314 float4 res;

315 res.x = v.x * xform[0] + v.y * xform[4] + v.z * xform[8] + xform[12];

316 res.y = v.x * xform[1] + v.y * xform[5] + v.z * xform[9] + xform[13];

317 res.z = v.x * xform[2] + v.y * xform[6] + v.z * xform[10] + xform[14];

318 res.w = 0.0;

319 return res;

320 }

321

322 void transform_ray(struct Ray *ray, global const float *xform, global const float *invtrans)

323 {

324 ray->origin = transform(ray->origin, xform);

325 ray->dir = transform(ray->dir, invtrans);

326 }

327

328 float4 calc_bary(float4 pt, global const struct Face *face, float4 norm)

329 {

330 float4 bc = (float4)(0, 0, 0, 0);

331

332 // calculate area of the whole triangle

333 float4 v1 = face->v[1].pos - face->v[0].pos;

334 float4 v2 = face->v[2].pos - face->v[0].pos;

335 float4 xv1v2 = cross(v1, v2);

336

337 float area = fabs(dot(xv1v2, norm)) * 0.5;

338 if(area < EPSILON) {

339 return bc;

340 }

341

342 float4 pv0 = face->v[0].pos - pt;

343 float4 pv1 = face->v[1].pos - pt;

344 float4 pv2 = face->v[2].pos - pt;

345

346 // calculate the area of each sub-triangle

347 float4 x12 = cross(pv1, pv2);

348 float4 x20 = cross(pv2, pv0);

349 float4 x01 = cross(pv0, pv1);

350

351 float a0 = fabs(dot(x12, norm)) * 0.5;

352 float a1 = fabs(dot(x20, norm)) * 0.5;

353 float a2 = fabs(dot(x01, norm)) * 0.5;

354

355 bc.x = native_divide(a0, area);

356 bc.y = native_divide(a1, area);

357 bc.z = native_divide(a2, area);

358 return bc;

359 }

360

361 float mean(float4 v)

362 {

363 return native_divide(v.x + v.y + v.z, 3.0);

364 }

365

366

367 const sampler_t kdsampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE | CLK_FILTER_NEAREST;

368

369 // read a KD-tree node from a texture scanline

370 void read_kdnode(int idx, struct KDNode *node, read_only image2d_t kdimg)

371 {

372 int2 tc;

373 tc.x = 0;

374 tc.y = idx;

375

376 node->aabb.min = read_imagef(kdimg, kdsampler, tc); tc.x++;

377 node->aabb.max = read_imagef(kdimg, kdsampler, tc);

378

379 tc.x = 2 + MAX_NODE_FACES / 4;

380 float4 pix = read_imagef(kdimg, kdsampler, tc);

381 node->num_faces = (int)pix.x;

382 node->left = (int)pix.y;

383 node->right = (int)pix.z;

384

385 tc.x = 2;

386 for(int i=0; i<node->num_faces; i+=4) {

387 float4 pix = read_imagef(kdimg, kdsampler, tc); tc.x++;

388 node->face_idx[i] = (int)pix.x;

389 node->face_idx[i + 1] = (int)pix.y;

390 node->face_idx[i + 2] = (int)pix.z;

391 node->face_idx[i + 3] = (int)pix.w;

392 }

393 }