comps(y, x) = root(comps, comps(y, x));
}
- void labelComponents(const DevMem2D& edges, DevMem2Di comps, cudaStream_t stream)
+ enum {CC_NO_COMPACT = 0, CC_COMPACT_LABELS = 1};
+
+ void labelComponents(const DevMem2D& edges, DevMem2Di comps, int flags, cudaStream_t stream)
{
dim3 block(CTA_SIZE_X, CTA_SIZE_Y);
dim3 grid(divUp(edges.cols, TILE_COLS), divUp(edges.rows, TILE_ROWS));
cudaSafeCall( cudaGetLastError() );
int tileSizeX = TILE_COLS, tileSizeY = TILE_ROWS;
-
- cudaSafeCall( cudaGetLastError() );
- // cudaSafeCall( cudaDeviceSynchronize() );
-
while (grid.x > 1 || grid.y > 1)
{
dim3 mergeGrid(ceilf(grid.x / 2.0), ceilf(grid.y / 2.0));
dim3 mergeBlock(STA_SIZE_MARGE_X, STA_SIZE_MARGE_Y);
- std::cout << "merging: " << grid.y << " x " << grid.x << " ---> " << mergeGrid.y << " x " << mergeGrid.x << " for tiles: " << tileSizeY << " x " << tileSizeX << std::endl;
+ // debug log
+ // std::cout << "merging: " << grid.y << " x " << grid.x << " ---> " << mergeGrid.y << " x " << mergeGrid.x << " for tiles: " << tileSizeY << " x " << tileSizeX << std::endl;
crossMerge<<<mergeGrid, mergeBlock, 0, stream>>>(2, 2, tileSizeY, tileSizeX, edges, comps, ceilf(grid.y / 2.0) - grid.y / 2, ceilf(grid.x / 2.0) - grid.x / 2);
tileSizeX <<= 1;
tileSizeY <<= 1;
grid.y = divUp(edges.rows, block.y);
flatten<<<grid, block, 0, stream>>>(edges, comps);
cudaSafeCall( cudaGetLastError() );
+
if (stream == 0)
cudaSafeCall( cudaDeviceSynchronize() );
}
{\r
namespace ccl\r
{\r
- void labelComponents(const DevMem2D& edges, DevMem2Di comps, cudaStream_t stream);\r
+ void labelComponents(const DevMem2D& edges, DevMem2Di comps, int flags, cudaStream_t stream);\r
\r
template<typename T>\r
void computeEdges(const DevMem2D& image, DevMem2D edges, const float4& lo, const float4& hi, cudaStream_t stream);\r
components.create(mask.size(), CV_32SC1);\r
\r
cudaStream_t stream = StreamAccessor::getStream(s);\r
- device::ccl::labelComponents(mask, components, stream);\r
+ device::ccl::labelComponents(mask, components, 0, stream);\r
}\r
\r
namespace\r