           activation(activation),
           fire(fire),
           inhibit(inhibit),
-          learning_rate(learning_rate)
+          learning_rate(learning_rate),
+          use_textures(true)
     {
         //cerr << "plan: num_layers = " << num_layers << endl;
 

    dim3 threads;
    
    size_t shared_mem_size;

    Plan(int num_layers,
         const int * architecture,

    }
}

__global__ void
train_examples_kernel_tex(const float * feature_vectors,  // feature vector [ni]
                          int feature_vector_width,
                          const int * labels,
                          const float * example_weights,
                          int num_layers,
                          const int * architecture,
                          const int * w_strides,
                          UpdateFloat * const * w_updates, // wt updates for each layer
                          UpdateFloat * const * b_updates, // bias upd for each layer
                          int activation,            // activation function
                          float fire,   // target value for firing neuron
                          float inhibit, // target value for inhibited neuron)
                          float learning_rate,
                          int num_threads_in_block,
                          int num_threads_on_multiprocessor,
                          int total_neurons,
                          float * layer_outputs,  // scratch space[total neurons]
                          int examples_per_block,
                          int total_num_examples,
                          int max_width)
{
    const unsigned block_num  = blockIdx.x;
    
    /* Where we accumulate our errors, layer by layer.  The size is that of
       the largest dimension. */
    extern __shared__ float scratch[];
    
    /* The layer outputs (activation of the neurons).  This is where the
       shared memory goes to.  Note that we store only the activated outputs,
       not the inputs.

       blockDim.x gives us the number of threads, which is also the size of
       the errors array, so that our layer outputs have to start at this
       offset.
    */

    // Get our private scratch memory for this block
    layer_outputs += block_num * total_neurons * 4;
    
    unsigned example_num_base = block_num * examples_per_block;
    unsigned last_example = min(total_num_examples, example_num_base + examples_per_block);

    unsigned example_num = example_num_base;

    // Do any others singly
    for (;  example_num < last_example;  ++example_num) {

        const float * input = feature_vectors + example_num * feature_vector_width;

        int label = labels[example_num];

        float example_weight = example_weights[example_num];

        train_example_tex(input, label, example_weight,
                      num_layers, scratch, architecture, w_strides,
                      w_updates, b_updates,
                      activation, fire, inhibit, learning_rate,
                      num_threads_in_block,
                      num_threads_on_multiprocessor,
                      total_neurons, layer_outputs);
    }
}


namespace ML {
namespace CUDA {
