|
|
|
@ -485,7 +485,9 @@ namespace cv { namespace gpu { namespace device |
|
|
|
|
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)); |
|
|
|
@ -494,15 +496,12 @@ namespace cv { namespace gpu { namespace device |
|
|
|
|
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; |
|
|
|
@ -515,6 +514,7 @@ namespace cv { namespace gpu { namespace device |
|
|
|
|
grid.y = divUp(edges.rows, block.y); |
|
|
|
|
flatten<<<grid, block, 0, stream>>>(edges, comps); |
|
|
|
|
cudaSafeCall( cudaGetLastError() ); |
|
|
|
|
|
|
|
|
|
if (stream == 0) |
|
|
|
|
cudaSafeCall( cudaDeviceSynchronize() ); |
|
|
|
|
} |
|
|
|
|