diff --git a/modules/cudastereo/src/cuda/stereosgm.cu b/modules/cudastereo/src/cuda/stereosgm.cu index 3ad7efb317..88c26450c9 100644 --- a/modules/cudastereo/src/cuda/stereosgm.cu +++ b/modules/cudastereo/src/cuda/stereosgm.cu @@ -1342,11 +1342,11 @@ void PathAggregation::operator() (const GpuMat& left, const GpuMat& right, GpuMa const Size size = left.size(); const int buffer_step = size.width * size.height * static_cast(MAX_DISPARITY); - CV_Assert(dest.rows == 1 && buffer_step * num_paths == dest.cols); + CV_Assert(dest.rows == num_paths && dest.cols == buffer_step); for (int i = 0; i < num_paths; ++i) { - subs[i] = dest.colRange(i * buffer_step, (i + 1) * buffer_step); + subs[i] = dest.row(i); } vertical::aggregateUp2DownPath(left, right, subs[0], p1, p2, min_disp, streams[0]); @@ -1447,12 +1447,10 @@ __global__ void winner_takes_all_kernel( ? REDUCTION_PER_THREAD : ACCUMULATION_INTERVAL; - const unsigned int cost_step = MAX_DISPARITY * width * height; const unsigned int warp_id = cudev::Warp::warpId(); const unsigned int lane_id = cudev::Warp::laneId(); const unsigned int y = blockIdx.x * WARPS_PER_BLOCK + warp_id; - const PtrStep src{ (uint8_t*)&_src(0, y * MAX_DISPARITY * width), height * width * MAX_DISPARITY * NUM_PATHS }; PtrStep left_dest{ _left_dest.ptr(y), _left_dest.step }; PtrStep right_dest{ _right_dest.ptr(y), _right_dest.step }; @@ -1492,7 +1490,7 @@ __global__ void winner_takes_all_kernel( { uint32_t load_buffer[ACCUMULATION_PER_THREAD]; load_uint8_vector( - load_buffer, &src(0, p * cost_step + offset)); + load_buffer, &_src(p, y * MAX_DISPARITY * width + offset)); for (unsigned int i = 0; i < ACCUMULATION_PER_THREAD; ++i) { sum[i] += load_buffer[i]; @@ -1583,7 +1581,7 @@ void winnerTakesAll(const GpuMat& src, GpuMat& left, GpuMat& right, float unique { cv::Size size = left.size(); int num_paths = mode == StereoSGBM::MODE_HH4 ? 4 : 8; - CV_Assert(src.rows == 1 && src.cols == size.width * size.height * static_cast(MAX_DISPARITY) * num_paths); + CV_Assert(src.rows == num_paths && src.cols == size.width * size.height * static_cast(MAX_DISPARITY)); CV_Assert(size == right.size()); CV_Assert(left.type() == right.type()); CV_Assert(src.type() == CV_8UC1); diff --git a/modules/cudastereo/src/stereosgm.cpp b/modules/cudastereo/src/stereosgm.cpp index 11f630a254..fe6765d4fc 100644 --- a/modules/cudastereo/src/stereosgm.cpp +++ b/modules/cudastereo/src/stereosgm.cpp @@ -116,7 +116,7 @@ void StereoSGMImpl::compute(InputArray _left, InputArray _right, OutputArray _di census_transform::censusTransform(left, censused_left, _stream); census_transform::censusTransform(right, censused_right, _stream); - ensureSizeIsEnough(1, size.width * size.height * params.numDisparities * num_paths, CV_8UC1, aggregated); + ensureSizeIsEnough(num_paths, size.width * size.height * params.numDisparities, CV_8UC1, aggregated); ensureSizeIsEnough(size, CV_16SC1, left_disp_tmp); ensureSizeIsEnough(size, CV_16SC1, right_disp_tmp); diff --git a/modules/cudastereo/test/test_sgm_funcs.cpp b/modules/cudastereo/test/test_sgm_funcs.cpp index 3f185ee857..440fe2d811 100644 --- a/modules/cudastereo/test/test_sgm_funcs.cpp +++ b/modules/cudastereo/test/test_sgm_funcs.cpp @@ -359,9 +359,8 @@ namespace opencv_test { namespace { for (int k = 0; k < disparity; ++k) { int cost_sum = 0; for (int p = 0; p < num_paths; ++p) { - cost_sum += static_cast(src.at(0, - p * disparity * width * height + - i * disparity * width + + cost_sum += static_cast(src.at(p, + i * disparity * width + j * disparity + k)); } @@ -418,7 +417,7 @@ namespace opencv_test { namespace { CUDA_TEST_P(StereoSGM_WinnerTakesAll, RandomLeft) { int num_paths = mode == cv::cuda::StereoSGM::MODE_HH4 ? 4 : 8; - cv::Mat aggregated = randomMat(cv::Size(size.width * size.height * DISPARITY * num_paths, 1), CV_8UC1, 0.0, 32.0); + cv::Mat aggregated = randomMat(cv::Size(size.width * size.height * DISPARITY, num_paths), CV_8UC1, 0.0, 32.0); cv::Mat dst_gold; winner_takes_all_left(aggregated, dst_gold, size.width, size.height, DISPARITY, num_paths, 0.95f, subpixel); diff --git a/modules/cudastereo/test/test_stereo.cpp b/modules/cudastereo/test/test_stereo.cpp index 59f2df9353..fa23ba81ab 100644 --- a/modules/cudastereo/test/test_stereo.cpp +++ b/modules/cudastereo/test/test_stereo.cpp @@ -1135,5 +1135,42 @@ INSTANTIATE_TEST_CASE_P(CUDA_Stereo, StereoSGM_NonAligned, testing::Combine( testing::Values(cv::Size(648, 488), cv::Size(320, 248)), testing::Values(cv::cuda::StereoSGM::MODE_HH4, cv::cuda::StereoSGM::MODE_HH))); +PARAM_TEST_CASE(StereoSGM_AggregatedBufferLayout, cv::cuda::DeviceInfo, cv::Size) +{ + cv::cuda::DeviceInfo devInfo; + cv::Size size; + + virtual void SetUp() + { + devInfo = GET_PARAM(0); + size = GET_PARAM(1); + cv::cuda::setDevice(devInfo.deviceID()); + } +}; + +CUDA_TEST_P(StereoSGM_AggregatedBufferLayout, MaxDisparityHH) +{ + cv::Mat left_cpu(size, CV_8UC1); + cv::Mat right_cpu(size, CV_8UC1); + cv::RNG& rng = cv::theRNG(); + rng.fill(left_cpu, cv::RNG::UNIFORM, 0, 256); + rng.fill(right_cpu, cv::RNG::UNIFORM, 0, 256); + + cv::cuda::GpuMat d_left, d_right, d_disp; + d_left.upload(left_cpu); + d_right.upload(right_cpu); + + cv::Ptr sgm = cv::cuda::createStereoSGM( + 0, 256, 10, 120, 5, cv::cuda::StereoSGM::MODE_HH); + sgm->compute(d_left, d_right, d_disp); + + EXPECT_EQ(d_disp.size(), size); + EXPECT_EQ(d_disp.type(), CV_16SC1); +} + +INSTANTIATE_TEST_CASE_P(CUDA_Stereo, StereoSGM_AggregatedBufferLayout, testing::Combine( + ALL_DEVICES, + testing::Values(cv::Size(320, 32)))); + }} // namespace #endif // HAVE_CUDA