 		boids[ix].y += height;
 }
 
-void apply_all_forces(boid* h_boids, int n, float dt, int width, int height) {
-	int blocksize = 64, boids_bytes = n * sizeof(boid);
+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);
-	static boid *d_boids = NULL;
-	if (!d_boids) {
-		cudaMalloc((void**) &d_boids, boids_bytes);
-		assert(d_boids);
-	}
-	check_cuda_error();
-	cudaMemcpy(d_boids, h_boids, boids_bytes, cudaMemcpyHostToDevice);
-	check_cuda_error();
 	apply_forces<<<blocks, threads>>>(d_boids, dt, width, height);
 	check_cuda_error();
-	cudaThreadSynchronize();
-	check_cuda_error();
-	cudaMemcpy(h_boids, d_boids, boids_bytes, cudaMemcpyDeviceToHost);
-	check_cuda_error();
-}
-
-static float* prepare_distance_cache(boid *_boids, int n) {
+}
+
+static float* prepare_distance_cache(int n) {
 	float *d_distance_cache = NULL;
 	assert(n > 0);
 	cudaMalloc((void**) &d_distance_cache, n * n * sizeof(float));

	}
}

void calculate_all_forces(boid* h_boids, int n, int eps, float *d_distance_cache) {
	int blocksize = 64, neighbours_bytes = n * n * sizeof(int),
		boids_bytes = n * sizeof(boid);
	dim3 threads(blocksize), blocks(n / blocksize);
	static boid *d_boids = NULL;
	static int *d_neighbours = NULL;
	if (!d_boids) {
		cudaMalloc((void**) &d_boids, boids_bytes);
		cudaMalloc((void**) &d_neighbours, neighbours_bytes);
		assert(d_boids && d_neighbours);
	}
	check_cuda_error();
	for (int i = 0; i < n; ++i)
		find_neighbours(d_neighbours + n * i, n, i, d_distance_cache, eps);
	check_cuda_error();
	cudaMemcpy(d_boids, h_boids, boids_bytes, cudaMemcpyHostToDevice);
	check_cuda_error();
	calculate_forces<<<blocks, threads>>>(d_boids,
			d_distance_cache, n, d_neighbours);
	check_cuda_error();
	cudaThreadSynchronize();
	check_cuda_error();
	cudaMemcpy(h_boids, d_boids, boids_bytes, cudaMemcpyDeviceToHost);
	check_cuda_error();
}


		   square(boids[iy].x - boids[ix].x);
}

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);
	count_distance<<<blocks, threads>>>(d_boids, d_cache, n);
}

__device__ void separation(boid *boids, int self, const int *neighbours,
