         {
         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;


    __syncthreads();
    
    body_force[threadIdx.x] = sum_force;
    body_torque[threadIdx.x] = sum_torque;
    
    // reduction within the current window
    unsigned int offset = min(window_size, nmax) >> 1;
    while (offset > 0)
        {

                                                 unsigned int n_bodies_per_block,
                                                 gpu_boxsize box)
    {
    // determine which body (0 ... n_bodies_per_block-1) this thread is working on
    // assign threads 0, 1, 2, ... to body 0, n, n+1, n+2, ... to body 1, and so on.
    unsigned int m = threadIdx.x / (blockDim.x / n_bodies_per_block);
    
    // body_force and body_torque are each shared memory arrays with 1 element per threads
    float3 *body_force = sum;
    float3 *body_torque = &sum[blockDim.x];
    
    // store ex_space, ey_space, ez_space, and the body index in shared memory. Up to 8 bodies per block can
    // be handled.
    __shared__ float4 ex_space[8], ey_space[8], ez_space[8];
    __shared__ int idx_body[8];

    // each thread makes partial sums of force and torque of all the particles that this thread loops over
    float3 sum_force = make_float3(0.0f, 0.0f, 0.0f);
    float3 sum_torque = make_float3(0.0f, 0.0f, 0.0f);
    
    // thread_mask is a bitmask that masks out the high bits in threadIdx.x.
    // threadIdx.x & thread_mask is an index from 0 to block_size/n_bodies_per_block-1 and determines what offset
    // this thread is to use when accessing the particles in the body
    if ((threadIdx.x & thread_mask) == 0)
        {
        // thread 0 for this body reads in the body id and orientation and stores them in shared memory
        int group_idx = blockIdx.x*n_bodies_per_block + m;
        if (group_idx < n_group_bodies)
            {
