     {
     unsigned int group_idx = blockIdx.x * blockDim.x + threadIdx.x;
     
+    if (group_idx >= n_group_bodies)
+        return;
+    
     // do velocity verlet update
     // v(t+deltaT/2) = v(t) + (1/2)a*deltaT
     // r(t+deltaT) = r(t) + v(t+deltaT/2)*deltaT
-    if (group_idx < n_group_bodies)
-        {
-        float body_mass;
-        float4 moment_inertia, com, vel, angmom, orientation, ex_space, ey_space, ez_space, force, torque;
-        int body_imagex, body_imagey, body_imagez;
-        float dt_half = 0.5 * deltaT;
-        
-        unsigned int idx_body = tex1Dfetch(rigid_data_body_indices_tex, group_idx);
-        if (idx_body < n_bodies)
-            {
-            body_mass = tex1Dfetch(rigid_data_body_mass_tex, idx_body);
-            moment_inertia = tex1Dfetch(rigid_data_moment_inertia_tex, idx_body);
-            com = tex1Dfetch(rigid_data_com_tex, idx_body);
-            vel = tex1Dfetch(rigid_data_vel_tex, idx_body);
-            angmom = tex1Dfetch(rigid_data_angmom_tex, idx_body);
-            orientation = tex1Dfetch(rigid_data_orientation_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);
-            force = tex1Dfetch(rigid_data_force_tex, idx_body);
-            torque = tex1Dfetch(rigid_data_torque_tex, idx_body);
-         
-            exyzFromQuaternion(orientation, ex_space, ey_space, ez_space);
-               
-            // update velocity
-            float dtfm = dt_half / body_mass;
-            
-            float4 vel2;
-            vel2.x = vel.x + dtfm * force.x;
-            vel2.y = vel.y + dtfm * force.y;
-            vel2.z = vel.z + dtfm * force.z;
-            vel2.w = vel.w;
-            
-            // update position
-            float4 pos2;
-            pos2.x = com.x + vel2.x * deltaT;
-            pos2.y = com.y + vel2.y * deltaT;
-            pos2.z = com.z + vel2.z * deltaT;
-            pos2.w = com.w;
-            
-            // read in body's image
-            // time to fix the periodic boundary conditions
-            float x_shift = rintf(pos2.x * box.Lxinv);
-            pos2.x -= box.Lx * x_shift;
-            body_imagex += (int)x_shift;
-            
-            float y_shift = rintf(pos2.y * box.Lyinv);
-            pos2.y -= box.Ly * y_shift;
-            body_imagey += (int)y_shift;
-            
-            float z_shift = rintf(pos2.z * box.Lzinv);
-            pos2.z -= box.Lz * z_shift;
-            body_imagez += (int)z_shift;
-       
-            // update the angular momentum and angular velocity
-            float4 angmom2;
-            angmom2.x = angmom.x + dt_half * torque.x;
-            angmom2.y = angmom.y + dt_half * torque.y;
-            angmom2.z = angmom.z + dt_half * torque.z;
-            angmom2.w = 0.0;
-            
-            float4 angvel2;
-            advanceQuaternion(angmom2, moment_inertia, angvel2, ex_space, ey_space, ez_space, deltaT, orientation);
-       
-            // write out the results (MEM_TRANSFER: ? bytes)
-            rdata_com[idx_body] = pos2;
-            rdata_vel[idx_body] = vel2;
-            rdata_angmom[idx_body] = angmom2;
-            rdata_angvel[idx_body] = angvel2;
-            rdata_orientation[idx_body] = orientation;
-            rdata_ex_space[idx_body] = ex_space;
-            rdata_ey_space[idx_body] = ey_space;
-            rdata_ez_space[idx_body] = ez_space;
-            rdata_body_imagex[idx_body] = body_imagex;
-            rdata_body_imagey[idx_body] = body_imagey;
-            rdata_body_imagez[idx_body] = body_imagez;
-            }
-        }
+    float body_mass;
+    float4 moment_inertia, com, vel, angmom, orientation, ex_space, ey_space, ez_space, force, torque;
+    int body_imagex, body_imagey, body_imagez;
+    float dt_half = 0.5 * deltaT;
+        
+    unsigned int idx_body = d_rigid_group[group_idx];
+    body_mass = d_rigid_mass[idx_body];
+    moment_inertia = d_rigid_mi[idx_body];
+    com = rdata_com[idx_body];
+    vel = rdata_vel[idx_body];
+    angmom = rdata_angmom[idx_body];
+    orientation = rdata_orientation[idx_body];
+    body_imagex = rdata_body_imagex[idx_body];
+    body_imagey = rdata_body_imagey[idx_body];
+    body_imagez = rdata_body_imagez[idx_body];
+    force = d_rigid_force[idx_body];
+    torque = d_rigid_torque[idx_body];
+    
+    exyzFromQuaternion(orientation, ex_space, ey_space, ez_space);
+        
+    // update velocity
+    float dtfm = dt_half / body_mass;
+    
+    float4 vel2;
+    vel2.x = vel.x + dtfm * force.x;
+    vel2.y = vel.y + dtfm * force.y;
+    vel2.z = vel.z + dtfm * force.z;
+    vel2.w = vel.w;
+    
+    // update position
+    float4 pos2;
+    pos2.x = com.x + vel2.x * deltaT;
+    pos2.y = com.y + vel2.y * deltaT;
+    pos2.z = com.z + vel2.z * deltaT;
+    pos2.w = com.w;
+    
+    // time to fix the periodic boundary conditions
+    float x_shift = rintf(pos2.x * box.Lxinv);
+    pos2.x -= box.Lx * x_shift;
+    body_imagex += (int)x_shift;
+    
+    float y_shift = rintf(pos2.y * box.Lyinv);
+    pos2.y -= box.Ly * y_shift;
+    body_imagey += (int)y_shift;
+    
+    float z_shift = rintf(pos2.z * box.Lzinv);
+    pos2.z -= box.Lz * z_shift;
+    body_imagez += (int)z_shift;
+
+    // update the angular momentum and angular velocity
+    float4 angmom2;
+    angmom2.x = angmom.x + dt_half * torque.x;
+    angmom2.y = angmom.y + dt_half * torque.y;
+    angmom2.z = angmom.z + dt_half * torque.z;
+    angmom2.w = 0.0;
+    
+    float4 angvel2;
+    advanceQuaternion(angmom2, moment_inertia, angvel2, ex_space, ey_space, ez_space, deltaT, orientation);
+
+    // write out the results
+    rdata_com[idx_body] = pos2;
+    rdata_vel[idx_body] = vel2;
+    rdata_angmom[idx_body] = angmom2;
+    rdata_angvel[idx_body] = angvel2;
+    rdata_orientation[idx_body] = orientation;
+    rdata_ex_space[idx_body] = ex_space;
+    rdata_ey_space[idx_body] = ey_space;
+    rdata_ez_space[idx_body] = ez_space;
+    rdata_body_imagex[idx_body] = body_imagex;
+    rdata_body_imagey[idx_body] = body_imagey;
+    rdata_body_imagez[idx_body] = body_imagez;
     }
 
 /*!

    gpu_nve_rigid_step_two_body_kernel<<< body_grid, body_threads >>>(rigid_data.vel, 
                                                                      rigid_data.angmom, 
                                                                      rigid_data.angvel,
                                                                      rigid_data.conjqm,
                                                                      n_group_bodies,
                                                                      n_bodies, 
                                                                      nmax, 
                                                                      box, 
                                                                      deltaT);
    block_size = 192;
    dim3 particle_grid(group_size/block_size+1, 1, 1);
    dim3 particle_threads(block_size, 1, 1);

                                                           rigid_data.body_imagey, 
                                                           rigid_data.body_imagez,
                                                           rigid_data.conjqm,
                                                           rigid_data.body_mass,
                                                           rigid_data.moment_inertia,
                                                           rigid_data.force,
                                                           rigid_data.torque,
                                                           rigid_data.body_indices,
                                                           n_group_bodies, 
                                                           n_bodies, 
                                                           box,
