 
             d_weight_updates_storage[l].sync(sync_to);
             std::copy(sync_to, sync_to + ni * w_stride, weight_updates[l]);
+
+#if 0
+            cerr << "first 10 weight updates for layer " << l << ": ";
+            for (unsigned i = 0;  i < 10;  ++i)
+                cerr << sync_to[i] << " ";
+            cerr << endl;
+#endif
         }
 
         for (unsigned l = 0;  l < plan.num_layers;  ++l) {

            float prev3 = (l == 0 ? input3[i] : last_layer_outputs3[i]); 
            float prev4 = (l == 0 ? input4[i] : last_layer_outputs4[i]); 

#if 1
            float update1 = prev1 * k1 * d1;
            float update2 = prev2 * k2 * d2;
            float update3 = prev3 * k3 * d3;
            float update4 = prev4 * k4 * d4;

            atomic_add(layer_updates[i * w_stride + tid], update1);
            atomic_add(layer_updates[i * w_stride + tid], update2);
            atomic_add(layer_updates[i * w_stride + tid], update3);
            atomic_add(layer_updates[i * w_stride + tid], update4);
#else
            float update
                = (k1 * d1 * prev1)
                + (k2 * d2 * prev2)
                + (k3 * d3 * prev3)
                + (k4 * d4 * prev4);

            atomic_add(layer_updates[i * w_stride + tid], update);
#endif
        }
#if 1
        float update1 = k1 * d1;
        float update2 = k2 * d2;
        float update3 = k3 * d3;
        float update4 = k4 * d4;

        atomic_add(layer_bias_updates[tid], update1);
        atomic_add(layer_bias_updates[tid], update2);
        atomic_add(layer_bias_updates[tid], update3);
        atomic_add(layer_bias_updates[tid], update4);
#else
        /* Update the bias */
        float update = ((k1 * d1) + (k2 * d2) + (k3 * d3) + (k4 * d4));

        atomic_add(layer_bias_updates[tid], update);
#endif
    }

#if 0
        for (unsigned i_ = start_at;  i_ < ni + start_at;  ++i_) {

            // 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);
        }
        
        /* Update the bias */
        float update = k * d;

        //layer_bias_updates[tid] += update;
        atomic_add(layer_bias_updates[tid], update);
#endif

}

__global__ void


        next_layer_outputs = this_layer_outputs + no;

#if defined(__DEVICE_EMULATION__) && 0
        if (tid == 0)
            fprintf(stderr, "fprop: tid %d layer %d ni %d no %d last_layer_outputs %p this_layer_outputs %p next_layer_outputs %p\n",
                    tid, l, ni, no, last_layer_outputs, this_layer_outputs,
