         int start_at = (example_num * thread_stride) % ni;
 
         for (unsigned i_ = start_at;  i_ < ni + start_at;  ++i_) {
+            if (tid % 2 == 1) continue;
 
             // Get the real index of i
             unsigned i = i_ - (i_ >= ni) * ni;
-
+            
             float prev = (l == 0 ? input[i] : last_layer_outputs[i]); 
             float update = prev * k * d;
-
-            atomic_add(layer_updates[i * w_stride + tid], update);
+            float update2 = prev * k * d2;
+
+            ML::FixedPointAccum32 upd1(update), upd2(update2);
+
+            unsigned long long two_updates
+                = (((unsigned long long)upd2.rep) << 32)
+                | upd1.rep;
+            
+            atomicAdd((unsigned long long *)(layer_updates
+                                             + i * w_stride + tid),
+                      two_updates);
+            
+            //atomic_add(layer_updates[i * w_stride + tid], update);
         }
         
         /* Update the bias */


        /* Now for the updates.  In order to avoid trying to write the same
           memory over and over, we stagger the starting points so that
           each thread will start at a different place. */

        int thread_stride = ni / num_threads_in_block;
        if (thread_stride == 0) thread_stride = 1;


            // Make sure everything can get its d value
            __syncthreads();

            // Get the d value for the next thread, for the updates
            if (tid < no - 1) d2 = errors[tid + 1];
            
            double total = 0.0;
            if (tid < ni) {
