                         + ez_space[m].y * particle_pos.z;
                 ri.z = ex_space[m].z * particle_pos.x + ey_space[m].z * particle_pos.y 
                         + ez_space[m].z * particle_pos.z;
-    
+                
+                // tally the torque in the per thread counter
                 sum_torque.x += ri.y * fi.z - ri.z * fi.y;
                 sum_torque.y += ri.z * fi.x - ri.x * fi.z;
                 sum_torque.z += ri.x * fi.y - ri.y * fi.x;

        __syncthreads();
        }
    
    if ((threadIdx.x & thread_mask) == 0 && idx_body[m] != -1)
        {
        rdata_force[idx_body[m]] = make_float4(body_force[threadIdx.x].x, body_force[threadIdx.x].y, body_force[threadIdx.x].z, 0.0f);

        {
        float4 fi = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
        
        // determine the index with this body that this particle should handle
        unsigned int k = start * window_size + (threadIdx.x & thread_mask);
        
        // if that index is in the body we are actually handling a real body
        if (k < nmax && idx_body[m] != -1)
            {
            // determine the particle idx of the particle
            int localidx = idx_body[m] * nmax + k;
            unsigned int pidx = d_rigid_particle_idx[localidx];
            
            // if this particle is actually in the body
            if (pidx != INVALID_INDEX)
                {
                // calculate body force and torques
                float4 particle_pos = d_rigid_particle_dis[localidx];
                fi = d_net_force[pidx];
                
                // tally the force in the per thread counter
                sum_force.x += fi.x;
                sum_force.y += fi.y;
                sum_force.z += fi.z;
