clray

annotate rt.cl @ 21:bd6c2b25f6e7

find changesets by author, revision, files, or words in the commit message
fixed, now we need to start with optimizations
author John Tsiombikas <nuclear@member.fsf.org>
date Tue, 10 Aug 2010 07:24:18 +0100
parents 63a6b46f58a0
children 6c44e4b1726d
rev   line source
nuclear@12 1 /* vim: set ft=opencl:ts=4:sw=4 */
nuclear@12 2
nuclear@2 3 struct RendInfo {
nuclear@2 4 int xsz, ysz;
nuclear@9 5 int num_faces, num_lights;
nuclear@2 6 int max_iter;
nuclear@16 7 float4 ambient;
nuclear@2 8 };
nuclear@2 9
nuclear@9 10 struct Vertex {
nuclear@2 11 float4 pos;
nuclear@9 12 float4 normal;
nuclear@12 13 float4 tex;
nuclear@12 14 float4 padding;
nuclear@9 15 };
nuclear@9 16
nuclear@9 17 struct Face {
nuclear@9 18 struct Vertex v[3];
nuclear@9 19 float4 normal;
nuclear@9 20 int matid;
nuclear@12 21 int padding[3];
nuclear@9 22 };
nuclear@9 23
nuclear@9 24 struct Material {
nuclear@5 25 float4 kd, ks;
nuclear@9 26 float kr, kt;
nuclear@9 27 float spow;
nuclear@12 28 float padding;
nuclear@2 29 };
nuclear@2 30
nuclear@3 31 struct Light {
nuclear@3 32 float4 pos, color;
nuclear@3 33 };
nuclear@3 34
nuclear@2 35 struct Ray {
nuclear@2 36 float4 origin, dir;
nuclear@2 37 };
nuclear@2 38
nuclear@2 39 struct SurfPoint {
nuclear@2 40 float t;
nuclear@12 41 float4 pos, norm, dbg;
nuclear@9 42 global const struct Face *obj;
nuclear@19 43 struct Material mat;
nuclear@2 44 };
nuclear@2 45
nuclear@16 46 struct Scene {
nuclear@16 47 float4 ambient;
nuclear@16 48 global const struct Face *faces;
nuclear@16 49 int num_faces;
nuclear@16 50 global const struct Light *lights;
nuclear@16 51 int num_lights;
nuclear@16 52 global const struct Material *matlib;
nuclear@16 53 };
nuclear@2 54
nuclear@16 55 #define MIN_ENERGY 0.001
nuclear@21 56 #define EPSILON 1e-5
nuclear@16 57
nuclear@16 58 float4 shade(struct Ray ray, struct Scene *scn, const struct SurfPoint *sp);
nuclear@16 59 bool find_intersection(struct Ray ray, const struct Scene *scn, struct SurfPoint *sp);
nuclear@9 60 bool intersect(struct Ray ray, global const struct Face *face, struct SurfPoint *sp);
nuclear@16 61
nuclear@8 62 float4 reflect(float4 v, float4 n);
nuclear@8 63 float4 transform(float4 v, global const float *xform);
nuclear@16 64 void transform_ray(struct Ray *ray, global const float *xform, global const float *invtrans);
nuclear@12 65 float4 calc_bary(float4 pt, global const struct Face *face, float4 norm);
nuclear@19 66 float mean(float4 v);
nuclear@4 67
nuclear@4 68 kernel void render(global float4 *fb,
nuclear@4 69 global const struct RendInfo *rinf,
nuclear@9 70 global const struct Face *faces,
nuclear@9 71 global const struct Material *matlib,
nuclear@4 72 global const struct Light *lights,
nuclear@7 73 global const struct Ray *primrays,
nuclear@12 74 global const float *xform,
nuclear@20 75 global const float *invtrans)
nuclear@2 76 {
nuclear@2 77 int idx = get_global_id(0);
nuclear@2 78
nuclear@16 79 struct Scene scn;
nuclear@16 80 scn.ambient = rinf->ambient;
nuclear@16 81 scn.faces = faces;
nuclear@16 82 scn.num_faces = rinf->num_faces;
nuclear@16 83 scn.lights = lights;
nuclear@16 84 scn.num_lights = rinf->num_lights;
nuclear@16 85 scn.matlib = matlib;
nuclear@8 86
nuclear@16 87 struct Ray ray = primrays[idx];
nuclear@16 88 transform_ray(&ray, xform, invtrans);
nuclear@4 89
nuclear@19 90 float4 pixel = (float4)(0, 0, 0, 0);
nuclear@19 91 float4 energy = (float4)(1.0, 1.0, 1.0, 1.0);
nuclear@19 92 int iter = 0;
nuclear@19 93
nuclear@19 94 while(iter++ < rinf->max_iter && mean(energy) > MIN_ENERGY) {
nuclear@19 95 struct SurfPoint sp;
nuclear@19 96 if(find_intersection(ray, &scn, &sp)) {
nuclear@19 97 pixel += shade(ray, &scn, &sp) * energy;
nuclear@19 98
nuclear@19 99 float4 refl_col = sp.mat.ks * sp.mat.kr;
nuclear@19 100
nuclear@19 101 ray.origin = sp.pos;
nuclear@19 102 ray.dir = reflect(-ray.dir, sp.norm);
nuclear@19 103
nuclear@19 104 energy *= sp.mat.ks * sp.mat.kr;
nuclear@19 105 } else {
nuclear@19 106 iter = INT_MAX - 1; // to break out of the loop
nuclear@19 107 }
nuclear@17 108 }
nuclear@19 109
nuclear@19 110 fb[idx] = pixel;
nuclear@4 111 }
nuclear@4 112
nuclear@16 113 float4 shade(struct Ray ray, struct Scene *scn, const struct SurfPoint *sp)
nuclear@16 114 {
nuclear@16 115 float4 norm = sp->norm;
nuclear@12 116 bool entering = true;
nuclear@12 117
nuclear@12 118 if(dot(ray.dir, norm) >= 0.0) {
nuclear@12 119 norm = -norm;
nuclear@12 120 entering = false;
nuclear@12 121 }
nuclear@12 122
nuclear@19 123 float4 dcol = scn->ambient * sp->mat.kd;
nuclear@8 124 float4 scol = (float4)(0, 0, 0, 0);
nuclear@5 125
nuclear@16 126 for(int i=0; i<scn->num_lights; i++) {
nuclear@16 127 float4 ldir = scn->lights[i].pos - sp->pos;
nuclear@5 128
nuclear@16 129 struct Ray shadowray;
nuclear@16 130 shadowray.origin = sp->pos;
nuclear@16 131 shadowray.dir = ldir;
nuclear@5 132
nuclear@16 133 if(!find_intersection(shadowray, scn, 0)) {
nuclear@16 134 ldir = normalize(ldir);
nuclear@16 135 float4 vdir = -normalize(ray.dir);
nuclear@16 136 float4 vref = reflect(vdir, norm);
nuclear@16 137
nuclear@16 138 float diff = fmax(dot(ldir, norm), 0.0f);
nuclear@19 139 dcol += sp->mat.kd * diff * scn->lights[i].color;
nuclear@16 140
nuclear@20 141 float spec = powr(fmax(dot(ldir, vref), 0.0f), sp->mat.spow);
nuclear@20 142 scol += sp->mat.ks * spec * scn->lights[i].color;
nuclear@16 143 }
nuclear@16 144 }
nuclear@16 145
nuclear@8 146 return dcol + scol;
nuclear@2 147 }
nuclear@2 148
nuclear@16 149
nuclear@16 150 bool find_intersection(struct Ray ray, const struct Scene *scn, struct SurfPoint *spres)
nuclear@12 151 {
nuclear@16 152 struct SurfPoint sp, sp0;
nuclear@16 153 sp0.t = 1.0;
nuclear@16 154 sp0.obj = 0;
nuclear@16 155
nuclear@16 156 for(int i=0; i<scn->num_faces; i++) {
nuclear@16 157 if(intersect(ray, scn->faces + i, &sp) && sp.t < sp0.t) {
nuclear@16 158 sp0 = sp;
nuclear@16 159 }
nuclear@16 160 }
nuclear@16 161
nuclear@16 162 if(!sp0.obj) {
nuclear@16 163 return false;
nuclear@16 164 }
nuclear@16 165
nuclear@16 166 if(spres) {
nuclear@16 167 *spres = sp0;
nuclear@19 168 spres->mat = scn->matlib[sp0.obj->matid];
nuclear@16 169 }
nuclear@16 170 return true;
nuclear@12 171 }
nuclear@12 172
nuclear@16 173 bool intersect(struct Ray ray, global const struct Face *face, struct SurfPoint *sp)
nuclear@2 174 {
nuclear@12 175 float4 origin = ray.origin;
nuclear@12 176 float4 dir = ray.dir;
nuclear@12 177 float4 norm = face->normal;
nuclear@12 178
nuclear@16 179 float ndotdir = dot(dir, norm);
nuclear@12 180
nuclear@9 181 if(fabs(ndotdir) <= EPSILON) {
nuclear@9 182 return false;
nuclear@9 183 }
nuclear@2 184
nuclear@9 185 float4 pt = face->v[0].pos;
nuclear@12 186 float4 vec = pt - origin;
nuclear@2 187
nuclear@16 188 float ndotvec = dot(norm, vec);
nuclear@9 189 float t = ndotvec / ndotdir;
nuclear@2 190
nuclear@2 191 if(t < EPSILON || t > 1.0) {
nuclear@2 192 return false;
nuclear@2 193 }
nuclear@12 194 pt = origin + dir * t;
nuclear@9 195
nuclear@12 196
nuclear@12 197 float4 bc = calc_bary(pt, face, norm);
nuclear@9 198 float bc_sum = bc.x + bc.y + bc.z;
nuclear@9 199
nuclear@20 200 if(bc_sum < 1.0 - EPSILON || bc_sum > 1.0 + EPSILON) {
nuclear@9 201 return false;
nuclear@12 202 bc *= 1.2;
nuclear@9 203 }
nuclear@2 204
nuclear@2 205 sp->t = t;
nuclear@9 206 sp->pos = pt;
nuclear@21 207 sp->norm = normalize(face->v[0].normal * bc.x + face->v[1].normal * bc.y + face->v[2].normal * bc.z);
nuclear@9 208 sp->obj = face;
nuclear@12 209 sp->dbg = bc;
nuclear@2 210 return true;
nuclear@2 211 }
nuclear@5 212
nuclear@8 213 float4 reflect(float4 v, float4 n)
nuclear@5 214 {
nuclear@12 215 float4 res = 2.0f * dot(v, n) * n - v;
nuclear@12 216 return res;
nuclear@5 217 }
nuclear@8 218
nuclear@8 219 float4 transform(float4 v, global const float *xform)
nuclear@8 220 {
nuclear@8 221 float4 res;
nuclear@8 222 res.x = v.x * xform[0] + v.y * xform[4] + v.z * xform[8] + xform[12];
nuclear@8 223 res.y = v.x * xform[1] + v.y * xform[5] + v.z * xform[9] + xform[13];
nuclear@8 224 res.z = v.x * xform[2] + v.y * xform[6] + v.z * xform[10] + xform[14];
nuclear@12 225 res.w = 0.0;
nuclear@8 226 return res;
nuclear@8 227 }
nuclear@8 228
nuclear@16 229 void transform_ray(struct Ray *ray, global const float *xform, global const float *invtrans)
nuclear@8 230 {
nuclear@16 231 ray->origin = transform(ray->origin, xform);
nuclear@16 232 ray->dir = transform(ray->dir, invtrans);
nuclear@8 233 }
nuclear@9 234
nuclear@12 235 float4 calc_bary(float4 pt, global const struct Face *face, float4 norm)
nuclear@9 236 {
nuclear@12 237 float4 bc = (float4)(0, 0, 0, 0);
nuclear@9 238
nuclear@12 239 // calculate area of the whole triangle
nuclear@12 240 float4 v1 = face->v[1].pos - face->v[0].pos;
nuclear@12 241 float4 v2 = face->v[2].pos - face->v[0].pos;
nuclear@12 242 float4 xv1v2 = cross(v1, v2);
nuclear@12 243
nuclear@16 244 float area = fabs(dot(xv1v2, norm)) * 0.5;
nuclear@9 245 if(area < EPSILON) {
nuclear@9 246 return bc;
nuclear@9 247 }
nuclear@9 248
nuclear@9 249 float4 pv0 = face->v[0].pos - pt;
nuclear@9 250 float4 pv1 = face->v[1].pos - pt;
nuclear@9 251 float4 pv2 = face->v[2].pos - pt;
nuclear@9 252
nuclear@12 253 // calculate the area of each sub-triangle
nuclear@12 254 float4 x12 = cross(pv1, pv2);
nuclear@12 255 float4 x20 = cross(pv2, pv0);
nuclear@12 256 float4 x01 = cross(pv0, pv1);
nuclear@12 257
nuclear@16 258 float a0 = fabs(dot(x12, norm)) * 0.5;
nuclear@16 259 float a1 = fabs(dot(x20, norm)) * 0.5;
nuclear@16 260 float a2 = fabs(dot(x01, norm)) * 0.5;
nuclear@9 261
nuclear@9 262 bc.x = a0 / area;
nuclear@9 263 bc.y = a1 / area;
nuclear@9 264 bc.z = a2 / area;
nuclear@9 265 return bc;
nuclear@9 266 }
nuclear@19 267
nuclear@19 268 float mean(float4 v)
nuclear@19 269 {
nuclear@19 270 return native_divide(v.x + v.y + v.z, 3.0);
nuclear@19 271 }