 
     int block_count_nodes = (num_movable_nodes + thread_count - 1) / thread_count;
     computeElectricForce<<<block_count_nodes, thread_count, 0, stream_movable>>>(
+        num_bins_x, num_bins_y,
+        num_movable_impacted_bins_x, num_movable_impacted_bins_y,
+        field_map_x_tensor, field_map_y_tensor,
+        x_tensor, y_tensor,
+        node_size_x_tensor, node_size_y_tensor,
+        bin_center_x_tensor, bin_center_y_tensor,
+        xl, yl, xh, yh,
+        bin_size_x, bin_size_y,
+        num_movable_nodes,
+        grad_x_tensor, grad_y_tensor);
+
+    if (num_filler_nodes)
+    {
+        status = cudaStreamCreate(&stream_filler);
+        if (status != cudaSuccess)
+        {
+            printf("cudaStreamCreate failed for stream_filler\n");
+            fflush(stdout);
+            return 1;
+        }
+
+        block_count_nodes = (num_filler_nodes + thread_count - 1) / thread_count;
+        int offset = num_nodes - num_filler_nodes;
+        computeElectricForce<<<block_count_nodes, thread_count, 0, stream_filler>>>(
             num_bins_x, num_bins_y,
-            num_movable_impacted_bins_x, num_movable_impacted_bins_y,
+            num_filler_impacted_bins_x, num_filler_impacted_bins_y,
             field_map_x_tensor, field_map_y_tensor,
-            x_tensor, y_tensor,
-            node_size_x_tensor, node_size_y_tensor,
+            x_tensor + offset, y_tensor + offset,
+            node_size_x_tensor + offset, node_size_y_tensor + offset,
             bin_center_x_tensor, bin_center_y_tensor,
             xl, yl, xh, yh,
             bin_size_x, bin_size_y,
-            num_movable_nodes,
-            grad_x_tensor, grad_y_tensor
-            );
-
-    if (num_filler_nodes)
-    {
-        status = cudaStreamCreate(&stream_filler);
-        if (status != cudaSuccess)
-        {
-            printf("cudaStreamCreate failed for stream_filler\n");
-            fflush(stdout);
-            return 1;
-        }
-
-        block_count_nodes = (num_filler_nodes + thread_count - 1) / thread_count;
-        computeElectricForce<<<block_count_nodes, thread_count, 0, stream_filler>>>(
-                num_bins_x, num_bins_y,
-                num_filler_impacted_bins_x, num_filler_impacted_bins_y,
-                field_map_x_tensor, field_map_y_tensor,
-                x_tensor+num_nodes-num_filler_nodes, y_tensor+num_nodes-num_filler_nodes,
-                node_size_x_tensor+num_nodes-num_filler_nodes, node_size_y_tensor+num_nodes-num_filler_nodes,
-                bin_center_x_tensor, bin_center_y_tensor,
-                xl, yl, xh, yh,
-                bin_size_x, bin_size_y,
-                num_filler_nodes,
-                grad_x_tensor+num_nodes-num_filler_nodes, grad_y_tensor+num_nodes-num_filler_nodes
-                );
+            num_filler_nodes,
+            grad_x_tensor + offset, grad_y_tensor + offset);
 
         status = cudaStreamDestroy(stream_filler);
         stream_filler = 0;

    return 0;
}

#define REGISTER_KERNEL_LAUNCHER(T) \
    int instantiateComputeElectricForceLauncher(\
            int num_bins_x, int num_bins_y, \
            int num_movable_impacted_bins_x, int num_movable_impacted_bins_y, \
            int num_filler_impacted_bins_x, int num_filler_impacted_bins_y, \
            const T* field_map_x_tensor, const T* field_map_y_tensor, \
            const T* x_tensor, const T* y_tensor, \
            const T* node_size_x_tensor, const T* node_size_y_tensor, \
            const T* bin_center_x_tensor, const T* bin_center_y_tensor, \
            T xl, T yl, T xh, T yh, \
            T bin_size_x, T bin_size_y, \
            int num_nodes, int num_movable_nodes, int num_filler_nodes, \
            T* grad_x_tensor, T* grad_y_tensor\
            )\
    { \
        return computeElectricForceCudaLauncher(\
                num_bins_x, num_bins_y, \
                num_movable_impacted_bins_x, num_movable_impacted_bins_y, \
                num_filler_impacted_bins_x, num_filler_impacted_bins_y, \
                field_map_x_tensor, field_map_y_tensor, \
                x_tensor, y_tensor, \
                node_size_x_tensor, node_size_y_tensor, \
                bin_center_x_tensor, bin_center_y_tensor, \
                xl, yl, xh, yh, \
                bin_size_x, bin_size_y, \
                num_nodes, num_movable_nodes, num_filler_nodes, \
                grad_x_tensor, grad_y_tensor\
                );\
    }

REGISTER_KERNEL_LAUNCHER(float);

        //node_y = min(node_y, yh-node_size_y);

        // x direction
        int bin_index_xl = int((node_x - xl) / bin_size_x);
        bin_index_xl = max(bin_index_xl, 0);
        int k = bin_index_xl + int(residual_index / num_impacted_bins_y);
        if (k + 1 > num_bins_x || node_x + node_size_x <= bin_center_x_tensor[k] - bin_size_x / 2)
        {
            return;
        }
        // y direction
        int bin_index_yl = int((node_y - yl) / bin_size_y);
        bin_index_yl = max(bin_index_yl, 0);
        int h = bin_index_yl + (residual_index % num_impacted_bins_y);
        if (h + 1 > num_bins_y || node_y + node_size_y <= bin_center_y_tensor[h] - bin_size_y / 2)
        {
            return;
        }
