         // We need our grid size to be exactly the maximum width of the output
         threads = dim3(max_width);
 
-        shared_mem_size = (max_width + total_neurons) * sizeof(float);
+        shared_mem_size = max_width * sizeof(float);
     }
 };
 

                    tid, l, ni, no, last_layer_outputs, this_layer_outputs);
#endif

#if 0
    std::copy(bias.begin(), bias.end(), output.begin());
    for (unsigned i = 0;  i < input.size();  ++i)
        SIMD::vec_add(&output[0], input[i], &weights[i][0], &output[0],
                      outputs());
        //for (unsigned o = 0;  o < output.size();  ++o)
        //    output[o] += input[i] * weights[i][o];
    transform(output);
#endif

        /* Add in the layer outputs.  We iterate with all threads */
        if (tid < no) {
            // Start off with the bias terms
            double accum = biases[l][tid];

            for (unsigned i = 0;  i < ni;  ++i) {
                float inval = (l == 0 ? input[i] : last_layer_outputs[i]);

                // Coalesced access; maybe texture would be better
                float weight = layer_weights[i * w_stride + tid];
                
                accum += weight * inval;
            }
            
            this_layer_outputs[tid] = transform(accum, activation);
        }

#if defined(__DEVICE_EMULATION__)

       the errors array, so that our layer outputs have to start at this
       offset.
    */
    //float * layer_outputs = scratch + blockDim.x;

    const float * input = feature_vectors + example_num * feature_vector_width;

