From b7c88ad3a1d6577ba1f815bc86e125263f69ab8c Mon Sep 17 00:00:00 2001 From: Eijiro SHIBUSAWA Date: Tue, 27 Jun 2023 23:28:51 +0900 Subject: [PATCH 1/2] * bug fix --- stereotgv/calcTensor.cu | 2 +- stereotgv/conversion.cu | 8 ++++---- stereotgv/medianFilter.cu | 6 +++--- stereotgv/stereotgv.cu | 4 ++-- stereotgv/updatePrimalVariables.cu | 2 +- 5 files changed, 11 insertions(+), 11 deletions(-) diff --git a/stereotgv/calcTensor.cu b/stereotgv/calcTensor.cu index 712de9a..94124bd 100644 --- a/stereotgv/calcTensor.cu +++ b/stereotgv/calcTensor.cu @@ -224,7 +224,7 @@ void StereoTgv::CalcTensorMasked(float* gray, float* mask, float beta, float gam float* a, float* b, float* c) { dim3 threads(BlockWidth, BlockHeight); - dim3 blocks(iDivUp(h, threads.x), iDivUp(h, threads.y)); + dim3 blocks(iDivUp(w, threads.x), iDivUp(h, threads.y)); // mirror if a coordinate value is out-of-range gray_img.addressMode[0] = cudaAddressModeMirror; diff --git a/stereotgv/conversion.cu b/stereotgv/conversion.cu index bd70f8b..49bc4dd 100644 --- a/stereotgv/conversion.cu +++ b/stereotgv/conversion.cu @@ -16,16 +16,16 @@ __global__ void TgvConvertKBKernel(float2* disparity, float u0 = (float)ix; float v0 = (float)iy; - float xprime0 = (u0 - focalx) / cx; - float yprime0 = (v0 - focaly) / cy; + float xprime0 = (u0 - cx) / focalx; + float yprime0 = (v0 - cy) / focaly; float u = disparity[pos].x; float v = disparity[pos].y; float u1 = u0 + u; float v1 = v0 + v; - float xprime1 = (u1 - focalx) / cx; - float yprime1 = (v1 - focaly) / cy; + float xprime1 = (u1 - cx) / focalx; + float yprime1 = (v1 - cy) / focaly; // Newton-Raphson Method Frame 0 float ru0 = sqrtf(xprime0 * xprime0 + yprime0 * yprime0); diff --git a/stereotgv/medianFilter.cu b/stereotgv/medianFilter.cu index a99f507..b941c09 100644 --- a/stereotgv/medianFilter.cu +++ b/stereotgv/medianFilter.cu @@ -121,7 +121,7 @@ __global__ void TgvMedianFilterKernel3(float* u, float* v, } float tmpu, tmpv; - for (int j = 0; j < 9; j++) { + for (int j = 0; j < 5; j++) { for (int i = j + 1; i < 9; i++) { if (mu[j] > mu[i]) { //Swap the variables. @@ -357,7 +357,7 @@ __global__ void TgvMedianFilterDisparityKernel3(float* u, } float tmpu, tmpv; - for (int j = 0; j < 9; j++) { + for (int j = 0; j < 5; j++) { for (int i = j + 1; i < 9; i++) { if (mu[j] > mu[i]) { //Swap the variables. @@ -490,7 +490,7 @@ __global__ void TgvMedianFilterDisparityMaskedKernel3(float* u, float* mask, } float tmpu, tmpv; - for (int j = 0; j < 9; j++) { + for (int j = 0; j < 5; j++) { for (int i = j + 1; i < 9; i++) { if (mu[j] > mu[i]) { //Swap the variables. diff --git a/stereotgv/stereotgv.cu b/stereotgv/stereotgv.cu index 8b4e460..39fe21c 100644 --- a/stereotgv/stereotgv.cu +++ b/stereotgv/stereotgv.cu @@ -181,9 +181,9 @@ int StereoTgv::loadVectorFields(cv::Mat translationVector, cv::Mat calibrationVe checkCudaErrors(cudaMemcpy(d_tvForward, (float2 *)translationVectorPad.ptr(), dataSize32fc2, cudaMemcpyHostToDevice)); - pTvForward[0] = d_tvForward; + Swap(pTvForward[0], d_tvForward); ScalarMultiply(d_tvForward, -1.0f, width, height, stride, d_tvBackward); - pTvBackward[0] = d_tvBackward; + Swap(pTvBackward[0], d_tvBackward); for (int level = 1; level < nLevels; level++) { //std::cout << "vectorfields " << pW[level] << " " << pH[level] << " " << pS[level] << std::endl; Downscale(pTvForward[level - 1], pW[level - 1], pH[level - 1], pS[level - 1], diff --git a/stereotgv/updatePrimalVariables.cu b/stereotgv/updatePrimalVariables.cu index e3164d5..1b48053 100644 --- a/stereotgv/updatePrimalVariables.cu +++ b/stereotgv/updatePrimalVariables.cu @@ -235,7 +235,7 @@ void TgvUpdatePrimalVariablesMaskedKernel(float* mask, float* u_, float2* v_, fl vpos.y = v_pos.y + (tau / eta_v2s) * (alpha1 * dq_tensor.y + alpha0 * div_q.y); } else { - vpos.x = v_pos.x + (tau) * (alpha1 * dq_tensor.x + alpha0 * div_q.x); + vpos.y = v_pos.y + (tau) * (alpha1 * dq_tensor.y + alpha0 * div_q.y); } v[pos] = vpos; From 5d75985a7856d18a3e18cbb9651968f3ff6017f5 Mon Sep 17 00:00:00 2001 From: Eijiro SHIBUSAWA Date: Wed, 28 Jun 2023 00:21:03 +0900 Subject: [PATCH 2/2] * fix agin --- stereotgv/stereotgv.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/stereotgv/stereotgv.cu b/stereotgv/stereotgv.cu index 39fe21c..b5895df 100644 --- a/stereotgv/stereotgv.cu +++ b/stereotgv/stereotgv.cu @@ -181,8 +181,8 @@ int StereoTgv::loadVectorFields(cv::Mat translationVector, cv::Mat calibrationVe checkCudaErrors(cudaMemcpy(d_tvForward, (float2 *)translationVectorPad.ptr(), dataSize32fc2, cudaMemcpyHostToDevice)); - Swap(pTvForward[0], d_tvForward); ScalarMultiply(d_tvForward, -1.0f, width, height, stride, d_tvBackward); + Swap(pTvForward[0], d_tvForward); Swap(pTvBackward[0], d_tvBackward); for (int level = 1; level < nLevels; level++) { //std::cout << "vectorfields " << pW[level] << " " << pH[level] << " " << pS[level] << std::endl;