             
             for (int x = 0;  x < N && x < valid_examples;  ++x)
                 scratch[x * scratch_stride + tid] = total[x];
-        }
-
-        // Again, threads indexed too low just leave
-        if (tid >= no) continue;
+
+            /* Scratch now contains the errors from the current layer. */
+        }
+
+        if (tid >= no) {
+            if (tid <= ni && l > 0) {
+                // Need to set up prev_outputs
+                for (int x = 0;  x < N && x < valid_examples;  ++x) {
+                    float prev = last_layer_outputs[x * total_neurons + tid]; 
+                    prev_outputs[x] = prev;
+                }
+            }
+            else {
+                //FOR_ALL_X(prev_outputs[x] = 0.0f);
+            }
+
+            // threads indexed too low just leave
+            continue;
+        }
         
         /* Update the weights. */
         float k[N];

    /* BPROP                                                                 */
    /*************************************************************************/

    /* How many output layers? */
    this_layer_outputs -= no;

    layer_weights -= ni * w_stride;

    /* First error calculation pass */
    float last_output[N];
    FOR_ALL_X(last_output[x] = scratch[x * scratch_stride + tid]);

    __syncthreads();

    for (int x = 0;  x < N && x < valid_examples;  ++x) {
        bool correct = (labels[x] == tid);
        float wanted = (correct ? fire : inhibit);
        scratch[x * scratch_stride + tid]
            = (tid < no ? wanted - last_output[x] : 0.0);
    }
    
    /* Let everything catch up */
    __syncthreads();

    /* Backpropegate. */

            double total_update = 0.0;

            for (int x = 0;  x < N && x < valid_examples;  ++x) {
                // All threads read the same value
                float prev
                    = (l == 0
                       ? input[x * input_stride + i]
                       : last_layer_outputs[x * total_neurons + i]); 
                float update = prev * k[x] * d[x];
                total_update += update;

                // If we've just read the value we need for the next
                // block, then store it.
                if (i == tid) prev_outputs[x] = prev;
            }
            
            atomic_add(layer_updates[i * w_stride + tid], total_update);
