 
 #define square(x) ((x)*(x))
 #define MAX_SPEED 4
+
+union boid_pack {
+	char buffer[sizeof(boid)];
+	boid my_boid;
+};
+
+texture<char, 1, cudaReadModeElementType> boids_texture;
+#define boid_tex_byte(offset) tex1Dfetch(boids_texture, (offset))
+#define boid_from_texture(offset) (((boid_pack) { { \
+	boid_tex_byte((offset) + 0), boid_tex_byte(offset + 1), \
+	boid_tex_byte((offset) + 2), boid_tex_byte(offset + 3), \
+	boid_tex_byte((offset) + 4), boid_tex_byte(offset + 5), \
+	boid_tex_byte((offset) + 6), boid_tex_byte(offset + 7), \
+	boid_tex_byte((offset) + 8), boid_tex_byte(offset + 9), \
+	boid_tex_byte((offset) + 10), boid_tex_byte(offset + 11), \
+	boid_tex_byte((offset) + 12), boid_tex_byte(offset + 13), \
+	boid_tex_byte((offset) + 14), boid_tex_byte(offset + 15), \
+	boid_tex_byte((offset) + 16), boid_tex_byte(offset + 17), \
+	boid_tex_byte((offset) + 18), boid_tex_byte(offset + 19), \
+	boid_tex_byte((offset) + 20), boid_tex_byte(offset + 21), \
+	boid_tex_byte((offset) + 22), boid_tex_byte(offset + 23) \
+	} } ).my_boid)
 
 #define check_cuda_error() {\
 	if (cudaError_t e = cudaGetLastError()) { \

	}
	cudaMemcpy(d_boids, sp->boids, boids_bytes, cudaMemcpyHostToDevice);
	check_cuda_error();
	reload_distance_cache(d_distance_cache, d_boids, sp->n);
	calculate_all_forces(d_boids, sp->n, sp->eps, d_distance_cache);
	apply_all_forces(d_boids, sp->n, sp->dt, sp->width, sp->height);
	cudaThreadSynchronize();

	assert(n > 0);
	cudaMalloc((void**) &d_boids, boids_bytes);
	assert(d_boids);
	cudaBindTexture(0, boids_texture, d_boids, n * sizeof(boid));
	check_cuda_error();
	return d_boids;
}

