         // no syncthreads here, we assume threads_per_particle < warp size
 
         // scan over flags
-        unsigned char n = 1;
         int k = 0;
         #if (__CUDA_ARCH__ >= 300)
+        unsigned char n = 1;
         k = warp_scan_sm30<threads_per_particle>::Scan(threadIdx.x % threads_per_particle, has_neighbor, &n);
         #endif
 
         if (has_neighbor && nneigh + k < nli.getH())
+            {
             d_nlist[nli(my_pidx, nneigh + k)] = neighbor;
+            if (my_pidx == 15000)
+                {
+                printf("Adding particle %d\n", neighbor);
+                }
+            }
 
         // increment total neighbor count
+        #if (__CUDA_ARCH__ >= 300)
         nneigh += n;
+        #else
+        if (has_neighbor)
+            nneigh++;
+        #endif
         } // end while
 
     if (threadIdx.x % threads_per_particle == 0)

                                     const Scalar3& ghost_width,
                                     const unsigned int compute_capability)
    {
    launcher<max_threads_per_particle>(d_nlist,
                                   d_n_neigh,
                                   d_last_updated_pos,


#include "NeighborListGPUBinned.cuh"
#include "TextureTools.h"
#include <stdio.h>

/*! \file NeighborListGPUBinned.cu
    \brief Defines GPU kernel code for O(N) neighbor list generation on the GPU
