From 8836624aba523607b49a27014991a932a695819f Mon Sep 17 00:00:00 2001 From: Nicolas Pope <nwpope@utu.fi> Date: Mon, 18 Mar 2019 13:30:08 +0200 Subject: [PATCH] Comments added --- cv-node/README.md | 7 ++++-- cv-node/src/algorithms/rtcensus.cu | 39 ++++++++++++++++++------------ 2 files changed, 29 insertions(+), 17 deletions(-) diff --git a/cv-node/README.md b/cv-node/README.md index 5e024ad8f..309950751 100644 --- a/cv-node/README.md +++ b/cv-node/README.md @@ -14,7 +14,8 @@ make ` ## Install -TODO. Currently, copy `<source_direction>/config/config.json` to `~/config/ftl`. +TODO. Currently, copy `<source_direction>/config/config.json` to +`~/.config/ftl/config.json` on Linux. ## Usage An optional command-line argument can be passed to specify a stereo video file @@ -30,7 +31,9 @@ JSON config option at `{ "disparity": { "algorithm": "sgbm" }}`. ### Calibration Cameras can be calibrated by using argument `--calibrate`. You either need to provide a calibration video as an argument or do it live. A checkerboard -grid pattern is required, the size can be configured in the json file. +grid pattern is required, the size can be configured in the json file. After +callibration the data is save and will be reloaded automatically next time you +run `cv-node`. Note: best calibration is not too close or far from the cameras, the board must be fully visible in each camera and move to cover as much of the visual field diff --git a/cv-node/src/algorithms/rtcensus.cu b/cv-node/src/algorithms/rtcensus.cu index ff0e489a5..bb130e0b5 100644 --- a/cv-node/src/algorithms/rtcensus.cu +++ b/cv-node/src/algorithms/rtcensus.cu @@ -26,29 +26,43 @@ using namespace cv; namespace ftl { namespace gpu { +// --- SUPPORT ----------------------------------------------------------------- + +/* + * Sparse 16x16 census (so 8x8) creating a 64bit mask + * (14) & (15), based upon (9) + */ __device__ uint64_t sparse_census(unsigned char *arr, size_t u, size_t v, size_t w) { uint64_t r = 0; unsigned char t = arr[v*w+u]; - for (int n=-7; n<=7; n+=2) { - auto u_ = u + n; for (int m=-7; m<=7; m+=2) { - auto v_ = v + m; - r <<= 1; - r |= XHI(t, arr[v_*w+u_]); - } + auto start_ix = (v + m)*w + u; + for (int n=-7; n<=7; n+=2) { + r <<= 1; + r |= XHI(t, arr[start_ix+n]); + } } return r; } +/* + * Parabolic interpolation between matched disparities either side. + * Results in subpixel disparity. (20). + */ __device__ float fit_parabola(size_t pi, uint16_t p, uint16_t pl, uint16_t pr) { float a = pr - pl; float b = 2 * (2 * p - pl - pr); return static_cast<float>(pi) + (a / b); } +// --- KERNELS ----------------------------------------------------------------- + +/* + * Calculate census mask for left and right images together. + */ __global__ void census_kernel(PtrStepSzb l, PtrStepSzb r, uint64_t *census) { //extern __shared__ uint64_t census[]; @@ -62,23 +76,18 @@ __global__ void census_kernel(PtrStepSzb l, PtrStepSzb r, uint64_t *census) { size_t width = l.cols; for (size_t v=v_start; v<v_end; v++) { - //for (size_t u=7; u<width-7; u++) { size_t ix = (u + v*width) * 2; uint64_t cenL = sparse_census(l.data, u, v, l.step); uint64_t cenR = sparse_census(r.data, u, v, r.step); census[ix] = cenL; census[ix + 1] = cenR; - - //disp(v,u) = (float)cenL; - //} } - - //__syncthreads(); - - return; } - + +/* + * Generate left and right disparity images from census data. (19) + */ __global__ void disp_kernel(float *disp_l, float *disp_r, size_t width, size_t height, uint64_t *census, size_t ds) { //extern __shared__ uint64_t cache[]; -- GitLab