clray

annotate rt.cl @ 10:2da57bedc550

find changesets by author, revision, files, or words in the commit message
added vs project
author John Tsiombikas
date Fri, 30 Jul 2010 19:28:18 +0100
parents deaf85acf6af
children 85fd61f374d9
rev   line source
nuclear@2 1 struct RendInfo {
nuclear@2 2 int xsz, ysz;
nuclear@9 3 int num_faces, num_lights;
nuclear@2 4 int max_iter;
nuclear@2 5 };
nuclear@2 6
nuclear@9 7 struct Vertex {
nuclear@2 8 float4 pos;
nuclear@9 9 float4 normal;
nuclear@9 10 float2 tex;
nuclear@9 11 };
nuclear@9 12
nuclear@9 13 struct Face {
nuclear@9 14 struct Vertex v[3];
nuclear@9 15 float4 normal;
nuclear@9 16 int matid;
nuclear@9 17 };
nuclear@9 18
nuclear@9 19 struct Material {
nuclear@5 20 float4 kd, ks;
nuclear@9 21 float kr, kt;
nuclear@9 22 float spow;
nuclear@2 23 };
nuclear@2 24
nuclear@3 25 struct Light {
nuclear@3 26 float4 pos, color;
nuclear@3 27 };
nuclear@3 28
nuclear@2 29 struct Ray {
nuclear@2 30 float4 origin, dir;
nuclear@2 31 };
nuclear@2 32
nuclear@2 33 struct SurfPoint {
nuclear@2 34 float t;
nuclear@8 35 float4 pos, norm;
nuclear@9 36 global const struct Face *obj;
nuclear@9 37 global const struct Material *mat;
nuclear@2 38 };
nuclear@2 39
nuclear@2 40 #define EPSILON 1e-6
nuclear@2 41
nuclear@5 42 float4 shade(struct Ray ray, struct SurfPoint sp,
nuclear@5 43 global const struct Light *lights, int num_lights);
nuclear@9 44 bool intersect(struct Ray ray, global const struct Face *face, struct SurfPoint *sp);
nuclear@8 45 float4 reflect(float4 v, float4 n);
nuclear@8 46 float4 transform(float4 v, global const float *xform);
nuclear@8 47 struct Ray transform_ray(global const struct Ray *ray, global const float *xform);
nuclear@9 48 float4 calc_bary(float4 pt, global const struct Face *face);
nuclear@2 49
nuclear@4 50
nuclear@4 51 kernel void render(global float4 *fb,
nuclear@4 52 global const struct RendInfo *rinf,
nuclear@9 53 global const struct Face *faces,
nuclear@9 54 global const struct Material *matlib,
nuclear@4 55 global const struct Light *lights,
nuclear@7 56 global const struct Ray *primrays,
nuclear@8 57 global const float *xform)
nuclear@2 58 {
nuclear@2 59 int idx = get_global_id(0);
nuclear@2 60
nuclear@8 61 struct Ray ray = transform_ray(primrays + idx, xform);
nuclear@8 62
nuclear@4 63 struct SurfPoint sp, sp0;
nuclear@4 64 sp0.t = FLT_MAX;
nuclear@5 65 sp0.obj = 0;
nuclear@4 66
nuclear@9 67 for(int i=0; i<rinf->num_faces; i++) {
nuclear@9 68 if(intersect(ray, faces + i, &sp) && sp.t < sp0.t) {
nuclear@4 69 sp0 = sp;
nuclear@4 70 }
nuclear@2 71 }
nuclear@3 72
nuclear@5 73 if(sp0.obj) {
nuclear@9 74 sp0.mat = matlib + sp0.obj->matid;
nuclear@5 75 fb[idx] = shade(ray, sp0, lights, rinf->num_lights);
nuclear@5 76 } else {
nuclear@5 77 fb[idx] = (float4)(0, 0, 0, 0);
nuclear@5 78 }
nuclear@4 79 }
nuclear@4 80
nuclear@5 81 float4 shade(struct Ray ray, struct SurfPoint sp,
nuclear@5 82 global const struct Light *lights, int num_lights)
nuclear@4 83 {
nuclear@8 84 float4 dcol = (float4)(0, 0, 0, 0);
nuclear@8 85 float4 scol = (float4)(0, 0, 0, 0);
nuclear@5 86
nuclear@5 87 for(int i=0; i<num_lights; i++) {
nuclear@8 88 float4 ldir = normalize(lights[i].pos - sp.pos);
nuclear@8 89 float4 vdir = -normalize(ray.dir);
nuclear@8 90 float4 vref = reflect(vdir, sp.norm);
nuclear@5 91
nuclear@5 92 float diff = fmax(dot(ldir, sp.norm), 0.0f);
nuclear@9 93 float spec = powr(fmax(dot(ldir, vref), 0.0f), sp.mat->spow);
nuclear@5 94
nuclear@9 95 dcol += sp.mat->kd * diff * lights[i].color;
nuclear@9 96 scol += sp.mat->ks * spec * lights[i].color;
nuclear@5 97 }
nuclear@5 98
nuclear@8 99 return dcol + scol;
nuclear@2 100 }
nuclear@2 101
nuclear@2 102 bool intersect(struct Ray ray,
nuclear@9 103 global const struct Face *face,
nuclear@2 104 struct SurfPoint *sp)
nuclear@2 105 {
nuclear@9 106 float ndotdir = dot(face->normal, ray.dir);
nuclear@9 107 if(fabs(ndotdir) <= EPSILON) {
nuclear@9 108 return false;
nuclear@9 109 }
nuclear@2 110
nuclear@9 111 float4 pt = face->v[0].pos;
nuclear@9 112 float4 vec = pt - ray.origin;
nuclear@2 113
nuclear@9 114 float ndotvec = dot(face->normal, vec);
nuclear@9 115 float t = ndotvec / ndotdir;
nuclear@2 116
nuclear@2 117 if(t < EPSILON || t > 1.0) {
nuclear@2 118 return false;
nuclear@2 119 }
nuclear@9 120 pt = ray.origin + ray.dir * t;
nuclear@9 121
nuclear@9 122 float4 bc = calc_bary(pt, face);
nuclear@9 123 float bc_sum = bc.x + bc.y + bc.z;
nuclear@9 124
nuclear@9 125 if(bc_sum < -EPSILON || bc_sum > 1.0) {
nuclear@9 126 return false;
nuclear@9 127 }
nuclear@2 128
nuclear@2 129 sp->t = t;
nuclear@9 130 sp->pos = pt;
nuclear@9 131 sp->norm = face->normal;
nuclear@9 132 sp->obj = face;
nuclear@2 133 return true;
nuclear@2 134 }
nuclear@5 135
nuclear@8 136 float4 reflect(float4 v, float4 n)
nuclear@5 137 {
nuclear@5 138 return 2.0f * dot(v, n) * n - v;
nuclear@5 139 }
nuclear@8 140
nuclear@8 141 float4 transform(float4 v, global const float *xform)
nuclear@8 142 {
nuclear@8 143 float4 res;
nuclear@8 144 res.x = v.x * xform[0] + v.y * xform[4] + v.z * xform[8] + xform[12];
nuclear@8 145 res.y = v.x * xform[1] + v.y * xform[5] + v.z * xform[9] + xform[13];
nuclear@8 146 res.z = v.x * xform[2] + v.y * xform[6] + v.z * xform[10] + xform[14];
nuclear@8 147 res.w = 1.0;
nuclear@8 148 return res;
nuclear@8 149 }
nuclear@8 150
nuclear@8 151 struct Ray transform_ray(global const struct Ray *ray, global const float *xform)
nuclear@8 152 {
nuclear@8 153 struct Ray res;
nuclear@8 154 float rot[16];
nuclear@8 155
nuclear@8 156 for(int i=0; i<16; i++) {
nuclear@8 157 rot[i] = xform[i];
nuclear@8 158 }
nuclear@8 159 rot[3] = rot[7] = rot[11] = rot[12] = rot[13] = rot[14] = 0.0f;
nuclear@8 160 rot[15] = 1.0f;
nuclear@8 161
nuclear@8 162 res.origin = transform(ray->origin, xform);
nuclear@8 163 res.dir = transform(ray->dir, xform);
nuclear@8 164 return res;
nuclear@8 165 }
nuclear@9 166
nuclear@9 167 float4 calc_bary(float4 pt, global const struct Face *face)
nuclear@9 168 {
nuclear@9 169 float4 bc = {0, 0, 0, 0};
nuclear@9 170
nuclear@9 171 float4 vi = face->v[1].pos - face->v[0].pos;
nuclear@9 172 float4 vj = face->v[2].pos - face->v[0].pos;
nuclear@9 173 float area = fabs(dot(cross(vi, vj), face->normal) / 2.0);
nuclear@9 174 if(area < EPSILON) {
nuclear@9 175 return bc;
nuclear@9 176 }
nuclear@9 177
nuclear@9 178 float4 pv0 = face->v[0].pos - pt;
nuclear@9 179 float4 pv1 = face->v[1].pos - pt;
nuclear@9 180 float4 pv2 = face->v[2].pos - pt;
nuclear@9 181
nuclear@9 182 // calculate the areas of each sub-triangle
nuclear@9 183 float a0 = fabs(dot(cross(pv1, pv2), face->normal) / 2.0);
nuclear@9 184 float a1 = fabs(dot(cross(pv2, pv0), face->normal) / 2.0);
nuclear@9 185 float a2 = fabs(dot(cross(pv0, pv1), face->normal) / 2.0);
nuclear@9 186
nuclear@9 187 bc.x = a0 / area;
nuclear@9 188 bc.y = a1 / area;
nuclear@9 189 bc.z = a2 / area;
nuclear@9 190 return bc;
nuclear@9 191 }