Merge remote-tracking branch 'origin/2.4' into merge-2.4
Conflicts: CMakeLists.txt cmake/OpenCVGenAndroidMK.cmake cmake/templates/OpenCV.mk.in doc/tutorials/viz/creating_widgets/creating_widgets.rst doc/tutorials/viz/launching_viz/launching_viz.rst doc/tutorials/viz/table_of_content_viz/images/image_effects.png doc/tutorials/viz/transformations/transformations.rst doc/tutorials/viz/widget_pose/widget_pose.rst modules/core/include/opencv2/core/affine.hpp modules/core/include/opencv2/core/core.hpp modules/gpu/perf/perf_imgproc.cpp modules/gpu/src/cuda/canny.cu modules/gpu/src/cuda/generalized_hough.cu modules/gpu/src/generalized_hough.cpp modules/gpu/src/imgproc.cpp modules/gpu/test/test_color.cpp modules/gpu/test/test_core.cpp modules/gpu/test/test_gpumat.cpp modules/gpu/test/test_hough.cpp modules/nonfree/CMakeLists.txt modules/nonfree/include/opencv2/nonfree/gpu.hpp modules/nonfree/perf/perf_gpu.cpp modules/nonfree/src/cuda/surf.cu modules/nonfree/src/precomp.hpp modules/nonfree/src/surf_gpu.cpp modules/nonfree/test/test_gpu.cpp modules/ocl/perf/perf_haar.cpp modules/stitching/CMakeLists.txt modules/stitching/include/opencv2/stitching/detail/matchers.hpp modules/stitching/include/opencv2/stitching/detail/seam_finders.hpp modules/stitching/include/opencv2/stitching/detail/warpers.hpp modules/stitching/include/opencv2/stitching/warpers.hpp modules/stitching/src/blenders.cpp modules/stitching/src/matchers.cpp modules/stitching/src/precomp.hpp modules/stitching/src/seam_finders.cpp modules/stitching/src/stitcher.cpp modules/stitching/src/warpers.cpp modules/viz/doc/widget.rst modules/viz/include/opencv2/viz/types.hpp modules/viz/include/opencv2/viz/viz3d.hpp modules/viz/include/opencv2/viz/widget_accessor.hpp modules/viz/src/precomp.hpp modules/viz/src/shapes.cpp modules/viz/src/vizcore.cpp modules/viz/src/vtk/vtkCloudMatSink.h modules/viz/src/vtk/vtkCloudMatSource.h modules/viz/test/test_precomp.hpp modules/viz/test/tests_simple.cpp samples/android/tutorial-4-cuda/CMakeLists.txt samples/android/tutorial-4-cuda/jni/Android.mk samples/android/tutorial-4-cuda/src/org/opencv/samples/tutorial4/Tutorial4Activity.java samples/cpp/stitching_detailed.cpp samples/cpp/tutorial_code/viz/creating_widgets.cpp samples/cpp/tutorial_code/viz/launching_viz.cpp samples/cpp/tutorial_code/viz/transformations.cpp samples/cpp/tutorial_code/viz/widget_pose.cpp
This commit is contained in:
@@ -70,7 +70,7 @@ namespace cv
|
||||
//Rodrigues vector
|
||||
Affine3(const Vec3& rvec, const Vec3& t = Vec3::all(0));
|
||||
|
||||
//Combines all contructors above. Supports 4x4, 3x4, 3x3, 1x3, 3x1 sizes of data matrix
|
||||
//Combines all contructors above. Supports 4x4, 4x3, 3x3, 1x3, 3x1 sizes of data matrix
|
||||
explicit Affine3(const Mat& data, const Vec3& t = Vec3::all(0));
|
||||
|
||||
//From 16th element array
|
||||
|
@@ -281,7 +281,7 @@ CUDA_TEST_P(ConvertTo, WithOutScaling)
|
||||
cv::Mat dst_gold;
|
||||
src.convertTo(dst_gold, depth2);
|
||||
|
||||
EXPECT_MAT_NEAR(dst_gold, dst, 1.0);
|
||||
EXPECT_MAT_NEAR(dst_gold, dst, depth2 < CV_32F ? 1.0 : 1e-4);
|
||||
}
|
||||
}
|
||||
|
||||
|
@@ -734,7 +734,7 @@ CUDA_TEST_P(Normalize, WithOutMask)
|
||||
cv::Mat dst_gold;
|
||||
cv::normalize(src, dst_gold, alpha, beta, norm_type, type);
|
||||
|
||||
EXPECT_MAT_NEAR(dst_gold, dst, 1.0);
|
||||
EXPECT_MAT_NEAR(dst_gold, dst, type < CV_32F ? 1.0 : 1e-4);
|
||||
}
|
||||
|
||||
CUDA_TEST_P(Normalize, WithMask)
|
||||
|
@@ -58,9 +58,9 @@ namespace canny
|
||||
|
||||
void calcMap(PtrStepSzi dx, PtrStepSzi dy, PtrStepSzf mag, PtrStepSzi map, float low_thresh, float high_thresh);
|
||||
|
||||
void edgesHysteresisLocal(PtrStepSzi map, ushort2* st1);
|
||||
void edgesHysteresisLocal(PtrStepSzi map, short2* st1);
|
||||
|
||||
void edgesHysteresisGlobal(PtrStepSzi map, ushort2* st1, ushort2* st2);
|
||||
void edgesHysteresisGlobal(PtrStepSzi map, short2* st1, short2* st2);
|
||||
|
||||
void getEdges(PtrStepSzi map, PtrStepSzb dst);
|
||||
}
|
||||
@@ -194,6 +194,8 @@ namespace
|
||||
|
||||
void CannyImpl::createBuf(Size image_size)
|
||||
{
|
||||
CV_Assert(image_size.width < std::numeric_limits<short>::max() && image_size.height < std::numeric_limits<short>::max());
|
||||
|
||||
ensureSizeIsEnough(image_size, CV_32SC1, dx_);
|
||||
ensureSizeIsEnough(image_size, CV_32SC1, dy_);
|
||||
|
||||
@@ -209,8 +211,8 @@ namespace
|
||||
ensureSizeIsEnough(image_size, CV_32FC1, mag_);
|
||||
ensureSizeIsEnough(image_size, CV_32SC1, map_);
|
||||
|
||||
ensureSizeIsEnough(1, image_size.area(), CV_16UC2, st1_);
|
||||
ensureSizeIsEnough(1, image_size.area(), CV_16UC2, st2_);
|
||||
ensureSizeIsEnough(1, image_size.area(), CV_16SC2, st1_);
|
||||
ensureSizeIsEnough(1, image_size.area(), CV_16SC2, st2_);
|
||||
}
|
||||
|
||||
void CannyImpl::CannyCaller(GpuMat& edges)
|
||||
@@ -218,9 +220,9 @@ namespace
|
||||
map_.setTo(Scalar::all(0));
|
||||
canny::calcMap(dx_, dy_, mag_, map_, static_cast<float>(low_thresh_), static_cast<float>(high_thresh_));
|
||||
|
||||
canny::edgesHysteresisLocal(map_, st1_.ptr<ushort2>());
|
||||
canny::edgesHysteresisLocal(map_, st1_.ptr<short2>());
|
||||
|
||||
canny::edgesHysteresisGlobal(map_, st1_.ptr<ushort2>(), st2_.ptr<ushort2>());
|
||||
canny::edgesHysteresisGlobal(map_, st1_.ptr<short2>(), st2_.ptr<short2>());
|
||||
|
||||
canny::getEdges(map_, edges);
|
||||
}
|
||||
|
@@ -239,30 +239,35 @@ namespace canny
|
||||
{
|
||||
__device__ int counter = 0;
|
||||
|
||||
__global__ void edgesHysteresisLocalKernel(PtrStepSzi map, ushort2* st)
|
||||
__device__ __forceinline__ bool checkIdx(int y, int x, int rows, int cols)
|
||||
{
|
||||
return (y >= 0) && (y < rows) && (x >= 0) && (x < cols);
|
||||
}
|
||||
|
||||
__global__ void edgesHysteresisLocalKernel(PtrStepSzi map, short2* st)
|
||||
{
|
||||
__shared__ volatile int smem[18][18];
|
||||
|
||||
const int x = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
const int y = blockIdx.y * blockDim.y + threadIdx.y;
|
||||
|
||||
smem[threadIdx.y + 1][threadIdx.x + 1] = x < map.cols && y < map.rows ? map(y, x) : 0;
|
||||
smem[threadIdx.y + 1][threadIdx.x + 1] = checkIdx(y, x, map.rows, map.cols) ? map(y, x) : 0;
|
||||
if (threadIdx.y == 0)
|
||||
smem[0][threadIdx.x + 1] = y > 0 ? map(y - 1, x) : 0;
|
||||
smem[0][threadIdx.x + 1] = checkIdx(y - 1, x, map.rows, map.cols) ? map(y - 1, x) : 0;
|
||||
if (threadIdx.y == blockDim.y - 1)
|
||||
smem[blockDim.y + 1][threadIdx.x + 1] = y + 1 < map.rows ? map(y + 1, x) : 0;
|
||||
smem[blockDim.y + 1][threadIdx.x + 1] = checkIdx(y + 1, x, map.rows, map.cols) ? map(y + 1, x) : 0;
|
||||
if (threadIdx.x == 0)
|
||||
smem[threadIdx.y + 1][0] = x > 0 ? map(y, x - 1) : 0;
|
||||
smem[threadIdx.y + 1][0] = checkIdx(y, x - 1, map.rows, map.cols) ? map(y, x - 1) : 0;
|
||||
if (threadIdx.x == blockDim.x - 1)
|
||||
smem[threadIdx.y + 1][blockDim.x + 1] = x + 1 < map.cols ? map(y, x + 1) : 0;
|
||||
smem[threadIdx.y + 1][blockDim.x + 1] = checkIdx(y, x + 1, map.rows, map.cols) ? map(y, x + 1) : 0;
|
||||
if (threadIdx.x == 0 && threadIdx.y == 0)
|
||||
smem[0][0] = y > 0 && x > 0 ? map(y - 1, x - 1) : 0;
|
||||
smem[0][0] = checkIdx(y - 1, x - 1, map.rows, map.cols) ? map(y - 1, x - 1) : 0;
|
||||
if (threadIdx.x == blockDim.x - 1 && threadIdx.y == 0)
|
||||
smem[0][blockDim.x + 1] = y > 0 && x + 1 < map.cols ? map(y - 1, x + 1) : 0;
|
||||
smem[0][blockDim.x + 1] = checkIdx(y - 1, x + 1, map.rows, map.cols) ? map(y - 1, x + 1) : 0;
|
||||
if (threadIdx.x == 0 && threadIdx.y == blockDim.y - 1)
|
||||
smem[blockDim.y + 1][0] = y + 1 < map.rows && x > 0 ? map(y + 1, x - 1) : 0;
|
||||
smem[blockDim.y + 1][0] = checkIdx(y + 1, x - 1, map.rows, map.cols) ? map(y + 1, x - 1) : 0;
|
||||
if (threadIdx.x == blockDim.x - 1 && threadIdx.y == blockDim.y - 1)
|
||||
smem[blockDim.y + 1][blockDim.x + 1] = y + 1 < map.rows && x + 1 < map.cols ? map(y + 1, x + 1) : 0;
|
||||
smem[blockDim.y + 1][blockDim.x + 1] = checkIdx(y + 1, x + 1, map.rows, map.cols) ? map(y + 1, x + 1) : 0;
|
||||
|
||||
__syncthreads();
|
||||
|
||||
@@ -317,11 +322,11 @@ namespace canny
|
||||
if (n > 0)
|
||||
{
|
||||
const int ind = ::atomicAdd(&counter, 1);
|
||||
st[ind] = make_ushort2(x, y);
|
||||
st[ind] = make_short2(x, y);
|
||||
}
|
||||
}
|
||||
|
||||
void edgesHysteresisLocal(PtrStepSzi map, ushort2* st1)
|
||||
void edgesHysteresisLocal(PtrStepSzi map, short2* st1)
|
||||
{
|
||||
void* counter_ptr;
|
||||
cudaSafeCall( cudaGetSymbolAddress(&counter_ptr, counter) );
|
||||
@@ -345,13 +350,13 @@ namespace canny
|
||||
__constant__ int c_dx[8] = {-1, 0, 1, -1, 1, -1, 0, 1};
|
||||
__constant__ int c_dy[8] = {-1, -1, -1, 0, 0, 1, 1, 1};
|
||||
|
||||
__global__ void edgesHysteresisGlobalKernel(PtrStepSzi map, ushort2* st1, ushort2* st2, const int count)
|
||||
__global__ void edgesHysteresisGlobalKernel(PtrStepSzi map, short2* st1, short2* st2, const int count)
|
||||
{
|
||||
const int stack_size = 512;
|
||||
|
||||
__shared__ int s_counter;
|
||||
__shared__ int s_ind;
|
||||
__shared__ ushort2 s_st[stack_size];
|
||||
__shared__ short2 s_st[stack_size];
|
||||
|
||||
if (threadIdx.x == 0)
|
||||
s_counter = 0;
|
||||
@@ -363,14 +368,14 @@ namespace canny
|
||||
if (ind >= count)
|
||||
return;
|
||||
|
||||
ushort2 pos = st1[ind];
|
||||
short2 pos = st1[ind];
|
||||
|
||||
if (threadIdx.x < 8)
|
||||
{
|
||||
pos.x += c_dx[threadIdx.x];
|
||||
pos.y += c_dy[threadIdx.x];
|
||||
|
||||
if (pos.x > 0 && pos.x < map.cols && pos.y > 0 && pos.y < map.rows && map(pos.y, pos.x) == 1)
|
||||
if (pos.x > 0 && pos.x < map.cols - 1 && pos.y > 0 && pos.y < map.rows - 1 && map(pos.y, pos.x) == 1)
|
||||
{
|
||||
map(pos.y, pos.x) = 2;
|
||||
|
||||
@@ -402,7 +407,7 @@ namespace canny
|
||||
pos.x += c_dx[threadIdx.x & 7];
|
||||
pos.y += c_dy[threadIdx.x & 7];
|
||||
|
||||
if (pos.x > 0 && pos.x < map.cols && pos.y > 0 && pos.y < map.rows && map(pos.y, pos.x) == 1)
|
||||
if (pos.x > 0 && pos.x < map.cols - 1 && pos.y > 0 && pos.y < map.rows - 1 && map(pos.y, pos.x) == 1)
|
||||
{
|
||||
map(pos.y, pos.x) = 2;
|
||||
|
||||
@@ -419,8 +424,10 @@ namespace canny
|
||||
{
|
||||
if (threadIdx.x == 0)
|
||||
{
|
||||
ind = ::atomicAdd(&counter, s_counter);
|
||||
s_ind = ind - s_counter;
|
||||
s_ind = ::atomicAdd(&counter, s_counter);
|
||||
|
||||
if (s_ind + s_counter > map.cols * map.rows)
|
||||
s_counter = 0;
|
||||
}
|
||||
|
||||
__syncthreads();
|
||||
@@ -432,7 +439,7 @@ namespace canny
|
||||
}
|
||||
}
|
||||
|
||||
void edgesHysteresisGlobal(PtrStepSzi map, ushort2* st1, ushort2* st2)
|
||||
void edgesHysteresisGlobal(PtrStepSzi map, short2* st1, short2* st2)
|
||||
{
|
||||
void* counter_ptr;
|
||||
cudaSafeCall( cudaGetSymbolAddress(&counter_ptr, canny::counter) );
|
||||
@@ -454,6 +461,8 @@ namespace canny
|
||||
|
||||
cudaSafeCall( cudaMemcpy(&count, counter_ptr, sizeof(int), cudaMemcpyDeviceToHost) );
|
||||
|
||||
count = min(count, map.cols * map.rows);
|
||||
|
||||
std::swap(st1, st2);
|
||||
}
|
||||
}
|
||||
|
@@ -715,7 +715,7 @@ CUDA_TEST_P(CvtColor, BGR2YCrCb)
|
||||
cv::Mat dst_gold;
|
||||
cv::cvtColor(src, dst_gold, cv::COLOR_BGR2YCrCb);
|
||||
|
||||
EXPECT_MAT_NEAR(dst_gold, dst, 1.0);
|
||||
EXPECT_MAT_NEAR(dst_gold, dst, depth == CV_32F ? 1e-2 : 1);
|
||||
}
|
||||
|
||||
CUDA_TEST_P(CvtColor, RGB2YCrCb)
|
||||
@@ -728,7 +728,7 @@ CUDA_TEST_P(CvtColor, RGB2YCrCb)
|
||||
cv::Mat dst_gold;
|
||||
cv::cvtColor(src, dst_gold, cv::COLOR_RGB2YCrCb);
|
||||
|
||||
EXPECT_MAT_NEAR(dst_gold, dst, 1.0);
|
||||
EXPECT_MAT_NEAR(dst_gold, dst, depth == CV_32F ? 1e-2 : 1);
|
||||
}
|
||||
|
||||
CUDA_TEST_P(CvtColor, BGR2YCrCb4)
|
||||
@@ -749,7 +749,7 @@ CUDA_TEST_P(CvtColor, BGR2YCrCb4)
|
||||
cv::split(h_dst, channels);
|
||||
cv::merge(channels, 3, h_dst);
|
||||
|
||||
EXPECT_MAT_NEAR(dst_gold, h_dst, 1.0);
|
||||
EXPECT_MAT_NEAR(dst_gold, h_dst, depth == CV_32F ? 1e-2 : 1);
|
||||
}
|
||||
|
||||
CUDA_TEST_P(CvtColor, RGBA2YCrCb4)
|
||||
@@ -771,7 +771,7 @@ CUDA_TEST_P(CvtColor, RGBA2YCrCb4)
|
||||
cv::split(h_dst, channels);
|
||||
cv::merge(channels, 3, h_dst);
|
||||
|
||||
EXPECT_MAT_NEAR(dst_gold, h_dst, 1.0);
|
||||
EXPECT_MAT_NEAR(dst_gold, h_dst, depth == CV_32F ? 1e-2 : 1);
|
||||
}
|
||||
|
||||
CUDA_TEST_P(CvtColor, YCrCb2BGR)
|
||||
|
@@ -48,7 +48,17 @@ public class OpenCVLoader
|
||||
*/
|
||||
public static boolean initDebug()
|
||||
{
|
||||
return StaticHelper.initOpenCV();
|
||||
return StaticHelper.initOpenCV(false);
|
||||
}
|
||||
|
||||
/**
|
||||
* Loads and initializes OpenCV library from current application package. Roughly, it's an analog of system.loadLibrary("opencv_java").
|
||||
* @param InitCuda load and initialize CUDA runtime libraries.
|
||||
* @return Returns true is initialization of OpenCV was successful.
|
||||
*/
|
||||
public static boolean initDebug(boolean InitCuda)
|
||||
{
|
||||
return StaticHelper.initOpenCV(InitCuda);
|
||||
}
|
||||
|
||||
/**
|
||||
|
@@ -7,11 +7,21 @@ import android.util.Log;
|
||||
|
||||
class StaticHelper {
|
||||
|
||||
public static boolean initOpenCV()
|
||||
public static boolean initOpenCV(boolean InitCuda)
|
||||
{
|
||||
boolean result;
|
||||
String libs = "";
|
||||
|
||||
if(InitCuda)
|
||||
{
|
||||
loadLibrary("cudart");
|
||||
loadLibrary("nppc");
|
||||
loadLibrary("nppi");
|
||||
loadLibrary("npps");
|
||||
loadLibrary("cufft");
|
||||
loadLibrary("cublas");
|
||||
}
|
||||
|
||||
Log.d(TAG, "Trying to get library list");
|
||||
|
||||
try
|
||||
@@ -52,7 +62,7 @@ class StaticHelper {
|
||||
try
|
||||
{
|
||||
System.loadLibrary(Name);
|
||||
Log.d(TAG, "OpenCV libs init was ok!");
|
||||
Log.d(TAG, "Library " + Name + " loaded");
|
||||
}
|
||||
catch(UnsatisfiedLinkError e)
|
||||
{
|
||||
|
@@ -512,6 +512,7 @@ void createLaplacePyrGpu(const Mat &img, int num_levels, std::vector<Mat> &pyr)
|
||||
(void)img;
|
||||
(void)num_levels;
|
||||
(void)pyr;
|
||||
CV_Error(Error::StsNotImplemented, "CUDA optimization is unavailable");
|
||||
#endif
|
||||
}
|
||||
|
||||
@@ -549,6 +550,7 @@ void restoreImageFromLaplacePyrGpu(std::vector<Mat> &pyr)
|
||||
gpu_pyr[0].download(pyr[0]);
|
||||
#else
|
||||
(void)pyr;
|
||||
CV_Error(Error::StsNotImplemented, "CUDA optimization is unavailable");
|
||||
#endif
|
||||
}
|
||||
|
||||
|
Reference in New Issue
Block a user