         
     // run the kernel: the shared memory size is used for dynamic memory allocation of extern __shared__ sum
     // tested with separate extern __shared__ body_force and body_torque each with size of nmax * sizeof(float4) but it did not work
-    unsigned int window_size = nmax;
+    unsigned int window_size = 64; // some fixed value
     unsigned int block_size = nmax; // each thread in a block takes care of a particle in a rigid body
     dim3 force_grid(n_bodies, 1, 1);
     dim3 force_threads(block_size, 1, 1); 
-    gpu_rigid_force_kernel<<< force_grid, force_threads, 2 * window_size * sizeof(float4) >>>(rigid_data.force, 
+    
+    if (nmax <= 256)   // 256 is some threshold for really big rigid bodies
+        gpu_rigid_force_kernel<<< force_grid, force_threads, 2 * nmax * sizeof(float4) >>>(rigid_data.force, 
                                                                                              rigid_data.torque, 
                                                                                              n_bodies, 
                                                                                              local_beg,
                                                                                              nmax,
                                                                                              window_size,
                                                                                              box);
-           
+    else gpu_rigid_force_sliding_kernel<<< force_grid, force_threads, 2 * window_size * sizeof(float4) >>>(rigid_data.force, 
+                                                                                             rigid_data.torque, 
+                                                                                             n_bodies, 
+                                                                                             local_beg,
+                                                                                             nmax,
+                                                                                             window_size,
+                                                                                             box);       
     if (!g_gpu_error_checking)
         {
         return cudaSuccess;

    c.w =   b.x * a.z + a.x * b.z - a.y * b.y;
    }

/* Normalize a quaternion
 */

__device__ void normalize(float4 &q)

    cudaError_t error = cudaBindTexture(0, rigid_data_com_tex, rigid_data.com, sizeof(float4) * n_bodies);
    if (error != cudaSuccess)
        return error;
    
    error = cudaBindTexture(0, rigid_data_particle_indices_tex, rigid_data.particle_indices, sizeof(unsigned int) * n_bodies * nmax);
    if (error != cudaSuccess)
        return error;    
    
    // bind the textures for particles
    error = cudaBindTexture(0, pdata_pos_tex, pdata.pos, sizeof(float4) * pdata.N);
    if (error != cudaSuccess)
