   {
     struct TrianglesGenerator : public CubeIndexEstimator
     {
-      enum { CTA_SIZE = 256 };
+      enum { CTA_SIZE = 256, MAX_GRID_SIZE_X = 65536 };
 
       const int* occupied_voxels;
       const int* vertex_ofssets;

  tg.cell_size.z = volume_size.z / VOLUME_Z;
  tg.output = output;

  dim3 block (TrianglesGenerator::CTA_SIZE);
  dim3 grid (divUp (tg.voxels_count, block.x));

  trianglesGeneratorKernel<<<grid, block>>>(tg);
  cudaSafeCall ( cudaGetLastError () );

      operator () () const
      {
        int tid = threadIdx.x;
        int idx = (blockIdx.y * MAX_GRID_SIZE_X + blockIdx.x) * CTA_SIZE + tid;
      

        if (idx >= voxels_count)
          return;
