clray

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