 	check_cuda_error();
 }
 
-__global__ static void count_distance(const boid *boids, float *distance,
-		const int n) {
+__global__ static void count_distance(float *distance, const int n) {
    int ix = blockIdx.x * blockDim.x + threadIdx.x,
-	   iy = blockIdx.y * blockDim.y + threadIdx.y;
+	   iy = blockIdx.y * blockDim.y + threadIdx.y,
+	   off_a = ix * sizeof(boid), off_b = iy * sizeof(boid);
+   boid b = boid_from_texture(off_b);
+   boid a = boid_from_texture(off_a);
    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) {
+	   distance[ix + n * iy] = square(b.y - a.y) + square(b.x - a.x);
+}
+
+static void reload_distance_cache(float *d_cache, 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);
+	count_distance<<<blocks, threads>>>(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);
	return d_boids;
}


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