                     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];
+        
+        // Start off with the bias terms
+        double accum = (tid < no ? biases[l][tid] : 0.0);
+
+        if (__any(tid < no)) {
 
             for (unsigned i = 0;  i < ni;  ++i) {
-                float inval = (l == 0 ? input[i] : last_layer_outputs[i]);
-
+                // No bank conflicts as all threads are accessing same value
+                float inval = scratch[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];
+                float weight
+                    = (tid < no ? layer_weights[i * w_stride + tid] : 0.0);
                 
                 accum += weight * inval;
             }
-            
-            this_layer_outputs[tid] = transform(accum, activation);
+        }         
+
+        // Let everything catch up so that we can write to scratch
+        __syncthreads();
+        
+        if (__any(tid < no)) {
+
+            if (tid < no)
+                /* this_layer_outputs[tid] = */ scratch[tid]
+                    = transform(accum, activation);
         }
 
 #if defined(__DEVICE_EMULATION__)

                example_num, label);
        for (unsigned i = 0;  i < no;  ++i) {
            fprintf(stderr, "output %d: value %f error %f correct %d\n",
                    i, this_layer_outputs[i], errors[i], (label == i));
        }
    }
#endif


        for (unsigned l = 0;  l < plan.num_layers;  ++l) {
            int no = plan.architecture[l + 1];
            d_bias_updates_storage[l].init_zeroed(no);
            bias_updates_vec[l] = d_bias_updates_storage[l];
        }

        d_bias_updates.init(&bias_updates_vec[0], plan.num_layers);

        // Get the scratch space
        //d_layer_outputs.init(plan.total_neurons);
        
        // Our grid size is one per example
        grid = dim3(num_feature_vectors);
