             body_force[ilocal].z = fi.z;
             body_force[ilocal].w = fi.w;
             
-            // project the position in the body frame to the space frame: ri = rotation_matrix * particle_pos
             ri.x = pos.x - com.x;
             ri.y = pos.y - com.y;
             ri.z = pos.z - com.z;

        body_force[threadIdx.x].z = fi.z;
        body_force[threadIdx.x].w = fi.w;
        
        // project the position in the body frame to the space frame: ri = rotation_matrix * particle_pos
        ri.x = pos.x - com.x;
        ri.y = pos.y - com.y;
        ri.z = pos.z - com.z;

//! Shared memory for body force and torque reduction, required allocation when the kernel is called
extern __shared__ float4 sum[];

//! Calculates the body forces and torques by summing the constituent particle forces
/*! \param rdata_force Particle data to step forward in time
    \param rdata_torque List of pointers to forces on each particle
    \param n_bodies Number of forces listed in \a force_data_ptrs
