diff --git a/src/kernels.cu b/src/kernels.cu
index b01bd38..2f3516a 100644
--- a/src/kernels.cu
+++ b/src/kernels.cu
@@ -66,11 +66,19 @@ void brute_force_kernel(double* x, double* y, double* z,
     double a_y_i = a_y[id];
     double a_z_i = a_z[id];
 
+    double temp_a_x = 0;
+    double temp_a_y = 0;
+    double temp_a_z = 0;
+
     calculate_force_device(x_old, y_old, z_old, m,
             x_old[id], y_old[id], z_old[id],
-            a_x[id], a_y[id], a_z[id],
+            temp_a_x, temp_a_y, temp_a_z,
             num_planets, eps2, G);
 
+    a_x[id] = temp_a_x;
+    a_y[id] = temp_a_y;
+    a_z[id] = temp_a_z;
+
     x[id] += v_x[id]*dt + a_x_i*0.5*dt*dt;
     y[id] += v_y[id]*dt + a_y_i*0.5*dt*dt;
     z[id] += v_z[id]*dt + a_z_i*0.5*dt*dt;
@@ -87,18 +95,6 @@ void brute_force_gpu_update(double* x, double* y, double* z,
         double* a_x, double* a_y, double* a_z,
         double* m, double G, double dt, int num_planets, double eps)
 {
-    double* x_old = (double*) malloc(num_planets*sizeof(double));
-    double* y_old = (double*) malloc(num_planets*sizeof(double));
-    double* z_old = (double*) malloc(num_planets*sizeof(double));
-
-    int i;
-    for(i=0; i<num_planets; i++)
-    {
-        x_old[i] = x[i];
-        y_old[i] = y[i];
-        z_old[i] = z[i];
-    }
-
     double *dev_x, *dev_y, *dev_z, *dev_x_old, *dev_y_old, *dev_z_old;
     double *dev_v_x, *dev_v_y, *dev_v_z, *dev_a_x, *dev_a_y, *dev_a_z;
     double *dev_m;
@@ -124,9 +120,9 @@ void brute_force_gpu_update(double* x, double* y, double* z,
     if( cudaMemcpy(dev_x, x, num_planets*sizeof(double), cudaMemcpyHostToDevice) != cudaSuccess ||
         cudaMemcpy(dev_y, y, num_planets*sizeof(double), cudaMemcpyHostToDevice) != cudaSuccess ||
         cudaMemcpy(dev_z, z, num_planets*sizeof(double), cudaMemcpyHostToDevice) != cudaSuccess ||
-        cudaMemcpy(dev_x_old, x_old, num_planets*sizeof(double), cudaMemcpyHostToDevice) != cudaSuccess ||
-        cudaMemcpy(dev_y_old, y_old, num_planets*sizeof(double), cudaMemcpyHostToDevice) != cudaSuccess ||
-        cudaMemcpy(dev_z_old, z_old, num_planets*sizeof(double), cudaMemcpyHostToDevice) != cudaSuccess ||
+        cudaMemcpy(dev_x_old, x, num_planets*sizeof(double), cudaMemcpyHostToDevice) != cudaSuccess ||
+        cudaMemcpy(dev_y_old, y, num_planets*sizeof(double), cudaMemcpyHostToDevice) != cudaSuccess ||
+        cudaMemcpy(dev_z_old, z, num_planets*sizeof(double), cudaMemcpyHostToDevice) != cudaSuccess ||
         cudaMemcpy(dev_v_x, v_x, num_planets*sizeof(double), cudaMemcpyHostToDevice) != cudaSuccess ||
         cudaMemcpy(dev_v_y, v_y, num_planets*sizeof(double), cudaMemcpyHostToDevice) != cudaSuccess ||
         cudaMemcpy(dev_v_z, v_z, num_planets*sizeof(double), cudaMemcpyHostToDevice) != cudaSuccess ||
@@ -139,7 +135,7 @@ void brute_force_gpu_update(double* x, double* y, double* z,
         exit(0);
     }
 
-    brute_force_kernel<<<ceil(num_planets/1024), 1024>>>(dev_x, dev_y, dev_z,
+    brute_force_kernel<<<ceil(num_planets/512), 512>>>(dev_x, dev_y, dev_z,
             dev_x_old, dev_y_old, dev_z_old,
             dev_v_x, dev_v_y, dev_v_z,
             dev_a_x, dev_a_y, dev_a_z,
@@ -182,8 +178,5 @@ void brute_force_gpu_update(double* x, double* y, double* z,
         exit(0);
     }
 
-    free(x_old);
-    free(y_old);
-    free(z_old);
 }