clray: 980bc07be868 rt.cl

clray

view rt.cl @ 39:980bc07be868

Implemented OpenGL/OpenCL interop, and removed the texture copy

author	John Tsiombikas <nuclear@member.fsf.org>
date	Fri, 27 Aug 2010 18:30:09 +0100
parents	7d77ded5f890
children	1bcbb53b3505

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 struct KDNode {

62 struct AABBox aabb;

63 int face_idx[32];

64 int num_faces;

65 int left, right;

66 int padding;

67 };

69 #define MIN_ENERGY 0.001

70 #define EPSILON 1e-5

72 float4 shade(struct Ray ray, struct Scene *scn, const struct SurfPoint *sp);

73 bool find_intersection(struct Ray ray, const struct Scene *scn, struct SurfPoint *sp);

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

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

77 float4 reflect(float4 v, float4 n);

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

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

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

81 float mean(float4 v);

84 kernel void render(write_only image2d_t fb,

85 global const struct RendInfo *rinf,

86 global const struct Face *faces,

87 global const struct Material *matlib,

88 global const struct Light *lights,

89 global const struct Ray *primrays,

90 global const float *xform,

91 global const float *invtrans,

92 global const struct KDNode *kdtree)

93 {

94 int idx = get_global_id(0);

96 struct Scene scn;

97 scn.ambient = rinf->ambient;

98 scn.faces = faces;

99 scn.num_faces = rinf->num_faces;

100 scn.lights = lights;

101 scn.num_lights = rinf->num_lights;

102 scn.matlib = matlib;

103 scn.kdtree = kdtree;

104

105 struct Ray ray = primrays[idx];

106 transform_ray(&ray, xform, invtrans);

107

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

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

110 int iter = 0;

111

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

113 struct SurfPoint sp;

114 if(find_intersection(ray, &scn, &sp)) {

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

116

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

118

119 ray.origin = sp.pos;

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

121

122 energy *= refl_col;

123 } else {

124 break;

125 }

126 }

127

128 int img_x = get_image_width(fb);

129

130 int2 coord;

131 coord.x = idx % img_x;

132 coord.y = idx / img_x;

133

134 write_imagef(fb, coord, pixel);

135 //fb[idx] = pixel;

136 }

137

138 float4 shade(struct Ray ray, struct Scene *scn, const struct SurfPoint *sp)

139 {

140 float4 norm = sp->norm;

141 bool entering = true;

142

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

144 norm = -norm;

145 entering = false;

146 }

147

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

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

150

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

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

153

154 struct Ray shadowray;

155 shadowray.origin = sp->pos;

156 shadowray.dir = ldir;

157

158 if(!find_intersection(shadowray, scn, 0)) {

159 ldir = normalize(ldir);

160 float4 vdir = -normalize(ray.dir);

161 float4 vref = reflect(vdir, norm);

162

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

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

165

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

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

168 }

169 }

170

171 return dcol + scol;

172 }

173

174 #define STACK_SIZE 64

175 bool find_intersection(struct Ray ray, const struct Scene *scn, struct SurfPoint *spres)

176 {

177 struct SurfPoint sp0;

178 sp0.t = 1.0;

179 sp0.obj = 0;

180

181 int idxstack[STACK_SIZE];

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

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

184

185 while(top > 0) {

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

187

188 global const struct KDNode *node = scn->kdtree + idx;

189

190 /*if(get_global_id(0) == 0) {

191 for(int i=0; i<top+1; i++) {

192 printf(" ");

193 }

194 printf("(%d) idx: %d (%p) num_faces: %d\n", top+1, idx, node, node->num_faces);

195 }*/

196

197 if(intersect_aabb(ray, node->aabb)) {

198 if(node->left == -1) {

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

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

201 struct SurfPoint spt;

202 int fidx = node->face_idx[i];

203

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

205 sp0 = spt;

206 }

207 }

208 } else {

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

210 /*if(get_global_id(0) == 0) {

211 printf("pushing %d's children %d and %d\n", idx, node->left, node->right);

212 }*/

213 idxstack[top++] = node->left;

214 idxstack[top++] = node->right;

215 }

216 }

217 }

218

219 if(!sp0.obj) {

220 return false;

221 }

222

223 if(spres) {

224 *spres = sp0;

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

226 }

227 return true;

228 }

229

230 /*bool find_intersection(struct Ray ray, const struct Scene *scn, struct SurfPoint *spres)

231 {

232 struct SurfPoint sp, sp0;

233 sp0.t = 1.0;

234 sp0.obj = 0;

235

236 for(int i=0; i<scn->num_faces; i++) {

237 if(intersect(ray, scn->faces + i, &sp) && sp.t < sp0.t) {

238 sp0 = sp;

239 }

240 }

241

242 if(!sp0.obj) {

243 return false;

244 }

245

246 if(spres) {

247 *spres = sp0;

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

249 }

250 return true;

251 }*/

252

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

254 {

255 float4 origin = ray.origin;

256 float4 dir = ray.dir;

257 float4 norm = face->normal;

258

259 float ndotdir = dot(dir, norm);

260

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

262 return false;

263 }

264

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

266 float4 vec = pt - origin;

267

268 float ndotvec = dot(norm, vec);

269 float t = ndotvec / ndotdir;

270

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

272 return false;

273 }

274 pt = origin + dir * t;

275

276

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

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

279

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

281 return false;

282 bc *= 1.2;

283 }

284

285 sp->t = t;

286 sp->pos = pt;

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

288 sp->obj = face;

289 sp->dbg = bc;

290 return true;

291 }

292

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

294 {

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

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

297 return true;

298 }

299

300 float4 bbox[2] = {

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

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

303 };

304

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

306 float invdirx = 1.0 / ray.dir.x;

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

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

309

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

311 float invdiry = 1.0 / ray.dir.y;

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

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

314

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

316 return false;

317 }

318

319 if(tymin > tmin) tmin = tymin;

320 if(tymax < tmax) tmax = tymax;

321

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

323 float invdirz = 1.0 / ray.dir.z;

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

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

326

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

328 return false;

329 }

330

331 return tmin < 1.0 && tmax > 0.0;

332 }

333

334 float4 reflect(float4 v, float4 n)

335 {

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

337 }

338

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

340 {

341 float4 res;

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

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

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

345 res.w = 0.0;

346 return res;

347 }

348

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

350 {

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

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

353 }

354

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

356 {

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

358

359 // calculate area of the whole triangle

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

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

362 float4 xv1v2 = cross(v1, v2);

363

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

365 if(area < EPSILON) {

366 return bc;

367 }

368

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

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

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

372

373 // calculate the area of each sub-triangle

374 float4 x12 = cross(pv1, pv2);

375 float4 x20 = cross(pv2, pv0);

376 float4 x01 = cross(pv0, pv1);

377

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

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

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

381

382 bc.x = a0 / area;

383 bc.y = a1 / area;

384 bc.z = a2 / area;

385 return bc;

386 }

387

388 float mean(float4 v)

389 {

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

391 }