clray
diff rt.cl @ 12:85fd61f374d9
fixed the bloody intersection bug
author | John Tsiombikas <nuclear@member.fsf.org> |
---|---|
date | Tue, 03 Aug 2010 13:06:59 +0100 |
parents | a09622aaa043 |
children | 754faf15ba36 |
line diff
1.1 --- a/rt.cl Sat Jul 31 22:23:57 2010 +0100 1.2 +++ b/rt.cl Tue Aug 03 13:06:59 2010 +0100 1.3 @@ -1,3 +1,5 @@ 1.4 +/* vim: set ft=opencl:ts=4:sw=4 */ 1.5 + 1.6 struct RendInfo { 1.7 int xsz, ysz; 1.8 int num_faces, num_lights; 1.9 @@ -7,19 +9,22 @@ 1.10 struct Vertex { 1.11 float4 pos; 1.12 float4 normal; 1.13 - float2 tex; 1.14 + float4 tex; 1.15 + float4 padding; 1.16 }; 1.17 1.18 struct Face { 1.19 struct Vertex v[3]; 1.20 float4 normal; 1.21 int matid; 1.22 + int padding[3]; 1.23 }; 1.24 1.25 struct Material { 1.26 float4 kd, ks; 1.27 float kr, kt; 1.28 float spow; 1.29 + float padding; 1.30 }; 1.31 1.32 struct Light { 1.33 @@ -32,7 +37,7 @@ 1.34 1.35 struct SurfPoint { 1.36 float t; 1.37 - float4 pos, norm; 1.38 + float4 pos, norm, dbg; 1.39 global const struct Face *obj; 1.40 global const struct Material *mat; 1.41 }; 1.42 @@ -44,9 +49,8 @@ 1.43 bool intersect(struct Ray ray, global const struct Face *face, struct SurfPoint *sp); 1.44 float4 reflect(float4 v, float4 n); 1.45 float4 transform(float4 v, global const float *xform); 1.46 -struct Ray transform_ray(global const struct Ray *ray, global const float *xform); 1.47 -float4 calc_bary(float4 pt, global const struct Face *face); 1.48 - 1.49 +struct Ray transform_ray(global const struct Ray *ray, global const float *xform, global const float *invtrans); 1.50 +float4 calc_bary(float4 pt, global const struct Face *face, float4 norm); 1.51 1.52 kernel void render(global float4 *fb, 1.53 global const struct RendInfo *rinf, 1.54 @@ -54,11 +58,12 @@ 1.55 global const struct Material *matlib, 1.56 global const struct Light *lights, 1.57 global const struct Ray *primrays, 1.58 - global const float *xform) 1.59 + global const float *xform, 1.60 + global const float *invtrans) 1.61 { 1.62 int idx = get_global_id(0); 1.63 1.64 - struct Ray ray = transform_ray(primrays + idx, xform); 1.65 + struct Ray ray = transform_ray(primrays + idx, xform, invtrans); 1.66 1.67 struct SurfPoint sp, sp0; 1.68 sp0.t = FLT_MAX; 1.69 @@ -81,61 +86,86 @@ 1.70 float4 shade(struct Ray ray, struct SurfPoint sp, 1.71 global const struct Light *lights, int num_lights) 1.72 { 1.73 + float4 norm = sp.norm; 1.74 + bool entering = true; 1.75 + 1.76 + if(dot(ray.dir, norm) >= 0.0) { 1.77 + norm = -norm; 1.78 + entering = false; 1.79 + } 1.80 + 1.81 float4 dcol = (float4)(0, 0, 0, 0); 1.82 float4 scol = (float4)(0, 0, 0, 0); 1.83 1.84 for(int i=0; i<num_lights; i++) { 1.85 float4 ldir = normalize(lights[i].pos - sp.pos); 1.86 float4 vdir = -normalize(ray.dir); 1.87 - float4 vref = reflect(vdir, sp.norm); 1.88 + float4 vref = reflect(vdir, norm); 1.89 1.90 - float diff = fmax(dot(ldir, sp.norm), 0.0f); 1.91 + float diff = fmax(dot(ldir, norm), 0.0f); 1.92 float spec = powr(fmax(dot(ldir, vref), 0.0f), sp.mat->spow); 1.93 1.94 dcol += sp.mat->kd * diff * lights[i].color; 1.95 - scol += sp.mat->ks * spec * lights[i].color; 1.96 + //scol += sp.mat->ks * spec * lights[i].color; 1.97 } 1.98 1.99 return dcol + scol; 1.100 } 1.101 1.102 +float dot3(float4 a, float4 b) 1.103 +{ 1.104 + return a.x * b.x + a.y * b.y + a.z * b.z; 1.105 +} 1.106 + 1.107 + 1.108 bool intersect(struct Ray ray, 1.109 global const struct Face *face, 1.110 struct SurfPoint *sp) 1.111 { 1.112 - float ndotdir = dot(face->normal, ray.dir); 1.113 + float4 origin = ray.origin; 1.114 + float4 dir = ray.dir; 1.115 + float4 norm = face->normal; 1.116 + 1.117 + float ndotdir = dot3(dir, norm); 1.118 + 1.119 if(fabs(ndotdir) <= EPSILON) { 1.120 return false; 1.121 } 1.122 1.123 float4 pt = face->v[0].pos; 1.124 - float4 vec = pt - ray.origin; 1.125 + float4 vec = pt - origin; 1.126 1.127 - float ndotvec = dot(face->normal, vec); 1.128 + float ndotvec = dot3(norm, vec); 1.129 float t = ndotvec / ndotdir; 1.130 1.131 if(t < EPSILON || t > 1.0) { 1.132 return false; 1.133 } 1.134 - pt = ray.origin + ray.dir * t; 1.135 + pt = origin + dir * t; 1.136 1.137 - float4 bc = calc_bary(pt, face); 1.138 + if(pt.w < 0.0) return false; 1.139 + 1.140 + 1.141 + float4 bc = calc_bary(pt, face, norm); 1.142 float bc_sum = bc.x + bc.y + bc.z; 1.143 1.144 - if(bc_sum < -EPSILON || bc_sum > 1.0) { 1.145 + if(bc_sum < 0.0 || bc_sum > 1.0 + EPSILON) { 1.146 return false; 1.147 + bc *= 1.2; 1.148 } 1.149 1.150 sp->t = t; 1.151 sp->pos = pt; 1.152 - sp->norm = face->normal; 1.153 + sp->norm = norm; 1.154 sp->obj = face; 1.155 + sp->dbg = bc; 1.156 return true; 1.157 } 1.158 1.159 float4 reflect(float4 v, float4 n) 1.160 { 1.161 - return 2.0f * dot(v, n) * n - v; 1.162 + float4 res = 2.0f * dot(v, n) * n - v; 1.163 + return res; 1.164 } 1.165 1.166 float4 transform(float4 v, global const float *xform) 1.167 @@ -144,33 +174,28 @@ 1.168 res.x = v.x * xform[0] + v.y * xform[4] + v.z * xform[8] + xform[12]; 1.169 res.y = v.x * xform[1] + v.y * xform[5] + v.z * xform[9] + xform[13]; 1.170 res.z = v.x * xform[2] + v.y * xform[6] + v.z * xform[10] + xform[14]; 1.171 - res.w = 1.0; 1.172 + res.w = 0.0; 1.173 return res; 1.174 } 1.175 1.176 -struct Ray transform_ray(global const struct Ray *ray, global const float *xform) 1.177 +struct Ray transform_ray(global const struct Ray *ray, global const float *xform, global const float *invtrans) 1.178 { 1.179 struct Ray res; 1.180 - float rot[16]; 1.181 - 1.182 - for(int i=0; i<16; i++) { 1.183 - rot[i] = xform[i]; 1.184 - } 1.185 - rot[3] = rot[7] = rot[11] = rot[12] = rot[13] = rot[14] = 0.0f; 1.186 - rot[15] = 1.0f; 1.187 - 1.188 res.origin = transform(ray->origin, xform); 1.189 - res.dir = transform(ray->dir, xform); 1.190 + res.dir = transform(ray->dir, invtrans); 1.191 return res; 1.192 } 1.193 1.194 -float4 calc_bary(float4 pt, global const struct Face *face) 1.195 +float4 calc_bary(float4 pt, global const struct Face *face, float4 norm) 1.196 { 1.197 - float4 bc = {0, 0, 0, 0}; 1.198 + float4 bc = (float4)(0, 0, 0, 0); 1.199 1.200 - float4 vi = face->v[1].pos - face->v[0].pos; 1.201 - float4 vj = face->v[2].pos - face->v[0].pos; 1.202 - float area = fabs(dot(cross(vi, vj), face->normal) / 2.0); 1.203 + // calculate area of the whole triangle 1.204 + float4 v1 = face->v[1].pos - face->v[0].pos; 1.205 + float4 v2 = face->v[2].pos - face->v[0].pos; 1.206 + float4 xv1v2 = cross(v1, v2); 1.207 + 1.208 + float area = fabs(dot3(xv1v2, norm)) * 0.5; 1.209 if(area < EPSILON) { 1.210 return bc; 1.211 } 1.212 @@ -179,10 +204,14 @@ 1.213 float4 pv1 = face->v[1].pos - pt; 1.214 float4 pv2 = face->v[2].pos - pt; 1.215 1.216 - // calculate the areas of each sub-triangle 1.217 - float a0 = fabs(dot(cross(pv1, pv2), face->normal) / 2.0); 1.218 - float a1 = fabs(dot(cross(pv2, pv0), face->normal) / 2.0); 1.219 - float a2 = fabs(dot(cross(pv0, pv1), face->normal) / 2.0); 1.220 + // calculate the area of each sub-triangle 1.221 + float4 x12 = cross(pv1, pv2); 1.222 + float4 x20 = cross(pv2, pv0); 1.223 + float4 x01 = cross(pv0, pv1); 1.224 + 1.225 + float a0 = fabs(dot3(x12, norm)) * 0.5; 1.226 + float a1 = fabs(dot3(x20, norm)) * 0.5; 1.227 + float a2 = fabs(dot3(x01, norm)) * 0.5; 1.228 1.229 bc.x = a0 / area; 1.230 bc.y = a1 / area;