       operator () () const
       {
         int tid = threadIdx.x;
-        int idx = blockIdx.x * CTA_SIZE + tid;
+        int idx = (blockIdx.y * MAX_GRID_SIZE_X + blockIdx.x) * CTA_SIZE + tid;
+      
 
         if (idx >= voxels_count)
           return;


void
pcl::device::generateTriangles (const PtrStep<short2>& volume, const DeviceArray2D<int>& occupied_voxels, const float3& volume_size, DeviceArray<PointType>& output)
{
  TrianglesGenerator tg;

  tg.volume = volume;
  tg.occupied_voxels = occupied_voxels.ptr (0);

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

  int blocks_num = divUp (tg.voxels_count, Tg::CTA_SIZE);

  dim3 block (Tg::CTA_SIZE);     
  dim3 grid(min(blocks_num, Tg::MAX_GRID_SIZE_X), divUp(blocks_num, Tg::MAX_GRID_SIZE_X));

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