     */
 
     // Get our private scratch memory for this block
-    layer_outputs += block_num * total_neurons;
+    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);
 
-    for (unsigned example_num = example_num_base;  example_num < last_example;  ++example_num) {
+    unsigned example_num = example_num_base;
+#if 0
+    for (;  example_num < last_example - 3;  example_num += 4) {
+        const float * input1 = feature_vectors + example_num * feature_vector_width;
+        const float * input2 = input1 + feature_vector_width;
+        const float * input3 = input2 + feature_vector_width;
+        const float * input4 = input3 + feature_vector_width;
+
+        int4 label;
+        label.x = labels[example_num];
+        label.y = labels[example_num + 1];
+        label.z = labels[example_num + 2];
+        label.w = labels[example_num + 3];
+
+        float4 example_weight;
+        example_weight.x = example_weights[example_num];
+        example_weight.y = example_weights[example_num + 1];
+        example_weight.z = example_weights[example_num + 2];
+        example_weight.w = example_weights[example_num + 3];
+
+        float * scratch1 = scratch;
+        float * scratch2 = scratch1 + max_width;
+        float * scratch3 = scratch2 + max_width;
+        float * scratch4 = scratch3 + max_width;
+
+        float * layer_outputs1 = layer_outputs;
+        float * layer_outputs2 = layer_outputs1 + total_neurons;
+        float * layer_outputs3 = layer_outputs2 + total_neurons;
+        float * layer_outputs4 = layer_outputs3 + total_neurons;
+        
+        train_4_examples(input1, input2, input3, input4,
+                         label, example_weight,
+                         num_layers, scratch1, scratch2, scratch3, scratch4,
+                         w, biases, architecture, w_strides,
+                         w_updates, b_updates,
+                         activation, fire, inhibit, learning_rate,
+                         num_threads_in_block, total_neurons,
+                         layer_outputs1, layer_outputs2, layer_outputs3,
+                         layer_outputs4);
+
+    }
+#endif
+
+#if 1
+    for (;  example_num < last_example;  example_num += 1) {
+        const float * input1 = feature_vectors + example_num * feature_vector_width;
+        const float * input2 = input1 + feature_vector_width;
+        const float * input3 = input1;
+        const float * input4 = input2;
+
+        int4 label;
+        label.x = labels[example_num];
+        label.y = labels[example_num + 1];
+        label.z = labels[example_num];
+        label.w = labels[example_num + 1];
+
+        float4 example_weight;
+        example_weight.x = 0.0;
+        example_weight.y = 0.0;
+        example_weight.z = example_weights[example_num];
+        example_weight.w = 0.0;//example_weights[example_num + 1];
+
+        float * scratch1 = scratch;
+        float * scratch2 = scratch1 + max_width;
+        float * scratch3 = scratch2 + max_width;
+        float * scratch4 = scratch3 + max_width;
+        
+        float * layer_outputs1 = layer_outputs;
+        float * layer_outputs2 = layer_outputs1 + total_neurons;
+        float * layer_outputs3 = layer_outputs2 + total_neurons;
+        float * layer_outputs4 = layer_outputs3 + total_neurons;
+        
+        train_4_examples(input1, input2, input3, input4,
+                         label, example_weight,
+                         num_layers, scratch1, scratch2, scratch3, scratch4,
+                         w, biases, architecture, w_strides,
+                         w_updates, b_updates,
+                         activation, fire, inhibit, learning_rate,
+                         num_threads_in_block, total_neurons,
+                         layer_outputs1, layer_outputs2, layer_outputs3,
+                         layer_outputs4);
+
+    }
+#endif
+
+    for (;  example_num < last_example;  ++example_num) {
 
         const float * input = feature_vectors + example_num * feature_vector_width;
 

    }
}

__global__ void
train_examples_kernel(const float * feature_vectors,  // feature vector [ni]
                      int feature_vector_width,

             plan.total_neurons,
             d_layer_outputs,
             num_examples_per_invocation,
             num_feature_vectors /* total num examples */,
             plan.max_width);

        //cerr << "launched" << endl;
    }
