         dim3 particle_threads(block_size, 1, 1);
         
         gpu_nve_rigid_step_one_particle_sliding_kernel<<< particle_grid, particle_threads >>>(pdata.pos, 
-                                                                     pdata.vel, 
+                                                                     pdata.vel,
                                                                      pdata.image,
+                                                                     pdata.mass,
+                                                                     d_net_force,
+                                                                     rigid_data.body_indices,
+                                                                     rigid_data.orientation,
+                                                                     rigid_data.com,
+                                                                     rigid_data.vel,
+                                                                     rigid_data.angvel,
+                                                                     rigid_data.body_imagex,
+                                                                     rigid_data.body_imagey,
+                                                                     rigid_data.body_imagez,
+                                                                     rigid_data.particle_indices,
+                                                                     rigid_data.particle_pos,
                                                                      rigid_data.particle_oldpos,
                                                                      rigid_data.particle_oldvel,
                                                                      rigid_data.virial,

    __shared__ float4 com, vel, angvel, ex_space, ey_space, ez_space;
    __shared__ int body_imagex, body_imagey, body_imagez;
    
    int idx_body = -1;
    
    float dt_half = 0.5 * deltaT; 
    
    if (group_idx < n_group_bodies)
        {
        idx_body = tex1Dfetch(rigid_data_body_indices_tex, group_idx);
       
        if (idx_body < n_bodies && threadIdx.x == 0)
            {
            com = tex1Dfetch(rigid_data_com_tex, idx_body);
            vel = tex1Dfetch(rigid_data_vel_tex, idx_body);
            angvel = tex1Dfetch(rigid_data_angvel_tex, idx_body);
            body_imagex = tex1Dfetch(rigid_data_body_imagex_tex, idx_body);
            body_imagey = tex1Dfetch(rigid_data_body_imagey_tex, idx_body);
            body_imagez = tex1Dfetch(rigid_data_body_imagez_tex, idx_body);
            
            Scalar4 orientation = tex1Dfetch(rigid_data_orientation_tex, idx_body);
            exyzFromQuaternion(orientation, ex_space, ey_space, ez_space);
            }
        }
        

extern "C" __global__ void gpu_nve_rigid_step_one_particle_sliding_kernel(float4* pdata_pos,
                                                        float4* pdata_vel,
                                                        int4* pdata_image,
                                                        float* d_particle_mass,
                                                        float4* d_net_force,
                                                        unsigned int *d_rigid_group,
                                                        float4* d_rigid_orientation,
                                                        float4* d_rigid_com,
                                                        float4* d_rigid_vel,
                                                        float4* d_rigid_angvel,
                                                        int* d_rigid_imagex,
                                                        int* d_rigid_imagey,
                                                        int* d_rigid_imagez,
                                                        unsigned int* d_rigid_particle_idx,
                                                        float4* d_rigid_particle_dis,
                                                        float4* rdata_oldpos,
                                                        float4* rdata_oldvel,
                                                        float *d_virial,
