Workaround to atomic float add

I'm trying to update the forces in a particle system after a nearest neighbors search.

The below example leads to a conflict when adding the forces to the neighbor particle.

On other systems I would use an atomic float add on these values but so far that is not available in Metal.

The addition operation of the forces and moments needs to be accurate.

I know one option is to do the force calculation twice but being an expensive calculation I'd prefer to do it only once.

Is there an alternate to accomplish this in Metal?

kernel void simpleExample(   device Particle *particle, 
                             device Contact *contact,
                             int home [thread_position_in_grid] ) {
        
        // int home; // Number of the home particle
        
     for (int i=0; i<contact[home].neighborcount; i++)="" {<br="">          particle[home].force = particcontact[home].ptr[j].force;
          particle[home].moment = particle[home].moment + contact[home].ptr[j].momentHome;


          int neighbor = contact[home].ptr[j].neighbor;


          /// These need to be an atomic float add. Needs to be accurate. But metal doesn't have an atomic float add
          particle[neighbor].force = particle[neighbor].force - contact[home].ptr[j].force;
          particle[neighbor].moment = particle[neighbor].moment + contact[home].ptr[j].momentNeighbor;
     }
}