 		   square(boids[iy].x - boids[ix].x);
 }
 
-void reload_distance_cache(float *d_cache, boid *boids, int n) {
-	int blocksize = 16, boids_bytes = n * sizeof(boid);
+void reload_distance_cache(float *d_cache, boid *d_boids, int n) {
+	const int blocksize = 16;
 	dim3 threads(blocksize, blocksize);
 	dim3 blocks(n / blocksize, n / blocksize);
-	static boid *d_boids = NULL;
-	if (!d_boids)
-		cudaMalloc((void**) &d_boids, boids_bytes);
-	cudaMemcpy(d_boids, boids, boids_bytes, cudaMemcpyHostToDevice);
 	count_distance<<<blocks, threads>>>(d_boids, d_cache, n);
-	cudaThreadSynchronize();
 }
 
 __device__ void separation(boid *boids, int self, const int *neighbours,

	return d_distance_cache;
}

void simulate(simulation_params *sp) {
	static float *d_distance_cache = NULL;
	if (!d_distance_cache)
		d_distance_cache = prepare_distance_cache(sp->boids, sp->n);
	reload_distance_cache(d_distance_cache, sp->boids, sp->n);
	calculate_all_forces(sp->boids, sp->n, sp->eps, d_distance_cache);
	apply_all_forces(sp->boids, sp->n, sp->dt, sp->width, sp->height);
}

		boids[ix].y += height;
}

void apply_all_forces(boid* d_boids, int n, float dt, int width, int height) {
	const int blocksize = 64;
	dim3 threads(blocksize), blocks(n / blocksize);
	apply_forces<<<blocks, threads>>>(d_boids, dt, width, height);
	check_cuda_error();
}

static float* prepare_distance_cache(int n) {
	float *d_distance_cache = NULL;
	assert(n > 0);
	cudaMalloc((void**) &d_distance_cache, n * n * sizeof(float));
