From 4376f7d409ea0bf6e364dabb0dc97c4bda7780d6 Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nwpope@utu.fi>
Date: Thu, 21 Mar 2019 07:28:25 +0200
Subject: [PATCH] Minor changes and comments
---
cv-node/src/algorithms/rtcensus.cu | 31 ++++++++++++++++--------------
cv-node/src/middlebury.cpp | 2 +-
2 files changed, 18 insertions(+), 15 deletions(-)
diff --git a/cv-node/src/algorithms/rtcensus.cu b/cv-node/src/algorithms/rtcensus.cu
index 39650ea5c..6d48a392a 100644
--- a/cv-node/src/algorithms/rtcensus.cu
+++ b/cv-node/src/algorithms/rtcensus.cu
@@ -84,20 +84,23 @@ __global__ void census_kernel(cudaTextureObject_t l, cudaTextureObject_t r,
}
}
-__forceinline__ __device__ unsigned long long int int2_as_longlong (uint2 a)
-{
- unsigned long long int res;
- asm ("mov.b64 %0, {%1,%2};" : "=l"(res) : "r"(a.x), "r"(a.y));
- return res;
+/* Convert vector uint2 (32bit x2) into a single uint64_t */
+__forceinline__ __device__ uint64_t uint2asull (uint2 a) {
+ uint64_t res;
+ asm ("mov.b64 %0, {%1,%2};" : "=l"(res) : "r"(a.x), "r"(a.y));
+ return res;
}
/*
* 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, cudaTextureObject_t censusL, cudaTextureObject_t censusR, size_t ds) {
+__global__ void disp_kernel(float *disp_l, float *disp_r,
+ size_t width, size_t height,
+ cudaTextureObject_t censusL, cudaTextureObject_t censusR,
+ size_t ds) {
//extern __shared__ uint64_t cache[];
- const int gamma = 5;
+ const int gamma = 10;
size_t u = (blockIdx.x * BLOCK_W) + threadIdx.x + RADIUS2;
size_t v_start = (blockIdx.y * ROWSperTHREAD) + RADIUS2;
@@ -147,11 +150,11 @@ __global__ void disp_kernel(float *disp_l, float *disp_r, size_t width, size_t h
- auto l2 = int2_as_longlong(tex2D<uint2>(censusL,u_,v_));
- auto l1 = int2_as_longlong(tex2D<uint2>(censusR,u_,v_));
+ auto l2 = uint2asull(tex2D<uint2>(censusL,u_,v_));
+ auto l1 = uint2asull(tex2D<uint2>(censusR,u_,v_));
- auto r1 = int2_as_longlong(tex2D<uint2>(censusL, u_+d, v_));
- auto r2 = int2_as_longlong(tex2D<uint2>(censusR, u_-d, v_));
+ auto r1 = uint2asull(tex2D<uint2>(censusL, u_+d, v_));
+ auto r2 = uint2asull(tex2D<uint2>(censusR, u_-d, v_));
hamming1 += __popcll(r1^l1);
hamming2 += __popcll(r2^l2);
@@ -190,8 +193,8 @@ __global__ void disp_kernel(float *disp_l, float *disp_r, size_t width, size_t h
// Confidence filter (25)
// TODO choice of gamma to depend on disparity variance
// Variance with next option, variance with neighbours, variance with past value
- disp_l[v*width+u] = ((min_disp2b - min_disp2) >= gamma) ? d2 : INFINITY;
- disp_r[v*width+u] = ((min_disp1b - min_disp1) >= gamma) ? d1 : INFINITY;
+ disp_l[v*width+u] = ((min_disp2b - min_disp2) >= gamma) ? d2 : NAN;
+ disp_r[v*width+u] = ((min_disp1b - min_disp1) >= gamma) ? d1 : NAN;
// TODO If disparity is 0.0f, perhaps
// Use previous value unless it conflicts with present
@@ -219,7 +222,7 @@ __global__ void consistency_kernel(float *d_sub_l, float *d_sub_r, PtrStepSz<flo
auto b = d_sub_r[v*w+u-a];
if (abs(a-b) <= 1.0) disp(v,u) = abs((a+b)/2); // was 1.0
- else disp(v,u) = INFINITY;
+ else disp(v,u) = NAN;
//}
}
diff --git a/cv-node/src/middlebury.cpp b/cv-node/src/middlebury.cpp
index b8b01e267..9252a4374 100644
--- a/cv-node/src/middlebury.cpp
+++ b/cv-node/src/middlebury.cpp
@@ -195,7 +195,7 @@ void ftl::middlebury::evaldisp(const Mat &disp, const Mat >disp, const Mat &ma
if (gt == INFINITY) // unknown
continue;
float d = scale * disp.at<float>(y / scale, x / scale, 0);
- int valid = (d != INFINITY);
+ int valid = (d != NAN); // NOTE: Is meant to be infinity in middlebury
if (valid) {
float maxd = scale * maxdisp; // max disp range
d = max(0.0f, min(maxd, d)); // clip disps to max disp range
--
GitLab