                  int max_width,
                  float * layer_outputs)  // global scratch space[total neurons]
 {
-#if 1
     // access thread id
     const unsigned tid = threadIdx.x;
 

        //float * d = (scratch + N * scratch_stride);
        float d[N];

        for (int x = 0;  x < N && x < valid_examples;  ++x) {
            float prev_output
                = (tid >= no
                   ? 0.0
                   : this_layer_outputs[x * total_neurons + tid]);
            
            float error = scratch[x * scratch_stride + tid];
        
            d[x] = (tid >= no ? 0.0 : delta(prev_output, error, activation));
        }
        
        if (l > 0) {

            float bias = layer_biases[o];

            for (int x = 0;  x < N && x < valid_examples;  ++x) {

                // XXX can get rid of this store for output layer if we don't
                // want to keep the data
                this_layer_outputs[x * total_neurons + o]
                    = scratch[x * scratch_stride + o]
                    = transform(accum[x] + bias, activation);
            }
        }
#endif
    }

    // layer_biases is no longer used

