 	}
 	cudaMemcpy(d_boids, sp->boids, boids_bytes, cudaMemcpyHostToDevice);
 	check_cuda_error();
-	reload_distance_cache(d_distance_cache, d_boids, sp->n);
+	reload_distance_cache(d_distance_cache, 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();

	check_cuda_error();
}

__global__ static void count_distance(const boid *boids, float *distance,
		const int n) {
   int ix = blockIdx.x * blockDim.x + threadIdx.x,
	   iy = blockIdx.y * blockDim.y + threadIdx.y;
   if (ix < n && iy < n)
	   distance[ix + n * iy] = square(boids[iy].y - boids[ix].y) +
		   square(boids[iy].x - boids[ix].x);
}

static 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__ static void separation(boid *boids, int self, const int *neighbours,

	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;
}

