clray: 97cfd9675310 rt.cl

clray

view rt.cl @ 28:97cfd9675310

trying to pass the kdtree to the kernel

author	John Tsiombikas <nuclear@member.fsf.org>
date	Sat, 21 Aug 2010 03:42:49 +0100
parents	51f115e337c2
children	353d80127627

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 AABBox aabb;

63 int face_idx[32];

64 int num_faces;

65 int padding[3];

66 };

68 #define MIN_ENERGY 0.001

69 #define EPSILON 1e-5

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

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

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

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

76 float4 reflect(float4 v, float4 n);

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

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

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

80 float mean(float4 v);

82 kernel void render(global float4 *fb,

83 global const struct RendInfo *rinf,

84 global const struct Face *faces,

85 global const struct Material *matlib,

86 global const struct Light *lights,

87 global const struct Ray *primrays,

88 global const float *xform,

89 global const float *invtrans,

90 global const struct KDNode *kdtree)

91 {

92 int idx = get_global_id(0);

94 struct Scene scn;

95 scn.ambient = rinf->ambient;

96 scn.faces = faces;

97 scn.num_faces = rinf->num_faces;

98 scn.lights = lights;

99 scn.num_lights = rinf->num_lights;

100 scn.matlib = matlib;

101

102 struct Ray ray = primrays[idx];

103 transform_ray(&ray, xform, invtrans);

104

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

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

107 int iter = 0;

108

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

110 struct SurfPoint sp;

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

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

113

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

115

116 ray.origin = sp.pos;

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

118

119 energy *= sp.mat.ks * sp.mat.kr;

120 } else {

121 iter = INT_MAX - 1; // to break out of the loop

122 }

123 }

124

125 fb[idx] = pixel;

126 }

127

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

129 {

130 float4 norm = sp->norm;

131 bool entering = true;

132

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

134 norm = -norm;

135 entering = false;

136 }

137

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

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

140

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

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

143

144 struct Ray shadowray;

145 shadowray.origin = sp->pos;

146 shadowray.dir = ldir;

147

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

149 ldir = normalize(ldir);

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

151 float4 vref = reflect(vdir, norm);

152

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

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

155

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

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

158 }

159 }

160

161 return dcol + scol;

162 }

163

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

165 {

166 return false;

167 }

168

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

170 {

171 struct SurfPoint sp, sp0;

172 sp0.t = 1.0;

173 sp0.obj = 0;

174

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

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

177 sp0 = sp;

178 }

179 }

180

181 if(!sp0.obj) {

182 return false;

183 }

184

185 if(spres) {

186 *spres = sp0;

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

188 }

189 return true;

190 }*/

191

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

193 {

194 float4 origin = ray.origin;

195 float4 dir = ray.dir;

196 float4 norm = face->normal;

197

198 float ndotdir = dot(dir, norm);

199

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

201 return false;

202 }

203

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

205 float4 vec = pt - origin;

206

207 float ndotvec = dot(norm, vec);

208 float t = ndotvec / ndotdir;

209

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

211 return false;

212 }

213 pt = origin + dir * t;

214

215

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

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

218

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

220 return false;

221 bc *= 1.2;

222 }

223

224 sp->t = t;

225 sp->pos = pt;

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

227 sp->obj = face;

228 sp->dbg = bc;

229 return true;

230 }

231

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

233 {

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

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

236 return true;

237 }

238

239 float4 bbox[2] = {aabb.min, aabb.max};

240

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

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

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

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

245

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

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

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

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

250

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

252 return false;

253 }

254

255 if(tymin > tmin) tmin = tymin;

256 if(tymax < tmax) tmax = tymax;

257

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

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

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

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

262

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

264 return false;

265 }

266

267 return tmin < t1 && tmax > t0;

268 }

269

270 float4 reflect(float4 v, float4 n)

271 {

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

273 }

274

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

276 {

277 float4 res;

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

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

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

281 res.w = 0.0;

282 return res;

283 }

284

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

286 {

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

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

289 }

290

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

292 {

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

294

295 // calculate area of the whole triangle

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

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

298 float4 xv1v2 = cross(v1, v2);

299

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

301 if(area < EPSILON) {

302 return bc;

303 }

304

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

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

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

308

309 // calculate the area of each sub-triangle

310 float4 x12 = cross(pv1, pv2);

311 float4 x20 = cross(pv2, pv0);

312 float4 x01 = cross(pv0, pv1);

313

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

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

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

317

318 bc.x = a0 / area;

319 bc.y = a1 / area;

320 bc.z = a2 / area;

321 return bc;

322 }

323

324 float mean(float4 v)

325 {

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

327 }