         /* 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 example will start at a different place, thus minimising
-           conflicting writes when we have multiple multiprocessors working
-           on the same thing. */
+           conflicting atomic updates when we have multiple multiprocessors
+           working on the same thing. */
 
         int thread_stride = ni / num_threads_in_block;
         if (thread_stride == 0) thread_stride = 1;

            
            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;
        
        /* Update the weights. */
        float k[N];

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

    /* Make this_layer_outputs point to the outputs of the layer before the
       output */
    this_layer_outputs -= no;

    layer_weights -= ni * w_stride;

    /* We keep the outputs of the next layer here to avoid needing to reload
       them. */
    float prev_outputs[N];
    FOR_ALL_X(prev_outputs[x] = scratch[x * scratch_stride + tid]);

    __syncthreads();

    /* Calculate the error terms for the output units.  Scratch contains the
       outputs of output layer, so we can access this instead of
       this_layer_outputs. */
    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 - prev_outputs[x] : 0.0);
    }
    
    /* Let everything catch up.  Scratch now contains the errors for the
       output layer. */
    __syncthreads();

    /* Backpropegate. */
