diff --git a/cv-node/README.md b/cv-node/README.md index 5e024ad8f0d2e4bd187528e33ff4b1a15e3aa234..309950751026663a24e0139d5731475b7e20299d 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 ff0e489a5ddcb32e5578da8f79032cac0e8191e7..bb130e0b5820df2114bbeb6393982a3a95237668 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[];