clray

changeset 16:9e4a28063394

cl compiler segfaults...
author John Tsiombikas <nuclear@member.fsf.org>
date Mon, 09 Aug 2010 04:18:21 +0100
parents 754faf15ba36
children 074a64b9d6bd
files rt.cl src/mesh.cc src/rt.cc test.cl
diffstat 4 files changed, 124 insertions(+), 68 deletions(-) [+]
line diff
     1.1 --- a/rt.cl	Sun Aug 08 09:51:45 2010 +0100
     1.2 +++ b/rt.cl	Mon Aug 09 04:18:21 2010 +0100
     1.3 @@ -4,6 +4,7 @@
     1.4  	int xsz, ysz;
     1.5  	int num_faces, num_lights;
     1.6  	int max_iter;
     1.7 +	float4 ambient;
     1.8  	int dbg;
     1.9  };
    1.10  
    1.11 @@ -34,6 +35,7 @@
    1.12  
    1.13  struct Ray {
    1.14  	float4 origin, dir;
    1.15 +	float energy, pad[3];
    1.16  };
    1.17  
    1.18  struct SurfPoint {
    1.19 @@ -43,14 +45,26 @@
    1.20  	global const struct Material *mat;
    1.21  };
    1.22  
    1.23 -#define EPSILON 1e-6
    1.24 +struct Scene {
    1.25 +	float4 ambient;
    1.26 +	global const struct Face *faces;
    1.27 +	int num_faces;
    1.28 +	global const struct Light *lights;
    1.29 +	int num_lights;
    1.30 +	global const struct Material *matlib;
    1.31 +};
    1.32  
    1.33 -float4 shade(struct Ray ray, struct SurfPoint sp,
    1.34 -		global const struct Light *lights, int num_lights);
    1.35 +#define MIN_ENERGY	0.001
    1.36 +#define EPSILON		1e-6
    1.37 +
    1.38 +float4 trace(struct Ray ray, struct Scene *scn);
    1.39 +float4 shade(struct Ray ray, struct Scene *scn, const struct SurfPoint *sp);
    1.40 +bool find_intersection(struct Ray ray, const struct Scene *scn, struct SurfPoint *sp);
    1.41  bool intersect(struct Ray ray, global const struct Face *face, struct SurfPoint *sp);
    1.42 +
    1.43  float4 reflect(float4 v, float4 n);
    1.44  float4 transform(float4 v, global const float *xform);
    1.45 -struct Ray transform_ray(global const struct Ray *ray, global const float *xform, global const float *invtrans);
    1.46 +void transform_ray(struct Ray *ray, global const float *xform, global const float *invtrans);
    1.47  float4 calc_bary(float4 pt, global const struct Face *face, float4 norm);
    1.48  
    1.49  kernel void render(global float4 *fb,
    1.50 @@ -64,79 +78,114 @@
    1.51  		global struct Face *outfaces)
    1.52  {
    1.53  	int idx = get_global_id(0);
    1.54 -	
    1.55 -	if(idx == 0) {
    1.56 -		for(int i=0; i<rinf->num_faces; i++) {
    1.57 -			outfaces[i] = faces[i];
    1.58 -		}
    1.59 -	}
    1.60  
    1.61 -	struct Ray ray = transform_ray(primrays + idx, xform, invtrans);
    1.62 +	struct Scene scn;
    1.63 +	scn.ambient = rinf->ambient;
    1.64 +	scn.faces = faces;
    1.65 +	scn.num_faces = rinf->num_faces;
    1.66 +	scn.lights = lights;
    1.67 +	scn.num_lights = rinf->num_lights;
    1.68 +	scn.matlib = matlib;
    1.69  
    1.70 -	struct SurfPoint sp, sp0;
    1.71 -	sp0.t = FLT_MAX;
    1.72 -	sp0.obj = 0;
    1.73 -	
    1.74 -	int max_faces = min(rinf->num_faces, rinf->dbg);
    1.75 +	struct Ray ray = primrays[idx];
    1.76 +	transform_ray(&ray, xform, invtrans);
    1.77  
    1.78 -	for(int i=0; i<max_faces; i++) {
    1.79 -		if(intersect(ray, faces + i, &sp) && sp.t < sp0.t) {
    1.80 -			sp0 = sp;
    1.81 -		}
    1.82 -	}
    1.83 -
    1.84 -	if(sp0.obj) {
    1.85 -		sp0.mat = matlib + sp0.obj->matid;
    1.86 -		fb[idx] = shade(ray, sp0, lights, rinf->num_lights);
    1.87 -	} else {
    1.88 -		fb[idx] = (float4)(0, 0, 0, 0);
    1.89 -	}
    1.90 +	fb[idx] = trace(ray, &scn);
    1.91  }
    1.92  
    1.93 -float4 shade(struct Ray ray, struct SurfPoint sp,
    1.94 -		global const struct Light *lights, int num_lights)
    1.95 +float4 trace(struct Ray ray, struct Scene *scn)
    1.96  {
    1.97 -	float4 norm = sp.norm;
    1.98 +	float4 color;
    1.99 +	struct SurfPoint sp;
   1.100 +
   1.101 +	if(find_intersection(ray, scn, &sp)) {
   1.102 +		color = shade(ray, scn, &sp);
   1.103 +	} else {
   1.104 +		color = (float4)(0, 0, 0, 0);
   1.105 +	}
   1.106 +	return color;
   1.107 +}
   1.108 +
   1.109 +float4 shade(struct Ray ray, struct Scene *scn, const struct SurfPoint *sp)
   1.110 +{
   1.111 +	float4 norm = sp->norm;
   1.112  	bool entering = true;
   1.113 +	struct Material mat = *sp->mat;
   1.114  
   1.115  	if(dot(ray.dir, norm) >= 0.0) {
   1.116  		norm = -norm;
   1.117  		entering = false;
   1.118  	}
   1.119  
   1.120 -	float4 dcol = (float4)(0.07, 0.07, 0.07, 0);
   1.121 +	float4 dcol = scn->ambient * mat.kd;
   1.122  	float4 scol = (float4)(0, 0, 0, 0);
   1.123  
   1.124 -	for(int i=0; i<num_lights; i++) {
   1.125 -		float4 ldir = normalize(lights[i].pos - sp.pos);
   1.126 -		float4 vdir = -normalize(ray.dir);
   1.127 -		float4 vref = reflect(vdir, norm);
   1.128 +	for(int i=0; i<scn->num_lights; i++) {
   1.129 +		float4 ldir = scn->lights[i].pos - sp->pos;
   1.130  
   1.131 -		float diff = fmax(dot(ldir, norm), 0.0f);
   1.132 -		float spec = powr(fmax(dot(ldir, vref), 0.0f), sp.mat->spow);
   1.133 +		struct Ray shadowray;
   1.134 +		shadowray.origin = sp->pos;
   1.135 +		shadowray.dir = ldir;
   1.136  
   1.137 -		dcol += sp.mat->kd * diff * lights[i].color;
   1.138 -		//scol += sp.mat->ks * spec * lights[i].color;
   1.139 +		if(!find_intersection(shadowray, scn, 0)) {
   1.140 +			ldir = normalize(ldir);
   1.141 +			float4 vdir = -normalize(ray.dir);
   1.142 +			float4 vref = reflect(vdir, norm);
   1.143 +
   1.144 +			float diff = fmax(dot(ldir, norm), 0.0f);
   1.145 +			dcol += mat.kd * diff * scn->lights[i].color;
   1.146 +
   1.147 +			//float spec = powr(fmax(dot(ldir, vref), 0.0f), mat.spow);
   1.148 +			//scol += mat.ks * spec * scn->lights[i].color;
   1.149 +		}
   1.150 +	}
   1.151 +
   1.152 +	float4 refl_col = mat.ks * mat.kr;
   1.153 +	float refl_coeff = (refl_col.x + refl_col.y + refl_col.z) / 3.0;
   1.154 +
   1.155 +	if(refl_coeff > MIN_ENERGY) {
   1.156 +		struct Ray refl_ray;
   1.157 +		refl_ray.origin = sp->pos;
   1.158 +		refl_ray.dir = reflect(-ray.dir, norm);
   1.159 +		refl_ray.energy *= refl_coeff;
   1.160 +
   1.161 +		scol += trace(refl_ray, scn) * refl_col;
   1.162  	}
   1.163  
   1.164  	return dcol + scol;
   1.165  }
   1.166  
   1.167 -float dot3(float4 a, float4 b)
   1.168 +
   1.169 +bool find_intersection(struct Ray ray, const struct Scene *scn, struct SurfPoint *spres)
   1.170  {
   1.171 -	return a.x * b.x + a.y * b.y + a.z * b.z;
   1.172 +	struct SurfPoint sp, sp0;
   1.173 +	sp0.t = 1.0;
   1.174 +	sp0.obj = 0;
   1.175 +
   1.176 +	for(int i=0; i<scn->num_faces; i++) {
   1.177 +		if(intersect(ray, scn->faces + i, &sp) && sp.t < sp0.t) {
   1.178 +			sp0 = sp;
   1.179 +		}
   1.180 +	}
   1.181 +
   1.182 +	if(!sp0.obj) {
   1.183 +		return false;
   1.184 +	}
   1.185 +
   1.186 +	if(spres) {
   1.187 +		*spres = sp0;
   1.188 +		spres->mat = scn->matlib + sp0.obj->matid;
   1.189 +	}
   1.190 +	return true;
   1.191  }
   1.192  
   1.193 -
   1.194 -bool intersect(struct Ray ray,
   1.195 -		global const struct Face *face,
   1.196 -		struct SurfPoint *sp)
   1.197 +bool intersect(struct Ray ray, global const struct Face *face, struct SurfPoint *sp)
   1.198  {
   1.199  	float4 origin = ray.origin;
   1.200  	float4 dir = ray.dir;
   1.201  	float4 norm = face->normal;
   1.202  
   1.203 -	float ndotdir = dot3(dir, norm);
   1.204 +	float ndotdir = dot(dir, norm);
   1.205  
   1.206  	if(fabs(ndotdir) <= EPSILON) {
   1.207  		return false;
   1.208 @@ -145,7 +194,7 @@
   1.209  	float4 pt = face->v[0].pos;
   1.210  	float4 vec = pt - origin;
   1.211  
   1.212 -	float ndotvec = dot3(norm, vec);
   1.213 +	float ndotvec = dot(norm, vec);
   1.214  	float t = ndotvec / ndotdir;
   1.215  
   1.216  	if(t < EPSILON || t > 1.0) {
   1.217 @@ -153,8 +202,6 @@
   1.218  	}
   1.219  	pt = origin + dir * t;
   1.220  
   1.221 -	if(pt.w < 0.0) return false;
   1.222 -
   1.223  
   1.224  	float4 bc = calc_bary(pt, face, norm);
   1.225  	float bc_sum = bc.x + bc.y + bc.z;
   1.226 @@ -188,12 +235,10 @@
   1.227  	return res;
   1.228  }
   1.229  
   1.230 -struct Ray transform_ray(global const struct Ray *ray, global const float *xform, global const float *invtrans)
   1.231 +void transform_ray(struct Ray *ray, global const float *xform, global const float *invtrans)
   1.232  {
   1.233 -	struct Ray res;
   1.234 -	res.origin = transform(ray->origin, xform);
   1.235 -	res.dir = transform(ray->dir, invtrans);
   1.236 -	return res;
   1.237 +	ray->origin = transform(ray->origin, xform);
   1.238 +	ray->dir = transform(ray->dir, invtrans);
   1.239  }
   1.240  
   1.241  float4 calc_bary(float4 pt, global const struct Face *face, float4 norm)
   1.242 @@ -205,7 +250,7 @@
   1.243  	float4 v2 = face->v[2].pos - face->v[0].pos;
   1.244  	float4 xv1v2 = cross(v1, v2);
   1.245  
   1.246 -	float area = fabs(dot3(xv1v2, norm)) * 0.5;
   1.247 +	float area = fabs(dot(xv1v2, norm)) * 0.5;
   1.248  	if(area < EPSILON) {
   1.249  		return bc;
   1.250  	}
   1.251 @@ -219,9 +264,9 @@
   1.252  	float4 x20 = cross(pv2, pv0);
   1.253  	float4 x01 = cross(pv0, pv1);
   1.254  
   1.255 -	float a0 = fabs(dot3(x12, norm)) * 0.5;
   1.256 -	float a1 = fabs(dot3(x20, norm)) * 0.5;
   1.257 -	float a2 = fabs(dot3(x01, norm)) * 0.5;
   1.258 +	float a0 = fabs(dot(x12, norm)) * 0.5;
   1.259 +	float a1 = fabs(dot(x20, norm)) * 0.5;
   1.260 +	float a2 = fabs(dot(x01, norm)) * 0.5;
   1.261  
   1.262  	bc.x = a0 / area;
   1.263  	bc.y = a1 / area;
     2.1 --- a/src/mesh.cc	Sun Aug 08 09:51:45 2010 +0100
     2.2 +++ b/src/mesh.cc	Mon Aug 09 04:18:21 2010 +0100
     2.3 @@ -28,6 +28,7 @@
     2.4  	CMD(KA),		\
     2.5  	CMD(KD),		\
     2.6  	CMD(KS),		\
     2.7 +	CMD(KR),		\
     2.8  	CMD(NS),		\
     2.9  	CMD(NI),		\
    2.10  	CMD(D),			\
    2.11 @@ -84,6 +85,7 @@
    2.12  	float shininess;	// Ns
    2.13  	float ior;			// Ni
    2.14  	float alpha;		// d, Tr
    2.15 +	float refl;			// Kr (my extesnsion)
    2.16  
    2.17  	string tex_dif, tex_spec, tex_shin, tex_alpha;	// map_Kd, map_Ks, map_Ns, map_d
    2.18  	string tex_refl;	// refl -type sphere|cube file
    2.19 @@ -316,7 +318,7 @@
    2.20  					mat.ks[2] = vmtl[i].specular.z;
    2.21  
    2.22  					mat.kt = 1.0 - vmtl[i].alpha;
    2.23 -					mat.kr = 0.0;	// TODO
    2.24 +					mat.kr = vmtl[i].refl;
    2.25  					mat.spow = vmtl[i].shininess;
    2.26  
    2.27  					matlib.push_back(mat);
    2.28 @@ -487,6 +489,12 @@
    2.29  			parse_color(&mat.specular);
    2.30  			break;
    2.31  
    2.32 +		case CMD_KR:
    2.33 +			if((tok = strtok(0, SEP)) && is_float(tok)) {
    2.34 +				mat.refl = atof(tok);
    2.35 +			}
    2.36 +			break;
    2.37 +
    2.38  		case CMD_NS:
    2.39  			if((tok = strtok(0, SEP)) && is_float(tok)) {
    2.40  				mat.shininess = atof(tok);
     3.1 --- a/src/rt.cc	Sun Aug 08 09:51:45 2010 +0100
     3.2 +++ b/src/rt.cc	Mon Aug 09 04:18:21 2010 +0100
     3.3 @@ -25,11 +25,13 @@
     3.4  	int xsz, ysz;
     3.5  	int num_faces, num_lights;
     3.6  	int max_iter;
     3.7 +	float ambient[4];
     3.8  	int dbg;
     3.9  };
    3.10  
    3.11  struct Ray {
    3.12  	float origin[4], dir[4];
    3.13 +	float energy;
    3.14  };
    3.15  
    3.16  struct Light {
    3.17 @@ -45,7 +47,7 @@
    3.18  static int global_size;
    3.19  
    3.20  static Light lightlist[] = {
    3.21 -	{{-10, 10, -20, 0}, {1, 1, 1, 1}}
    3.22 +	{{-10, 13, -20, 0}, {1, 1, 1, 1}}
    3.23  };
    3.24  
    3.25  
    3.26 @@ -55,6 +57,9 @@
    3.27  bool init_renderer(int xsz, int ysz, Scene *scn)
    3.28  {
    3.29  	// render info
    3.30 +	rinf.ambient[0] = rinf.ambient[1] = rinf.ambient[2] = 0.075;
    3.31 +	rinf.ambient[3] = 0.0;
    3.32 +
    3.33  	rinf.xsz = xsz;
    3.34  	rinf.ysz = ysz;
    3.35  	rinf.num_faces = scn->get_num_faces();
    3.36 @@ -98,6 +103,10 @@
    3.37  		return false;
    3.38  	}
    3.39  
    3.40 +	if(!prog->build()) {
    3.41 +		return false;
    3.42 +	}
    3.43 +
    3.44  	delete [] prim_rays;
    3.45  
    3.46  	global_size = xsz * ysz;
    3.47 @@ -250,7 +259,7 @@
    3.48  	py *= 100.0;
    3.49  	pz *= 100.0;
    3.50  
    3.51 -	Ray ray = {{0, 0, 0, 1}, {px, py, -pz, 1}};
    3.52 +	Ray ray = {{0, 0, 0, 1}, {px, py, -pz, 1}, 1.0};
    3.53  	return ray;
    3.54  }
    3.55  
     4.1 --- a/test.cl	Sun Aug 08 09:51:45 2010 +0100
     4.2 +++ /dev/null	Thu Jan 01 00:00:00 1970 +0000
     4.3 @@ -1,6 +0,0 @@
     4.4 -__kernel void test(__global int *dst, __global const int *src, __global const int foo)
     4.5 -{
     4.6 -	int idx = get_global_id(0);
     4.7 -
     4.8 -	dst[idx] = src[idx] * foo;
     4.9 -}