     }
     __syncthreads();
     if (gid < i_to) {
+      unsigned int self_id = gid + 1;
       unsigned int cluster_id = clustering[gid];
       if (cluster_id == 0) {
-        cluster_id = gid + 1;
+        cluster_id = self_id;
       }
       unsigned int ref_id = tid+bsize;
       // load reference coordinates for re-use into shared memory

      cudaFree(d_coords_sorted[i_gpu]);
      cudaFree(d_clustering[i_gpu]);
    }

//TODO
    std::cout << "#  " << clustering_sorted[0] << std::endl;

//    return clustering;
    return normalized_cluster_names(first_frame_above_threshold
                                  , clustering
                                  , fe_sorted);

               , tmp_coords_sorted.data()
               , sizeof(float) * n_rows * n_cols
               , cudaMemcpyHostToDevice);


      //TODO change kernel to use two lists: own state + min next-neighbor state


      // copy prev results to GPU-buffer (and set the rest to zero)
      cudaMemset(d_clustering[i_gpu]
               , 0
               , sizeof(unsigned int) * n_rows);
      cudaMemcpy(d_clustering[i_gpu]
               , prev_clustering_sorted.data()
               , sizeof(unsigned int) * n_rows
