 
             __syncthreads();
 
-            if (tid < no) scratch[tid] = biases[l][tid];
-
+            if (tid < no)
+                scratch[tid] = tex1Dfetch(biases_tex,
+                                          layer_bias_offset
+                                          + tid);
+            
             __syncthreads();
             
             /* Now we accumulate them, allowing each thread to increment in its

            if (tid < ni) {
                for (unsigned o = 0;  o < no;  ++o) {
                    float d = scratch[o];  // may be the d from another thread
                    float update = d * layer_weights[tid * w_stride + o];
                    total += update;
                }
            }


        // Since we do 4 examples per loop, we need enough memory for all of
        // the four outputs for a single layer
        shared_mem_size = max_width * sizeof(float)* 4;
    }
};

