mirror of https://github.com/opencv/opencv.git
Alexander Mordvintesv
13 years ago
commit
4fb15ae1f0
121 changed files with 12554 additions and 8455 deletions
File diff suppressed because it is too large
Load Diff
File diff suppressed because it is too large
Load Diff
@ -0,0 +1,125 @@ |
||||
#include "perf_precomp.hpp" |
||||
|
||||
using namespace std; |
||||
using namespace cv; |
||||
using namespace cv::gpu; |
||||
using namespace cvtest; |
||||
using namespace testing; |
||||
|
||||
void printOsInfo() |
||||
{ |
||||
#if defined _WIN32 |
||||
# if defined _WIN64 |
||||
cout << "OS: Windows x64 \n" << endl; |
||||
# else |
||||
cout << "OS: Windows x32 \n" << endl; |
||||
# endif |
||||
#elif defined linux |
||||
# if defined _LP64 |
||||
cout << "OS: Linux x64 \n" << endl; |
||||
# else |
||||
cout << "OS: Linux x32 \n" << endl; |
||||
# endif |
||||
#elif defined __APPLE__ |
||||
# if defined _LP64 |
||||
cout << "OS: Apple x64 \n" << endl; |
||||
# else |
||||
cout << "OS: Apple x32 \n" << endl; |
||||
# endif |
||||
#endif |
||||
} |
||||
|
||||
void printCudaInfo() |
||||
{ |
||||
#ifndef HAVE_CUDA |
||||
cout << "OpenCV was built without CUDA support \n" << endl; |
||||
#else |
||||
int driver; |
||||
cudaDriverGetVersion(&driver); |
||||
|
||||
cout << "CUDA Driver version: " << driver << '\n'; |
||||
cout << "CUDA Runtime version: " << CUDART_VERSION << '\n'; |
||||
|
||||
cout << endl; |
||||
|
||||
cout << "GPU module was compiled for the following GPU archs:" << endl; |
||||
cout << " BIN: " << CUDA_ARCH_BIN << '\n'; |
||||
cout << " PTX: " << CUDA_ARCH_PTX << '\n'; |
||||
|
||||
cout << endl; |
||||
|
||||
int deviceCount = getCudaEnabledDeviceCount(); |
||||
cout << "CUDA device count: " << deviceCount << '\n'; |
||||
|
||||
cout << endl; |
||||
|
||||
for (int i = 0; i < deviceCount; ++i) |
||||
{ |
||||
DeviceInfo info(i); |
||||
|
||||
cout << "Device [" << i << "] \n"; |
||||
cout << "\t Name: " << info.name() << '\n'; |
||||
cout << "\t Compute capability: " << info.majorVersion() << '.' << info.minorVersion()<< '\n'; |
||||
cout << "\t Multi Processor Count: " << info.multiProcessorCount() << '\n'; |
||||
cout << "\t Total memory: " << static_cast<int>(static_cast<int>(info.totalMemory() / 1024.0) / 1024.0) << " Mb \n"; |
||||
cout << "\t Free memory: " << static_cast<int>(static_cast<int>(info.freeMemory() / 1024.0) / 1024.0) << " Mb \n"; |
||||
if (!info.isCompatible()) |
||||
cout << "\t !!! This device is NOT compatible with current GPU module build \n"; |
||||
|
||||
cout << endl; |
||||
} |
||||
#endif |
||||
} |
||||
|
||||
int main(int argc, char** argv) |
||||
{ |
||||
CommandLineParser cmd(argc, (const char**) argv, |
||||
"{ print_info_only | print_info_only | false | Print information about system and exit }" |
||||
"{ device | device | 0 | Device on which tests will be executed }" |
||||
"{ cpu | cpu | false | Run tests on cpu }" |
||||
); |
||||
|
||||
printOsInfo(); |
||||
printCudaInfo(); |
||||
|
||||
if (cmd.get<bool>("print_info_only")) |
||||
return 0; |
||||
|
||||
int device = cmd.get<int>("device"); |
||||
bool cpu = cmd.get<bool>("cpu"); |
||||
#ifndef HAVE_CUDA |
||||
cpu = true; |
||||
#endif |
||||
|
||||
if (cpu) |
||||
{ |
||||
runOnGpu = false; |
||||
|
||||
cout << "Run tests on CPU \n" << endl; |
||||
} |
||||
else |
||||
{ |
||||
runOnGpu = true; |
||||
|
||||
if (device < 0 || device >= getCudaEnabledDeviceCount()) |
||||
{ |
||||
cerr << "Incorrect device index - " << device << endl; |
||||
return -1; |
||||
} |
||||
|
||||
DeviceInfo info(device); |
||||
if (!info.isCompatible()) |
||||
{ |
||||
cerr << "Device " << device << " [" << info.name() << "] is NOT compatible with current GPU module build" << endl; |
||||
return -1; |
||||
} |
||||
|
||||
setDevice(device); |
||||
|
||||
cout << "Run tests on device " << device << " [" << info.name() << "] \n" << endl; |
||||
} |
||||
|
||||
InitGoogleTest(&argc, argv); |
||||
perf::TestBase::Init(argc, argv); |
||||
return RUN_ALL_TESTS(); |
||||
} |
||||
File diff suppressed because it is too large
Load Diff
@ -1,209 +1,278 @@ |
||||
#include "perf_precomp.hpp" |
||||
|
||||
#ifdef HAVE_CUDA |
||||
using namespace std; |
||||
using namespace testing; |
||||
|
||||
namespace { |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// SURF
|
||||
|
||||
GPU_PERF_TEST_1(SURF, cv::gpu::DeviceInfo) |
||||
{ |
||||
cv::gpu::DeviceInfo devInfo = GetParam(); |
||||
cv::gpu::setDevice(devInfo.deviceID()); |
||||
DEF_PARAM_TEST_1(Image, string); |
||||
|
||||
cv::Mat img_host = readImage("gpu/perf/aloe.jpg", cv::IMREAD_GRAYSCALE); |
||||
ASSERT_FALSE(img_host.empty()); |
||||
PERF_TEST_P(Image, Features2D_SURF, Values<string>("gpu/perf/aloe.jpg")) |
||||
{ |
||||
declare.time(50.0); |
||||
|
||||
cv::gpu::SURF_GPU surf; |
||||
cv::Mat img = readImage(GetParam(), cv::IMREAD_GRAYSCALE); |
||||
ASSERT_FALSE(img.empty()); |
||||
|
||||
cv::gpu::GpuMat img(img_host); |
||||
cv::gpu::GpuMat keypoints, descriptors; |
||||
if (runOnGpu) |
||||
{ |
||||
cv::gpu::SURF_GPU d_surf; |
||||
|
||||
surf(img, cv::gpu::GpuMat(), keypoints, descriptors); |
||||
cv::gpu::GpuMat d_img(img); |
||||
cv::gpu::GpuMat d_keypoints, d_descriptors; |
||||
|
||||
declare.time(2.0); |
||||
d_surf(d_img, cv::gpu::GpuMat(), d_keypoints, d_descriptors); |
||||
|
||||
TEST_CYCLE() |
||||
TEST_CYCLE() |
||||
{ |
||||
d_surf(d_img, cv::gpu::GpuMat(), d_keypoints, d_descriptors); |
||||
} |
||||
} |
||||
else |
||||
{ |
||||
surf(img, cv::gpu::GpuMat(), keypoints, descriptors); |
||||
cv::SURF surf; |
||||
|
||||
std::vector<cv::KeyPoint> keypoints; |
||||
cv::Mat descriptors; |
||||
|
||||
surf(img, cv::noArray(), keypoints, descriptors); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
keypoints.clear(); |
||||
surf(img, cv::noArray(), keypoints, descriptors); |
||||
} |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Features2D, SURF, ALL_DEVICES); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// FAST
|
||||
|
||||
GPU_PERF_TEST_1(FAST, cv::gpu::DeviceInfo) |
||||
PERF_TEST_P(Image, Features2D_FAST, Values<string>("gpu/perf/aloe.jpg")) |
||||
{ |
||||
cv::gpu::DeviceInfo devInfo = GetParam(); |
||||
cv::gpu::setDevice(devInfo.deviceID()); |
||||
cv::Mat img = readImage(GetParam(), cv::IMREAD_GRAYSCALE); |
||||
ASSERT_FALSE(img.empty()); |
||||
|
||||
cv::Mat img_host = readImage("gpu/perf/aloe.jpg", cv::IMREAD_GRAYSCALE); |
||||
ASSERT_FALSE(img_host.empty()); |
||||
|
||||
cv::gpu::FAST_GPU fast(20); |
||||
if (runOnGpu) |
||||
{ |
||||
cv::gpu::FAST_GPU d_fast(20); |
||||
|
||||
cv::gpu::GpuMat img(img_host); |
||||
cv::gpu::GpuMat keypoints; |
||||
cv::gpu::GpuMat d_img(img); |
||||
cv::gpu::GpuMat d_keypoints; |
||||
|
||||
fast(img, cv::gpu::GpuMat(), keypoints); |
||||
d_fast(d_img, cv::gpu::GpuMat(), d_keypoints); |
||||
|
||||
TEST_CYCLE() |
||||
TEST_CYCLE() |
||||
{ |
||||
d_fast(d_img, cv::gpu::GpuMat(), d_keypoints); |
||||
} |
||||
} |
||||
else |
||||
{ |
||||
fast(img, cv::gpu::GpuMat(), keypoints); |
||||
std::vector<cv::KeyPoint> keypoints; |
||||
|
||||
cv::FAST(img, keypoints, 20); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
keypoints.clear(); |
||||
cv::FAST(img, keypoints, 20); |
||||
} |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Features2D, FAST, ALL_DEVICES); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// ORB
|
||||
|
||||
GPU_PERF_TEST_1(ORB, cv::gpu::DeviceInfo) |
||||
PERF_TEST_P(Image, Features2D_ORB, Values<string>("gpu/perf/aloe.jpg")) |
||||
{ |
||||
cv::gpu::DeviceInfo devInfo = GetParam(); |
||||
cv::gpu::setDevice(devInfo.deviceID()); |
||||
cv::Mat img = readImage(GetParam(), cv::IMREAD_GRAYSCALE); |
||||
ASSERT_FALSE(img.empty()); |
||||
|
||||
cv::Mat img_host = readImage("gpu/perf/aloe.jpg", cv::IMREAD_GRAYSCALE); |
||||
ASSERT_FALSE(img_host.empty()); |
||||
if (runOnGpu) |
||||
{ |
||||
cv::gpu::ORB_GPU d_orb(4000); |
||||
|
||||
cv::gpu::ORB_GPU orb(4000); |
||||
cv::gpu::GpuMat d_img(img); |
||||
cv::gpu::GpuMat d_keypoints, d_descriptors; |
||||
|
||||
cv::gpu::GpuMat img(img_host); |
||||
cv::gpu::GpuMat keypoints, descriptors; |
||||
d_orb(d_img, cv::gpu::GpuMat(), d_keypoints, d_descriptors); |
||||
|
||||
TEST_CYCLE() |
||||
TEST_CYCLE() |
||||
{ |
||||
d_orb(d_img, cv::gpu::GpuMat(), d_keypoints, d_descriptors); |
||||
} |
||||
} |
||||
else |
||||
{ |
||||
orb(img, cv::gpu::GpuMat(), keypoints, descriptors); |
||||
cv::ORB orb(4000); |
||||
|
||||
std::vector<cv::KeyPoint> keypoints; |
||||
cv::Mat descriptors; |
||||
|
||||
orb(img, cv::noArray(), keypoints, descriptors); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
keypoints.clear(); |
||||
orb(img, cv::noArray(), keypoints, descriptors); |
||||
} |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Features2D, ORB, ALL_DEVICES); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// BruteForceMatcher_match
|
||||
// BFMatch
|
||||
|
||||
IMPLEMENT_PARAM_CLASS(DescriptorSize, int) |
||||
DEF_PARAM_TEST(DescSize_Norm, int, NormType); |
||||
|
||||
GPU_PERF_TEST(BruteForceMatcher_match, cv::gpu::DeviceInfo, DescriptorSize, NormType) |
||||
PERF_TEST_P(DescSize_Norm, Features2D_BFMatch, Combine(Values(64, 128, 256), Values(NormType(cv::NORM_L1), NormType(cv::NORM_L2), NormType(cv::NORM_HAMMING)))) |
||||
{ |
||||
cv::gpu::DeviceInfo devInfo = GET_PARAM(0); |
||||
cv::gpu::setDevice(devInfo.deviceID()); |
||||
declare.time(20.0); |
||||
|
||||
int desc_size = GET_PARAM(1); |
||||
int normType = GET_PARAM(2); |
||||
int desc_size = GET_PARAM(0); |
||||
int normType = GET_PARAM(1); |
||||
|
||||
int type = normType == cv::NORM_HAMMING ? CV_8U : CV_32F; |
||||
|
||||
cv::Mat query_host(3000, desc_size, type); |
||||
fill(query_host, 0.0, 10.0); |
||||
|
||||
cv::Mat train_host(3000, desc_size, type); |
||||
fill(train_host, 0.0, 10.0); |
||||
cv::Mat query(3000, desc_size, type); |
||||
fillRandom(query); |
||||
|
||||
cv::gpu::BFMatcher_GPU matcher(normType); |
||||
cv::Mat train(3000, desc_size, type); |
||||
fillRandom(train); |
||||
|
||||
cv::gpu::GpuMat query(query_host); |
||||
cv::gpu::GpuMat train(train_host); |
||||
cv::gpu::GpuMat trainIdx, distance; |
||||
if (runOnGpu) |
||||
{ |
||||
cv::gpu::BFMatcher_GPU d_matcher(normType); |
||||
|
||||
matcher.matchSingle(query, train, trainIdx, distance); |
||||
cv::gpu::GpuMat d_query(query); |
||||
cv::gpu::GpuMat d_train(train); |
||||
cv::gpu::GpuMat d_trainIdx, d_distance; |
||||
|
||||
declare.time(3.0); |
||||
d_matcher.matchSingle(d_query, d_train, d_trainIdx, d_distance); |
||||
|
||||
TEST_CYCLE() |
||||
TEST_CYCLE() |
||||
{ |
||||
d_matcher.matchSingle(d_query, d_train, d_trainIdx, d_distance); |
||||
} |
||||
} |
||||
else |
||||
{ |
||||
matcher.matchSingle(query, train, trainIdx, distance); |
||||
cv::BFMatcher matcher(normType); |
||||
|
||||
std::vector<cv::DMatch> matches; |
||||
|
||||
matcher.match(query, train, matches); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
matcher.match(query, train, matches); |
||||
} |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Features2D, BruteForceMatcher_match, testing::Combine( |
||||
ALL_DEVICES, |
||||
testing::Values(DescriptorSize(64), DescriptorSize(128), DescriptorSize(256)), |
||||
testing::Values(NormType(cv::NORM_L1), NormType(cv::NORM_L2), NormType(cv::NORM_HAMMING)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// BruteForceMatcher_knnMatch
|
||||
// BFKnnMatch
|
||||
|
||||
IMPLEMENT_PARAM_CLASS(K, int) |
||||
DEF_PARAM_TEST(DescSize_K_Norm, int, int, NormType); |
||||
|
||||
GPU_PERF_TEST(BruteForceMatcher_knnMatch, cv::gpu::DeviceInfo, DescriptorSize, K, NormType) |
||||
PERF_TEST_P(DescSize_K_Norm, Features2D_BFKnnMatch, Combine( |
||||
Values(64, 128, 256), |
||||
Values(2, 3), |
||||
Values(NormType(cv::NORM_L1), NormType(cv::NORM_L2), NormType(cv::NORM_HAMMING)))) |
||||
{ |
||||
cv::gpu::DeviceInfo devInfo = GET_PARAM(0); |
||||
cv::gpu::setDevice(devInfo.deviceID()); |
||||
declare.time(30.0); |
||||
|
||||
int desc_size = GET_PARAM(1); |
||||
int k = GET_PARAM(2); |
||||
int normType = GET_PARAM(3); |
||||
int desc_size = GET_PARAM(0); |
||||
int k = GET_PARAM(1); |
||||
int normType = GET_PARAM(2); |
||||
|
||||
int type = normType == cv::NORM_HAMMING ? CV_8U : CV_32F; |
||||
|
||||
cv::Mat query_host(3000, desc_size, type); |
||||
fill(query_host, 0.0, 10.0); |
||||
|
||||
cv::Mat train_host(3000, desc_size, type); |
||||
fill(train_host, 0.0, 10.0); |
||||
cv::Mat query(3000, desc_size, type); |
||||
fillRandom(query); |
||||
|
||||
cv::gpu::BFMatcher_GPU matcher(normType); |
||||
cv::Mat train(3000, desc_size, type); |
||||
fillRandom(train); |
||||
|
||||
cv::gpu::GpuMat query(query_host); |
||||
cv::gpu::GpuMat train(train_host); |
||||
cv::gpu::GpuMat trainIdx, distance, allDist; |
||||
if (runOnGpu) |
||||
{ |
||||
cv::gpu::BFMatcher_GPU d_matcher(normType); |
||||
|
||||
matcher.knnMatchSingle(query, train, trainIdx, distance, allDist, k); |
||||
cv::gpu::GpuMat d_query(query); |
||||
cv::gpu::GpuMat d_train(train); |
||||
cv::gpu::GpuMat d_trainIdx, d_distance, d_allDist; |
||||
|
||||
declare.time(3.0); |
||||
d_matcher.knnMatchSingle(d_query, d_train, d_trainIdx, d_distance, d_allDist, k); |
||||
|
||||
TEST_CYCLE() |
||||
TEST_CYCLE() |
||||
{ |
||||
d_matcher.knnMatchSingle(d_query, d_train, d_trainIdx, d_distance, d_allDist, k); |
||||
} |
||||
} |
||||
else |
||||
{ |
||||
matcher.knnMatchSingle(query, train, trainIdx, distance, allDist, k); |
||||
cv::BFMatcher matcher(normType); |
||||
|
||||
std::vector< std::vector<cv::DMatch> > matches; |
||||
|
||||
matcher.knnMatch(query, train, matches, k); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
matcher.knnMatch(query, train, matches, k); |
||||
} |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Features2D, BruteForceMatcher_knnMatch, testing::Combine( |
||||
ALL_DEVICES, |
||||
testing::Values(DescriptorSize(64), DescriptorSize(128), DescriptorSize(256)), |
||||
testing::Values(K(2), K(3)), |
||||
testing::Values(NormType(cv::NORM_L1), NormType(cv::NORM_L2), NormType(cv::NORM_HAMMING)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// BruteForceMatcher_radiusMatch
|
||||
// BFRadiusMatch
|
||||
|
||||
GPU_PERF_TEST(BruteForceMatcher_radiusMatch, cv::gpu::DeviceInfo, DescriptorSize, NormType) |
||||
PERF_TEST_P(DescSize_Norm, Features2D_BFRadiusMatch, Combine(Values(64, 128, 256), Values(NormType(cv::NORM_L1), NormType(cv::NORM_L2), NormType(cv::NORM_HAMMING)))) |
||||
{ |
||||
cv::gpu::DeviceInfo devInfo = GET_PARAM(0); |
||||
cv::gpu::setDevice(devInfo.deviceID()); |
||||
declare.time(30.0); |
||||
|
||||
int desc_size = GET_PARAM(1); |
||||
int normType = GET_PARAM(2); |
||||
int desc_size = GET_PARAM(0); |
||||
int normType = GET_PARAM(1); |
||||
|
||||
int type = normType == cv::NORM_HAMMING ? CV_8U : CV_32F; |
||||
|
||||
cv::Mat query_host(3000, desc_size, type); |
||||
fill(query_host, 0.0, 1.0); |
||||
|
||||
cv::Mat train_host(3000, desc_size, type); |
||||
fill(train_host, 0.0, 1.0); |
||||
cv::Mat query(3000, desc_size, type); |
||||
fillRandom(query, 0.0, 1.0); |
||||
|
||||
cv::gpu::BFMatcher_GPU matcher(normType); |
||||
cv::Mat train(3000, desc_size, type); |
||||
fillRandom(train, 0.0, 1.0); |
||||
|
||||
cv::gpu::GpuMat query(query_host); |
||||
cv::gpu::GpuMat train(train_host); |
||||
cv::gpu::GpuMat trainIdx, nMatches, distance; |
||||
if (runOnGpu) |
||||
{ |
||||
cv::gpu::BFMatcher_GPU d_matcher(normType); |
||||
|
||||
matcher.radiusMatchSingle(query, train, trainIdx, distance, nMatches, 2.0); |
||||
cv::gpu::GpuMat d_query(query); |
||||
cv::gpu::GpuMat d_train(train); |
||||
cv::gpu::GpuMat d_trainIdx, d_nMatches, d_distance; |
||||
|
||||
declare.time(3.0); |
||||
d_matcher.radiusMatchSingle(d_query, d_train, d_trainIdx, d_distance, d_nMatches, 2.0); |
||||
|
||||
TEST_CYCLE() |
||||
TEST_CYCLE() |
||||
{ |
||||
d_matcher.radiusMatchSingle(d_query, d_train, d_trainIdx, d_distance, d_nMatches, 2.0); |
||||
} |
||||
} |
||||
else |
||||
{ |
||||
matcher.radiusMatchSingle(query, train, trainIdx, distance, nMatches, 2.0); |
||||
cv::BFMatcher matcher(normType); |
||||
|
||||
std::vector< std::vector<cv::DMatch> > matches; |
||||
|
||||
matcher.radiusMatch(query, train, matches, 2.0); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
matcher.radiusMatch(query, train, matches, 2.0); |
||||
} |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Features2D, BruteForceMatcher_radiusMatch, testing::Combine( |
||||
ALL_DEVICES, |
||||
testing::Values(DescriptorSize(64), DescriptorSize(128), DescriptorSize(256)), |
||||
testing::Values(NormType(cv::NORM_L1), NormType(cv::NORM_L2), NormType(cv::NORM_HAMMING)))); |
||||
|
||||
#endif |
||||
} // namespace
|
||||
|
||||
@ -1,308 +1,379 @@ |
||||
#include "perf_precomp.hpp" |
||||
|
||||
#ifdef HAVE_CUDA |
||||
using namespace std; |
||||
using namespace testing; |
||||
|
||||
namespace { |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// Blur
|
||||
|
||||
IMPLEMENT_PARAM_CLASS(KernelSize, int) |
||||
DEF_PARAM_TEST(Sz_Type_KernelSz, cv::Size, MatType, int); |
||||
|
||||
GPU_PERF_TEST(Blur, cv::gpu::DeviceInfo, cv::Size, MatType, KernelSize) |
||||
PERF_TEST_P(Sz_Type_KernelSz, Filters_Blur, Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8UC1, CV_8UC4), Values(3, 5, 7))) |
||||
{ |
||||
cv::gpu::DeviceInfo devInfo = GET_PARAM(0); |
||||
cv::gpu::setDevice(devInfo.deviceID()); |
||||
declare.time(20.0); |
||||
|
||||
cv::Size size = GET_PARAM(1); |
||||
int type = GET_PARAM(2); |
||||
int ksize = GET_PARAM(3); |
||||
cv::Size size = GET_PARAM(0); |
||||
int type = GET_PARAM(1); |
||||
int ksize = GET_PARAM(2); |
||||
|
||||
cv::Mat src_host(size, type); |
||||
fill(src_host, 0.0, 255.0); |
||||
cv::Mat src(size, type); |
||||
fillRandom(src); |
||||
|
||||
cv::gpu::GpuMat src(src_host); |
||||
cv::gpu::GpuMat dst; |
||||
if (runOnGpu) |
||||
{ |
||||
cv::gpu::GpuMat d_src(src); |
||||
cv::gpu::GpuMat d_dst; |
||||
|
||||
cv::gpu::blur(src, dst, cv::Size(ksize, ksize)); |
||||
cv::gpu::blur(d_src, d_dst, cv::Size(ksize, ksize)); |
||||
|
||||
TEST_CYCLE() |
||||
TEST_CYCLE() |
||||
{ |
||||
cv::gpu::blur(d_src, d_dst, cv::Size(ksize, ksize)); |
||||
} |
||||
} |
||||
else |
||||
{ |
||||
cv::gpu::blur(src, dst, cv::Size(ksize, ksize)); |
||||
cv::Mat dst; |
||||
|
||||
cv::blur(src, dst, cv::Size(ksize, ksize)); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::blur(src, dst, cv::Size(ksize, ksize)); |
||||
} |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Filters, Blur, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(MatType(CV_8UC1), MatType(CV_8UC4)), |
||||
testing::Values(KernelSize(3), KernelSize(5), KernelSize(7)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// Sobel
|
||||
|
||||
GPU_PERF_TEST(Sobel, cv::gpu::DeviceInfo, cv::Size, MatType, KernelSize) |
||||
PERF_TEST_P(Sz_Type_KernelSz, Filters_Sobel, Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8UC1, CV_8UC4, CV_32FC1), Values(3, 5, 7, 9, 11, 13, 15))) |
||||
{ |
||||
cv::gpu::DeviceInfo devInfo = GET_PARAM(0); |
||||
cv::gpu::setDevice(devInfo.deviceID()); |
||||
declare.time(20.0); |
||||
|
||||
cv::Size size = GET_PARAM(1); |
||||
int type = GET_PARAM(2); |
||||
int ksize = GET_PARAM(3); |
||||
cv::Size size = GET_PARAM(0); |
||||
int type = GET_PARAM(1); |
||||
int ksize = GET_PARAM(2); |
||||
|
||||
cv::Mat src_host(size, type); |
||||
fill(src_host, 0.0, 255.0); |
||||
cv::Mat src(size, type); |
||||
fillRandom(src); |
||||
|
||||
cv::gpu::GpuMat src(src_host); |
||||
cv::gpu::GpuMat dst; |
||||
cv::gpu::GpuMat buf; |
||||
if (runOnGpu) |
||||
{ |
||||
cv::gpu::GpuMat d_src(src); |
||||
cv::gpu::GpuMat d_dst; |
||||
cv::gpu::GpuMat d_buf; |
||||
|
||||
cv::gpu::Sobel(src, dst, -1, 1, 1, buf, ksize); |
||||
cv::gpu::Sobel(d_src, d_dst, -1, 1, 1, d_buf, ksize); |
||||
|
||||
TEST_CYCLE() |
||||
TEST_CYCLE() |
||||
{ |
||||
cv::gpu::Sobel(d_src, d_dst, -1, 1, 1, d_buf, ksize); |
||||
} |
||||
} |
||||
else |
||||
{ |
||||
cv::gpu::Sobel(src, dst, -1, 1, 1, buf, ksize); |
||||
cv::Mat dst; |
||||
|
||||
cv::Sobel(src, dst, -1, 1, 1, ksize); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::Sobel(src, dst, -1, 1, 1, ksize); |
||||
} |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Filters, Sobel, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(MatType(CV_8UC1), MatType(CV_8UC4), MatType(CV_32FC1)), |
||||
testing::Values(KernelSize(3), KernelSize(5), KernelSize(7), KernelSize(9), KernelSize(11), KernelSize(13), KernelSize(15)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// Scharr
|
||||
|
||||
GPU_PERF_TEST(Scharr, cv::gpu::DeviceInfo, cv::Size, MatType) |
||||
PERF_TEST_P(Sz_Type, Filters_Scharr, Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8UC1, CV_8UC4, CV_32FC1))) |
||||
{ |
||||
cv::gpu::DeviceInfo devInfo = GET_PARAM(0); |
||||
cv::gpu::setDevice(devInfo.deviceID()); |
||||
declare.time(20.0); |
||||
|
||||
cv::Size size = GET_PARAM(1); |
||||
int type = GET_PARAM(2); |
||||
cv::Size size = GET_PARAM(0); |
||||
int type = GET_PARAM(1); |
||||
|
||||
cv::Mat src_host(size, type); |
||||
fill(src_host, 0.0, 255.0); |
||||
cv::Mat src(size, type); |
||||
fillRandom(src); |
||||
|
||||
cv::gpu::GpuMat src(src_host); |
||||
cv::gpu::GpuMat dst; |
||||
cv::gpu::GpuMat buf; |
||||
if (runOnGpu) |
||||
{ |
||||
cv::gpu::GpuMat d_src(src); |
||||
cv::gpu::GpuMat d_dst; |
||||
cv::gpu::GpuMat d_buf; |
||||
|
||||
cv::gpu::Scharr(src, dst, -1, 1, 0, buf); |
||||
cv::gpu::Scharr(d_src, d_dst, -1, 1, 0, d_buf); |
||||
|
||||
TEST_CYCLE() |
||||
TEST_CYCLE() |
||||
{ |
||||
cv::gpu::Scharr(d_src, d_dst, -1, 1, 0, d_buf); |
||||
} |
||||
} |
||||
else |
||||
{ |
||||
cv::gpu::Scharr(src, dst, -1, 1, 0, buf); |
||||
cv::Mat dst; |
||||
|
||||
cv::Scharr(src, dst, -1, 1, 0); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::Scharr(src, dst, -1, 1, 0); |
||||
} |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Filters, Scharr, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(MatType(CV_8UC1), MatType(CV_8UC4), MatType(CV_32FC1)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// GaussianBlur
|
||||
|
||||
GPU_PERF_TEST(GaussianBlur, cv::gpu::DeviceInfo, cv::Size, MatType, KernelSize) |
||||
PERF_TEST_P(Sz_Type_KernelSz, Filters_GaussianBlur, Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8UC1, CV_8UC4, CV_32FC1), Values(3, 5, 7, 9, 11, 13, 15))) |
||||
{ |
||||
cv::gpu::DeviceInfo devInfo = GET_PARAM(0); |
||||
cv::gpu::setDevice(devInfo.deviceID()); |
||||
declare.time(20.0); |
||||
|
||||
cv::Size size = GET_PARAM(1); |
||||
int type = GET_PARAM(2); |
||||
int ksize = GET_PARAM(3); |
||||
cv::Size size = GET_PARAM(0); |
||||
int type = GET_PARAM(1); |
||||
int ksize = GET_PARAM(2); |
||||
|
||||
cv::Mat src_host(size, type); |
||||
fill(src_host, 0.0, 255.0); |
||||
cv::Mat src(size, type); |
||||
fillRandom(src); |
||||
|
||||
cv::gpu::GpuMat src(src_host); |
||||
cv::gpu::GpuMat dst; |
||||
cv::gpu::GpuMat buf; |
||||
if (runOnGpu) |
||||
{ |
||||
cv::gpu::GpuMat d_src(src); |
||||
cv::gpu::GpuMat d_dst; |
||||
cv::gpu::GpuMat d_buf; |
||||
|
||||
cv::gpu::GaussianBlur(src, dst, cv::Size(ksize, ksize), buf, 0.5); |
||||
cv::gpu::GaussianBlur(d_src, d_dst, cv::Size(ksize, ksize), d_buf, 0.5); |
||||
|
||||
TEST_CYCLE() |
||||
TEST_CYCLE() |
||||
{ |
||||
cv::gpu::GaussianBlur(d_src, d_dst, cv::Size(ksize, ksize), d_buf, 0.5); |
||||
} |
||||
} |
||||
else |
||||
{ |
||||
cv::gpu::GaussianBlur(src, dst, cv::Size(ksize, ksize), buf, 0.5); |
||||
cv::Mat dst; |
||||
|
||||
cv::GaussianBlur(src, dst, cv::Size(ksize, ksize), 0.5); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::GaussianBlur(src, dst, cv::Size(ksize, ksize), 0.5); |
||||
} |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Filters, GaussianBlur, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(MatType(CV_8UC1), MatType(CV_8UC4), MatType(CV_32FC1)), |
||||
testing::Values(KernelSize(3), KernelSize(5), KernelSize(7), KernelSize(9), KernelSize(11), KernelSize(13), KernelSize(15)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// Laplacian
|
||||
|
||||
GPU_PERF_TEST(Laplacian, cv::gpu::DeviceInfo, cv::Size, MatType, KernelSize) |
||||
PERF_TEST_P(Sz_Type_KernelSz, Filters_Laplacian, Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4), Values(1, 3))) |
||||
{ |
||||
cv::gpu::DeviceInfo devInfo = GET_PARAM(0); |
||||
cv::gpu::setDevice(devInfo.deviceID()); |
||||
declare.time(20.0); |
||||
|
||||
cv::Size size = GET_PARAM(1); |
||||
int type = GET_PARAM(2); |
||||
int ksize = GET_PARAM(3); |
||||
cv::Size size = GET_PARAM(0); |
||||
int type = GET_PARAM(1); |
||||
int ksize = GET_PARAM(2); |
||||
|
||||
cv::Mat src_host(size, type); |
||||
fill(src_host, 0.0, 255.0); |
||||
cv::Mat src(size, type); |
||||
fillRandom(src); |
||||
|
||||
cv::gpu::GpuMat src(src_host); |
||||
cv::gpu::GpuMat dst; |
||||
if (runOnGpu) |
||||
{ |
||||
cv::gpu::GpuMat d_src(src); |
||||
cv::gpu::GpuMat d_dst; |
||||
|
||||
cv::gpu::Laplacian(src, dst, -1, ksize); |
||||
cv::gpu::Laplacian(d_src, d_dst, -1, ksize); |
||||
|
||||
TEST_CYCLE() |
||||
TEST_CYCLE() |
||||
{ |
||||
cv::gpu::Laplacian(d_src, d_dst, -1, ksize); |
||||
} |
||||
} |
||||
else |
||||
{ |
||||
cv::gpu::Laplacian(src, dst, -1, ksize); |
||||
cv::Mat dst; |
||||
|
||||
cv::Laplacian(src, dst, -1, ksize); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::Laplacian(src, dst, -1, ksize); |
||||
} |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Filters, Laplacian, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(MatType(CV_8UC1), MatType(CV_8UC4), MatType(CV_32FC1), MatType(CV_32FC4)), |
||||
testing::Values(KernelSize(1), KernelSize(3)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// Erode
|
||||
|
||||
GPU_PERF_TEST(Erode, cv::gpu::DeviceInfo, cv::Size, MatType) |
||||
PERF_TEST_P(Sz_Type, Filters_Erode, Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8UC1, CV_8UC4))) |
||||
{ |
||||
cv::gpu::DeviceInfo devInfo = GET_PARAM(0); |
||||
cv::gpu::setDevice(devInfo.deviceID()); |
||||
declare.time(20.0); |
||||
|
||||
cv::Size size = GET_PARAM(1); |
||||
int type = GET_PARAM(2); |
||||
cv::Size size = GET_PARAM(0); |
||||
int type = GET_PARAM(1); |
||||
|
||||
cv::Mat src_host(size, type); |
||||
fill(src_host, 0.0, 255.0); |
||||
cv::Mat src(size, type); |
||||
fillRandom(src); |
||||
|
||||
cv::Mat ker = cv::getStructuringElement(cv::MORPH_RECT, cv::Size(3, 3)); |
||||
|
||||
cv::gpu::GpuMat src(src_host); |
||||
cv::gpu::GpuMat dst; |
||||
cv::gpu::GpuMat buf; |
||||
if (runOnGpu) |
||||
{ |
||||
cv::gpu::GpuMat d_src(src); |
||||
cv::gpu::GpuMat d_dst; |
||||
cv::gpu::GpuMat d_buf; |
||||
|
||||
cv::gpu::erode(src, dst, ker, buf); |
||||
cv::gpu::erode(d_src, d_dst, ker, d_buf); |
||||
|
||||
TEST_CYCLE() |
||||
TEST_CYCLE() |
||||
{ |
||||
cv::gpu::erode(d_src, d_dst, ker, d_buf); |
||||
} |
||||
} |
||||
else |
||||
{ |
||||
cv::gpu::erode(src, dst, ker, buf); |
||||
cv::Mat dst; |
||||
|
||||
cv::erode(src, dst, ker); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::erode(src, dst, ker); |
||||
} |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Filters, Erode, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(MatType(CV_8UC1), MatType(CV_8UC4)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// Dilate
|
||||
|
||||
GPU_PERF_TEST(Dilate, cv::gpu::DeviceInfo, cv::Size, MatType) |
||||
PERF_TEST_P(Sz_Type, Filters_Dilate, Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8UC1, CV_8UC4))) |
||||
{ |
||||
cv::gpu::DeviceInfo devInfo = GET_PARAM(0); |
||||
cv::gpu::setDevice(devInfo.deviceID()); |
||||
declare.time(20.0); |
||||
|
||||
cv::Size size = GET_PARAM(1); |
||||
int type = GET_PARAM(2); |
||||
cv::Size size = GET_PARAM(0); |
||||
int type = GET_PARAM(1); |
||||
|
||||
cv::Mat src_host(size, type); |
||||
fill(src_host, 0.0, 255.0); |
||||
cv::Mat src(size, type); |
||||
fillRandom(src); |
||||
|
||||
cv::Mat ker = cv::getStructuringElement(cv::MORPH_RECT, cv::Size(3, 3)); |
||||
|
||||
cv::gpu::GpuMat src(src_host); |
||||
cv::gpu::GpuMat dst; |
||||
cv::gpu::GpuMat buf; |
||||
if (runOnGpu) |
||||
{ |
||||
cv::gpu::GpuMat d_src(src); |
||||
cv::gpu::GpuMat d_dst; |
||||
cv::gpu::GpuMat d_buf; |
||||
|
||||
cv::gpu::dilate(src, dst, ker, buf); |
||||
cv::gpu::dilate(d_src, d_dst, ker, d_buf); |
||||
|
||||
TEST_CYCLE() |
||||
TEST_CYCLE() |
||||
{ |
||||
cv::gpu::dilate(d_src, d_dst, ker, d_buf); |
||||
} |
||||
} |
||||
else |
||||
{ |
||||
cv::gpu::dilate(src, dst, ker, buf); |
||||
cv::Mat dst; |
||||
|
||||
cv::dilate(src, dst, ker); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::dilate(src, dst, ker); |
||||
} |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Filters, Dilate, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(MatType(CV_8UC1), MatType(CV_8UC4)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// MorphologyEx
|
||||
|
||||
CV_ENUM(MorphOp, cv::MORPH_OPEN, cv::MORPH_CLOSE, cv::MORPH_GRADIENT, cv::MORPH_TOPHAT, cv::MORPH_BLACKHAT) |
||||
#define ALL_MORPH_OPS testing::Values(MorphOp(cv::MORPH_OPEN), MorphOp(cv::MORPH_CLOSE), MorphOp(cv::MORPH_GRADIENT), MorphOp(cv::MORPH_TOPHAT), MorphOp(cv::MORPH_BLACKHAT)) |
||||
#define ALL_MORPH_OPS ValuesIn(MorphOp::all()) |
||||
|
||||
GPU_PERF_TEST(MorphologyEx, cv::gpu::DeviceInfo, cv::Size, MatType, MorphOp) |
||||
DEF_PARAM_TEST(Sz_Type_Op, cv::Size, MatType, MorphOp); |
||||
|
||||
PERF_TEST_P(Sz_Type_Op, Filters_MorphologyEx, Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8UC1, CV_8UC4), ALL_MORPH_OPS)) |
||||
{ |
||||
cv::gpu::DeviceInfo devInfo = GET_PARAM(0); |
||||
cv::gpu::setDevice(devInfo.deviceID()); |
||||
declare.time(20.0); |
||||
|
||||
cv::Size size = GET_PARAM(1); |
||||
int type = GET_PARAM(2); |
||||
int morphOp = GET_PARAM(3); |
||||
cv::Size size = GET_PARAM(0); |
||||
int type = GET_PARAM(1); |
||||
int morphOp = GET_PARAM(2); |
||||
|
||||
cv::Mat src_host(size, type); |
||||
fill(src_host, 0.0, 255.0); |
||||
cv::Mat src(size, type); |
||||
fillRandom(src); |
||||
|
||||
cv::Mat ker = cv::getStructuringElement(cv::MORPH_RECT, cv::Size(3, 3)); |
||||
|
||||
cv::gpu::GpuMat src(src_host); |
||||
cv::gpu::GpuMat dst; |
||||
cv::gpu::GpuMat buf1; |
||||
cv::gpu::GpuMat buf2; |
||||
if (runOnGpu) |
||||
{ |
||||
cv::gpu::GpuMat d_src(src); |
||||
cv::gpu::GpuMat d_dst; |
||||
cv::gpu::GpuMat d_buf1; |
||||
cv::gpu::GpuMat d_buf2; |
||||
|
||||
cv::gpu::morphologyEx(src, dst, morphOp, ker, buf1, buf2); |
||||
cv::gpu::morphologyEx(d_src, d_dst, morphOp, ker, d_buf1, d_buf2); |
||||
|
||||
TEST_CYCLE() |
||||
TEST_CYCLE() |
||||
{ |
||||
cv::gpu::morphologyEx(d_src, d_dst, morphOp, ker, d_buf1, d_buf2); |
||||
} |
||||
} |
||||
else |
||||
{ |
||||
cv::gpu::morphologyEx(src, dst, morphOp, ker, buf1, buf2); |
||||
cv::Mat dst; |
||||
|
||||
cv::morphologyEx(src, dst, morphOp, ker); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::morphologyEx(src, dst, morphOp, ker); |
||||
} |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Filters, MorphologyEx, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(MatType(CV_8UC1), MatType(CV_8UC4)), |
||||
ALL_MORPH_OPS)); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// Filter2D
|
||||
|
||||
GPU_PERF_TEST(Filter2D, cv::gpu::DeviceInfo, cv::Size, MatType, KernelSize) |
||||
PERF_TEST_P(Sz_Type_KernelSz, Filters_Filter2D, Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4), Values(3, 5, 7, 9, 11, 13, 15))) |
||||
{ |
||||
cv::gpu::DeviceInfo devInfo = GET_PARAM(0); |
||||
cv::gpu::setDevice(devInfo.deviceID()); |
||||
declare.time(20.0); |
||||
|
||||
cv::Size size = GET_PARAM(1); |
||||
int type = GET_PARAM(2); |
||||
int ksize = GET_PARAM(3); |
||||
cv::Size size = GET_PARAM(0); |
||||
int type = GET_PARAM(1); |
||||
int ksize = GET_PARAM(2); |
||||
|
||||
cv::Mat src_host(size, type); |
||||
fill(src_host, 0.0, 255.0); |
||||
cv::Mat src(size, type); |
||||
fillRandom(src); |
||||
|
||||
cv::Mat kernel(ksize, ksize, CV_32FC1); |
||||
fill(kernel, 0.0, 1.0); |
||||
fillRandom(kernel, 0.0, 1.0); |
||||
|
||||
cv::gpu::GpuMat src(src_host); |
||||
cv::gpu::GpuMat dst; |
||||
if (runOnGpu) |
||||
{ |
||||
cv::gpu::GpuMat d_src(src); |
||||
cv::gpu::GpuMat d_dst; |
||||
|
||||
cv::gpu::filter2D(src, dst, -1, kernel); |
||||
cv::gpu::filter2D(d_src, d_dst, -1, kernel); |
||||
|
||||
TEST_CYCLE() |
||||
TEST_CYCLE() |
||||
{ |
||||
cv::gpu::filter2D(d_src, d_dst, -1, kernel); |
||||
} |
||||
} |
||||
else |
||||
{ |
||||
cv::gpu::filter2D(src, dst, -1, kernel); |
||||
cv::Mat dst; |
||||
|
||||
cv::filter2D(src, dst, -1, kernel); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::filter2D(src, dst, -1, kernel); |
||||
} |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Filters, Filter2D, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(MatType(CV_8UC1), MatType(CV_8UC4), MatType(CV_32FC1), MatType(CV_32FC4)), |
||||
testing::Values(KernelSize(3), KernelSize(5), KernelSize(7), KernelSize(9), KernelSize(11), KernelSize(13), KernelSize(15)))); |
||||
|
||||
#endif |
||||
} // namespace
|
||||
|
||||
File diff suppressed because it is too large
Load Diff
@ -1,75 +1,141 @@ |
||||
/*M///////////////////////////////////////////////////////////////////////////////////////
|
||||
//
|
||||
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
|
||||
//
|
||||
// By downloading, copying, installing or using the software you agree to this license.
|
||||
// If you do not agree to this license, do not download, install,
|
||||
// copy or use the software.
|
||||
//
|
||||
//
|
||||
// License Agreement
|
||||
// For Open Source Computer Vision Library
|
||||
//
|
||||
// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
|
||||
// Copyright (C) 2008-2011, Willow Garage Inc., all rights reserved.
|
||||
// Third party copyrights are property of their respective owners.
|
||||
//
|
||||
// Redistribution and use in source and binary forms, with or without modification,
|
||||
// are permitted provided that the following conditions are met:
|
||||
//
|
||||
// * Redistributions of source code must retain the above copyright notice,
|
||||
// this list of conditions and the following disclaimer.
|
||||
//
|
||||
// * Redistributions in binary form must reproduce the above copyright notice,
|
||||
// this list of conditions and the following disclaimer in the documentation
|
||||
// and/or other materials provided with the distribution.
|
||||
//
|
||||
// * The name of the copyright holders may not be used to endorse or promote products
|
||||
// derived from this software without specific prior written permission.
|
||||
//
|
||||
// This software is provided by the copyright holders and contributors "as is" and
|
||||
// any express or implied warranties, including, but not limited to, the implied
|
||||
// warranties of merchantability and fitness for a particular purpose are disclaimed.
|
||||
// In no event shall the Intel Corporation or contributors be liable for any direct,
|
||||
// indirect, incidental, special, exemplary, or consequential damages
|
||||
// (including, but not limited to, procurement of substitute goods or services;
|
||||
// loss of use, data, or profits; or business interruption) however caused
|
||||
// and on any theory of liability, whether in contract, strict liability,
|
||||
// or tort (including negligence or otherwise) arising in any way out of
|
||||
// the use of this software, even if advised of the possibility of such damage.
|
||||
//M*/
|
||||
|
||||
#include "perf_precomp.hpp" |
||||
|
||||
#ifdef HAVE_CUDA |
||||
using namespace std; |
||||
using namespace testing; |
||||
|
||||
namespace { |
||||
|
||||
DEF_PARAM_TEST_1(Image, string); |
||||
|
||||
GPU_PERF_TEST(ConnectedComponents, cv::gpu::DeviceInfo, cv::Size) |
||||
struct GreedyLabeling |
||||
{ |
||||
cv::gpu::DeviceInfo devInfo = GET_PARAM(0); |
||||
cv::gpu::setDevice(devInfo.deviceID()); |
||||
struct dot |
||||
{ |
||||
int x; |
||||
int y; |
||||
|
||||
static dot make(int i, int j) |
||||
{ |
||||
dot d; d.x = i; d.y = j; |
||||
return d; |
||||
} |
||||
}; |
||||
|
||||
struct InInterval |
||||
{ |
||||
InInterval(const int& _lo, const int& _hi) : lo(-_lo), hi(_hi) {}; |
||||
const int lo, hi; |
||||
|
||||
bool operator() (const unsigned char a, const unsigned char b) const |
||||
{ |
||||
int d = a - b; |
||||
return lo <= d && d <= hi; |
||||
} |
||||
|
||||
private: |
||||
InInterval& operator=(const InInterval&); |
||||
|
||||
|
||||
}; |
||||
|
||||
GreedyLabeling(cv::Mat img) |
||||
: image(img), _labels(image.size(), CV_32SC1, cv::Scalar::all(-1)) {stack = new dot[image.cols * image.rows];} |
||||
|
||||
~GreedyLabeling(){delete[] stack;} |
||||
|
||||
void operator() (cv::Mat labels) const |
||||
{ |
||||
labels.setTo(cv::Scalar::all(-1)); |
||||
InInterval inInt(0, 2); |
||||
int cc = -1; |
||||
|
||||
int* dist_labels = (int*)labels.data; |
||||
int pitch = static_cast<int>(labels.step1()); |
||||
|
||||
unsigned char* source = (unsigned char*)image.data; |
||||
int width = image.cols; |
||||
int height = image.rows; |
||||
|
||||
for (int j = 0; j < image.rows; ++j) |
||||
for (int i = 0; i < image.cols; ++i) |
||||
{ |
||||
if (dist_labels[j * pitch + i] != -1) continue; |
||||
|
||||
cv::Mat image = readImage("gpu/labeling/aloe-disp.png", cv::IMREAD_GRAYSCALE); |
||||
dot* top = stack; |
||||
dot p = dot::make(i, j); |
||||
cc++; |
||||
|
||||
// cv::threshold(image, image, 150, 255, CV_THRESH_BINARY);
|
||||
dist_labels[j * pitch + i] = cc; |
||||
|
||||
cv::gpu::GpuMat mask; |
||||
mask.create(image.rows, image.cols, CV_8UC1); |
||||
while (top >= stack) |
||||
{ |
||||
int* dl = &dist_labels[p.y * pitch + p.x]; |
||||
unsigned char* sp = &source[p.y * image.step1() + p.x]; |
||||
|
||||
cv::gpu::GpuMat components; |
||||
components.create(image.rows, image.cols, CV_32SC1); |
||||
dl[0] = cc; |
||||
|
||||
cv::gpu::connectivityMask(cv::gpu::GpuMat(image), mask, cv::Scalar::all(0), cv::Scalar::all(2)); |
||||
//right
|
||||
if( p.x < (width - 1) && dl[ +1] == -1 && inInt(sp[0], sp[+1])) |
||||
*top++ = dot::make(p.x + 1, p.y); |
||||
|
||||
ASSERT_NO_THROW(cv::gpu::labelComponents(mask, components)); |
||||
//left
|
||||
if( p.x > 0 && dl[-1] == -1 && inInt(sp[0], sp[-1])) |
||||
*top++ = dot::make(p.x - 1, p.y); |
||||
|
||||
//bottom
|
||||
if( p.y < (height - 1) && dl[+pitch] == -1 && inInt(sp[0], sp[+image.step1()])) |
||||
*top++ = dot::make(p.x, p.y + 1); |
||||
|
||||
//top
|
||||
if( p.y > 0 && dl[-pitch] == -1 && inInt(sp[0], sp[-static_cast<int>(image.step1())])) |
||||
*top++ = dot::make(p.x, p.y - 1); |
||||
|
||||
p = *--top; |
||||
} |
||||
} |
||||
} |
||||
|
||||
cv::Mat image; |
||||
cv::Mat _labels; |
||||
dot* stack; |
||||
}; |
||||
|
||||
PERF_TEST_P(Image, Labeling_ConnectedComponents, Values<string>("gpu/labeling/aloe-disp.png")) |
||||
{ |
||||
declare.time(1.0); |
||||
|
||||
TEST_CYCLE() |
||||
cv::Mat image = readImage(GetParam(), cv::IMREAD_GRAYSCALE); |
||||
|
||||
if (runOnGpu) |
||||
{ |
||||
cv::gpu::labelComponents(mask, components); |
||||
cv::gpu::GpuMat mask; |
||||
mask.create(image.rows, image.cols, CV_8UC1); |
||||
|
||||
cv::gpu::GpuMat components; |
||||
components.create(image.rows, image.cols, CV_32SC1); |
||||
|
||||
cv::gpu::connectivityMask(cv::gpu::GpuMat(image), mask, cv::Scalar::all(0), cv::Scalar::all(2)); |
||||
|
||||
ASSERT_NO_THROW(cv::gpu::labelComponents(mask, components)); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::gpu::labelComponents(mask, components); |
||||
} |
||||
} |
||||
} |
||||
else |
||||
{ |
||||
GreedyLabeling host(image); |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Labeling, ConnectedComponents, testing::Combine(ALL_DEVICES, testing::Values(cv::Size(261, 262)))); |
||||
host(host._labels); |
||||
|
||||
declare.time(1.0); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
host(host._labels); |
||||
} |
||||
} |
||||
} |
||||
|
||||
#endif |
||||
} // namespace
|
||||
|
||||
@ -1,20 +0,0 @@ |
||||
#include "perf_precomp.hpp" |
||||
|
||||
#ifdef HAVE_CUDA |
||||
|
||||
int main(int argc, char **argv) |
||||
{ |
||||
testing::InitGoogleTest(&argc, argv); |
||||
perf::TestBase::Init(argc, argv); |
||||
return RUN_ALL_TESTS(); |
||||
} |
||||
|
||||
#else |
||||
|
||||
int main() |
||||
{ |
||||
printf("OpenCV was built without CUDA support\n"); |
||||
return 0; |
||||
} |
||||
|
||||
#endif |
||||
@ -1,141 +1,169 @@ |
||||
#include "perf_precomp.hpp" |
||||
|
||||
#ifdef HAVE_CUDA |
||||
using namespace std; |
||||
using namespace testing; |
||||
|
||||
namespace { |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// SetTo
|
||||
|
||||
GPU_PERF_TEST(SetTo, cv::gpu::DeviceInfo, cv::Size, MatType) |
||||
PERF_TEST_P(Sz_Depth_Cn, MatOp_SetTo, Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8U, CV_16U, CV_32F, CV_64F), Values(1, 3, 4))) |
||||
{ |
||||
cv::gpu::DeviceInfo devInfo = GET_PARAM(0); |
||||
cv::gpu::setDevice(devInfo.deviceID()); |
||||
cv::Size size = GET_PARAM(0); |
||||
int depth = GET_PARAM(1); |
||||
int channels = GET_PARAM(2); |
||||
|
||||
cv::Size size = GET_PARAM(1); |
||||
int type = GET_PARAM(2); |
||||
int type = CV_MAKE_TYPE(depth, channels); |
||||
|
||||
cv::gpu::GpuMat src(size, type); |
||||
cv::Scalar val(1, 2, 3, 4); |
||||
|
||||
src.setTo(val); |
||||
if (runOnGpu) |
||||
{ |
||||
cv::gpu::GpuMat d_src(size, type); |
||||
|
||||
d_src.setTo(val); |
||||
|
||||
TEST_CYCLE() |
||||
TEST_CYCLE() |
||||
{ |
||||
d_src.setTo(val); |
||||
} |
||||
} |
||||
else |
||||
{ |
||||
cv::Mat src(size, type); |
||||
|
||||
src.setTo(val); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
src.setTo(val); |
||||
} |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(MatOp, SetTo, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(MatType(CV_8UC1), MatType(CV_8UC3), MatType(CV_8UC4), |
||||
MatType(CV_16UC1), MatType(CV_16UC3), MatType(CV_16UC4), |
||||
MatType(CV_32FC1), MatType(CV_32FC3), MatType(CV_32FC4), |
||||
MatType(CV_64FC1), MatType(CV_64FC3), MatType(CV_64FC4)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// SetToMasked
|
||||
|
||||
GPU_PERF_TEST(SetToMasked, cv::gpu::DeviceInfo, cv::Size, MatType) |
||||
PERF_TEST_P(Sz_Depth_Cn, MatOp_SetToMasked, Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8U, CV_16U, CV_32F, CV_64F), Values(1, 3, 4))) |
||||
{ |
||||
cv::gpu::DeviceInfo devInfo = GET_PARAM(0); |
||||
cv::gpu::setDevice(devInfo.deviceID()); |
||||
cv::Size size = GET_PARAM(0); |
||||
int depth = GET_PARAM(1); |
||||
int channels = GET_PARAM(2); |
||||
|
||||
cv::Size size = GET_PARAM(1); |
||||
int type = GET_PARAM(2); |
||||
int type = CV_MAKE_TYPE(depth, channels); |
||||
|
||||
cv::Mat src_host(size, type); |
||||
fill(src_host, 0, 255); |
||||
cv::Mat src(size, type); |
||||
fillRandom(src); |
||||
|
||||
cv::Mat mask_host(size, CV_8UC1); |
||||
fill(mask_host, 0, 2); |
||||
cv::Mat mask(size, CV_8UC1); |
||||
fillRandom(mask, 0, 2); |
||||
|
||||
cv::gpu::GpuMat src(src_host); |
||||
cv::Scalar val(1, 2, 3, 4); |
||||
cv::gpu::GpuMat mask(mask_host); |
||||
|
||||
src.setTo(val, mask); |
||||
if (runOnGpu) |
||||
{ |
||||
cv::gpu::GpuMat d_src(src); |
||||
cv::gpu::GpuMat d_mask(mask); |
||||
|
||||
d_src.setTo(val, d_mask); |
||||
|
||||
TEST_CYCLE() |
||||
TEST_CYCLE() |
||||
{ |
||||
d_src.setTo(val, d_mask); |
||||
} |
||||
} |
||||
else |
||||
{ |
||||
src.setTo(val, mask); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
src.setTo(val, mask); |
||||
} |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(MatOp, SetToMasked, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(MatType(CV_8UC1), MatType(CV_8UC3), MatType(CV_8UC4), |
||||
MatType(CV_16UC1), MatType(CV_16UC3), MatType(CV_16UC4), |
||||
MatType(CV_32FC1), MatType(CV_32FC3), MatType(CV_32FC4), |
||||
MatType(CV_64FC1), MatType(CV_64FC3), MatType(CV_64FC4)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// CopyToMasked
|
||||
|
||||
GPU_PERF_TEST(CopyToMasked, cv::gpu::DeviceInfo, cv::Size, MatType) |
||||
PERF_TEST_P(Sz_Depth_Cn, MatOp_CopyToMasked, Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8U, CV_16U, CV_32F, CV_64F), Values(1, 3, 4))) |
||||
{ |
||||
cv::gpu::DeviceInfo devInfo = GET_PARAM(0); |
||||
cv::gpu::setDevice(devInfo.deviceID()); |
||||
cv::Size size = GET_PARAM(0); |
||||
int depth = GET_PARAM(1); |
||||
int channels = GET_PARAM(2); |
||||
|
||||
cv::Size size = GET_PARAM(1); |
||||
int type = GET_PARAM(2); |
||||
int type = CV_MAKE_TYPE(depth, channels); |
||||
|
||||
cv::Mat src_host(size, type); |
||||
fill(src_host, 0, 255); |
||||
cv::Mat src(size, type); |
||||
fillRandom(src); |
||||
|
||||
cv::Mat mask_host(size, CV_8UC1); |
||||
fill(mask_host, 0, 2); |
||||
cv::Mat mask(size, CV_8UC1); |
||||
fillRandom(mask, 0, 2); |
||||
|
||||
cv::gpu::GpuMat src(src_host); |
||||
cv::gpu::GpuMat mask(mask_host); |
||||
cv::gpu::GpuMat dst; |
||||
if (runOnGpu) |
||||
{ |
||||
cv::gpu::GpuMat d_src(src); |
||||
cv::gpu::GpuMat d_mask(mask); |
||||
cv::gpu::GpuMat d_dst; |
||||
|
||||
src.copyTo(dst, mask); |
||||
d_src.copyTo(d_dst, d_mask); |
||||
|
||||
TEST_CYCLE() |
||||
TEST_CYCLE() |
||||
{ |
||||
d_src.copyTo(d_dst, d_mask); |
||||
} |
||||
} |
||||
else |
||||
{ |
||||
cv::Mat dst; |
||||
|
||||
src.copyTo(dst, mask); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
src.copyTo(dst, mask); |
||||
} |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(MatOp, CopyToMasked, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(MatType(CV_8UC1), MatType(CV_8UC3), MatType(CV_8UC4), |
||||
MatType(CV_16UC1), MatType(CV_16UC3), MatType(CV_16UC4), |
||||
MatType(CV_32FC1), MatType(CV_32FC3), MatType(CV_32FC4), |
||||
MatType(CV_64FC1), MatType(CV_64FC3), MatType(CV_64FC4)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// ConvertTo
|
||||
|
||||
GPU_PERF_TEST(ConvertTo, cv::gpu::DeviceInfo, cv::Size, MatDepth, MatDepth) |
||||
{ |
||||
cv::gpu::DeviceInfo devInfo = GET_PARAM(0); |
||||
cv::gpu::setDevice(devInfo.deviceID()); |
||||
DEF_PARAM_TEST(Sz_2Depth, cv::Size, MatDepth, MatDepth); |
||||
|
||||
cv::Size size = GET_PARAM(1); |
||||
int depth1 = GET_PARAM(2); |
||||
int depth2 = GET_PARAM(3); |
||||
PERF_TEST_P(Sz_2Depth, MatOp_ConvertTo, Combine(GPU_TYPICAL_MAT_SIZES, Values(CV_8U, CV_16U, CV_32F, CV_64F), Values(CV_8U, CV_16U, CV_32F, CV_64F))) |
||||
{ |
||||
cv::Size size = GET_PARAM(0); |
||||
int depth1 = GET_PARAM(1); |
||||
int depth2 = GET_PARAM(2); |
||||
|
||||
cv::Mat src_host(size, depth1); |
||||
fill(src_host, 0, 255); |
||||
cv::Mat src(size, depth1); |
||||
fillRandom(src); |
||||
|
||||
cv::gpu::GpuMat src(src_host); |
||||
cv::gpu::GpuMat dst; |
||||
if (runOnGpu) |
||||
{ |
||||
cv::gpu::GpuMat d_src(src); |
||||
cv::gpu::GpuMat d_dst; |
||||
|
||||
src.convertTo(dst, depth2, 0.5, 1.0); |
||||
d_src.convertTo(d_dst, depth2, 0.5, 1.0); |
||||
|
||||
TEST_CYCLE() |
||||
TEST_CYCLE() |
||||
{ |
||||
d_src.convertTo(d_dst, depth2, 0.5, 1.0); |
||||
} |
||||
} |
||||
else |
||||
{ |
||||
cv::Mat dst; |
||||
|
||||
src.convertTo(dst, depth2, 0.5, 1.0); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
src.convertTo(dst, depth2, 0.5, 1.0); |
||||
} |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(MatOp, ConvertTo, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(MatDepth(CV_8U), MatDepth(CV_16U), MatDepth(CV_32F), MatDepth(CV_64F)), |
||||
testing::Values(MatDepth(CV_8U), MatDepth(CV_16U), MatDepth(CV_32F), MatDepth(CV_64F)))); |
||||
|
||||
#endif |
||||
} // namespace
|
||||
|
||||
@ -1,85 +1,131 @@ |
||||
#include "perf_precomp.hpp" |
||||
|
||||
#ifdef HAVE_CUDA |
||||
using namespace std; |
||||
using namespace testing; |
||||
|
||||
namespace { |
||||
|
||||
///////////////////////////////////////////////////////////////
|
||||
// HOG
|
||||
|
||||
GPU_PERF_TEST_1(HOG, cv::gpu::DeviceInfo) |
||||
{ |
||||
cv::gpu::DeviceInfo devInfo = GetParam(); |
||||
cv::gpu::setDevice(devInfo.deviceID()); |
||||
DEF_PARAM_TEST_1(Image, string); |
||||
|
||||
cv::Mat img_host = readImage("gpu/hog/road.png", cv::IMREAD_GRAYSCALE); |
||||
ASSERT_FALSE(img_host.empty()); |
||||
PERF_TEST_P(Image, ObjDetect_HOG, Values<string>("gpu/hog/road.png")) |
||||
{ |
||||
cv::Mat img = readImage(GetParam(), cv::IMREAD_GRAYSCALE); |
||||
ASSERT_FALSE(img.empty()); |
||||
|
||||
cv::gpu::GpuMat img(img_host); |
||||
std::vector<cv::Rect> found_locations; |
||||
|
||||
cv::gpu::HOGDescriptor hog; |
||||
hog.setSVMDetector(cv::gpu::HOGDescriptor::getDefaultPeopleDetector()); |
||||
if (runOnGpu) |
||||
{ |
||||
cv::gpu::GpuMat d_img(img); |
||||
|
||||
cv::gpu::HOGDescriptor d_hog; |
||||
d_hog.setSVMDetector(cv::gpu::HOGDescriptor::getDefaultPeopleDetector()); |
||||
|
||||
hog.detectMultiScale(img, found_locations); |
||||
d_hog.detectMultiScale(d_img, found_locations); |
||||
|
||||
TEST_CYCLE() |
||||
TEST_CYCLE() |
||||
{ |
||||
d_hog.detectMultiScale(d_img, found_locations); |
||||
} |
||||
} |
||||
else |
||||
{ |
||||
cv::HOGDescriptor hog; |
||||
hog.setSVMDetector(cv::gpu::HOGDescriptor::getDefaultPeopleDetector()); |
||||
|
||||
hog.detectMultiScale(img, found_locations); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
hog.detectMultiScale(img, found_locations); |
||||
} |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(ObjDetect, HOG, ALL_DEVICES); |
||||
|
||||
///////////////////////////////////////////////////////////////
|
||||
// HaarClassifier
|
||||
|
||||
GPU_PERF_TEST_1(HaarClassifier, cv::gpu::DeviceInfo) |
||||
typedef pair<string, string> pair_string; |
||||
DEF_PARAM_TEST_1(ImageAndCascade, pair_string); |
||||
|
||||
PERF_TEST_P(ImageAndCascade, ObjDetect_HaarClassifier, |
||||
Values<pair_string>(make_pair("gpu/haarcascade/group_1_640x480_VGA.pgm", "gpu/perf/haarcascade_frontalface_alt.xml"))) |
||||
{ |
||||
cv::gpu::DeviceInfo devInfo = GetParam(); |
||||
cv::gpu::setDevice(devInfo.deviceID()); |
||||
cv::Mat img = readImage(GetParam().first, cv::IMREAD_GRAYSCALE); |
||||
ASSERT_FALSE(img.empty()); |
||||
|
||||
cv::Mat img_host = readImage("gpu/haarcascade/group_1_640x480_VGA.pgm", cv::IMREAD_GRAYSCALE); |
||||
ASSERT_FALSE(img_host.empty()); |
||||
if (runOnGpu) |
||||
{ |
||||
cv::gpu::CascadeClassifier_GPU d_cascade; |
||||
ASSERT_TRUE(d_cascade.load(perf::TestBase::getDataPath(GetParam().second))); |
||||
|
||||
cv::gpu::CascadeClassifier_GPU cascade; |
||||
cv::gpu::GpuMat d_img(img); |
||||
cv::gpu::GpuMat d_objects_buffer; |
||||
|
||||
ASSERT_TRUE(cascade.load(perf::TestBase::getDataPath("gpu/perf/haarcascade_frontalface_alt.xml"))); |
||||
d_cascade.detectMultiScale(d_img, d_objects_buffer); |
||||
|
||||
cv::gpu::GpuMat img(img_host); |
||||
cv::gpu::GpuMat objects_buffer; |
||||
TEST_CYCLE() |
||||
{ |
||||
d_cascade.detectMultiScale(d_img, d_objects_buffer); |
||||
} |
||||
} |
||||
else |
||||
{ |
||||
cv::CascadeClassifier cascade; |
||||
ASSERT_TRUE(cascade.load(perf::TestBase::getDataPath("gpu/perf/haarcascade_frontalface_alt.xml"))); |
||||
|
||||
cascade.detectMultiScale(img, objects_buffer); |
||||
std::vector<cv::Rect> rects; |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cascade.detectMultiScale(img, objects_buffer); |
||||
cascade.detectMultiScale(img, rects); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cascade.detectMultiScale(img, rects); |
||||
} |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(ObjDetect, HaarClassifier, ALL_DEVICES); |
||||
///////////////////////////////////////////////////////////////
|
||||
// LBP cascade
|
||||
|
||||
//===================== LBP cascade ==========================//
|
||||
GPU_PERF_TEST_1(LBPClassifier, cv::gpu::DeviceInfo) |
||||
PERF_TEST_P(ImageAndCascade, ObjDetect_LBPClassifier, |
||||
Values<pair_string>(make_pair("gpu/haarcascade/group_1_640x480_VGA.pgm", "gpu/lbpcascade/lbpcascade_frontalface.xml"))) |
||||
{ |
||||
cv::gpu::DeviceInfo devInfo = GetParam(); |
||||
cv::gpu::setDevice(devInfo.deviceID()); |
||||
|
||||
cv::Mat img_host = readImage("gpu/haarcascade/group_1_640x480_VGA.pgm", cv::IMREAD_GRAYSCALE); |
||||
ASSERT_FALSE(img_host.empty()); |
||||
cv::Mat img = readImage(GetParam().first, cv::IMREAD_GRAYSCALE); |
||||
ASSERT_FALSE(img.empty()); |
||||
|
||||
if (runOnGpu) |
||||
{ |
||||
cv::gpu::CascadeClassifier_GPU d_cascade; |
||||
ASSERT_TRUE(d_cascade.load(perf::TestBase::getDataPath(GetParam().second))); |
||||
|
||||
cv::gpu::GpuMat d_img(img); |
||||
cv::gpu::GpuMat d_gpu_rects; |
||||
|
||||
cv::gpu::GpuMat img(img_host); |
||||
cv::gpu::GpuMat gpu_rects; |
||||
cv::gpu::CascadeClassifier_GPU cascade; |
||||
ASSERT_TRUE(cascade.load(perf::TestBase::getDataPath("gpu/lbpcascade/lbpcascade_frontalface.xml"))); |
||||
d_cascade.detectMultiScale(d_img, d_gpu_rects); |
||||
|
||||
cascade.detectMultiScale(img, gpu_rects); |
||||
TEST_CYCLE() |
||||
TEST_CYCLE() |
||||
{ |
||||
d_cascade.detectMultiScale(d_img, d_gpu_rects); |
||||
} |
||||
} |
||||
else |
||||
{ |
||||
cascade.detectMultiScale(img, gpu_rects); |
||||
cv::CascadeClassifier cascade; |
||||
ASSERT_TRUE(cascade.load(perf::TestBase::getDataPath("gpu/lbpcascade/lbpcascade_frontalface.xml"))); |
||||
|
||||
std::vector<cv::Rect> rects; |
||||
|
||||
cascade.detectMultiScale(img, rects); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cascade.detectMultiScale(img, rects); |
||||
} |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(ObjDetect, LBPClassifier, ALL_DEVICES); |
||||
|
||||
#endif |
||||
} // namespace
|
||||
|
||||
@ -1,77 +0,0 @@ |
||||
#ifndef __OPENCV_PERF_GPU_UTILITY_HPP__ |
||||
#define __OPENCV_PERF_GPU_UTILITY_HPP__ |
||||
|
||||
void fill(cv::Mat& m, double a, double b); |
||||
|
||||
using perf::MatType; |
||||
using perf::MatDepth; |
||||
|
||||
CV_ENUM(BorderMode, cv::BORDER_REFLECT101, cv::BORDER_REPLICATE, cv::BORDER_CONSTANT, cv::BORDER_REFLECT, cv::BORDER_WRAP) |
||||
CV_ENUM(Interpolation, cv::INTER_NEAREST, cv::INTER_LINEAR, cv::INTER_CUBIC, cv::INTER_AREA) |
||||
CV_ENUM(NormType, cv::NORM_INF, cv::NORM_L1, cv::NORM_L2, cv::NORM_HAMMING) |
||||
|
||||
struct CvtColorInfo |
||||
{ |
||||
int scn; |
||||
int dcn; |
||||
int code; |
||||
|
||||
explicit CvtColorInfo(int scn_=0, int dcn_=0, int code_=0) : scn(scn_), dcn(dcn_), code(code_) {} |
||||
}; |
||||
|
||||
void PrintTo(const CvtColorInfo& info, std::ostream* os); |
||||
|
||||
#define IMPLEMENT_PARAM_CLASS(name, type) \ |
||||
class name \
|
||||
{ \
|
||||
public: \
|
||||
name ( type arg = type ()) : val_(arg) {} \
|
||||
operator type () const {return val_;} \
|
||||
private: \
|
||||
type val_; \
|
||||
}; \
|
||||
inline void PrintTo( name param, std::ostream* os) \
|
||||
{ \
|
||||
*os << #name << " = " << testing::PrintToString(static_cast< type >(param)); \
|
||||
} |
||||
|
||||
IMPLEMENT_PARAM_CLASS(Channels, int) |
||||
|
||||
namespace cv { namespace gpu |
||||
{ |
||||
void PrintTo(const cv::gpu::DeviceInfo& info, std::ostream* os); |
||||
}} |
||||
|
||||
#define GPU_PERF_TEST(name, ...) \ |
||||
struct name : perf::TestBaseWithParam< std::tr1::tuple< __VA_ARGS__ > > \
|
||||
{ \
|
||||
public: \
|
||||
name() {} \
|
||||
protected: \
|
||||
void PerfTestBody(); \
|
||||
}; \
|
||||
TEST_P(name, perf){ RunPerfTestBody(); } \
|
||||
void name :: PerfTestBody() |
||||
|
||||
#define GPU_PERF_TEST_1(name, param_type) \ |
||||
struct name : perf::TestBaseWithParam< param_type > \
|
||||
{ \
|
||||
public: \
|
||||
name() {} \
|
||||
protected: \
|
||||
void PerfTestBody(); \
|
||||
}; \
|
||||
TEST_P(name, perf){ RunPerfTestBody(); } \
|
||||
void name :: PerfTestBody() |
||||
|
||||
#define GPU_TYPICAL_MAT_SIZES testing::Values(perf::szSXGA, perf::sz1080p, cv::Size(1800, 1500)) |
||||
|
||||
cv::Mat readImage(const std::string& fileName, int flags = cv::IMREAD_COLOR); |
||||
|
||||
const std::vector<cv::gpu::DeviceInfo>& devices(); |
||||
|
||||
#define ALL_DEVICES testing::ValuesIn(devices()) |
||||
|
||||
#define GET_PARAM(k) std::tr1::get< k >(GetParam()) |
||||
|
||||
#endif // __OPENCV_PERF_GPU_UTILITY_HPP__
|
||||
File diff suppressed because it is too large
Load Diff
@ -0,0 +1,45 @@ |
||||
#ifndef __OPENCV_PERF_GPU_UTILITY_HPP__ |
||||
#define __OPENCV_PERF_GPU_UTILITY_HPP__ |
||||
|
||||
#include "opencv2/core/core.hpp" |
||||
#include "opencv2/core/gpumat.hpp" |
||||
#include "opencv2/imgproc/imgproc.hpp" |
||||
#include "opencv2/ts/ts_perf.hpp" |
||||
|
||||
extern bool runOnGpu; |
||||
|
||||
void fillRandom(cv::Mat& m, double a = 0.0, double b = 255.0); |
||||
cv::Mat readImage(const std::string& fileName, int flags = cv::IMREAD_COLOR); |
||||
|
||||
using perf::MatType; |
||||
using perf::MatDepth; |
||||
|
||||
CV_ENUM(BorderMode, cv::BORDER_REFLECT101, cv::BORDER_REPLICATE, cv::BORDER_CONSTANT, cv::BORDER_REFLECT, cv::BORDER_WRAP) |
||||
#define ALL_BORDER_MODES testing::ValuesIn(BorderMode::all()) |
||||
CV_ENUM(Interpolation, cv::INTER_NEAREST, cv::INTER_LINEAR, cv::INTER_CUBIC, cv::INTER_AREA) |
||||
#define ALL_INTERPOLATIONS testing::ValuesIn(Interpolation::all()) |
||||
CV_ENUM(NormType, cv::NORM_INF, cv::NORM_L1, cv::NORM_L2, cv::NORM_HAMMING) |
||||
|
||||
struct CvtColorInfo |
||||
{ |
||||
int scn; |
||||
int dcn; |
||||
int code; |
||||
|
||||
explicit CvtColorInfo(int scn_=0, int dcn_=0, int code_=0) : scn(scn_), dcn(dcn_), code(code_) {} |
||||
}; |
||||
void PrintTo(const CvtColorInfo& info, std::ostream* os); |
||||
|
||||
#define GET_PARAM(k) std::tr1::get< k >(GetParam()) |
||||
|
||||
#define DEF_PARAM_TEST(name, ...) typedef ::perf::TestBaseWithParam< std::tr1::tuple< __VA_ARGS__ > > name |
||||
#define DEF_PARAM_TEST_1(name, param_type) typedef ::perf::TestBaseWithParam< param_type > name |
||||
|
||||
DEF_PARAM_TEST_1(Sz, cv::Size); |
||||
typedef perf::Size_MatType Sz_Type; |
||||
DEF_PARAM_TEST(Sz_Depth, cv::Size, MatDepth); |
||||
DEF_PARAM_TEST(Sz_Depth_Cn, cv::Size, MatDepth, int); |
||||
|
||||
#define GPU_TYPICAL_MAT_SIZES testing::Values(perf::szSXGA, perf::sz720p, perf::sz1080p) |
||||
|
||||
#endif // __OPENCV_PERF_GPU_UTILITY_HPP__
|
||||
@ -1,136 +0,0 @@ |
||||
#include "perf_cpu_precomp.hpp" |
||||
|
||||
#ifdef HAVE_CUDA |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// StereoBM
|
||||
|
||||
GPU_PERF_TEST_1(StereoBM, cv::gpu::DeviceInfo) |
||||
{ |
||||
cv::Mat img_l = readImage("gpu/perf/aloe.jpg", cv::IMREAD_GRAYSCALE); |
||||
ASSERT_FALSE(img_l.empty()); |
||||
|
||||
cv::Mat img_r = readImage("gpu/perf/aloeR.jpg", cv::IMREAD_GRAYSCALE); |
||||
ASSERT_FALSE(img_r.empty()); |
||||
|
||||
cv::StereoBM bm(0, 256); |
||||
|
||||
cv::Mat dst; |
||||
|
||||
bm(img_l, img_r, dst); |
||||
|
||||
declare.time(5.0); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
bm(img_l, img_r, dst); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Calib3D, StereoBM, ALL_DEVICES); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// ProjectPoints
|
||||
|
||||
IMPLEMENT_PARAM_CLASS(Count, int) |
||||
|
||||
GPU_PERF_TEST(ProjectPoints, cv::gpu::DeviceInfo, Count) |
||||
{ |
||||
int count = GET_PARAM(1); |
||||
|
||||
cv::Mat src(1, count, CV_32FC3); |
||||
fill(src, -100, 100); |
||||
|
||||
cv::Mat rvec = cv::Mat::ones(1, 3, CV_32FC1); |
||||
cv::Mat tvec = cv::Mat::ones(1, 3, CV_32FC1); |
||||
cv::Mat camera_mat = cv::Mat::ones(3, 3, CV_32FC1); |
||||
cv::Mat dst; |
||||
|
||||
cv::projectPoints(src, rvec, tvec, camera_mat, cv::noArray(), dst); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::projectPoints(src, rvec, tvec, camera_mat, cv::noArray(), dst); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Calib3D, ProjectPoints, testing::Combine( |
||||
ALL_DEVICES, |
||||
testing::Values<Count>(5000, 10000, 20000))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// SolvePnPRansac
|
||||
|
||||
GPU_PERF_TEST(SolvePnPRansac, cv::gpu::DeviceInfo, Count) |
||||
{ |
||||
int count = GET_PARAM(1); |
||||
|
||||
cv::Mat object(1, count, CV_32FC3); |
||||
fill(object, -100, 100); |
||||
|
||||
cv::Mat camera_mat(3, 3, CV_32FC1); |
||||
fill(camera_mat, 0.5, 1); |
||||
camera_mat.at<float>(0, 1) = 0.f; |
||||
camera_mat.at<float>(1, 0) = 0.f; |
||||
camera_mat.at<float>(2, 0) = 0.f; |
||||
camera_mat.at<float>(2, 1) = 0.f; |
||||
|
||||
cv::Mat dist_coef(1, 8, CV_32F, cv::Scalar::all(0)); |
||||
|
||||
std::vector<cv::Point2f> image_vec; |
||||
cv::Mat rvec_gold(1, 3, CV_32FC1); |
||||
fill(rvec_gold, 0, 1); |
||||
cv::Mat tvec_gold(1, 3, CV_32FC1); |
||||
fill(tvec_gold, 0, 1); |
||||
cv::projectPoints(object, rvec_gold, tvec_gold, camera_mat, dist_coef, image_vec); |
||||
|
||||
cv::Mat image(1, count, CV_32FC2, &image_vec[0]); |
||||
|
||||
cv::Mat rvec; |
||||
cv::Mat tvec; |
||||
|
||||
cv::solvePnPRansac(object, image, camera_mat, dist_coef, rvec, tvec); |
||||
|
||||
declare.time(10.0); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::solvePnPRansac(object, image, camera_mat, dist_coef, rvec, tvec); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Calib3D, SolvePnPRansac, testing::Combine( |
||||
ALL_DEVICES, |
||||
testing::Values<Count>(5000, 10000, 20000))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// ReprojectImageTo3D
|
||||
|
||||
GPU_PERF_TEST(ReprojectImageTo3D, cv::gpu::DeviceInfo, cv::Size, MatDepth) |
||||
{ |
||||
cv::Size size = GET_PARAM(1); |
||||
int depth = GET_PARAM(2); |
||||
|
||||
cv::Mat src(size, depth); |
||||
fill(src, 5.0, 30.0); |
||||
|
||||
cv::Mat Q(4, 4, CV_32FC1); |
||||
fill(Q, 0.1, 1.0); |
||||
|
||||
cv::Mat dst; |
||||
|
||||
cv::reprojectImageTo3D(src, dst, Q); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::reprojectImageTo3D(src, dst, Q); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Calib3D, ReprojectImageTo3D, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values<MatDepth>(CV_8U, CV_16S))); |
||||
|
||||
#endif |
||||
|
||||
File diff suppressed because it is too large
Load Diff
@ -1 +0,0 @@ |
||||
#include "perf_cpu_precomp.hpp" |
||||
@ -1,32 +0,0 @@ |
||||
#ifdef __GNUC__ |
||||
# pragma GCC diagnostic ignored "-Wmissing-declarations" |
||||
# pragma GCC diagnostic ignored "-Wmissing-prototypes" //OSX
|
||||
#endif |
||||
|
||||
#ifndef __OPENCV_PERF_CPU_PRECOMP_HPP__ |
||||
#define __OPENCV_PERF_CPU_PRECOMP_HPP__ |
||||
|
||||
#include <cstdio> |
||||
#include <iostream> |
||||
|
||||
#include "cvconfig.h" |
||||
|
||||
#include "opencv2/ts/ts.hpp" |
||||
#include "opencv2/ts/ts_perf.hpp" |
||||
|
||||
#include "opencv2/core/core.hpp" |
||||
#include "opencv2/highgui/highgui.hpp" |
||||
#include "opencv2/gpu/gpu.hpp" |
||||
#include "opencv2/imgproc/imgproc.hpp" |
||||
#include "opencv2/video/video.hpp" |
||||
#include "opencv2/calib3d/calib3d.hpp" |
||||
#include "opencv2/nonfree/nonfree.hpp" |
||||
#include "opencv2/legacy/legacy.hpp" |
||||
|
||||
#include "perf_utility.hpp" |
||||
|
||||
#ifdef GTEST_CREATE_SHARED_LIBRARY |
||||
#error no modules except ts should have GTEST_CREATE_SHARED_LIBRARY defined |
||||
#endif |
||||
|
||||
#endif |
||||
@ -1,187 +0,0 @@ |
||||
#include "perf_cpu_precomp.hpp" |
||||
|
||||
#ifdef HAVE_CUDA |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// SURF
|
||||
|
||||
GPU_PERF_TEST_1(SURF, cv::gpu::DeviceInfo) |
||||
{ |
||||
cv::Mat img = readImage("gpu/perf/aloe.jpg", cv::IMREAD_GRAYSCALE); |
||||
ASSERT_FALSE(img.empty()); |
||||
|
||||
cv::SURF surf; |
||||
|
||||
std::vector<cv::KeyPoint> keypoints; |
||||
cv::Mat descriptors; |
||||
|
||||
surf(img, cv::noArray(), keypoints, descriptors); |
||||
|
||||
declare.time(50.0); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
keypoints.clear(); |
||||
surf(img, cv::noArray(), keypoints, descriptors); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Features2D, SURF, ALL_DEVICES); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// FAST
|
||||
|
||||
GPU_PERF_TEST_1(FAST, cv::gpu::DeviceInfo) |
||||
{ |
||||
cv::Mat img = readImage("gpu/perf/aloe.jpg", cv::IMREAD_GRAYSCALE); |
||||
ASSERT_FALSE(img.empty()); |
||||
|
||||
std::vector<cv::KeyPoint> keypoints; |
||||
|
||||
cv::FAST(img, keypoints, 20); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
keypoints.clear(); |
||||
cv::FAST(img, keypoints, 20); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Features2D, FAST, ALL_DEVICES); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// ORB
|
||||
|
||||
GPU_PERF_TEST_1(ORB, cv::gpu::DeviceInfo) |
||||
{ |
||||
cv::Mat img = readImage("gpu/perf/aloe.jpg", cv::IMREAD_GRAYSCALE); |
||||
ASSERT_FALSE(img.empty()); |
||||
|
||||
cv::ORB orb(4000); |
||||
|
||||
std::vector<cv::KeyPoint> keypoints; |
||||
cv::Mat descriptors; |
||||
|
||||
orb(img, cv::noArray(), keypoints, descriptors); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
keypoints.clear(); |
||||
orb(img, cv::noArray(), keypoints, descriptors); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Features2D, ORB, ALL_DEVICES); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// BruteForceMatcher_match
|
||||
|
||||
IMPLEMENT_PARAM_CLASS(DescriptorSize, int) |
||||
|
||||
GPU_PERF_TEST(BruteForceMatcher_match, cv::gpu::DeviceInfo, DescriptorSize, NormType) |
||||
{ |
||||
int desc_size = GET_PARAM(1); |
||||
int normType = GET_PARAM(2); |
||||
|
||||
int type = normType == cv::NORM_HAMMING ? CV_8U : CV_32F; |
||||
|
||||
cv::Mat query(3000, desc_size, type); |
||||
fill(query, 0.0, 10.0); |
||||
|
||||
cv::Mat train(3000, desc_size, type); |
||||
fill(train, 0.0, 10.0); |
||||
|
||||
cv::BFMatcher matcher(normType); |
||||
|
||||
std::vector<cv::DMatch> matches; |
||||
|
||||
matcher.match(query, train, matches); |
||||
|
||||
declare.time(20.0); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
matcher.match(query, train, matches); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Features2D, BruteForceMatcher_match, testing::Combine( |
||||
ALL_DEVICES, |
||||
testing::Values(DescriptorSize(64), DescriptorSize(128), DescriptorSize(256)), |
||||
testing::Values(NormType(cv::NORM_L1), NormType(cv::NORM_L2), NormType(cv::NORM_HAMMING)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// BruteForceMatcher_knnMatch
|
||||
|
||||
IMPLEMENT_PARAM_CLASS(K, int) |
||||
|
||||
GPU_PERF_TEST(BruteForceMatcher_knnMatch, cv::gpu::DeviceInfo, DescriptorSize, K, NormType) |
||||
{ |
||||
int desc_size = GET_PARAM(1); |
||||
int k = GET_PARAM(2); |
||||
int normType = GET_PARAM(3); |
||||
|
||||
int type = normType == cv::NORM_HAMMING ? CV_8U : CV_32F; |
||||
|
||||
cv::Mat query(3000, desc_size, type); |
||||
fill(query, 0.0, 10.0); |
||||
|
||||
cv::Mat train(3000, desc_size, type); |
||||
fill(train, 0.0, 10.0); |
||||
|
||||
cv::BFMatcher matcher(normType); |
||||
|
||||
std::vector< std::vector<cv::DMatch> > matches; |
||||
|
||||
matcher.knnMatch(query, train, matches, k); |
||||
|
||||
declare.time(30.0); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
matcher.knnMatch(query, train, matches, k); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Features2D, BruteForceMatcher_knnMatch, testing::Combine( |
||||
ALL_DEVICES, |
||||
testing::Values(DescriptorSize(64), DescriptorSize(128), DescriptorSize(256)), |
||||
testing::Values(K(2), K(3)), |
||||
testing::Values(NormType(cv::NORM_L1), NormType(cv::NORM_L2), NormType(cv::NORM_HAMMING)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// BruteForceMatcher_radiusMatch
|
||||
|
||||
GPU_PERF_TEST(BruteForceMatcher_radiusMatch, cv::gpu::DeviceInfo, DescriptorSize, NormType) |
||||
{ |
||||
int desc_size = GET_PARAM(1); |
||||
int normType = GET_PARAM(2); |
||||
|
||||
int type = normType == cv::NORM_HAMMING ? CV_8U : CV_32F; |
||||
|
||||
cv::Mat query(3000, desc_size, type); |
||||
fill(query, 0.0, 1.0); |
||||
|
||||
cv::Mat train(3000, desc_size, type); |
||||
fill(train, 0.0, 1.0); |
||||
|
||||
cv::BFMatcher matcher(normType); |
||||
|
||||
std::vector< std::vector<cv::DMatch> > matches; |
||||
|
||||
matcher.radiusMatch(query, train, matches, 2.0); |
||||
|
||||
declare.time(30.0); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
matcher.radiusMatch(query, train, matches, 2.0); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Features2D, BruteForceMatcher_radiusMatch, testing::Combine( |
||||
ALL_DEVICES, |
||||
testing::Values(DescriptorSize(64), DescriptorSize(128), DescriptorSize(256)), |
||||
testing::Values(NormType(cv::NORM_L1), NormType(cv::NORM_L2), NormType(cv::NORM_HAMMING)))); |
||||
|
||||
#endif |
||||
@ -1,283 +0,0 @@ |
||||
#include "perf_cpu_precomp.hpp" |
||||
|
||||
#ifdef HAVE_CUDA |
||||
|
||||
IMPLEMENT_PARAM_CLASS(KernelSize, int) |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// Blur
|
||||
|
||||
GPU_PERF_TEST(Blur, cv::gpu::DeviceInfo, cv::Size, MatType, KernelSize) |
||||
{ |
||||
cv::Size size = GET_PARAM(1); |
||||
int type = GET_PARAM(2); |
||||
int ksize = GET_PARAM(3); |
||||
|
||||
cv::Mat src(size, type); |
||||
fill(src, 0.0, 255.0); |
||||
|
||||
cv::Mat dst; |
||||
|
||||
cv::blur(src, dst, cv::Size(ksize, ksize)); |
||||
|
||||
declare.time(20.0); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::blur(src, dst, cv::Size(ksize, ksize)); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Filters, Blur, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(MatType(CV_8UC1), MatType(CV_8UC4)), |
||||
testing::Values(KernelSize(3), KernelSize(5), KernelSize(7)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// Sobel
|
||||
|
||||
GPU_PERF_TEST(Sobel, cv::gpu::DeviceInfo, cv::Size, MatType, KernelSize) |
||||
{ |
||||
cv::Size size = GET_PARAM(1); |
||||
int type = GET_PARAM(2); |
||||
int ksize = GET_PARAM(3); |
||||
|
||||
cv::Mat src(size, type); |
||||
fill(src, 0.0, 255.0); |
||||
|
||||
cv::Mat dst; |
||||
|
||||
cv::Sobel(src, dst, -1, 1, 1, ksize); |
||||
|
||||
declare.time(20.0); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::Sobel(src, dst, -1, 1, 1, ksize); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Filters, Sobel, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(MatType(CV_8UC1), MatType(CV_8UC4), MatType(CV_32FC1)), |
||||
testing::Values(KernelSize(3), KernelSize(5), KernelSize(7), KernelSize(9), KernelSize(11), KernelSize(13), KernelSize(15)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// Scharr
|
||||
|
||||
GPU_PERF_TEST(Scharr, cv::gpu::DeviceInfo, cv::Size, MatType) |
||||
{ |
||||
cv::Size size = GET_PARAM(1); |
||||
int type = GET_PARAM(2); |
||||
|
||||
cv::Mat src(size, type); |
||||
fill(src, 0.0, 255.0); |
||||
|
||||
cv::Mat dst; |
||||
|
||||
cv::Scharr(src, dst, -1, 1, 0); |
||||
|
||||
declare.time(20.0); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::Scharr(src, dst, -1, 1, 0); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Filters, Scharr, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(MatType(CV_8UC1), MatType(CV_8UC4), MatType(CV_32FC1)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// GaussianBlur
|
||||
|
||||
GPU_PERF_TEST(GaussianBlur, cv::gpu::DeviceInfo, cv::Size, MatType, KernelSize) |
||||
{ |
||||
cv::Size size = GET_PARAM(1); |
||||
int type = GET_PARAM(2); |
||||
int ksize = GET_PARAM(3); |
||||
|
||||
cv::Mat src(size, type); |
||||
fill(src, 0.0, 255.0); |
||||
|
||||
cv::Mat dst; |
||||
|
||||
cv::GaussianBlur(src, dst, cv::Size(ksize, ksize), 0.5); |
||||
|
||||
declare.time(20.0); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::GaussianBlur(src, dst, cv::Size(ksize, ksize), 0.5); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Filters, GaussianBlur, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(MatType(CV_8UC1), MatType(CV_8UC4), MatType(CV_32FC1)), |
||||
testing::Values(KernelSize(3), KernelSize(5), KernelSize(7), KernelSize(9), KernelSize(11), KernelSize(13), KernelSize(15)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// Laplacian
|
||||
|
||||
GPU_PERF_TEST(Laplacian, cv::gpu::DeviceInfo, cv::Size, MatType, KernelSize) |
||||
{ |
||||
cv::Size size = GET_PARAM(1); |
||||
int type = GET_PARAM(2); |
||||
int ksize = GET_PARAM(3); |
||||
|
||||
cv::Mat src(size, type); |
||||
fill(src, 0.0, 255.0); |
||||
|
||||
cv::Mat dst; |
||||
|
||||
cv::Laplacian(src, dst, -1, ksize); |
||||
|
||||
declare.time(20.0); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::Laplacian(src, dst, -1, ksize); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Filters, Laplacian, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(MatType(CV_8UC1), MatType(CV_8UC4), MatType(CV_32FC1), MatType(CV_32FC4)), |
||||
testing::Values(KernelSize(1), KernelSize(3)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// Erode
|
||||
|
||||
GPU_PERF_TEST(Erode, cv::gpu::DeviceInfo, cv::Size, MatType) |
||||
{ |
||||
cv::Size size = GET_PARAM(1); |
||||
int type = GET_PARAM(2); |
||||
|
||||
cv::Mat src(size, type); |
||||
fill(src, 0.0, 255.0); |
||||
|
||||
cv::Mat ker = cv::getStructuringElement(cv::MORPH_RECT, cv::Size(3, 3)); |
||||
|
||||
cv::Mat dst; |
||||
|
||||
cv::erode(src, dst, ker); |
||||
|
||||
declare.time(20.0); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::erode(src, dst, ker); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Filters, Erode, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(MatType(CV_8UC1), MatType(CV_8UC4)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// Dilate
|
||||
|
||||
GPU_PERF_TEST(Dilate, cv::gpu::DeviceInfo, cv::Size, MatType) |
||||
{ |
||||
cv::Size size = GET_PARAM(1); |
||||
int type = GET_PARAM(2); |
||||
|
||||
cv::Mat src(size, type); |
||||
fill(src, 0.0, 255.0); |
||||
|
||||
cv::Mat ker = cv::getStructuringElement(cv::MORPH_RECT, cv::Size(3, 3)); |
||||
|
||||
cv::Mat dst; |
||||
|
||||
cv::dilate(src, dst, ker); |
||||
|
||||
declare.time(20.0); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::dilate(src, dst, ker); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Filters, Dilate, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(MatType(CV_8UC1), MatType(CV_8UC4)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// MorphologyEx
|
||||
|
||||
CV_ENUM(MorphOp, cv::MORPH_OPEN, cv::MORPH_CLOSE, cv::MORPH_GRADIENT, cv::MORPH_TOPHAT, cv::MORPH_BLACKHAT) |
||||
#define ALL_MORPH_OPS testing::Values(MorphOp(cv::MORPH_OPEN), MorphOp(cv::MORPH_CLOSE), MorphOp(cv::MORPH_GRADIENT), MorphOp(cv::MORPH_TOPHAT), MorphOp(cv::MORPH_BLACKHAT)) |
||||
|
||||
GPU_PERF_TEST(MorphologyEx, cv::gpu::DeviceInfo, cv::Size, MatType, MorphOp) |
||||
{ |
||||
cv::Size size = GET_PARAM(1); |
||||
int type = GET_PARAM(2); |
||||
int morphOp = GET_PARAM(3); |
||||
|
||||
cv::Mat src(size, type); |
||||
fill(src, 0.0, 255.0); |
||||
|
||||
cv::Mat dst; |
||||
|
||||
cv::Mat ker = cv::getStructuringElement(cv::MORPH_RECT, cv::Size(3, 3)); |
||||
|
||||
cv::morphologyEx(src, dst, morphOp, ker); |
||||
|
||||
declare.time(20.0); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::morphologyEx(src, dst, morphOp, ker); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Filters, MorphologyEx, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(MatType(CV_8UC1), MatType(CV_8UC4)), |
||||
ALL_MORPH_OPS)); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// Filter2D
|
||||
|
||||
GPU_PERF_TEST(Filter2D, cv::gpu::DeviceInfo, cv::Size, MatType, KernelSize) |
||||
{ |
||||
cv::Size size = GET_PARAM(1); |
||||
int type = GET_PARAM(2); |
||||
int ksize = GET_PARAM(3); |
||||
|
||||
cv::Mat src(size, type); |
||||
fill(src, 0.0, 255.0); |
||||
|
||||
cv::Mat kernel(ksize, ksize, CV_32FC1); |
||||
fill(kernel, 0.0, 1.0); |
||||
|
||||
cv::Mat dst; |
||||
|
||||
cv::filter2D(src, dst, -1, kernel); |
||||
|
||||
declare.time(20.0); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::filter2D(src, dst, -1, kernel); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Filters, Filter2D, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(MatType(CV_8UC1), MatType(CV_8UC4), MatType(CV_32FC1), MatType(CV_32FC4)), |
||||
testing::Values(KernelSize(3), KernelSize(5), KernelSize(7), KernelSize(9), KernelSize(11), KernelSize(13), KernelSize(15)))); |
||||
|
||||
#endif |
||||
@ -1,771 +0,0 @@ |
||||
#include "perf_cpu_precomp.hpp" |
||||
|
||||
#ifdef HAVE_CUDA |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// Remap
|
||||
|
||||
GPU_PERF_TEST(Remap, cv::gpu::DeviceInfo, cv::Size, MatType, Interpolation, BorderMode) |
||||
{ |
||||
cv::Size size = GET_PARAM(1); |
||||
int type = GET_PARAM(2); |
||||
int interpolation = GET_PARAM(3); |
||||
int borderMode = GET_PARAM(4); |
||||
|
||||
cv::Mat src(size, type); |
||||
fill(src, 0, 255); |
||||
|
||||
cv::Mat xmap(size, CV_32FC1); |
||||
fill(xmap, 0, size.width); |
||||
|
||||
cv::Mat ymap(size, CV_32FC1); |
||||
fill(ymap, 0, size.height); |
||||
|
||||
cv::Mat dst; |
||||
|
||||
cv::remap(src, dst, xmap, ymap, interpolation, borderMode); |
||||
|
||||
declare.time(20.0); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::remap(src, dst, xmap, ymap, interpolation, borderMode); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(ImgProc, Remap, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(MatType(CV_8UC1), MatType(CV_8UC3), MatType(CV_8UC4), |
||||
MatType(CV_16UC1), MatType(CV_16UC3), MatType(CV_16UC4), |
||||
MatType(CV_32FC1), MatType(CV_32FC3), MatType(CV_32FC4)), |
||||
testing::Values(Interpolation(cv::INTER_NEAREST), Interpolation(cv::INTER_LINEAR), Interpolation(cv::INTER_CUBIC)), |
||||
testing::Values(BorderMode(cv::BORDER_REFLECT101), BorderMode(cv::BORDER_REPLICATE), BorderMode(cv::BORDER_CONSTANT), BorderMode(cv::BORDER_REFLECT), BorderMode(cv::BORDER_WRAP)))); |
||||
|
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// Resize
|
||||
|
||||
IMPLEMENT_PARAM_CLASS(Scale, double) |
||||
|
||||
GPU_PERF_TEST(Resize, cv::gpu::DeviceInfo, cv::Size, MatType, Interpolation, Scale) |
||||
{ |
||||
cv::Size size = GET_PARAM(1); |
||||
int type = GET_PARAM(2); |
||||
int interpolation = GET_PARAM(3); |
||||
double f = GET_PARAM(4); |
||||
|
||||
cv::Mat src(size, type); |
||||
fill(src, 0, 255); |
||||
|
||||
cv::Mat dst; |
||||
|
||||
cv::resize(src, dst, cv::Size(), f, f, interpolation); |
||||
|
||||
declare.time(20.0); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::resize(src, dst, cv::Size(), f, f, interpolation); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(ImgProc, Resize, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(MatType(CV_8UC1), MatType(CV_8UC3), MatType(CV_8UC4), |
||||
MatType(CV_16UC1), MatType(CV_16UC3), MatType(CV_16UC4), |
||||
MatType(CV_32FC1), MatType(CV_32FC3), MatType(CV_32FC4)), |
||||
testing::Values(Interpolation(cv::INTER_NEAREST), Interpolation(cv::INTER_LINEAR), |
||||
Interpolation(cv::INTER_CUBIC), Interpolation(cv::INTER_AREA)), |
||||
testing::Values(Scale(0.5), Scale(0.3), Scale(2.0)))); |
||||
|
||||
GPU_PERF_TEST(ResizeArea, cv::gpu::DeviceInfo, cv::Size, MatType, Scale) |
||||
{ |
||||
cv::Size size = GET_PARAM(1); |
||||
int type = GET_PARAM(2); |
||||
int interpolation = cv::INTER_AREA; |
||||
double f = GET_PARAM(3); |
||||
|
||||
cv::Mat src_host(size, type); |
||||
fill(src_host, 0, 255); |
||||
|
||||
cv::Mat src(src_host); |
||||
cv::Mat dst; |
||||
|
||||
cv::resize(src, dst, cv::Size(), f, f, interpolation); |
||||
|
||||
declare.time(1.0); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::resize(src, dst, cv::Size(), f, f, interpolation); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(ImgProc, ResizeArea, testing::Combine( |
||||
ALL_DEVICES, |
||||
testing::Values(perf::sz1080p, cv::Size(4096, 2048)), |
||||
testing::Values(MatType(CV_8UC1)/*, MatType(CV_8UC3), MatType(CV_8UC4),
|
||||
MatType(CV_16UC1), MatType(CV_16UC3), MatType(CV_16UC4), |
||||
MatType(CV_32FC1), MatType(CV_32FC3), MatType(CV_32FC4)*/), |
||||
testing::Values(Scale(0.2),Scale(0.1),Scale(0.05)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// WarpAffine
|
||||
|
||||
GPU_PERF_TEST(WarpAffine, cv::gpu::DeviceInfo, cv::Size, MatType, Interpolation, BorderMode) |
||||
{ |
||||
cv::Size size = GET_PARAM(1); |
||||
int type = GET_PARAM(2); |
||||
int interpolation = GET_PARAM(3); |
||||
int borderMode = GET_PARAM(4); |
||||
|
||||
cv::Mat src(size, type); |
||||
fill(src, 0, 255); |
||||
|
||||
cv::Mat dst; |
||||
|
||||
const double aplha = CV_PI / 4; |
||||
double mat[2][3] = { {std::cos(aplha), -std::sin(aplha), src.cols / 2}, |
||||
{std::sin(aplha), std::cos(aplha), 0}}; |
||||
cv::Mat M(2, 3, CV_64F, (void*) mat); |
||||
|
||||
cv::warpAffine(src, dst, M, size, interpolation, borderMode); |
||||
|
||||
declare.time(20.0); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::warpAffine(src, dst, M, size, interpolation, borderMode); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(ImgProc, WarpAffine, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(MatType(CV_8UC1), MatType(CV_8UC3), MatType(CV_8UC4), |
||||
MatType(CV_16UC1), MatType(CV_16UC3), MatType(CV_16UC4), |
||||
MatType(CV_32FC1), MatType(CV_32FC3), MatType(CV_32FC4)), |
||||
testing::Values(Interpolation(cv::INTER_NEAREST), Interpolation(cv::INTER_LINEAR), Interpolation(cv::INTER_CUBIC)), |
||||
testing::Values(BorderMode(cv::BORDER_REFLECT101), BorderMode(cv::BORDER_REPLICATE), BorderMode(cv::BORDER_CONSTANT), BorderMode(cv::BORDER_REFLECT), BorderMode(cv::BORDER_WRAP)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// WarpPerspective
|
||||
|
||||
GPU_PERF_TEST(WarpPerspective, cv::gpu::DeviceInfo, cv::Size, MatType, Interpolation, BorderMode) |
||||
{ |
||||
cv::Size size = GET_PARAM(1); |
||||
int type = GET_PARAM(2); |
||||
int interpolation = GET_PARAM(3); |
||||
int borderMode = GET_PARAM(4); |
||||
|
||||
cv::Mat src(size, type); |
||||
fill(src, 0, 255); |
||||
|
||||
cv::Mat dst; |
||||
|
||||
const double aplha = CV_PI / 4; |
||||
double mat[3][3] = { {std::cos(aplha), -std::sin(aplha), src.cols / 2}, |
||||
{std::sin(aplha), std::cos(aplha), 0}, |
||||
{0.0, 0.0, 1.0}}; |
||||
cv::Mat M(3, 3, CV_64F, (void*) mat); |
||||
|
||||
cv::warpPerspective(src, dst, M, size, interpolation, borderMode); |
||||
|
||||
declare.time(20.0); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::warpPerspective(src, dst, M, size, interpolation, borderMode); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(ImgProc, WarpPerspective, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(MatType(CV_8UC1), MatType(CV_8UC3), MatType(CV_8UC4), |
||||
MatType(CV_16UC1), MatType(CV_16UC3), MatType(CV_16UC4), |
||||
MatType(CV_32FC1), MatType(CV_32FC3), MatType(CV_32FC4)), |
||||
testing::Values(Interpolation(cv::INTER_NEAREST), Interpolation(cv::INTER_LINEAR), Interpolation(cv::INTER_CUBIC)), |
||||
testing::Values(BorderMode(cv::BORDER_REFLECT101), BorderMode(cv::BORDER_REPLICATE), BorderMode(cv::BORDER_CONSTANT), BorderMode(cv::BORDER_REFLECT), BorderMode(cv::BORDER_WRAP)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// CopyMakeBorder
|
||||
|
||||
GPU_PERF_TEST(CopyMakeBorder, cv::gpu::DeviceInfo, cv::Size, MatType, BorderMode) |
||||
{ |
||||
cv::Size size = GET_PARAM(1); |
||||
int type = GET_PARAM(2); |
||||
int borderType = GET_PARAM(3); |
||||
|
||||
cv::Mat src(size, type); |
||||
fill(src, 0, 255); |
||||
|
||||
cv::Mat dst; |
||||
|
||||
cv::copyMakeBorder(src, dst, 5, 5, 5, 5, borderType); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::copyMakeBorder(src, dst, 5, 5, 5, 5, borderType); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(ImgProc, CopyMakeBorder, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(MatType(CV_8UC1), MatType(CV_8UC3), MatType(CV_8UC4), |
||||
MatType(CV_16UC1), MatType(CV_16UC3), MatType(CV_16UC4), |
||||
MatType(CV_32FC1), MatType(CV_32FC3), MatType(CV_32FC4)), |
||||
testing::Values(BorderMode(cv::BORDER_REFLECT101), BorderMode(cv::BORDER_REPLICATE), BorderMode(cv::BORDER_CONSTANT), BorderMode(cv::BORDER_REFLECT), BorderMode(cv::BORDER_WRAP)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// Threshold
|
||||
|
||||
CV_ENUM(ThreshOp, cv::THRESH_BINARY, cv::THRESH_BINARY_INV, cv::THRESH_TRUNC, cv::THRESH_TOZERO, cv::THRESH_TOZERO_INV) |
||||
#define ALL_THRESH_OPS testing::Values(ThreshOp(cv::THRESH_BINARY), ThreshOp(cv::THRESH_BINARY_INV), ThreshOp(cv::THRESH_TRUNC), ThreshOp(cv::THRESH_TOZERO), ThreshOp(cv::THRESH_TOZERO_INV)) |
||||
|
||||
GPU_PERF_TEST(Threshold, cv::gpu::DeviceInfo, cv::Size, MatDepth, ThreshOp) |
||||
{ |
||||
cv::Size size = GET_PARAM(1); |
||||
int depth = GET_PARAM(2); |
||||
int threshOp = GET_PARAM(3); |
||||
|
||||
cv::Mat src(size, depth); |
||||
fill(src, 0, 255); |
||||
|
||||
cv::Mat dst; |
||||
|
||||
cv::threshold(src, dst, 100.0, 255.0, threshOp); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::threshold(src, dst, 100.0, 255.0, threshOp); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(ImgProc, Threshold, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(MatDepth(CV_8U), MatDepth(CV_16U), MatDepth(CV_32F), MatDepth(CV_64F)), |
||||
ALL_THRESH_OPS)); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// Integral
|
||||
|
||||
GPU_PERF_TEST(Integral, cv::gpu::DeviceInfo, cv::Size) |
||||
{ |
||||
cv::Size size = GET_PARAM(1); |
||||
|
||||
cv::Mat src(size, CV_8UC1); |
||||
fill(src, 0, 255); |
||||
|
||||
cv::Mat dst; |
||||
|
||||
cv::integral(src, dst); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::integral(src, dst); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(ImgProc, Integral, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES)); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// HistEven_OneChannel
|
||||
|
||||
GPU_PERF_TEST(HistEven_OneChannel, cv::gpu::DeviceInfo, cv::Size, MatDepth) |
||||
{ |
||||
cv::Size size = GET_PARAM(1); |
||||
int depth = GET_PARAM(2); |
||||
|
||||
cv::Mat src(size, depth); |
||||
fill(src, 0, 255); |
||||
|
||||
int hbins = 30; |
||||
float hranges[] = {0.0f, 180.0f}; |
||||
cv::Mat hist; |
||||
int histSize[] = {hbins}; |
||||
const float* ranges[] = {hranges}; |
||||
int channels[] = {0}; |
||||
|
||||
cv::calcHist(&src, 1, channels, cv::Mat(), hist, 1, histSize, ranges); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::calcHist(&src, 1, channels, cv::Mat(), hist, 1, histSize, ranges); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(ImgProc, HistEven_OneChannel, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(MatDepth(CV_8U), MatDepth(CV_16U), MatDepth(CV_16S)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// EqualizeHist
|
||||
|
||||
GPU_PERF_TEST(EqualizeHist, cv::gpu::DeviceInfo, cv::Size) |
||||
{ |
||||
cv::Size size = GET_PARAM(1); |
||||
|
||||
cv::Mat src(size, CV_8UC1); |
||||
fill(src, 0, 255); |
||||
|
||||
cv::Mat dst; |
||||
|
||||
cv::equalizeHist(src, dst); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::equalizeHist(src, dst); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(ImgProc, EqualizeHist, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES)); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// Canny
|
||||
|
||||
IMPLEMENT_PARAM_CLASS(AppertureSize, int) |
||||
IMPLEMENT_PARAM_CLASS(L2gradient, bool) |
||||
|
||||
GPU_PERF_TEST(Canny, cv::gpu::DeviceInfo, AppertureSize, L2gradient) |
||||
{ |
||||
int apperture_size = GET_PARAM(1); |
||||
bool useL2gradient = GET_PARAM(2); |
||||
|
||||
cv::Mat image = readImage("perf/1280x1024.jpg", cv::IMREAD_GRAYSCALE); |
||||
ASSERT_FALSE(image.empty()); |
||||
|
||||
cv::Mat dst; |
||||
|
||||
cv::Canny(image, dst, 50.0, 100.0, apperture_size, useL2gradient); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::Canny(image, dst, 50.0, 100.0, apperture_size, useL2gradient); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(ImgProc, Canny, testing::Combine( |
||||
ALL_DEVICES, |
||||
testing::Values(AppertureSize(3), AppertureSize(5)), |
||||
testing::Values(L2gradient(false), L2gradient(true)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// MeanShiftFiltering
|
||||
|
||||
GPU_PERF_TEST_1(MeanShiftFiltering, cv::gpu::DeviceInfo) |
||||
{ |
||||
cv::Mat img = readImage("gpu/meanshift/cones.png"); |
||||
ASSERT_FALSE(img.empty()); |
||||
|
||||
cv::Mat dst; |
||||
|
||||
cv::pyrMeanShiftFiltering(img, dst, 50, 50); |
||||
|
||||
declare.time(15.0); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::pyrMeanShiftFiltering(img, dst, 50, 50); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(ImgProc, MeanShiftFiltering, ALL_DEVICES); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// Convolve
|
||||
|
||||
IMPLEMENT_PARAM_CLASS(KSize, int) |
||||
IMPLEMENT_PARAM_CLASS(Ccorr, bool) |
||||
|
||||
GPU_PERF_TEST(Convolve, cv::gpu::DeviceInfo, cv::Size, KSize, Ccorr) |
||||
{ |
||||
cv::Size size = GET_PARAM(1); |
||||
int templ_size = GET_PARAM(2); |
||||
bool ccorr = GET_PARAM(3); |
||||
|
||||
ASSERT_FALSE(ccorr); |
||||
|
||||
cv::Mat image(size, CV_32FC1); |
||||
image.setTo(1.0); |
||||
|
||||
cv::Mat templ(templ_size, templ_size, CV_32FC1); |
||||
templ.setTo(1.0); |
||||
|
||||
cv::Mat dst; |
||||
|
||||
cv::filter2D(image, dst, image.depth(), templ); |
||||
|
||||
declare.time(10.0); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::filter2D(image, dst, image.depth(), templ); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(ImgProc, Convolve, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(KSize(3), KSize(9), KSize(17), KSize(27), KSize(32), KSize(64)), |
||||
testing::Values(Ccorr(false), Ccorr(true)))); |
||||
|
||||
////////////////////////////////////////////////////////////////////////////////
|
||||
// MatchTemplate_8U
|
||||
|
||||
CV_ENUM(TemplateMethod, cv::TM_SQDIFF, cv::TM_SQDIFF_NORMED, cv::TM_CCORR, cv::TM_CCORR_NORMED, cv::TM_CCOEFF, cv::TM_CCOEFF_NORMED) |
||||
#define ALL_TEMPLATE_METHODS testing::Values(TemplateMethod(cv::TM_SQDIFF), TemplateMethod(cv::TM_SQDIFF_NORMED), TemplateMethod(cv::TM_CCORR), TemplateMethod(cv::TM_CCORR_NORMED), TemplateMethod(cv::TM_CCOEFF), TemplateMethod(cv::TM_CCOEFF_NORMED)) |
||||
|
||||
IMPLEMENT_PARAM_CLASS(TemplateSize, cv::Size) |
||||
|
||||
GPU_PERF_TEST(MatchTemplate_8U, cv::gpu::DeviceInfo, cv::Size, TemplateSize, Channels, TemplateMethod) |
||||
{ |
||||
cv::Size size = GET_PARAM(1); |
||||
cv::Size templ_size = GET_PARAM(2); |
||||
int cn = GET_PARAM(3); |
||||
int method = GET_PARAM(4); |
||||
|
||||
cv::Mat image(size, CV_MAKE_TYPE(CV_8U, cn)); |
||||
fill(image, 0, 255); |
||||
|
||||
cv::Mat templ(templ_size, CV_MAKE_TYPE(CV_8U, cn)); |
||||
fill(templ, 0, 255); |
||||
|
||||
cv::Mat dst; |
||||
|
||||
cv::matchTemplate(image, templ, dst, method); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::matchTemplate(image, templ, dst, method); |
||||
} |
||||
}; |
||||
|
||||
INSTANTIATE_TEST_CASE_P(ImgProc, MatchTemplate_8U, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(TemplateSize(cv::Size(5, 5)), TemplateSize(cv::Size(16, 16)), TemplateSize(cv::Size(30, 30))), |
||||
testing::Values(Channels(1), Channels(3), Channels(4)), |
||||
ALL_TEMPLATE_METHODS)); |
||||
|
||||
////////////////////////////////////////////////////////////////////////////////
|
||||
// MatchTemplate_32F
|
||||
|
||||
GPU_PERF_TEST(MatchTemplate_32F, cv::gpu::DeviceInfo, cv::Size, TemplateSize, Channels, TemplateMethod) |
||||
{ |
||||
cv::Size size = GET_PARAM(1); |
||||
cv::Size templ_size = GET_PARAM(2); |
||||
int cn = GET_PARAM(3); |
||||
int method = GET_PARAM(4); |
||||
|
||||
cv::Mat image(size, CV_MAKE_TYPE(CV_32F, cn)); |
||||
fill(image, 0, 255); |
||||
|
||||
cv::Mat templ(templ_size, CV_MAKE_TYPE(CV_32F, cn)); |
||||
fill(templ, 0, 255); |
||||
|
||||
cv::Mat dst; |
||||
|
||||
cv::matchTemplate(image, templ, dst, method); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::matchTemplate(image, templ, dst, method); |
||||
} |
||||
}; |
||||
|
||||
INSTANTIATE_TEST_CASE_P(ImgProc, MatchTemplate_32F, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(TemplateSize(cv::Size(5, 5)), TemplateSize(cv::Size(16, 16)), TemplateSize(cv::Size(30, 30))), |
||||
testing::Values(Channels(1), Channels(3), Channels(4)), |
||||
testing::Values(TemplateMethod(cv::TM_SQDIFF), TemplateMethod(cv::TM_CCORR)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// MulSpectrums
|
||||
|
||||
CV_FLAGS(DftFlags, 0, cv::DFT_INVERSE, cv::DFT_SCALE, cv::DFT_ROWS, cv::DFT_COMPLEX_OUTPUT, cv::DFT_REAL_OUTPUT) |
||||
|
||||
GPU_PERF_TEST(MulSpectrums, cv::gpu::DeviceInfo, cv::Size, DftFlags) |
||||
{ |
||||
cv::Size size = GET_PARAM(1); |
||||
int flag = GET_PARAM(2); |
||||
|
||||
cv::Mat a(size, CV_32FC2); |
||||
fill(a, 0, 100); |
||||
|
||||
cv::Mat b(size, CV_32FC2); |
||||
fill(b, 0, 100); |
||||
|
||||
cv::Mat dst; |
||||
|
||||
cv::mulSpectrums(a, b, dst, flag); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::mulSpectrums(a, b, dst, flag); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(ImgProc, MulSpectrums, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(DftFlags(0), DftFlags(cv::DFT_ROWS)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// Dft
|
||||
|
||||
GPU_PERF_TEST(Dft, cv::gpu::DeviceInfo, cv::Size, DftFlags) |
||||
{ |
||||
cv::Size size = GET_PARAM(1); |
||||
int flag = GET_PARAM(2); |
||||
|
||||
cv::Mat src(size, CV_32FC2); |
||||
fill(src, 0, 100); |
||||
|
||||
cv::Mat dst; |
||||
|
||||
cv::dft(src, dst, flag); |
||||
|
||||
declare.time(10.0); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::dft(src, dst, flag); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(ImgProc, Dft, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(DftFlags(0), DftFlags(cv::DFT_ROWS), DftFlags(cv::DFT_INVERSE)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// CornerHarris
|
||||
|
||||
IMPLEMENT_PARAM_CLASS(BlockSize, int) |
||||
IMPLEMENT_PARAM_CLASS(ApertureSize, int) |
||||
|
||||
GPU_PERF_TEST(CornerHarris, cv::gpu::DeviceInfo, MatType, BorderMode, BlockSize, ApertureSize) |
||||
{ |
||||
int type = GET_PARAM(1); |
||||
int borderType = GET_PARAM(2); |
||||
int blockSize = GET_PARAM(3); |
||||
int apertureSize = GET_PARAM(4); |
||||
|
||||
cv::Mat img = readImage("gpu/stereobm/aloe-L.png", cv::IMREAD_GRAYSCALE); |
||||
ASSERT_FALSE(img.empty()); |
||||
|
||||
img.convertTo(img, type, type == CV_32F ? 1.0 / 255.0 : 1.0); |
||||
|
||||
cv::Mat dst; |
||||
|
||||
double k = 0.5; |
||||
|
||||
cv::cornerHarris(img, dst, blockSize, apertureSize, k, borderType); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::cornerHarris(img, dst, blockSize, apertureSize, k, borderType); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(ImgProc, CornerHarris, testing::Combine( |
||||
ALL_DEVICES, |
||||
testing::Values(MatType(CV_8UC1), MatType(CV_32FC1)), |
||||
testing::Values(BorderMode(cv::BORDER_REFLECT101), BorderMode(cv::BORDER_REPLICATE), BorderMode(cv::BORDER_REFLECT)), |
||||
testing::Values(BlockSize(3), BlockSize(5), BlockSize(7)), |
||||
testing::Values(ApertureSize(0), ApertureSize(3), ApertureSize(5), ApertureSize(7)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// CornerMinEigenVal
|
||||
|
||||
GPU_PERF_TEST(CornerMinEigenVal, cv::gpu::DeviceInfo, MatType, BorderMode, BlockSize, ApertureSize) |
||||
{ |
||||
int type = GET_PARAM(1); |
||||
int borderType = GET_PARAM(2); |
||||
int blockSize = GET_PARAM(3); |
||||
int apertureSize = GET_PARAM(4); |
||||
|
||||
cv::Mat img = readImage("gpu/stereobm/aloe-L.png", cv::IMREAD_GRAYSCALE); |
||||
ASSERT_FALSE(img.empty()); |
||||
|
||||
img.convertTo(img, type, type == CV_32F ? 1.0 / 255.0 : 1.0); |
||||
|
||||
cv::Mat dst; |
||||
|
||||
cv::cornerMinEigenVal(img, dst, blockSize, apertureSize, borderType); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::cornerMinEigenVal(img, dst, blockSize, apertureSize, borderType); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(ImgProc, CornerMinEigenVal, testing::Combine( |
||||
ALL_DEVICES, |
||||
testing::Values(MatType(CV_8UC1), MatType(CV_32FC1)), |
||||
testing::Values(BorderMode(cv::BORDER_REFLECT101), BorderMode(cv::BORDER_REPLICATE), BorderMode(cv::BORDER_REFLECT)), |
||||
testing::Values(BlockSize(3), BlockSize(5), BlockSize(7)), |
||||
testing::Values(ApertureSize(0), ApertureSize(3), ApertureSize(5), ApertureSize(7)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// PyrDown
|
||||
|
||||
GPU_PERF_TEST(PyrDown, cv::gpu::DeviceInfo, cv::Size, MatType) |
||||
{ |
||||
cv::Size size = GET_PARAM(1); |
||||
int type = GET_PARAM(2); |
||||
|
||||
cv::Mat src(size, type); |
||||
fill(src, 0, 255); |
||||
|
||||
cv::Mat dst; |
||||
|
||||
cv::pyrDown(src, dst); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::pyrDown(src, dst); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(ImgProc, PyrDown, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(MatType(CV_8UC1), MatType(CV_8UC3), MatType(CV_8UC4), |
||||
MatType(CV_16UC1), MatType(CV_16UC3), MatType(CV_16UC4), |
||||
MatType(CV_32FC1), MatType(CV_32FC3), MatType(CV_32FC4)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// PyrUp
|
||||
|
||||
GPU_PERF_TEST(PyrUp, cv::gpu::DeviceInfo, cv::Size, MatType) |
||||
{ |
||||
cv::Size size = GET_PARAM(1); |
||||
int type = GET_PARAM(2); |
||||
|
||||
cv::Mat src(size, type); |
||||
fill(src, 0, 255); |
||||
|
||||
cv::Mat dst; |
||||
|
||||
cv::pyrUp(src, dst); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::pyrUp(src, dst); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(ImgProc, PyrUp, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(MatType(CV_8UC1), MatType(CV_8UC3), MatType(CV_8UC4), |
||||
MatType(CV_16UC1), MatType(CV_16UC3), MatType(CV_16UC4), |
||||
MatType(CV_32FC1), MatType(CV_32FC3), MatType(CV_32FC4)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// CvtColor
|
||||
|
||||
GPU_PERF_TEST(CvtColor, cv::gpu::DeviceInfo, cv::Size, MatDepth, CvtColorInfo) |
||||
{ |
||||
cv::Size size = GET_PARAM(1); |
||||
int depth = GET_PARAM(2); |
||||
CvtColorInfo info = GET_PARAM(3); |
||||
|
||||
cv::Mat src(size, CV_MAKETYPE(depth, info.scn)); |
||||
fill(src, 0, 255); |
||||
|
||||
cv::Mat dst; |
||||
|
||||
cv::cvtColor(src, dst, info.code, info.dcn); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::cvtColor(src, dst, info.code, info.dcn); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(ImgProc, CvtColor, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(MatDepth(CV_8U), MatDepth(CV_16U), MatDepth(CV_32F)), |
||||
testing::Values(CvtColorInfo(4, 4, cv::COLOR_RGBA2BGRA), |
||||
CvtColorInfo(4, 1, cv::COLOR_BGRA2GRAY), |
||||
CvtColorInfo(1, 4, cv::COLOR_GRAY2BGRA), |
||||
CvtColorInfo(3, 3, cv::COLOR_BGR2XYZ), |
||||
CvtColorInfo(3, 3, cv::COLOR_XYZ2BGR), |
||||
CvtColorInfo(3, 3, cv::COLOR_BGR2YCrCb), |
||||
CvtColorInfo(3, 3, cv::COLOR_YCrCb2BGR), |
||||
CvtColorInfo(3, 3, cv::COLOR_BGR2YUV), |
||||
CvtColorInfo(3, 3, cv::COLOR_YUV2BGR), |
||||
CvtColorInfo(3, 3, cv::COLOR_BGR2HSV), |
||||
CvtColorInfo(3, 3, cv::COLOR_HSV2BGR), |
||||
CvtColorInfo(3, 3, cv::COLOR_BGR2HLS), |
||||
CvtColorInfo(3, 3, cv::COLOR_HLS2BGR), |
||||
CvtColorInfo(3, 3, cv::COLOR_BGR2Lab), |
||||
CvtColorInfo(3, 3, cv::COLOR_RGB2Lab), |
||||
CvtColorInfo(3, 3, cv::COLOR_BGR2Luv), |
||||
CvtColorInfo(3, 3, cv::COLOR_RGB2Luv), |
||||
CvtColorInfo(3, 3, cv::COLOR_Lab2BGR), |
||||
CvtColorInfo(3, 3, cv::COLOR_Lab2RGB), |
||||
CvtColorInfo(3, 3, cv::COLOR_Luv2BGR), |
||||
CvtColorInfo(3, 3, cv::COLOR_Luv2RGB), |
||||
CvtColorInfo(1, 3, cv::COLOR_BayerBG2BGR), |
||||
CvtColorInfo(1, 3, cv::COLOR_BayerGB2BGR), |
||||
CvtColorInfo(1, 3, cv::COLOR_BayerRG2BGR), |
||||
CvtColorInfo(1, 3, cv::COLOR_BayerGR2BGR), |
||||
CvtColorInfo(4, 4, cv::COLOR_RGBA2mRGBA)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// HoughLines
|
||||
|
||||
IMPLEMENT_PARAM_CLASS(DoSort, bool) |
||||
|
||||
GPU_PERF_TEST(HoughLines, cv::gpu::DeviceInfo, cv::Size, DoSort) |
||||
{ |
||||
declare.time(30.0); |
||||
|
||||
const cv::Size size = GET_PARAM(1); |
||||
|
||||
const float rho = 1.0f; |
||||
const float theta = CV_PI / 180.0f; |
||||
const int threshold = 300; |
||||
|
||||
cv::RNG rng(123456789); |
||||
|
||||
cv::Mat src(size, CV_8UC1, cv::Scalar::all(0)); |
||||
|
||||
const int numLines = rng.uniform(500, 2000); |
||||
for (int i = 0; i < numLines; ++i) |
||||
{ |
||||
cv::Point p1(rng.uniform(0, src.cols), rng.uniform(0, src.rows)); |
||||
cv::Point p2(rng.uniform(0, src.cols), rng.uniform(0, src.rows)); |
||||
cv::line(src, p1, p2, cv::Scalar::all(255), 2); |
||||
} |
||||
|
||||
std::vector<cv::Vec2f> lines; |
||||
cv::HoughLines(src, lines, rho, theta, threshold); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::HoughLines(src, lines, rho, theta, threshold); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(ImgProc, HoughLines, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(DoSort(false), DoSort(true)))); |
||||
|
||||
#endif |
||||
@ -1,158 +0,0 @@ |
||||
/*M///////////////////////////////////////////////////////////////////////////////////////
|
||||
//
|
||||
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
|
||||
//
|
||||
// By downloading, copying, installing or using the software you agree to this license.
|
||||
// If you do not agree to this license, do not download, install,
|
||||
// copy or use the software.
|
||||
//
|
||||
//
|
||||
// License Agreement
|
||||
// For Open Source Computer Vision Library
|
||||
//
|
||||
// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
|
||||
// Copyright (C) 2008-2011, Willow Garage Inc., all rights reserved.
|
||||
// Third party copyrights are property of their respective owners.
|
||||
//
|
||||
// Redistribution and use in source and binary forms, with or without modification,
|
||||
// are permitted provided that the following conditions are met:
|
||||
//
|
||||
// * Redistributions of source code must retain the above copyright notice,
|
||||
// this list of conditions and the following disclaimer.
|
||||
//
|
||||
// * Redistributions in binary form must reproduce the above copyright notice,
|
||||
// this list of conditions and the following disclaimer in the documentation
|
||||
// and/or other materials provided with the distribution.
|
||||
//
|
||||
// * The name of the copyright holders may not be used to endorse or promote products
|
||||
// derived from this software without specific prior written permission.
|
||||
//
|
||||
// This software is provided by the copyright holders and contributors "as is" and
|
||||
// any express or implied warranties, including, but not limited to, the implied
|
||||
// warranties of merchantability and fitness for a particular purpose are disclaimed.
|
||||
// In no event shall the Intel Corporation or contributors be liable for any direct,
|
||||
// indirect, incidental, special, exemplary, or consequential damages
|
||||
// (including, but not limited to, procurement of substitute goods or services;
|
||||
// loss of use, data, or profits; or business interruption) however caused
|
||||
// and on any theory of liability, whether in contract, strict liability,
|
||||
// or tort (including negligence or otherwise) arising in any way out of
|
||||
// the use of this software, even if advised of the possibility of such damage.
|
||||
//M*/
|
||||
|
||||
#include "perf_precomp.hpp" |
||||
|
||||
#ifdef HAVE_CUDA |
||||
|
||||
namespace { |
||||
|
||||
struct GreedyLabeling |
||||
{ |
||||
struct dot |
||||
{ |
||||
int x; |
||||
int y; |
||||
|
||||
static dot make(int i, int j) |
||||
{ |
||||
dot d; d.x = i; d.y = j; |
||||
return d; |
||||
} |
||||
}; |
||||
|
||||
struct InInterval |
||||
{ |
||||
InInterval(const int& _lo, const int& _hi) : lo(-_lo), hi(_hi) {}; |
||||
const int lo, hi; |
||||
|
||||
bool operator() (const unsigned char a, const unsigned char b) const |
||||
{ |
||||
int d = a - b; |
||||
return lo <= d && d <= hi; |
||||
} |
||||
}; |
||||
|
||||
GreedyLabeling(cv::Mat img) |
||||
: image(img), _labels(image.size(), CV_32SC1, cv::Scalar::all(-1)) {stack = new dot[image.cols * image.rows];} |
||||
|
||||
~GreedyLabeling(){delete[] stack;} |
||||
|
||||
void operator() (cv::Mat labels) const |
||||
{ |
||||
labels.setTo(cv::Scalar::all(-1)); |
||||
InInterval inInt(0, 2); |
||||
int cc = -1; |
||||
|
||||
int* dist_labels = (int*)labels.data; |
||||
int pitch = labels.step1(); |
||||
|
||||
unsigned char* source = (unsigned char*)image.data; |
||||
int width = image.cols; |
||||
int height = image.rows; |
||||
|
||||
for (int j = 0; j < image.rows; ++j) |
||||
for (int i = 0; i < image.cols; ++i) |
||||
{ |
||||
if (dist_labels[j * pitch + i] != -1) continue; |
||||
|
||||
dot* top = stack; |
||||
dot p = dot::make(i, j); |
||||
cc++; |
||||
|
||||
dist_labels[j * pitch + i] = cc; |
||||
|
||||
while (top >= stack) |
||||
{ |
||||
int* dl = &dist_labels[p.y * pitch + p.x]; |
||||
unsigned char* sp = &source[p.y * image.step1() + p.x]; |
||||
|
||||
dl[0] = cc; |
||||
|
||||
//right
|
||||
if( p.x < (width - 1) && dl[ +1] == -1 && inInt(sp[0], sp[+1])) |
||||
*top++ = dot::make(p.x + 1, p.y); |
||||
|
||||
//left
|
||||
if( p.x > 0 && dl[-1] == -1 && inInt(sp[0], sp[-1])) |
||||
*top++ = dot::make(p.x - 1, p.y); |
||||
|
||||
//bottom
|
||||
if( p.y < (height - 1) && dl[+pitch] == -1 && inInt(sp[0], sp[+image.step1()])) |
||||
*top++ = dot::make(p.x, p.y + 1); |
||||
|
||||
//top
|
||||
if( p.y > 0 && dl[-pitch] == -1 && inInt(sp[0], sp[-image.step1()])) |
||||
*top++ = dot::make(p.x, p.y - 1); |
||||
|
||||
p = *--top; |
||||
} |
||||
} |
||||
} |
||||
|
||||
cv::Mat image; |
||||
cv::Mat _labels; |
||||
dot* stack; |
||||
}; |
||||
} |
||||
|
||||
GPU_PERF_TEST(ConnectedComponents, cv::gpu::DeviceInfo, cv::Size) |
||||
{ |
||||
cv::gpu::DeviceInfo devInfo = GET_PARAM(0); |
||||
cv::gpu::setDevice(devInfo.deviceID()); |
||||
|
||||
cv::Mat image = readImage("gpu/labeling/aloe-disp.png", cv::IMREAD_GRAYSCALE); |
||||
|
||||
GreedyLabeling host(image); |
||||
|
||||
host(host._labels); |
||||
|
||||
declare.time(1.0); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
host(host._labels); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Labeling, ConnectedComponents, testing::Combine(ALL_DEVICES, testing::Values(cv::Size(261, 262)))); |
||||
|
||||
#endif |
||||
@ -1,20 +0,0 @@ |
||||
#include "perf_cpu_precomp.hpp" |
||||
|
||||
#ifdef HAVE_CUDA |
||||
|
||||
int main(int argc, char **argv) |
||||
{ |
||||
testing::InitGoogleTest(&argc, argv); |
||||
perf::TestBase::Init(argc, argv); |
||||
return RUN_ALL_TESTS(); |
||||
} |
||||
|
||||
#else |
||||
|
||||
int main() |
||||
{ |
||||
printf("OpenCV was built without CUDA support\n"); |
||||
return 0; |
||||
} |
||||
|
||||
#endif |
||||
@ -1,124 +0,0 @@ |
||||
#include "perf_cpu_precomp.hpp" |
||||
|
||||
#ifdef HAVE_CUDA |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// SetTo
|
||||
|
||||
GPU_PERF_TEST(SetTo, cv::gpu::DeviceInfo, cv::Size, MatType) |
||||
{ |
||||
cv::Size size = GET_PARAM(1); |
||||
int type = GET_PARAM(2); |
||||
|
||||
cv::Mat src(size, type); |
||||
cv::Scalar val(1, 2, 3, 4); |
||||
|
||||
src.setTo(val); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
src.setTo(val); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(MatOp, SetTo, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(MatType(CV_8UC1), MatType(CV_8UC3), MatType(CV_8UC4), |
||||
MatType(CV_16UC1), MatType(CV_16UC3), MatType(CV_16UC4), |
||||
MatType(CV_32FC1), MatType(CV_32FC3), MatType(CV_32FC4), |
||||
MatType(CV_64FC1), MatType(CV_64FC3), MatType(CV_64FC4)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// SetToMasked
|
||||
|
||||
GPU_PERF_TEST(SetToMasked, cv::gpu::DeviceInfo, cv::Size, MatType) |
||||
{ |
||||
cv::Size size = GET_PARAM(1); |
||||
int type = GET_PARAM(2); |
||||
|
||||
cv::Mat src(size, type); |
||||
fill(src, 0, 255); |
||||
|
||||
cv::Mat mask(size, CV_8UC1); |
||||
fill(mask, 0, 2); |
||||
|
||||
cv::Scalar val(1, 2, 3, 4); |
||||
|
||||
src.setTo(val, mask); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
src.setTo(val, mask); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(MatOp, SetToMasked, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(MatType(CV_8UC1), MatType(CV_8UC3), MatType(CV_8UC4), |
||||
MatType(CV_16UC1), MatType(CV_16UC3), MatType(CV_16UC4), |
||||
MatType(CV_32FC1), MatType(CV_32FC3), MatType(CV_32FC4), |
||||
MatType(CV_64FC1), MatType(CV_64FC3), MatType(CV_64FC4)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// CopyToMasked
|
||||
|
||||
GPU_PERF_TEST(CopyToMasked, cv::gpu::DeviceInfo, cv::Size, MatType) |
||||
{ |
||||
cv::Size size = GET_PARAM(1); |
||||
int type = GET_PARAM(2); |
||||
|
||||
cv::Mat src(size, type); |
||||
fill(src, 0, 255); |
||||
|
||||
cv::Mat mask(size, CV_8UC1); |
||||
fill(mask, 0, 2); |
||||
|
||||
cv::Mat dst; |
||||
|
||||
src.copyTo(dst, mask); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
src.copyTo(dst, mask); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(MatOp, CopyToMasked, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(MatType(CV_8UC1), MatType(CV_8UC3), MatType(CV_8UC4), |
||||
MatType(CV_16UC1), MatType(CV_16UC3), MatType(CV_16UC4), |
||||
MatType(CV_32FC1), MatType(CV_32FC3), MatType(CV_32FC4), |
||||
MatType(CV_64FC1), MatType(CV_64FC3), MatType(CV_64FC4)))); |
||||
|
||||
//////////////////////////////////////////////////////////////////////
|
||||
// ConvertTo
|
||||
|
||||
GPU_PERF_TEST(ConvertTo, cv::gpu::DeviceInfo, cv::Size, MatDepth, MatDepth) |
||||
{ |
||||
cv::Size size = GET_PARAM(1); |
||||
int depth1 = GET_PARAM(2); |
||||
int depth2 = GET_PARAM(3); |
||||
|
||||
cv::Mat src(size, depth1); |
||||
fill(src, 0, 255); |
||||
|
||||
cv::Mat dst; |
||||
|
||||
src.convertTo(dst, depth2, 0.5, 1.0); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
src.convertTo(dst, depth2, 0.5, 1.0); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(MatOp, ConvertTo, testing::Combine( |
||||
ALL_DEVICES, |
||||
GPU_TYPICAL_MAT_SIZES, |
||||
testing::Values(MatDepth(CV_8U), MatDepth(CV_16U), MatDepth(CV_32F), MatDepth(CV_64F)), |
||||
testing::Values(MatDepth(CV_8U), MatDepth(CV_16U), MatDepth(CV_32F), MatDepth(CV_64F)))); |
||||
|
||||
#endif |
||||
@ -1,74 +0,0 @@ |
||||
#include "perf_cpu_precomp.hpp" |
||||
|
||||
#ifdef HAVE_CUDA |
||||
|
||||
///////////////////////////////////////////////////////////////
|
||||
// HOG
|
||||
|
||||
GPU_PERF_TEST_1(HOG, cv::gpu::DeviceInfo) |
||||
{ |
||||
cv::Mat img = readImage("gpu/hog/road.png", cv::IMREAD_GRAYSCALE); |
||||
ASSERT_FALSE(img.empty()); |
||||
|
||||
std::vector<cv::Rect> found_locations; |
||||
|
||||
cv::HOGDescriptor hog; |
||||
hog.setSVMDetector(cv::gpu::HOGDescriptor::getDefaultPeopleDetector()); |
||||
|
||||
hog.detectMultiScale(img, found_locations); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
hog.detectMultiScale(img, found_locations); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(ObjDetect, HOG, ALL_DEVICES); |
||||
|
||||
///////////////////////////////////////////////////////////////
|
||||
// HaarClassifier
|
||||
|
||||
GPU_PERF_TEST_1(HaarClassifier, cv::gpu::DeviceInfo) |
||||
{ |
||||
cv::Mat img = readImage("gpu/haarcascade/group_1_640x480_VGA.pgm", cv::IMREAD_GRAYSCALE); |
||||
ASSERT_FALSE(img.empty()); |
||||
|
||||
cv::CascadeClassifier cascade; |
||||
|
||||
ASSERT_TRUE(cascade.load(perf::TestBase::getDataPath("gpu/perf/haarcascade_frontalface_alt.xml"))); |
||||
|
||||
std::vector<cv::Rect> rects; |
||||
|
||||
cascade.detectMultiScale(img, rects); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cascade.detectMultiScale(img, rects); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(ObjDetect, HaarClassifier, ALL_DEVICES); |
||||
|
||||
//===================== LBP cascade ==========================//
|
||||
GPU_PERF_TEST_1(LBPClassifier, cv::gpu::DeviceInfo) |
||||
{ |
||||
cv::Mat img = readImage("gpu/haarcascade/group_1_640x480_VGA.pgm", cv::IMREAD_GRAYSCALE); |
||||
ASSERT_FALSE(img.empty()); |
||||
|
||||
cv::CascadeClassifier cascade; |
||||
|
||||
ASSERT_TRUE(cascade.load(perf::TestBase::getDataPath("gpu/lbpcascade/lbpcascade_frontalface.xml"))); |
||||
|
||||
std::vector<cv::Rect> rects; |
||||
|
||||
cascade.detectMultiScale(img, rects); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cascade.detectMultiScale(img, rects); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(ObjDetect, LBPClassifier, ALL_DEVICES); |
||||
|
||||
#endif |
||||
@ -1,220 +0,0 @@ |
||||
#include "perf_cpu_precomp.hpp" |
||||
|
||||
using namespace std; |
||||
using namespace cv; |
||||
using namespace cv::gpu; |
||||
|
||||
void fill(Mat& m, double a, double b) |
||||
{ |
||||
RNG rng(123456789); |
||||
rng.fill(m, RNG::UNIFORM, a, b); |
||||
} |
||||
|
||||
void PrintTo(const CvtColorInfo& info, ostream* os) |
||||
{ |
||||
static const char* str[] = |
||||
{ |
||||
"BGR2BGRA", |
||||
"BGRA2BGR", |
||||
"BGR2RGBA", |
||||
"RGBA2BGR", |
||||
"BGR2RGB", |
||||
"BGRA2RGBA", |
||||
|
||||
"BGR2GRAY", |
||||
"RGB2GRAY", |
||||
"GRAY2BGR", |
||||
"GRAY2BGRA", |
||||
"BGRA2GRAY", |
||||
"RGBA2GRAY", |
||||
|
||||
"BGR2BGR565", |
||||
"RGB2BGR565", |
||||
"BGR5652BGR", |
||||
"BGR5652RGB", |
||||
"BGRA2BGR565", |
||||
"RGBA2BGR565", |
||||
"BGR5652BGRA", |
||||
"BGR5652RGBA", |
||||
|
||||
"GRAY2BGR565", |
||||
"BGR5652GRAY", |
||||
|
||||
"BGR2BGR555", |
||||
"RGB2BGR555", |
||||
"BGR5552BGR", |
||||
"BGR5552RGB", |
||||
"BGRA2BGR555", |
||||
"RGBA2BGR555", |
||||
"BGR5552BGRA", |
||||
"BGR5552RGBA", |
||||
|
||||
"GRAY2BGR555", |
||||
"BGR5552GRAY", |
||||
|
||||
"BGR2XYZ", |
||||
"RGB2XYZ", |
||||
"XYZ2BGR", |
||||
"XYZ2RGB", |
||||
|
||||
"BGR2YCrCb", |
||||
"RGB2YCrCb", |
||||
"YCrCb2BGR", |
||||
"YCrCb2RGB", |
||||
|
||||
"BGR2HSV", |
||||
"RGB2HSV", |
||||
|
||||
"", |
||||
"", |
||||
|
||||
"BGR2Lab", |
||||
"RGB2Lab", |
||||
|
||||
"BayerBG2BGR", |
||||
"BayerGB2BGR", |
||||
"BayerRG2BGR", |
||||
"BayerGR2BGR", |
||||
|
||||
"BGR2Luv", |
||||
"RGB2Luv", |
||||
|
||||
"BGR2HLS", |
||||
"RGB2HLS", |
||||
|
||||
"HSV2BGR", |
||||
"HSV2RGB", |
||||
|
||||
"Lab2BGR", |
||||
"Lab2RGB", |
||||
"Luv2BGR", |
||||
"Luv2RGB", |
||||
|
||||
"HLS2BGR", |
||||
"HLS2RGB", |
||||
|
||||
"BayerBG2BGR_VNG", |
||||
"BayerGB2BGR_VNG", |
||||
"BayerRG2BGR_VNG", |
||||
"BayerGR2BGR_VNG", |
||||
|
||||
"BGR2HSV_FULL", |
||||
"RGB2HSV_FULL", |
||||
"BGR2HLS_FULL", |
||||
"RGB2HLS_FULL", |
||||
|
||||
"HSV2BGR_FULL", |
||||
"HSV2RGB_FULL", |
||||
"HLS2BGR_FULL", |
||||
"HLS2RGB_FULL", |
||||
|
||||
"LBGR2Lab", |
||||
"LRGB2Lab", |
||||
"LBGR2Luv", |
||||
"LRGB2Luv", |
||||
|
||||
"Lab2LBGR", |
||||
"Lab2LRGB", |
||||
"Luv2LBGR", |
||||
"Luv2LRGB", |
||||
|
||||
"BGR2YUV", |
||||
"RGB2YUV", |
||||
"YUV2BGR", |
||||
"YUV2RGB", |
||||
|
||||
"BayerBG2GRAY", |
||||
"BayerGB2GRAY", |
||||
"BayerRG2GRAY", |
||||
"BayerGR2GRAY", |
||||
|
||||
//YUV 4:2:0 formats family
|
||||
"YUV2RGB_NV12", |
||||
"YUV2BGR_NV12", |
||||
"YUV2RGB_NV21", |
||||
"YUV2BGR_NV21", |
||||
|
||||
"YUV2RGBA_NV12", |
||||
"YUV2BGRA_NV12", |
||||
"YUV2RGBA_NV21", |
||||
"YUV2BGRA_NV21", |
||||
|
||||
"YUV2RGB_YV12", |
||||
"YUV2BGR_YV12", |
||||
"YUV2RGB_IYUV", |
||||
"YUV2BGR_IYUV", |
||||
|
||||
"YUV2RGBA_YV12", |
||||
"YUV2BGRA_YV12", |
||||
"YUV2RGBA_IYUV", |
||||
"YUV2BGRA_IYUV", |
||||
|
||||
"YUV2GRAY_420", |
||||
|
||||
//YUV 4:2:2 formats family
|
||||
"YUV2RGB_UYVY", |
||||
"YUV2BGR_UYVY", |
||||
"YUV2RGB_VYUY", |
||||
"YUV2BGR_VYUY", |
||||
|
||||
"YUV2RGBA_UYVY", |
||||
"YUV2BGRA_UYVY", |
||||
"YUV2RGBA_VYUY", |
||||
"YUV2BGRA_VYUY", |
||||
|
||||
"YUV2RGB_YUY2", |
||||
"YUV2BGR_YUY2", |
||||
"YUV2RGB_YVYU", |
||||
"YUV2BGR_YVYU", |
||||
|
||||
"YUV2RGBA_YUY2", |
||||
"YUV2BGRA_YUY2", |
||||
"YUV2RGBA_YVYU", |
||||
"YUV2BGRA_YVYU", |
||||
|
||||
"YUV2GRAY_UYVY", |
||||
"YUV2GRAY_YUY2", |
||||
|
||||
// alpha premultiplication
|
||||
"RGBA2mRGBA", |
||||
"mRGBA2RGBA", |
||||
|
||||
"COLORCVT_MAX" |
||||
}; |
||||
|
||||
*os << str[info.code]; |
||||
} |
||||
|
||||
void cv::gpu::PrintTo(const DeviceInfo& info, ostream* os) |
||||
{ |
||||
*os << info.name(); |
||||
} |
||||
|
||||
Mat readImage(const string& fileName, int flags) |
||||
{ |
||||
return imread(perf::TestBase::getDataPath(fileName), flags); |
||||
} |
||||
|
||||
const vector<DeviceInfo>& devices() |
||||
{ |
||||
static vector<DeviceInfo> devs; |
||||
static bool first = true; |
||||
|
||||
if (first) |
||||
{ |
||||
int deviceCount = getCudaEnabledDeviceCount(); |
||||
|
||||
devs.reserve(deviceCount); |
||||
|
||||
for (int i = 0; i < deviceCount; ++i) |
||||
{ |
||||
DeviceInfo info(i); |
||||
if (info.isCompatible()) |
||||
devs.push_back(info); |
||||
} |
||||
|
||||
first = false; |
||||
} |
||||
|
||||
return devs; |
||||
} |
||||
@ -1,77 +0,0 @@ |
||||
#ifndef __OPENCV_PERF_GPU_UTILITY_HPP__ |
||||
#define __OPENCV_PERF_GPU_UTILITY_HPP__ |
||||
|
||||
void fill(cv::Mat& m, double a, double b); |
||||
|
||||
using perf::MatType; |
||||
using perf::MatDepth; |
||||
|
||||
CV_ENUM(BorderMode, cv::BORDER_REFLECT101, cv::BORDER_REPLICATE, cv::BORDER_CONSTANT, cv::BORDER_REFLECT, cv::BORDER_WRAP) |
||||
CV_ENUM(Interpolation, cv::INTER_NEAREST, cv::INTER_LINEAR, cv::INTER_CUBIC, cv::INTER_AREA) |
||||
CV_ENUM(NormType, cv::NORM_INF, cv::NORM_L1, cv::NORM_L2, cv::NORM_HAMMING) |
||||
|
||||
struct CvtColorInfo |
||||
{ |
||||
int scn; |
||||
int dcn; |
||||
int code; |
||||
|
||||
explicit CvtColorInfo(int scn_=0, int dcn_=0, int code_=0) : scn(scn_), dcn(dcn_), code(code_) {} |
||||
}; |
||||
|
||||
void PrintTo(const CvtColorInfo& info, std::ostream* os); |
||||
|
||||
#define IMPLEMENT_PARAM_CLASS(name, type) \ |
||||
class name \
|
||||
{ \
|
||||
public: \
|
||||
name ( type arg = type ()) : val_(arg) {} \
|
||||
operator type () const {return val_;} \
|
||||
private: \
|
||||
type val_; \
|
||||
}; \
|
||||
inline void PrintTo( name param, std::ostream* os) \
|
||||
{ \
|
||||
*os << #name << " = " << testing::PrintToString(static_cast< type >(param)); \
|
||||
} |
||||
|
||||
IMPLEMENT_PARAM_CLASS(Channels, int) |
||||
|
||||
namespace cv { namespace gpu |
||||
{ |
||||
void PrintTo(const cv::gpu::DeviceInfo& info, std::ostream* os); |
||||
}} |
||||
|
||||
#define GPU_PERF_TEST(name, ...) \ |
||||
struct name : perf::TestBaseWithParam< std::tr1::tuple< __VA_ARGS__ > > \
|
||||
{ \
|
||||
public: \
|
||||
name() {} \
|
||||
protected: \
|
||||
void PerfTestBody(); \
|
||||
}; \
|
||||
TEST_P(name, perf){ RunPerfTestBody(); } \
|
||||
void name :: PerfTestBody() |
||||
|
||||
#define GPU_PERF_TEST_1(name, param_type) \ |
||||
struct name : perf::TestBaseWithParam< param_type > \
|
||||
{ \
|
||||
public: \
|
||||
name() {} \
|
||||
protected: \
|
||||
void PerfTestBody(); \
|
||||
}; \
|
||||
TEST_P(name, perf){ RunPerfTestBody(); } \
|
||||
void name :: PerfTestBody() |
||||
|
||||
#define GPU_TYPICAL_MAT_SIZES testing::Values(perf::szSXGA, perf::sz1080p, cv::Size(1800, 1500)) |
||||
|
||||
cv::Mat readImage(const std::string& fileName, int flags = cv::IMREAD_COLOR); |
||||
|
||||
const std::vector<cv::gpu::DeviceInfo>& devices(); |
||||
|
||||
#define ALL_DEVICES testing::ValuesIn(devices()) |
||||
|
||||
#define GET_PARAM(k) std::tr1::get< k >(GetParam()) |
||||
|
||||
#endif // __OPENCV_PERF_GPU_UTILITY_HPP__
|
||||
@ -1,466 +0,0 @@ |
||||
#include "perf_cpu_precomp.hpp" |
||||
|
||||
#ifdef HAVE_CUDA |
||||
|
||||
//////////////////////////////////////////////////////
|
||||
// GoodFeaturesToTrack
|
||||
|
||||
IMPLEMENT_PARAM_CLASS(MinDistance, double) |
||||
|
||||
GPU_PERF_TEST(GoodFeaturesToTrack, cv::gpu::DeviceInfo, MinDistance) |
||||
{ |
||||
double minDistance = GET_PARAM(1); |
||||
|
||||
cv::Mat image = readImage("gpu/perf/aloe.jpg", cv::IMREAD_GRAYSCALE); |
||||
ASSERT_FALSE(image.empty()); |
||||
|
||||
cv::Mat corners; |
||||
|
||||
cv::goodFeaturesToTrack(image, corners, 8000, 0.01, minDistance); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::goodFeaturesToTrack(image, corners, 8000, 0.01, minDistance); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Video, GoodFeaturesToTrack, testing::Combine( |
||||
ALL_DEVICES, |
||||
testing::Values(MinDistance(0.0), MinDistance(3.0)))); |
||||
|
||||
//////////////////////////////////////////////////////
|
||||
// PyrLKOpticalFlowSparse
|
||||
|
||||
IMPLEMENT_PARAM_CLASS(GraySource, bool) |
||||
IMPLEMENT_PARAM_CLASS(Points, int) |
||||
IMPLEMENT_PARAM_CLASS(WinSize, int) |
||||
IMPLEMENT_PARAM_CLASS(Levels, int) |
||||
IMPLEMENT_PARAM_CLASS(Iters, int) |
||||
|
||||
GPU_PERF_TEST(PyrLKOpticalFlowSparse, cv::gpu::DeviceInfo, GraySource, Points, WinSize, Levels, Iters) |
||||
{ |
||||
bool useGray = GET_PARAM(1); |
||||
int points = GET_PARAM(2); |
||||
int win_size = GET_PARAM(3); |
||||
int levels = GET_PARAM(4); |
||||
int iters = GET_PARAM(5); |
||||
|
||||
cv::Mat frame0 = readImage("gpu/opticalflow/frame0.png", useGray ? cv::IMREAD_GRAYSCALE : cv::IMREAD_COLOR); |
||||
ASSERT_FALSE(frame0.empty()); |
||||
|
||||
cv::Mat frame1 = readImage("gpu/opticalflow/frame1.png", useGray ? cv::IMREAD_GRAYSCALE : cv::IMREAD_COLOR); |
||||
ASSERT_FALSE(frame1.empty()); |
||||
|
||||
cv::Mat gray_frame; |
||||
if (useGray) |
||||
gray_frame = frame0; |
||||
else |
||||
cv::cvtColor(frame0, gray_frame, cv::COLOR_BGR2GRAY); |
||||
|
||||
cv::Mat pts; |
||||
cv::goodFeaturesToTrack(gray_frame, pts, points, 0.01, 0.0); |
||||
|
||||
cv::Mat nextPts; |
||||
cv::Mat status; |
||||
|
||||
cv::calcOpticalFlowPyrLK(frame0, frame1, pts, nextPts, status, cv::noArray(), |
||||
cv::Size(win_size, win_size), levels - 1, |
||||
cv::TermCriteria(cv::TermCriteria::COUNT + cv::TermCriteria::EPS, iters, 0.01)); |
||||
|
||||
declare.time(20.0); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::calcOpticalFlowPyrLK(frame0, frame1, pts, nextPts, status, cv::noArray(), |
||||
cv::Size(win_size, win_size), levels - 1, |
||||
cv::TermCriteria(cv::TermCriteria::COUNT + cv::TermCriteria::EPS, iters, 0.01)); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Video, PyrLKOpticalFlowSparse, testing::Combine( |
||||
ALL_DEVICES, |
||||
testing::Values(GraySource(true), GraySource(false)), |
||||
testing::Values(Points(1000), Points(2000), Points(4000), Points(8000)), |
||||
testing::Values(WinSize(9), WinSize(13), WinSize(17), WinSize(21)), |
||||
testing::Values(Levels(1), Levels(2), Levels(3)), |
||||
testing::Values(Iters(1), Iters(10), Iters(30)))); |
||||
|
||||
//////////////////////////////////////////////////////
|
||||
// FarnebackOpticalFlowTest
|
||||
|
||||
GPU_PERF_TEST_1(FarnebackOpticalFlowTest, cv::gpu::DeviceInfo) |
||||
{ |
||||
cv::Mat frame0 = readImage("gpu/opticalflow/frame0.png", cv::IMREAD_GRAYSCALE); |
||||
ASSERT_FALSE(frame0.empty()); |
||||
|
||||
cv::Mat frame1 = readImage("gpu/opticalflow/frame1.png", cv::IMREAD_GRAYSCALE); |
||||
ASSERT_FALSE(frame1.empty()); |
||||
|
||||
cv::Mat flow; |
||||
|
||||
int numLevels = 5; |
||||
double pyrScale = 0.5; |
||||
int winSize = 13; |
||||
int numIters = 10; |
||||
int polyN = 5; |
||||
double polySigma = 1.1; |
||||
int flags = 0; |
||||
|
||||
cv::calcOpticalFlowFarneback(frame0, frame1, flow, pyrScale, numLevels, winSize, numIters, polyN, polySigma, flags); |
||||
|
||||
declare.time(10); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
cv::calcOpticalFlowFarneback(frame0, frame1, flow, pyrScale, numLevels, winSize, numIters, polyN, polySigma, flags); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Video, FarnebackOpticalFlowTest, ALL_DEVICES); |
||||
|
||||
//////////////////////////////////////////////////////
|
||||
// FGDStatModel
|
||||
|
||||
namespace cv |
||||
{ |
||||
template<> void Ptr<CvBGStatModel>::delete_obj() |
||||
{ |
||||
cvReleaseBGStatModel(&obj); |
||||
} |
||||
} |
||||
|
||||
GPU_PERF_TEST(FGDStatModel, cv::gpu::DeviceInfo, std::string) |
||||
{ |
||||
std::string inputFile = perf::TestBase::getDataPath(std::string("gpu/video/") + GET_PARAM(1)); |
||||
|
||||
cv::VideoCapture cap(inputFile); |
||||
ASSERT_TRUE(cap.isOpened()); |
||||
|
||||
cv::Mat frame; |
||||
cap >> frame; |
||||
ASSERT_FALSE(frame.empty()); |
||||
|
||||
IplImage ipl_frame = frame; |
||||
cv::Ptr<CvBGStatModel> model(cvCreateFGDStatModel(&ipl_frame)); |
||||
|
||||
declare.time(60); |
||||
|
||||
for (int i = 0; i < 10; ++i) |
||||
{ |
||||
cap >> frame; |
||||
ASSERT_FALSE(frame.empty()); |
||||
|
||||
ipl_frame = frame; |
||||
|
||||
startTimer(); |
||||
next(); |
||||
|
||||
cvUpdateBGStatModel(&ipl_frame, model); |
||||
|
||||
stopTimer(); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Video, FGDStatModel, testing::Combine( |
||||
ALL_DEVICES, |
||||
testing::Values(std::string("768x576.avi"), std::string("1920x1080.avi")))); |
||||
|
||||
//////////////////////////////////////////////////////
|
||||
// MOG
|
||||
|
||||
IMPLEMENT_PARAM_CLASS(LearningRate, double) |
||||
|
||||
GPU_PERF_TEST(MOG, cv::gpu::DeviceInfo, std::string, Channels, LearningRate) |
||||
{ |
||||
std::string inputFile = perf::TestBase::getDataPath(std::string("gpu/video/") + GET_PARAM(1)); |
||||
int cn = GET_PARAM(2); |
||||
double learningRate = GET_PARAM(3); |
||||
|
||||
cv::VideoCapture cap(inputFile); |
||||
ASSERT_TRUE(cap.isOpened()); |
||||
|
||||
cv::Mat frame; |
||||
|
||||
cv::BackgroundSubtractorMOG mog; |
||||
cv::Mat foreground; |
||||
|
||||
cap >> frame; |
||||
ASSERT_FALSE(frame.empty()); |
||||
|
||||
if (cn != 3) |
||||
{ |
||||
cv::Mat temp; |
||||
if (cn == 1) |
||||
cv::cvtColor(frame, temp, cv::COLOR_BGR2GRAY); |
||||
else |
||||
cv::cvtColor(frame, temp, cv::COLOR_BGR2BGRA); |
||||
cv::swap(temp, frame); |
||||
} |
||||
|
||||
mog(frame, foreground, learningRate); |
||||
|
||||
for (int i = 0; i < 10; ++i) |
||||
{ |
||||
cap >> frame; |
||||
ASSERT_FALSE(frame.empty()); |
||||
|
||||
if (cn != 3) |
||||
{ |
||||
cv::Mat temp; |
||||
if (cn == 1) |
||||
cv::cvtColor(frame, temp, cv::COLOR_BGR2GRAY); |
||||
else |
||||
cv::cvtColor(frame, temp, cv::COLOR_BGR2BGRA); |
||||
cv::swap(temp, frame); |
||||
} |
||||
|
||||
startTimer(); next(); |
||||
mog(frame, foreground, learningRate); |
||||
stopTimer(); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Video, MOG, testing::Combine( |
||||
ALL_DEVICES, |
||||
testing::Values(std::string("768x576.avi"), std::string("1920x1080.avi")), |
||||
testing::Values(Channels(1), Channels(3)/*, Channels(4)*/), |
||||
testing::Values(LearningRate(0.0), LearningRate(0.01)))); |
||||
|
||||
//////////////////////////////////////////////////////
|
||||
// MOG2
|
||||
|
||||
GPU_PERF_TEST(MOG2_update, cv::gpu::DeviceInfo, std::string, Channels) |
||||
{ |
||||
std::string inputFile = perf::TestBase::getDataPath(std::string("gpu/video/") + GET_PARAM(1)); |
||||
int cn = GET_PARAM(2); |
||||
|
||||
cv::VideoCapture cap(inputFile); |
||||
ASSERT_TRUE(cap.isOpened()); |
||||
|
||||
cv::Mat frame; |
||||
|
||||
cv::BackgroundSubtractorMOG2 mog2; |
||||
cv::Mat foreground; |
||||
|
||||
cap >> frame; |
||||
ASSERT_FALSE(frame.empty()); |
||||
|
||||
if (cn != 3) |
||||
{ |
||||
cv::Mat temp; |
||||
if (cn == 1) |
||||
cv::cvtColor(frame, temp, cv::COLOR_BGR2GRAY); |
||||
else |
||||
cv::cvtColor(frame, temp, cv::COLOR_BGR2BGRA); |
||||
cv::swap(temp, frame); |
||||
} |
||||
|
||||
mog2(frame, foreground); |
||||
|
||||
for (int i = 0; i < 10; ++i) |
||||
{ |
||||
cap >> frame; |
||||
ASSERT_FALSE(frame.empty()); |
||||
|
||||
if (cn != 3) |
||||
{ |
||||
cv::Mat temp; |
||||
if (cn == 1) |
||||
cv::cvtColor(frame, temp, cv::COLOR_BGR2GRAY); |
||||
else |
||||
cv::cvtColor(frame, temp, cv::COLOR_BGR2BGRA); |
||||
cv::swap(temp, frame); |
||||
} |
||||
|
||||
startTimer(); next(); |
||||
mog2(frame, foreground); |
||||
stopTimer(); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Video, MOG2_update, testing::Combine( |
||||
ALL_DEVICES, |
||||
testing::Values(std::string("768x576.avi"), std::string("1920x1080.avi")), |
||||
testing::Values(Channels(1), Channels(3)/*, Channels(4)*/))); |
||||
|
||||
GPU_PERF_TEST(MOG2_getBackgroundImage, cv::gpu::DeviceInfo, std::string, Channels) |
||||
{ |
||||
std::string inputFile = perf::TestBase::getDataPath(std::string("gpu/video/") + GET_PARAM(1)); |
||||
int cn = GET_PARAM(2); |
||||
|
||||
cv::VideoCapture cap(inputFile); |
||||
ASSERT_TRUE(cap.isOpened()); |
||||
|
||||
cv::Mat frame; |
||||
|
||||
cv::BackgroundSubtractorMOG2 mog2; |
||||
cv::Mat foreground; |
||||
|
||||
for (int i = 0; i < 10; ++i) |
||||
{ |
||||
cap >> frame; |
||||
ASSERT_FALSE(frame.empty()); |
||||
|
||||
if (cn != 3) |
||||
{ |
||||
cv::Mat temp; |
||||
if (cn == 1) |
||||
cv::cvtColor(frame, temp, cv::COLOR_BGR2GRAY); |
||||
else |
||||
cv::cvtColor(frame, temp, cv::COLOR_BGR2BGRA); |
||||
cv::swap(temp, frame); |
||||
} |
||||
|
||||
mog2(frame, foreground); |
||||
} |
||||
|
||||
cv::Mat background; |
||||
mog2.getBackgroundImage(background); |
||||
|
||||
TEST_CYCLE() |
||||
{ |
||||
mog2.getBackgroundImage(background); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Video, MOG2_getBackgroundImage, testing::Combine( |
||||
ALL_DEVICES, |
||||
testing::Values(std::string("768x576.avi"), std::string("1920x1080.avi")), |
||||
testing::Values(/*Channels(1),*/ Channels(3)/*, Channels(4)*/))); |
||||
|
||||
//////////////////////////////////////////////////////
|
||||
// GMG
|
||||
|
||||
IMPLEMENT_PARAM_CLASS(MaxFeatures, int) |
||||
|
||||
GPU_PERF_TEST(GMG, cv::gpu::DeviceInfo, std::string, Channels, MaxFeatures) |
||||
{ |
||||
std::string inputFile = perf::TestBase::getDataPath(std::string("gpu/video/") + GET_PARAM(1)); |
||||
int cn = GET_PARAM(2); |
||||
int maxFeatures = GET_PARAM(3); |
||||
|
||||
cv::VideoCapture cap(inputFile); |
||||
ASSERT_TRUE(cap.isOpened()); |
||||
|
||||
cv::Mat frame; |
||||
cap >> frame; |
||||
ASSERT_FALSE(frame.empty()); |
||||
|
||||
if (cn != 3) |
||||
{ |
||||
cv::Mat temp; |
||||
if (cn == 1) |
||||
cv::cvtColor(frame, temp, cv::COLOR_BGR2GRAY); |
||||
else |
||||
cv::cvtColor(frame, temp, cv::COLOR_BGR2BGRA); |
||||
cv::swap(temp, frame); |
||||
} |
||||
|
||||
cv::Mat fgmask; |
||||
cv::Mat zeros(frame.size(), CV_8UC1, cv::Scalar::all(0)); |
||||
|
||||
cv::BackgroundSubtractorGMG gmg; |
||||
gmg.set("maxFeatures", maxFeatures); |
||||
gmg.initialize(frame.size(), 0.0, 255.0); |
||||
|
||||
gmg(frame, fgmask); |
||||
|
||||
for (int i = 0; i < 150; ++i) |
||||
{ |
||||
cap >> frame; |
||||
if (frame.empty()) |
||||
{ |
||||
cap.open(inputFile); |
||||
cap >> frame; |
||||
} |
||||
|
||||
if (cn != 3) |
||||
{ |
||||
cv::Mat temp; |
||||
if (cn == 1) |
||||
cv::cvtColor(frame, temp, cv::COLOR_BGR2GRAY); |
||||
else |
||||
cv::cvtColor(frame, temp, cv::COLOR_BGR2BGRA); |
||||
cv::swap(temp, frame); |
||||
} |
||||
|
||||
startTimer(); next(); |
||||
gmg(frame, fgmask); |
||||
stopTimer(); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Video, GMG, testing::Combine( |
||||
ALL_DEVICES, |
||||
testing::Values(std::string("768x576.avi"), std::string("1920x1080.avi")), |
||||
testing::Values(Channels(1), Channels(3), Channels(4)), |
||||
testing::Values(MaxFeatures(20), MaxFeatures(40), MaxFeatures(60)))); |
||||
|
||||
//////////////////////////////////////////////////////
|
||||
// VideoWriter
|
||||
|
||||
#ifdef WIN32 |
||||
|
||||
GPU_PERF_TEST(VideoWriter, cv::gpu::DeviceInfo, std::string) |
||||
{ |
||||
const double FPS = 25.0; |
||||
|
||||
std::string inputFile = perf::TestBase::getDataPath(std::string("gpu/video/") + GET_PARAM(1)); |
||||
std::string outputFile = cv::tempfile(".avi"); |
||||
|
||||
cv::VideoCapture reader(inputFile); |
||||
ASSERT_TRUE( reader.isOpened() ); |
||||
|
||||
cv::VideoWriter writer; |
||||
|
||||
cv::Mat frame; |
||||
|
||||
declare.time(30); |
||||
|
||||
for (int i = 0; i < 10; ++i) |
||||
{ |
||||
reader >> frame; |
||||
ASSERT_FALSE(frame.empty()); |
||||
|
||||
if (!writer.isOpened()) |
||||
writer.open(outputFile, CV_FOURCC('X', 'V', 'I', 'D'), FPS, frame.size()); |
||||
|
||||
startTimer(); next(); |
||||
writer.write(frame); |
||||
stopTimer(); |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Video, VideoWriter, testing::Combine( |
||||
ALL_DEVICES, |
||||
testing::Values(std::string("768x576.avi"), std::string("1920x1080.avi")))); |
||||
|
||||
#endif // WIN32
|
||||
|
||||
//////////////////////////////////////////////////////
|
||||
// VideoReader
|
||||
|
||||
GPU_PERF_TEST(VideoReader, cv::gpu::DeviceInfo, std::string) |
||||
{ |
||||
std::string inputFile = perf::TestBase::getDataPath(std::string("gpu/video/") + GET_PARAM(1)); |
||||
|
||||
cv::VideoCapture reader(inputFile); |
||||
ASSERT_TRUE( reader.isOpened() ); |
||||
|
||||
cv::Mat frame; |
||||
|
||||
reader >> frame; |
||||
|
||||
declare.time(20); |
||||
|
||||
TEST_CYCLE_N(10) |
||||
{ |
||||
reader >> frame; |
||||
} |
||||
} |
||||
|
||||
INSTANTIATE_TEST_CASE_P(Video, VideoReader, testing::Combine( |
||||
ALL_DEVICES, |
||||
testing::Values(std::string("768x576.avi"), std::string("1920x1080.avi")))); |
||||
|
||||
#endif |
||||
@ -0,0 +1,385 @@ |
||||
/*M/////////////////////////////////////////////////////////////////////////////////////// |
||||
// |
||||
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. |
||||
// |
||||
// By downloading, copying, installing or using the software you agree to this license. |
||||
// If you do not agree to this license, do not download, install, |
||||
// copy or use the software. |
||||
// |
||||
// |
||||
// License Agreement |
||||
// For Open Source Computer Vision Library |
||||
// |
||||
// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. |
||||
// Copyright (C) 2009, Willow Garage Inc., all rights reserved. |
||||
// Third party copyrights are property of their respective owners. |
||||
// |
||||
// Redistribution and use in source and binary forms, with or without modification, |
||||
// are permitted provided that the following conditions are met: |
||||
// |
||||
// * Redistribution's of source code must retain the above copyright notice, |
||||
// this list of conditions and the following disclaimer. |
||||
// |
||||
// * Redistribution's in binary form must reproduce the above copyright notice, |
||||
// this list of conditions and the following disclaimer in the documentation |
||||
// and/or other materials provided with the distribution. |
||||
// |
||||
// * The name of the copyright holders may not be used to endorse or promote products |
||||
// derived from this software without specific prior written permission. |
||||
// |
||||
// This software is provided by the copyright holders and contributors "as is" and |
||||
// any express or bpied warranties, including, but not limited to, the bpied |
||||
// warranties of merchantability and fitness for a particular purpose are disclaimed. |
||||
// In no event shall the Intel Corporation or contributors be liable for any direct, |
||||
// indirect, incidental, special, exemplary, or consequential damages |
||||
// (including, but not limited to, procurement of substitute goods or services; |
||||
// loss of use, data, or profits; or business interruption) however caused |
||||
// and on any theory of liability, whether in contract, strict liability, |
||||
// or tort (including negligence or otherwise) arising in any way out of |
||||
// the use of this software, even if advised of the possibility of such damage. |
||||
// |
||||
//M*/ |
||||
|
||||
#include "opencv2/gpu/device/common.hpp" |
||||
|
||||
namespace cv { namespace gpu { namespace device |
||||
{ |
||||
namespace imgproc |
||||
{ |
||||
// Utility function to extract unsigned chars from an unsigned integer |
||||
__device__ uchar4 int_to_uchar4(unsigned int in) |
||||
{ |
||||
uchar4 bytes; |
||||
bytes.x = (in && 0x000000ff) >> 0; |
||||
bytes.y = (in && 0x0000ff00) >> 8; |
||||
bytes.z = (in && 0x00ff0000) >> 16; |
||||
bytes.w = (in && 0xff000000) >> 24; |
||||
return bytes; |
||||
} |
||||
|
||||
__global__ void shfl_integral_horizontal(const PtrStep_<uint4> img, PtrStep_<uint4> integral) |
||||
{ |
||||
#if __CUDA_ARCH__ >= 300 |
||||
__shared__ int sums[128]; |
||||
|
||||
const int id = threadIdx.x; |
||||
const int lane_id = id % warpSize; |
||||
const int warp_id = id / warpSize; |
||||
|
||||
const uint4 data = img(blockIdx.x, id); |
||||
|
||||
const uchar4 a = int_to_uchar4(data.x); |
||||
const uchar4 b = int_to_uchar4(data.y); |
||||
const uchar4 c = int_to_uchar4(data.z); |
||||
const uchar4 d = int_to_uchar4(data.w); |
||||
|
||||
int result[16]; |
||||
|
||||
result[0] = a.x; |
||||
result[1] = result[0] + a.y; |
||||
result[2] = result[1] + a.z; |
||||
result[3] = result[2] + a.w; |
||||
|
||||
result[4] = result[3] + b.x; |
||||
result[5] = result[4] + b.y; |
||||
result[6] = result[5] + b.z; |
||||
result[7] = result[6] + b.w; |
||||
|
||||
result[8] = result[7] + c.x; |
||||
result[9] = result[8] + c.y; |
||||
result[10] = result[9] + c.z; |
||||
result[11] = result[10] + c.w; |
||||
|
||||
result[12] = result[11] + d.x; |
||||
result[13] = result[12] + d.y; |
||||
result[14] = result[13] + d.z; |
||||
result[15] = result[14] + d.w; |
||||
|
||||
int sum = result[15]; |
||||
|
||||
// the prefix sum for each thread's 16 value is computed, |
||||
// now the final sums (result[15]) need to be shared |
||||
// with the other threads and add. To do this, |
||||
// the __shfl_up() instruction is used and a shuffle scan |
||||
// operation is performed to distribute the sums to the correct |
||||
// threads |
||||
#pragma unroll |
||||
for (int i = 1; i < 32; i *= 2) |
||||
{ |
||||
const int n = __shfl_up(sum, i, 32); |
||||
|
||||
if (lane_id >= i) |
||||
{ |
||||
#pragma unroll |
||||
for (int i = 0; i < 16; ++i) |
||||
result[i] += n; |
||||
|
||||
sum += n; |
||||
} |
||||
} |
||||
|
||||
// Now the final sum for the warp must be shared |
||||
// between warps. This is done by each warp |
||||
// having a thread store to shared memory, then |
||||
// having some other warp load the values and |
||||
// compute a prefix sum, again by using __shfl_up. |
||||
// The results are uniformly added back to the warps. |
||||
// last thread in the warp holding sum of the warp |
||||
// places that in shared |
||||
if (threadIdx.x % warpSize == warpSize - 1) |
||||
sums[warp_id] = result[15]; |
||||
|
||||
__syncthreads(); |
||||
|
||||
if (warp_id == 0) |
||||
{ |
||||
int warp_sum = sums[lane_id]; |
||||
|
||||
#pragma unroll |
||||
for (int i = 1; i <= 32; i *= 2) |
||||
{ |
||||
const int n = __shfl_up(warp_sum, i, 32); |
||||
|
||||
if (lane_id >= i) |
||||
warp_sum += n; |
||||
} |
||||
|
||||
sums[lane_id] = warp_sum; |
||||
} |
||||
|
||||
__syncthreads(); |
||||
|
||||
int blockSum = 0; |
||||
|
||||
// fold in unused warp |
||||
if (warp_id > 0) |
||||
{ |
||||
blockSum = sums[warp_id - 1]; |
||||
|
||||
#pragma unroll |
||||
for (int i = 0; i < 16; ++i) |
||||
result[i] += blockSum; |
||||
} |
||||
|
||||
// assemble result |
||||
// Each thread has 16 values to write, which are |
||||
// now integer data (to avoid overflow). Instead of |
||||
// each thread writing consecutive uint4s, the |
||||
// approach shown here experiments using |
||||
// the shuffle command to reformat the data |
||||
// inside the registers so that each thread holds |
||||
// consecutive data to be written so larger contiguous |
||||
// segments can be assembled for writing. |
||||
|
||||
/* |
||||
For example data that needs to be written as |
||||
|
||||
GMEM[16] <- x0 x1 x2 x3 y0 y1 y2 y3 z0 z1 z2 z3 w0 w1 w2 w3 |
||||
but is stored in registers (r0..r3), in four threads (0..3) as: |
||||
|
||||
threadId 0 1 2 3 |
||||
r0 x0 y0 z0 w0 |
||||
r1 x1 y1 z1 w1 |
||||
r2 x2 y2 z2 w2 |
||||
r3 x3 y3 z3 w3 |
||||
|
||||
after apply __shfl_xor operations to move data between registers r1..r3: |
||||
|
||||
threadId 00 01 10 11 |
||||
x0 y0 z0 w0 |
||||
xor(01)->y1 x1 w1 z1 |
||||
xor(10)->z2 w2 x2 y2 |
||||
xor(11)->w3 z3 y3 x3 |
||||
|
||||
and now x0..x3, and z0..z3 can be written out in order by all threads. |
||||
|
||||
In the current code, each register above is actually representing |
||||
four integers to be written as uint4's to GMEM. |
||||
*/ |
||||
|
||||
result[4] = __shfl_xor(result[4] , 1, 32); |
||||
result[5] = __shfl_xor(result[5] , 1, 32); |
||||
result[6] = __shfl_xor(result[6] , 1, 32); |
||||
result[7] = __shfl_xor(result[7] , 1, 32); |
||||
|
||||
result[8] = __shfl_xor(result[8] , 2, 32); |
||||
result[9] = __shfl_xor(result[9] , 2, 32); |
||||
result[10] = __shfl_xor(result[10], 2, 32); |
||||
result[11] = __shfl_xor(result[11], 2, 32); |
||||
|
||||
result[12] = __shfl_xor(result[12], 3, 32); |
||||
result[13] = __shfl_xor(result[13], 3, 32); |
||||
result[14] = __shfl_xor(result[14], 3, 32); |
||||
result[15] = __shfl_xor(result[15], 3, 32); |
||||
|
||||
uint4* integral_row = integral.ptr(blockIdx.x); |
||||
uint4 output; |
||||
|
||||
/////// |
||||
|
||||
if (threadIdx.x % 4 == 0) |
||||
output = make_uint4(result[0], result[1], result[2], result[3]); |
||||
|
||||
if (threadIdx.x % 4 == 1) |
||||
output = make_uint4(result[4], result[5], result[6], result[7]); |
||||
|
||||
if (threadIdx.x % 4 == 2) |
||||
output = make_uint4(result[8], result[9], result[10], result[11]); |
||||
|
||||
if (threadIdx.x % 4 == 3) |
||||
output = make_uint4(result[12], result[13], result[14], result[15]); |
||||
|
||||
integral_row[threadIdx.x % 4 + (threadIdx.x / 4) * 16] = output; |
||||
|
||||
/////// |
||||
|
||||
if (threadIdx.x % 4 == 2) |
||||
output = make_uint4(result[0], result[1], result[2], result[3]); |
||||
|
||||
if (threadIdx.x % 4 == 3) |
||||
output = make_uint4(result[4], result[5], result[6], result[7]); |
||||
|
||||
if (threadIdx.x % 4 == 0) |
||||
output = make_uint4(result[8], result[9], result[10], result[11]); |
||||
|
||||
if (threadIdx.x % 4 == 1) |
||||
output = make_uint4(result[12], result[13], result[14], result[15]); |
||||
|
||||
integral_row[(threadIdx.x + 2) % 4 + (threadIdx.x / 4) * 16 + 8] = output; |
||||
|
||||
// continuning from the above example, |
||||
// this use of __shfl_xor() places the y0..y3 and w0..w3 data |
||||
// in order. |
||||
|
||||
#pragma unroll |
||||
for (int i = 0; i < 16; ++i) |
||||
result[i] = __shfl_xor(result[i], 1, 32); |
||||
|
||||
if (threadIdx.x % 4 == 0) |
||||
output = make_uint4(result[0], result[1], result[2], result[3]); |
||||
|
||||
if (threadIdx.x % 4 == 1) |
||||
output = make_uint4(result[4], result[5], result[6], result[7]); |
||||
|
||||
if (threadIdx.x % 4 == 2) |
||||
output = make_uint4(result[8], result[9], result[10], result[11]); |
||||
|
||||
if (threadIdx.x % 4 == 3) |
||||
output = make_uint4(result[12], result[13], result[14], result[15]); |
||||
|
||||
integral_row[threadIdx.x % 4 + (threadIdx.x / 4) * 16 + 4] = output; |
||||
|
||||
/////// |
||||
|
||||
if (threadIdx.x % 4 == 2) |
||||
output = make_uint4(result[0], result[1], result[2], result[3]); |
||||
|
||||
if (threadIdx.x % 4 == 3) |
||||
output = make_uint4(result[4], result[5], result[6], result[7]); |
||||
|
||||
if (threadIdx.x % 4 == 0) |
||||
output = make_uint4(result[8], result[9], result[10], result[11]); |
||||
|
||||
if (threadIdx.x % 4 == 1) |
||||
output = make_uint4(result[12], result[13], result[14], result[15]); |
||||
|
||||
integral_row[(threadIdx.x + 2) % 4 + (threadIdx.x / 4) * 16 + 12] = output; |
||||
#endif |
||||
} |
||||
|
||||
// This kernel computes columnwise prefix sums. When the data input is |
||||
// the row sums from above, this completes the integral image. |
||||
// The approach here is to have each block compute a local set of sums. |
||||
// First , the data covered by the block is loaded into shared memory, |
||||
// then instead of performing a sum in shared memory using __syncthreads |
||||
// between stages, the data is reformatted so that the necessary sums |
||||
// occur inside warps and the shuffle scan operation is used. |
||||
// The final set of sums from the block is then propgated, with the block |
||||
// computing "down" the image and adding the running sum to the local |
||||
// block sums. |
||||
__global__ void shfl_integral_vertical(DevMem2D_<unsigned int> integral) |
||||
{ |
||||
#if __CUDA_ARCH__ >= 300 |
||||
__shared__ unsigned int sums[32][9]; |
||||
|
||||
const int tidx = blockIdx.x * blockDim.x + threadIdx.x; |
||||
const int lane_id = tidx % 8; |
||||
|
||||
if (tidx >= integral.cols) |
||||
return; |
||||
|
||||
sums[threadIdx.x][threadIdx.y] = 0; |
||||
__syncthreads(); |
||||
|
||||
unsigned int stepSum = 0; |
||||
|
||||
for (int y = threadIdx.y; y < integral.rows; y += blockDim.y) |
||||
{ |
||||
unsigned int* p = integral.ptr(y) + tidx; |
||||
|
||||
unsigned int sum = *p; |
||||
|
||||
sums[threadIdx.x][threadIdx.y] = sum; |
||||
__syncthreads(); |
||||
|
||||
// place into SMEM |
||||
// shfl scan reduce the SMEM, reformating so the column |
||||
// sums are computed in a warp |
||||
// then read out properly |
||||
const int j = threadIdx.x % 8; |
||||
const int k = threadIdx.x / 8 + threadIdx.y * 4; |
||||
|
||||
int partial_sum = sums[k][j]; |
||||
|
||||
for (int i = 1; i <= 8; i *= 2) |
||||
{ |
||||
int n = __shfl_up(partial_sum, i, 32); |
||||
|
||||
if (lane_id >= i) |
||||
partial_sum += n; |
||||
} |
||||
|
||||
sums[k][j] = partial_sum; |
||||
__syncthreads(); |
||||
|
||||
if (threadIdx.y > 0) |
||||
sum += sums[threadIdx.x][threadIdx.y - 1]; |
||||
|
||||
sum += stepSum; |
||||
stepSum += sums[threadIdx.x][blockDim.y - 1]; |
||||
|
||||
__syncthreads(); |
||||
|
||||
*p = sum; |
||||
} |
||||
#endif |
||||
} |
||||
|
||||
void shfl_integral_gpu(DevMem2Db img, DevMem2D_<unsigned int> integral, cudaStream_t stream) |
||||
{ |
||||
{ |
||||
// each thread handles 16 values, use 1 block/row |
||||
const int block = img.cols / 16; |
||||
|
||||
// launch 1 block / row |
||||
const int grid = img.rows; |
||||
|
||||
cudaSafeCall( cudaFuncSetCacheConfig(shfl_integral_horizontal, cudaFuncCachePreferL1) ); |
||||
|
||||
shfl_integral_horizontal<<<grid, block, 0, stream>>>((DevMem2D_<uint4>) img, (DevMem2D_<uint4>) integral); |
||||
cudaSafeCall( cudaGetLastError() ); |
||||
} |
||||
|
||||
{ |
||||
const dim3 block(32, 8); |
||||
const dim3 grid(divUp(integral.cols, block.x), 1); |
||||
|
||||
shfl_integral_vertical<<<grid, block, 0, stream>>>(integral); |
||||
cudaSafeCall( cudaGetLastError() ); |
||||
} |
||||
|
||||
if (stream == 0) |
||||
cudaSafeCall( cudaDeviceSynchronize() ); |
||||
} |
||||
} |
||||
}}} |
||||
@ -0,0 +1,163 @@ |
||||
/*
|
||||
* cap_ios.h |
||||
* For iOS video I/O |
||||
* by Eduard Feicho on 29/07/12 |
||||
* Copyright 2012. All rights reserved. |
||||
* |
||||
* Redistribution and use in source and binary forms, with or without |
||||
* modification, are permitted provided that the following conditions are met: |
||||
* |
||||
* 1. Redistributions of source code must retain the above copyright notice, |
||||
* this list of conditions and the following disclaimer. |
||||
* 2. Redistributions in binary form must reproduce the above copyright notice, |
||||
* this list of conditions and the following disclaimer in the documentation |
||||
* and/or other materials provided with the distribution. |
||||
* 3. The name of the author may not be used to endorse or promote products |
||||
* derived from this software without specific prior written permission. |
||||
* |
||||
* THIS SOFTWARE IS PROVIDED BY THE AUTHOR "AS IS" AND ANY EXPRESS OR IMPLIED |
||||
* WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF |
||||
* MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO |
||||
* EVENT SHALL THE AUTHOR BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, |
||||
* SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, |
||||
* PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; |
||||
* OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, |
||||
* WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR |
||||
* OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF |
||||
* ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. |
||||
* |
||||
*/ |
||||
|
||||
#import <UIKit/UIKit.h> |
||||
#import <Accelerate/Accelerate.h> |
||||
#import <AVFoundation/AVFoundation.h> |
||||
#import <ImageIO/ImageIO.h> |
||||
#include "opencv2/core/core.hpp" |
||||
|
||||
/////////////////////////////////////// CvAbstractCamera /////////////////////////////////////
|
||||
|
||||
@class CvAbstractCamera; |
||||
|
||||
@interface CvAbstractCamera : NSObject |
||||
{ |
||||
AVCaptureSession* captureSession; |
||||
AVCaptureConnection* videoCaptureConnection; |
||||
AVCaptureVideoPreviewLayer *captureVideoPreviewLayer; |
||||
|
||||
UIDeviceOrientation currentDeviceOrientation; |
||||
|
||||
BOOL cameraAvailable; |
||||
BOOL captureSessionLoaded; |
||||
BOOL running; |
||||
BOOL useAVCaptureVideoPreviewLayer; |
||||
|
||||
AVCaptureDevicePosition defaultAVCaptureDevicePosition; |
||||
AVCaptureVideoOrientation defaultAVCaptureVideoOrientation; |
||||
NSString *const defaultAVCaptureSessionPreset; |
||||
|
||||
int defaultFPS; |
||||
|
||||
UIView* parentView; |
||||
|
||||
int imageWidth; |
||||
int imageHeight; |
||||
} |
||||
|
||||
@property (nonatomic, retain) AVCaptureSession* captureSession; |
||||
@property (nonatomic, retain) AVCaptureConnection* videoCaptureConnection; |
||||
|
||||
@property (nonatomic, readonly) BOOL running; |
||||
@property (nonatomic, readonly) BOOL captureSessionLoaded; |
||||
|
||||
@property (nonatomic, assign) int defaultFPS; |
||||
@property (nonatomic, assign) AVCaptureDevicePosition defaultAVCaptureDevicePosition; |
||||
@property (nonatomic, assign) AVCaptureVideoOrientation defaultAVCaptureVideoOrientation; |
||||
@property (nonatomic, assign) BOOL useAVCaptureVideoPreviewLayer; |
||||
@property (nonatomic, strong) NSString *const defaultAVCaptureSessionPreset; |
||||
|
||||
@property (nonatomic, assign) int imageWidth; |
||||
@property (nonatomic, assign) int imageHeight; |
||||
|
||||
@property (nonatomic, retain) UIView* parentView; |
||||
|
||||
- (void)start; |
||||
- (void)stop; |
||||
- (void)switchCameras; |
||||
|
||||
- (id)initWithParentView:(UIView*)parent; |
||||
|
||||
- (void)createCaptureOutput; |
||||
- (void)createVideoPreviewLayer; |
||||
- (void)updateOrientation; |
||||
|
||||
|
||||
@end |
||||
|
||||
///////////////////////////////// CvVideoCamera ///////////////////////////////////////////
|
||||
|
||||
@class CvVideoCamera; |
||||
|
||||
@protocol CvVideoCameraDelegate <NSObject> |
||||
|
||||
#ifdef __cplusplus |
||||
// delegate method for processing image frames
|
||||
- (void)processImage:(cv::Mat&)image; |
||||
#endif |
||||
|
||||
@end |
||||
|
||||
@interface CvVideoCamera : CvAbstractCamera<AVCaptureVideoDataOutputSampleBufferDelegate> |
||||
{ |
||||
AVCaptureVideoDataOutput *videoDataOutput; |
||||
|
||||
dispatch_queue_t videoDataOutputQueue; |
||||
CALayer *customPreviewLayer; |
||||
|
||||
BOOL grayscaleMode; |
||||
|
||||
BOOL recordVideo; |
||||
AVAssetWriterInput* recordAssetWriterInput; |
||||
AVAssetWriterInputPixelBufferAdaptor* recordPixelBufferAdaptor; |
||||
AVAssetWriter* recordAssetWriter; |
||||
|
||||
CMTime lastSampleTime; |
||||
|
||||
} |
||||
|
||||
@property (nonatomic, assign) id<CvVideoCameraDelegate> delegate; |
||||
@property (nonatomic, assign) BOOL grayscaleMode; |
||||
|
||||
@property (nonatomic, assign) BOOL recordVideo; |
||||
@property (nonatomic, retain) AVAssetWriterInput* recordAssetWriterInput; |
||||
@property (nonatomic, retain) AVAssetWriterInputPixelBufferAdaptor* recordPixelBufferAdaptor; |
||||
@property (nonatomic, retain) AVAssetWriter* recordAssetWriter; |
||||
|
||||
- (void)adjustLayoutToInterfaceOrientation:(UIInterfaceOrientation)interfaceOrientation; |
||||
- (void)layoutPreviewLayer; |
||||
- (void)saveVideo; |
||||
- (NSURL *)videoFileURL; |
||||
|
||||
|
||||
@end |
||||
|
||||
///////////////////////////////// CvPhotoCamera ///////////////////////////////////////////
|
||||
|
||||
@class CvPhotoCamera; |
||||
|
||||
@protocol CvPhotoCameraDelegate <NSObject> |
||||
|
||||
- (void)photoCamera:(CvPhotoCamera*)photoCamera capturedImage:(UIImage *)image; |
||||
- (void)photoCameraCancel:(CvPhotoCamera*)photoCamera; |
||||
|
||||
@end |
||||
|
||||
@interface CvPhotoCamera : CvAbstractCamera |
||||
{ |
||||
AVCaptureStillImageOutput *stillImageOutput; |
||||
} |
||||
|
||||
@property (nonatomic, assign) id<CvPhotoCameraDelegate> delegate; |
||||
|
||||
- (void)takePicture; |
||||
|
||||
@end |
||||
@ -0,0 +1,408 @@ |
||||
/* |
||||
* cap_ios_abstract_camera.mm |
||||
* For iOS video I/O |
||||
* by Eduard Feicho on 29/07/12 |
||||
* Copyright 2012. All rights reserved. |
||||
* |
||||
* Redistribution and use in source and binary forms, with or without |
||||
* modification, are permitted provided that the following conditions are met: |
||||
* |
||||
* 1. Redistributions of source code must retain the above copyright notice, |
||||
* this list of conditions and the following disclaimer. |
||||
* 2. Redistributions in binary form must reproduce the above copyright notice, |
||||
* this list of conditions and the following disclaimer in the documentation |
||||
* and/or other materials provided with the distribution. |
||||
* 3. The name of the author may not be used to endorse or promote products |
||||
* derived from this software without specific prior written permission. |
||||
* |
||||
* THIS SOFTWARE IS PROVIDED BY THE AUTHOR "AS IS" AND ANY EXPRESS OR IMPLIED |
||||
* WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF |
||||
* MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO |
||||
* EVENT SHALL THE AUTHOR BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, |
||||
* SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, |
||||
* PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; |
||||
* OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, |
||||
* WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR |
||||
* OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF |
||||
* ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. |
||||
* |
||||
*/ |
||||
|
||||
|
||||
#import "opencv2/highgui/cap_ios.h" |
||||
#include "precomp.hpp" |
||||
|
||||
#pragma mark - Private Interface |
||||
|
||||
@interface CvAbstractCamera () |
||||
|
||||
@property (nonatomic, retain) AVCaptureVideoPreviewLayer* captureVideoPreviewLayer; |
||||
|
||||
- (void)deviceOrientationDidChange:(NSNotification*)notification; |
||||
- (void)startCaptureSession; |
||||
|
||||
- (void)setDesiredCameraPosition:(AVCaptureDevicePosition)desiredPosition; |
||||
|
||||
- (void)updateSize; |
||||
|
||||
@end |
||||
|
||||
|
||||
#pragma mark - Implementation |
||||
|
||||
|
||||
@implementation CvAbstractCamera |
||||
|
||||
|
||||
|
||||
#pragma mark Public |
||||
|
||||
@synthesize imageWidth; |
||||
@synthesize imageHeight; |
||||
|
||||
|
||||
@synthesize defaultFPS; |
||||
@synthesize defaultAVCaptureDevicePosition; |
||||
@synthesize defaultAVCaptureVideoOrientation; |
||||
@synthesize defaultAVCaptureSessionPreset; |
||||
|
||||
|
||||
|
||||
@synthesize captureSession; |
||||
@synthesize captureVideoPreviewLayer; |
||||
@synthesize videoCaptureConnection; |
||||
@synthesize running; |
||||
@synthesize captureSessionLoaded; |
||||
@synthesize useAVCaptureVideoPreviewLayer; |
||||
|
||||
@synthesize parentView; |
||||
|
||||
#pragma mark - Constructors |
||||
|
||||
- (id)init; |
||||
{ |
||||
self = [super init]; |
||||
if (self) { |
||||
// react to device orientation notifications |
||||
[[NSNotificationCenter defaultCenter] addObserver:self |
||||
selector:@selector(deviceOrientationDidChange:) |
||||
name:UIDeviceOrientationDidChangeNotification |
||||
object:nil]; |
||||
[[UIDevice currentDevice] beginGeneratingDeviceOrientationNotifications]; |
||||
currentDeviceOrientation = [[UIDevice currentDevice] orientation]; |
||||
|
||||
|
||||
// check if camera available |
||||
cameraAvailable = [UIImagePickerController isSourceTypeAvailable:UIImagePickerControllerSourceTypeCamera]; |
||||
NSLog(@"camera available: %@", (cameraAvailable == YES ? @"YES" : @"NO") ); |
||||
|
||||
running = NO; |
||||
|
||||
// set camera default configuration |
||||
self.defaultAVCaptureDevicePosition = AVCaptureDevicePositionFront; |
||||
self.defaultAVCaptureVideoOrientation = AVCaptureVideoOrientationLandscapeLeft; |
||||
self.defaultFPS = 15; |
||||
self.defaultAVCaptureSessionPreset = AVCaptureSessionPreset352x288; |
||||
|
||||
self.parentView = nil; |
||||
self.useAVCaptureVideoPreviewLayer = NO; |
||||
} |
||||
return self; |
||||
} |
||||
|
||||
|
||||
|
||||
- (id)initWithParentView:(UIView*)parent; |
||||
{ |
||||
self = [super init]; |
||||
if (self) { |
||||
// react to device orientation notifications |
||||
[[NSNotificationCenter defaultCenter] addObserver:self |
||||
selector:@selector(deviceOrientationDidChange:) |
||||
name:UIDeviceOrientationDidChangeNotification |
||||
object:nil]; |
||||
[[UIDevice currentDevice] beginGeneratingDeviceOrientationNotifications]; |
||||
currentDeviceOrientation = [[UIDevice currentDevice] orientation]; |
||||
|
||||
|
||||
// check if camera available |
||||
cameraAvailable = [UIImagePickerController isSourceTypeAvailable:UIImagePickerControllerSourceTypeCamera]; |
||||
NSLog(@"camera available: %@", (cameraAvailable == YES ? @"YES" : @"NO") ); |
||||
|
||||
running = NO; |
||||
|
||||
// set camera default configuration |
||||
self.defaultAVCaptureDevicePosition = AVCaptureDevicePositionFront; |
||||
self.defaultAVCaptureVideoOrientation = AVCaptureVideoOrientationLandscapeLeft; |
||||
self.defaultFPS = 15; |
||||
self.defaultAVCaptureSessionPreset = AVCaptureSessionPreset640x480; |
||||
|
||||
self.parentView = parent; |
||||
self.useAVCaptureVideoPreviewLayer = YES; |
||||
} |
||||
return self; |
||||
} |
||||
|
||||
|
||||
|
||||
- (void)dealloc; |
||||
{ |
||||
[[NSNotificationCenter defaultCenter] removeObserver:self]; |
||||
[[UIDevice currentDevice] endGeneratingDeviceOrientationNotifications]; |
||||
} |
||||
|
||||
|
||||
#pragma mark - Public interface |
||||
|
||||
|
||||
- (void)start; |
||||
{ |
||||
if (![NSThread isMainThread]) { |
||||
NSLog(@"[Camera] Warning: Call start only from main thread"); |
||||
[self performSelectorOnMainThread:@selector(start) withObject:nil waitUntilDone:NO]; |
||||
return; |
||||
} |
||||
|
||||
if (running == YES) { |
||||
return; |
||||
} |
||||
running = YES; |
||||
|
||||
// TOOD update image size data before actually starting (needed for recording) |
||||
[self updateSize]; |
||||
|
||||
if (cameraAvailable) { |
||||
[self startCaptureSession]; |
||||
} |
||||
} |
||||
|
||||
|
||||
- (void)pause; |
||||
{ |
||||
running = NO; |
||||
[self.captureSession stopRunning]; |
||||
} |
||||
|
||||
|
||||
|
||||
- (void)stop; |
||||
{ |
||||
running = NO; |
||||
|
||||
// Release any retained subviews of the main view. |
||||
// e.g. self.myOutlet = nil; |
||||
|
||||
[self.captureSession stopRunning]; |
||||
self.captureSession = nil; |
||||
self.captureVideoPreviewLayer = nil; |
||||
self.videoCaptureConnection = nil; |
||||
captureSessionLoaded = NO; |
||||
} |
||||
|
||||
|
||||
|
||||
// use front/back camera |
||||
- (void)switchCameras; |
||||
{ |
||||
BOOL was_running = self.running; |
||||
if (was_running) { |
||||
[self stop]; |
||||
} |
||||
if (self.defaultAVCaptureDevicePosition == AVCaptureDevicePositionFront) { |
||||
self.defaultAVCaptureDevicePosition = AVCaptureDevicePositionBack; |
||||
} else { |
||||
self.defaultAVCaptureDevicePosition = AVCaptureDevicePositionFront; |
||||
} |
||||
if (was_running) { |
||||
[self start]; |
||||
} |
||||
} |
||||
|
||||
|
||||
|
||||
#pragma mark - Device Orientation Changes |
||||
|
||||
|
||||
- (void)deviceOrientationDidChange:(NSNotification*)notification |
||||
{ |
||||
UIDeviceOrientation orientation = [UIDevice currentDevice].orientation; |
||||
|
||||
switch (orientation) |
||||
{ |
||||
case UIDeviceOrientationPortrait: |
||||
case UIDeviceOrientationPortraitUpsideDown: |
||||
case UIDeviceOrientationLandscapeLeft: |
||||
case UIDeviceOrientationLandscapeRight: |
||||
currentDeviceOrientation = orientation; |
||||
break; |
||||
|
||||
case UIDeviceOrientationFaceUp: |
||||
case UIDeviceOrientationFaceDown: |
||||
default: |
||||
break; |
||||
} |
||||
NSLog(@"deviceOrientationDidChange: %d", orientation); |
||||
|
||||
[self updateOrientation]; |
||||
} |
||||
|
||||
|
||||
|
||||
#pragma mark - Private Interface |
||||
|
||||
- (void)createCaptureSession; |
||||
{ |
||||
// set a av capture session preset |
||||
self.captureSession = [[AVCaptureSession alloc] init]; |
||||
if ([self.captureSession canSetSessionPreset:self.defaultAVCaptureSessionPreset]) { |
||||
[self.captureSession setSessionPreset:self.defaultAVCaptureSessionPreset]; |
||||
} else if ([self.captureSession canSetSessionPreset:AVCaptureSessionPresetLow]) { |
||||
[self.captureSession setSessionPreset:AVCaptureSessionPresetLow]; |
||||
} else { |
||||
NSLog(@"[Camera] Error: could not set session preset"); |
||||
} |
||||
} |
||||
|
||||
- (void)createCaptureDevice; |
||||
{ |
||||
// setup the device |
||||
AVCaptureDevice *device = [AVCaptureDevice defaultDeviceWithMediaType:AVMediaTypeVideo]; |
||||
[self setDesiredCameraPosition:self.defaultAVCaptureDevicePosition]; |
||||
NSLog(@"[Camera] device connected? %@", device.connected ? @"YES" : @"NO"); |
||||
NSLog(@"[Camera] device position %@", (device.position == AVCaptureDevicePositionBack) ? @"back" : @"front"); |
||||
} |
||||
|
||||
|
||||
- (void)createVideoPreviewLayer; |
||||
{ |
||||
self.captureVideoPreviewLayer = [[AVCaptureVideoPreviewLayer alloc] initWithSession:self.captureSession]; |
||||
|
||||
if ([self.captureVideoPreviewLayer isOrientationSupported]) { |
||||
[self.captureVideoPreviewLayer setOrientation:self.defaultAVCaptureVideoOrientation]; |
||||
} |
||||
|
||||
if (parentView != nil) { |
||||
self.captureVideoPreviewLayer.frame = self.parentView.bounds; |
||||
self.captureVideoPreviewLayer.videoGravity = AVLayerVideoGravityResizeAspectFill; |
||||
[self.parentView.layer addSublayer:self.captureVideoPreviewLayer]; |
||||
} |
||||
NSLog(@"[Camera] created AVCaptureVideoPreviewLayer"); |
||||
} |
||||
|
||||
|
||||
|
||||
|
||||
- (void)setDesiredCameraPosition:(AVCaptureDevicePosition)desiredPosition; |
||||
{ |
||||
for (AVCaptureDevice *device in [AVCaptureDevice devicesWithMediaType:AVMediaTypeVideo]) { |
||||
if ([device position] == desiredPosition) { |
||||
[self.captureSession beginConfiguration]; |
||||
|
||||
NSError* error; |
||||
AVCaptureDeviceInput *input = [AVCaptureDeviceInput deviceInputWithDevice:device error:&error]; |
||||
if (!input) { |
||||
NSLog(@"error creating input %@", [error localizedDescription]); |
||||
} |
||||
|
||||
// support for autofocus |
||||
if ([device isFocusModeSupported:AVCaptureFocusModeContinuousAutoFocus]) { |
||||
NSError *error = nil; |
||||
if ([device lockForConfiguration:&error]) { |
||||
device.focusMode = AVCaptureFocusModeContinuousAutoFocus; |
||||
[device unlockForConfiguration]; |
||||
} else { |
||||
NSLog(@"unable to lock device for autofocos configuration %@", [error localizedDescription]); |
||||
} |
||||
} |
||||
[self.captureSession addInput:input]; |
||||
|
||||
for (AVCaptureInput *oldInput in self.captureSession.inputs) { |
||||
[self.captureSession removeInput:oldInput]; |
||||
} |
||||
[self.captureSession addInput:input]; |
||||
[self.captureSession commitConfiguration]; |
||||
|
||||
break; |
||||
} |
||||
} |
||||
} |
||||
|
||||
|
||||
|
||||
- (void)startCaptureSession |
||||
{ |
||||
if (!cameraAvailable) { |
||||
return; |
||||
} |
||||
|
||||
if (self.captureSessionLoaded == NO) { |
||||
[self createCaptureSession]; |
||||
[self createCaptureDevice]; |
||||
[self createCaptureOutput]; |
||||
|
||||
// setup preview layer |
||||
if (self.useAVCaptureVideoPreviewLayer) { |
||||
[self createVideoPreviewLayer]; |
||||
} else { |
||||
[self createCustomVideoPreview]; |
||||
} |
||||
|
||||
captureSessionLoaded = YES; |
||||
} |
||||
|
||||
[self.captureSession startRunning]; |
||||
} |
||||
|
||||
|
||||
- (void)createCaptureOutput; |
||||
{ |
||||
[NSException raise:NSInternalInconsistencyException |
||||
format:@"You must override %@ in a subclass", NSStringFromSelector(_cmd)]; |
||||
} |
||||
|
||||
- (void)createCustomVideoPreview; |
||||
{ |
||||
[NSException raise:NSInternalInconsistencyException |
||||
format:@"You must override %@ in a subclass", NSStringFromSelector(_cmd)]; |
||||
} |
||||
|
||||
- (void)updateOrientation; |
||||
{ |
||||
// nothing to do here |
||||
} |
||||
|
||||
|
||||
- (void)updateSize; |
||||
{ |
||||
if ([self.defaultAVCaptureSessionPreset isEqualToString:AVCaptureSessionPresetPhoto]) { |
||||
//TODO: find the correct resolution |
||||
self.imageWidth = 640; |
||||
self.imageHeight = 480; |
||||
} else if ([self.defaultAVCaptureSessionPreset isEqualToString:AVCaptureSessionPresetHigh]) { |
||||
//TODO: find the correct resolution |
||||
self.imageWidth = 640; |
||||
self.imageHeight = 480; |
||||
} else if ([self.defaultAVCaptureSessionPreset isEqualToString:AVCaptureSessionPresetMedium]) { |
||||
//TODO: find the correct resolution |
||||
self.imageWidth = 640; |
||||
self.imageHeight = 480; |
||||
} else if ([self.defaultAVCaptureSessionPreset isEqualToString:AVCaptureSessionPresetLow]) { |
||||
//TODO: find the correct resolution |
||||
self.imageWidth = 640; |
||||
self.imageHeight = 480; |
||||
} else if ([self.defaultAVCaptureSessionPreset isEqualToString:AVCaptureSessionPreset352x288]) { |
||||
self.imageWidth = 352; |
||||
self.imageHeight = 288; |
||||
} else if ([self.defaultAVCaptureSessionPreset isEqualToString:AVCaptureSessionPreset640x480]) { |
||||
self.imageWidth = 640; |
||||
self.imageHeight = 480; |
||||
} else if ([self.defaultAVCaptureSessionPreset isEqualToString:AVCaptureSessionPreset1280x720]) { |
||||
self.imageWidth = 1280; |
||||
self.imageHeight = 720; |
||||
} else { |
||||
self.imageWidth = 640; |
||||
self.imageHeight = 480; |
||||
} |
||||
} |
||||
|
||||
@end |
||||
@ -0,0 +1,165 @@ |
||||
/* |
||||
* cap_ios_photo_camera.mm |
||||
* For iOS video I/O |
||||
* by Eduard Feicho on 29/07/12 |
||||
* Copyright 2012. All rights reserved. |
||||
* |
||||
* Redistribution and use in source and binary forms, with or without |
||||
* modification, are permitted provided that the following conditions are met: |
||||
* |
||||
* 1. Redistributions of source code must retain the above copyright notice, |
||||
* this list of conditions and the following disclaimer. |
||||
* 2. Redistributions in binary form must reproduce the above copyright notice, |
||||
* this list of conditions and the following disclaimer in the documentation |
||||
* and/or other materials provided with the distribution. |
||||
* 3. The name of the author may not be used to endorse or promote products |
||||
* derived from this software without specific prior written permission. |
||||
* |
||||
* THIS SOFTWARE IS PROVIDED BY THE AUTHOR "AS IS" AND ANY EXPRESS OR IMPLIED |
||||
* WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF |
||||
* MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO |
||||
* EVENT SHALL THE AUTHOR BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, |
||||
* SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, |
||||
* PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; |
||||
* OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, |
||||
* WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR |
||||
* OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF |
||||
* ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. |
||||
* |
||||
*/ |
||||
|
||||
|
||||
#import "opencv2/highgui/cap_ios.h" |
||||
#include "precomp.hpp" |
||||
|
||||
#pragma mark - Private Interface |
||||
|
||||
|
||||
@interface CvPhotoCamera () |
||||
|
||||
@property (nonatomic, retain) AVCaptureStillImageOutput* stillImageOutput; |
||||
|
||||
@end |
||||
|
||||
|
||||
|
||||
#pragma mark - Implementation |
||||
|
||||
|
||||
@implementation CvPhotoCamera |
||||
|
||||
|
||||
|
||||
#pragma mark Public |
||||
|
||||
@synthesize stillImageOutput; |
||||
@synthesize delegate; |
||||
|
||||
|
||||
#pragma mark - Public interface |
||||
|
||||
|
||||
- (void)takePicture |
||||
{ |
||||
if (cameraAvailable == NO) { |
||||
return; |
||||
} |
||||
cameraAvailable = NO; |
||||
|
||||
|
||||
[self.stillImageOutput captureStillImageAsynchronouslyFromConnection:self.videoCaptureConnection |
||||
completionHandler: |
||||
^(CMSampleBufferRef imageSampleBuffer, NSError *error) |
||||
{ |
||||
if (error == nil && imageSampleBuffer != NULL) |
||||
{ |
||||
// TODO check |
||||
// NSNumber* imageOrientation = [UIImage cgImageOrientationForUIDeviceOrientation:currentDeviceOrientation]; |
||||
// CMSetAttachment(imageSampleBuffer, kCGImagePropertyOrientation, imageOrientation, 1); |
||||
|
||||
NSData *jpegData = [AVCaptureStillImageOutput jpegStillImageNSDataRepresentation:imageSampleBuffer]; |
||||
|
||||
dispatch_async(dispatch_get_main_queue(), ^{ |
||||
[self.captureSession stopRunning]; |
||||
|
||||
// Make sure we create objects on the main thread in the main context |
||||
UIImage* newImage = [UIImage imageWithData:jpegData]; |
||||
|
||||
//UIImageOrientation orientation = [newImage imageOrientation]; |
||||
|
||||
// TODO: only apply rotation, don't scale, since we can set this directly in the camera |
||||
/* |
||||
switch (orientation) { |
||||
case UIImageOrientationUp: |
||||
case UIImageOrientationDown: |
||||
newImage = [newImage imageWithAppliedRotationAndMaxSize:CGSizeMake(640.0, 480.0)]; |
||||
break; |
||||
case UIImageOrientationLeft: |
||||
case UIImageOrientationRight: |
||||
newImage = [newImage imageWithMaxSize:CGSizeMake(640.0, 480.0)]; |
||||
default: |
||||
break; |
||||
} |
||||
*/ |
||||
|
||||
// We have captured the image, we can allow the user to take another picture |
||||
cameraAvailable = YES; |
||||
|
||||
NSLog(@"CvPhotoCamera captured image"); |
||||
if (self.delegate) { |
||||
[self.delegate photoCamera:self capturedImage:newImage]; |
||||
} |
||||
|
||||
[self.captureSession startRunning]; |
||||
}); |
||||
} |
||||
}]; |
||||
|
||||
|
||||
} |
||||
|
||||
- (void)stop; |
||||
{ |
||||
[super stop]; |
||||
self.stillImageOutput = nil; |
||||
} |
||||
|
||||
|
||||
#pragma mark - Private Interface |
||||
|
||||
|
||||
- (void)createStillImageOutput; |
||||
{ |
||||
// setup still image output with jpeg codec |
||||
self.stillImageOutput = [[AVCaptureStillImageOutput alloc] init]; |
||||
NSDictionary *outputSettings = [NSDictionary dictionaryWithObjectsAndKeys:AVVideoCodecJPEG, AVVideoCodecKey, nil]; |
||||
[self.stillImageOutput setOutputSettings:outputSettings]; |
||||
[self.captureSession addOutput:self.stillImageOutput]; |
||||
|
||||
for (AVCaptureConnection *connection in self.stillImageOutput.connections) { |
||||
for (AVCaptureInputPort *port in [connection inputPorts]) { |
||||
if ([port.mediaType isEqual:AVMediaTypeVideo]) { |
||||
self.videoCaptureConnection = connection; |
||||
break; |
||||
} |
||||
} |
||||
if (self.videoCaptureConnection) { |
||||
break; |
||||
} |
||||
} |
||||
NSLog(@"[Camera] still image output created"); |
||||
} |
||||
|
||||
|
||||
- (void)createCaptureOutput; |
||||
{ |
||||
[self createStillImageOutput]; |
||||
} |
||||
|
||||
- (void)createCustomVideoPreview; |
||||
{ |
||||
//do nothing, always use AVCaptureVideoPreviewLayer |
||||
} |
||||
|
||||
|
||||
@end |
||||
@ -0,0 +1,585 @@ |
||||
/* |
||||
* cap_ios_video_camera.mm |
||||
* For iOS video I/O |
||||
* by Eduard Feicho on 29/07/12 |
||||
* Copyright 2012. All rights reserved. |
||||
* |
||||
* Redistribution and use in source and binary forms, with or without |
||||
* modification, are permitted provided that the following conditions are met: |
||||
* |
||||
* 1. Redistributions of source code must retain the above copyright notice, |
||||
* this list of conditions and the following disclaimer. |
||||
* 2. Redistributions in binary form must reproduce the above copyright notice, |
||||
* this list of conditions and the following disclaimer in the documentation |
||||
* and/or other materials provided with the distribution. |
||||
* 3. The name of the author may not be used to endorse or promote products |
||||
* derived from this software without specific prior written permission. |
||||
* |
||||
* THIS SOFTWARE IS PROVIDED BY THE AUTHOR "AS IS" AND ANY EXPRESS OR IMPLIED |
||||
* WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF |
||||
* MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO |
||||
* EVENT SHALL THE AUTHOR BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, |
||||
* SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, |
||||
* PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; |
||||
* OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, |
||||
* WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR |
||||
* OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF |
||||
* ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. |
||||
* |
||||
*/ |
||||
|
||||
#import "opencv2/highgui/cap_ios.h" |
||||
#include "precomp.hpp" |
||||
|
||||
#import <AssetsLibrary/AssetsLibrary.h> |
||||
|
||||
|
||||
static CGFloat DegreesToRadians(CGFloat degrees) {return degrees * M_PI / 180;}; |
||||
|
||||
#pragma mark - Private Interface |
||||
|
||||
|
||||
|
||||
|
||||
@interface CvVideoCamera () |
||||
|
||||
- (void)createVideoDataOutput; |
||||
- (void)createVideoFileOutput; |
||||
|
||||
|
||||
@property (nonatomic, retain) CALayer *customPreviewLayer; |
||||
@property (nonatomic, retain) AVCaptureVideoDataOutput *videoDataOutput; |
||||
|
||||
@end |
||||
|
||||
|
||||
|
||||
#pragma mark - Implementation |
||||
|
||||
|
||||
|
||||
@implementation CvVideoCamera |
||||
|
||||
|
||||
|
||||
|
||||
@synthesize delegate; |
||||
@synthesize grayscaleMode; |
||||
|
||||
@synthesize customPreviewLayer; |
||||
@synthesize videoDataOutput; |
||||
|
||||
@synthesize recordVideo; |
||||
//@synthesize videoFileOutput; |
||||
@synthesize recordAssetWriterInput; |
||||
@synthesize recordPixelBufferAdaptor; |
||||
@synthesize recordAssetWriter; |
||||
|
||||
|
||||
|
||||
#pragma mark - Constructors |
||||
|
||||
- (id)initWithParentView:(UIView*)parent; |
||||
{ |
||||
self = [super initWithParentView:parent]; |
||||
if (self) { |
||||
self.useAVCaptureVideoPreviewLayer = NO; |
||||
self.recordVideo = NO; |
||||
} |
||||
return self; |
||||
} |
||||
|
||||
|
||||
|
||||
#pragma mark - Public interface |
||||
|
||||
|
||||
- (void)start; |
||||
{ |
||||
[super start]; |
||||
|
||||
if (self.recordVideo == YES) { |
||||
NSError* error; |
||||
if ([[NSFileManager defaultManager] fileExistsAtPath:[self videoFileString]]) { |
||||
[[NSFileManager defaultManager] removeItemAtPath:[self videoFileString] error:&error]; |
||||
} |
||||
if (error == nil) { |
||||
NSLog(@"[Camera] Delete file %@", [self videoFileString]); |
||||
} |
||||
} |
||||
} |
||||
|
||||
|
||||
|
||||
- (void)stop; |
||||
{ |
||||
[super stop]; |
||||
|
||||
self.videoDataOutput = nil; |
||||
if (videoDataOutputQueue) { |
||||
dispatch_release(videoDataOutputQueue); |
||||
} |
||||
|
||||
if (self.recordVideo == YES) { |
||||
|
||||
if (self.recordAssetWriter.status == AVAssetWriterStatusWriting) { |
||||
[self.recordAssetWriter finishWriting]; |
||||
NSLog(@"[Camera] recording stopped"); |
||||
} else { |
||||
NSLog(@"[Camera] Recording Error: asset writer status is not writing"); |
||||
} |
||||
|
||||
self.recordAssetWriter = nil; |
||||
self.recordAssetWriterInput = nil; |
||||
self.recordPixelBufferAdaptor = nil; |
||||
} |
||||
|
||||
[self.customPreviewLayer removeFromSuperlayer]; |
||||
self.customPreviewLayer = nil; |
||||
} |
||||
|
||||
// TODO fix |
||||
- (void)adjustLayoutToInterfaceOrientation:(UIInterfaceOrientation)interfaceOrientation; |
||||
{ |
||||
|
||||
NSLog(@"layout preview layer"); |
||||
if (self.parentView != nil) { |
||||
|
||||
CALayer* layer = self.customPreviewLayer; |
||||
CGRect bounds = self.customPreviewLayer.bounds; |
||||
int rotation_angle = 0; |
||||
bool flip_bounds = false; |
||||
|
||||
switch (interfaceOrientation) { |
||||
case UIInterfaceOrientationPortrait: |
||||
NSLog(@"to Portrait"); |
||||
rotation_angle = 270; |
||||
break; |
||||
case UIInterfaceOrientationPortraitUpsideDown: |
||||
rotation_angle = 90; |
||||
NSLog(@"to UpsideDown"); |
||||
break; |
||||
case UIInterfaceOrientationLandscapeLeft: |
||||
rotation_angle = 0; |
||||
NSLog(@"to LandscapeLeft"); |
||||
break; |
||||
case UIInterfaceOrientationLandscapeRight: |
||||
rotation_angle = 180; |
||||
NSLog(@"to LandscapeRight"); |
||||
break; |
||||
default: |
||||
break; // leave the layer in its last known orientation |
||||
} |
||||
|
||||
switch (defaultAVCaptureVideoOrientation) { |
||||
case AVCaptureVideoOrientationLandscapeRight: |
||||
rotation_angle += 180; |
||||
break; |
||||
case AVCaptureVideoOrientationPortraitUpsideDown: |
||||
rotation_angle += 270; |
||||
break; |
||||
case AVCaptureVideoOrientationPortrait: |
||||
rotation_angle += 90; |
||||
case AVCaptureVideoOrientationLandscapeLeft: |
||||
break; |
||||
default: |
||||
break; |
||||
} |
||||
rotation_angle = rotation_angle % 360; |
||||
|
||||
if (rotation_angle == 90 || rotation_angle == 270) { |
||||
flip_bounds = true; |
||||
} |
||||
|
||||
if (flip_bounds) { |
||||
NSLog(@"flip bounds"); |
||||
bounds = CGRectMake(0, 0, bounds.size.height, bounds.size.width); |
||||
} |
||||
|
||||
layer.position = CGPointMake(self.parentView.frame.size.width/2., self.parentView.frame.size.height/2.); |
||||
self.customPreviewLayer.bounds = CGRectMake(0, 0, self.parentView.frame.size.width, self.parentView.frame.size.height); |
||||
|
||||
layer.affineTransform = CGAffineTransformMakeRotation( DegreesToRadians(rotation_angle) ); |
||||
layer.bounds = bounds; |
||||
} |
||||
|
||||
} |
||||
|
||||
// TODO fix |
||||
- (void)layoutPreviewLayer; |
||||
{ |
||||
NSLog(@"layout preview layer"); |
||||
if (self.parentView != nil) { |
||||
|
||||
CALayer* layer = self.customPreviewLayer; |
||||
CGRect bounds = self.customPreviewLayer.bounds; |
||||
int rotation_angle = 0; |
||||
bool flip_bounds = false; |
||||
|
||||
switch (currentDeviceOrientation) { |
||||
case UIDeviceOrientationPortrait: |
||||
rotation_angle = 270; |
||||
break; |
||||
case UIDeviceOrientationPortraitUpsideDown: |
||||
rotation_angle = 90; |
||||
break; |
||||
case UIDeviceOrientationLandscapeLeft: |
||||
NSLog(@"left"); |
||||
rotation_angle = 180; |
||||
break; |
||||
case UIDeviceOrientationLandscapeRight: |
||||
NSLog(@"right"); |
||||
rotation_angle = 0; |
||||
break; |
||||
case UIDeviceOrientationFaceUp: |
||||
case UIDeviceOrientationFaceDown: |
||||
default: |
||||
break; // leave the layer in its last known orientation |
||||
} |
||||
|
||||
switch (defaultAVCaptureVideoOrientation) { |
||||
case AVCaptureVideoOrientationLandscapeRight: |
||||
rotation_angle += 180; |
||||
break; |
||||
case AVCaptureVideoOrientationPortraitUpsideDown: |
||||
rotation_angle += 270; |
||||
break; |
||||
case AVCaptureVideoOrientationPortrait: |
||||
rotation_angle += 90; |
||||
case AVCaptureVideoOrientationLandscapeLeft: |
||||
break; |
||||
default: |
||||
break; |
||||
} |
||||
rotation_angle = rotation_angle % 360; |
||||
|
||||
if (rotation_angle == 90 || rotation_angle == 270) { |
||||
flip_bounds = true; |
||||
} |
||||
|
||||
if (flip_bounds) { |
||||
NSLog(@"flip bounds"); |
||||
bounds = CGRectMake(0, 0, bounds.size.height, bounds.size.width); |
||||
} |
||||
|
||||
layer.position = CGPointMake(self.parentView.frame.size.width/2., self.parentView.frame.size.height/2.); |
||||
layer.affineTransform = CGAffineTransformMakeRotation( DegreesToRadians(rotation_angle) ); |
||||
layer.bounds = bounds; |
||||
} |
||||
|
||||
} |
||||
|
||||
|
||||
|
||||
|
||||
#pragma mark - Private Interface |
||||
|
||||
|
||||
|
||||
- (void)createVideoDataOutput; |
||||
{ |
||||
// Make a video data output |
||||
self.videoDataOutput = [AVCaptureVideoDataOutput new]; |
||||
|
||||
// In grayscale mode we want YUV (YpCbCr 4:2:0) so we can directly access the graylevel intensity values (Y component) |
||||
// In color mode we, BGRA format is used |
||||
OSType format = self.grayscaleMode ? kCVPixelFormatType_420YpCbCr8BiPlanarFullRange : kCVPixelFormatType_32BGRA; |
||||
|
||||
self.videoDataOutput.videoSettings = [NSDictionary dictionaryWithObject:[NSNumber numberWithUnsignedInt:format] |
||||
forKey:(id)kCVPixelBufferPixelFormatTypeKey]; |
||||
|
||||
// discard if the data output queue is blocked (as we process the still image) |
||||
[self.videoDataOutput setAlwaysDiscardsLateVideoFrames:YES]; |
||||
|
||||
if ( [self.captureSession canAddOutput:self.videoDataOutput] ) { |
||||
[self.captureSession addOutput:self.videoDataOutput]; |
||||
} |
||||
[[self.videoDataOutput connectionWithMediaType:AVMediaTypeVideo] setEnabled:YES]; |
||||
|
||||
|
||||
// set default FPS |
||||
if ([self.videoDataOutput connectionWithMediaType:AVMediaTypeVideo].supportsVideoMinFrameDuration) { |
||||
[self.videoDataOutput connectionWithMediaType:AVMediaTypeVideo].videoMinFrameDuration = CMTimeMake(1, self.defaultFPS); |
||||
} |
||||
if ([self.videoDataOutput connectionWithMediaType:AVMediaTypeVideo].supportsVideoMaxFrameDuration) { |
||||
[self.videoDataOutput connectionWithMediaType:AVMediaTypeVideo].videoMaxFrameDuration = CMTimeMake(1, self.defaultFPS); |
||||
} |
||||
|
||||
// set video mirroring for front camera (more intuitive) |
||||
if ([self.videoDataOutput connectionWithMediaType:AVMediaTypeVideo].supportsVideoMirroring) { |
||||
if (self.defaultAVCaptureDevicePosition == AVCaptureDevicePositionFront) { |
||||
[self.videoDataOutput connectionWithMediaType:AVMediaTypeVideo].videoMirrored = YES; |
||||
} else { |
||||
[self.videoDataOutput connectionWithMediaType:AVMediaTypeVideo].videoMirrored = NO; |
||||
} |
||||
} |
||||
|
||||
// set default video orientation |
||||
if ([self.videoDataOutput connectionWithMediaType:AVMediaTypeVideo].supportsVideoOrientation) { |
||||
[self.videoDataOutput connectionWithMediaType:AVMediaTypeVideo].videoOrientation = self.defaultAVCaptureVideoOrientation; |
||||
} |
||||
|
||||
|
||||
// create a custom preview layer |
||||
self.customPreviewLayer = [CALayer layer]; |
||||
self.customPreviewLayer.bounds = CGRectMake(0, 0, self.parentView.frame.size.width, self.parentView.frame.size.height); |
||||
[self layoutPreviewLayer]; |
||||
|
||||
// create a serial dispatch queue used for the sample buffer delegate as well as when a still image is captured |
||||
// a serial dispatch queue must be used to guarantee that video frames will be delivered in order |
||||
// see the header doc for setSampleBufferDelegate:queue: for more information |
||||
videoDataOutputQueue = dispatch_queue_create("VideoDataOutputQueue", DISPATCH_QUEUE_SERIAL); |
||||
[self.videoDataOutput setSampleBufferDelegate:self queue:videoDataOutputQueue]; |
||||
|
||||
|
||||
NSLog(@"[Camera] created AVCaptureVideoDataOutput at %d FPS", self.defaultFPS); |
||||
} |
||||
|
||||
|
||||
|
||||
- (void)createVideoFileOutput; |
||||
{ |
||||
/* Video File Output in H.264, via AVAsserWriter */ |
||||
NSLog(@"Create Video with dimensions %dx%d", self.imageWidth, self.imageHeight); |
||||
|
||||
NSDictionary *outputSettings |
||||
= [NSDictionary dictionaryWithObjectsAndKeys:[NSNumber numberWithInt:self.imageWidth], AVVideoWidthKey, |
||||
[NSNumber numberWithInt:self.imageHeight], AVVideoHeightKey, |
||||
AVVideoCodecH264, AVVideoCodecKey, |
||||
nil |
||||
]; |
||||
|
||||
|
||||
self.recordAssetWriterInput = [AVAssetWriterInput assetWriterInputWithMediaType:AVMediaTypeVideo outputSettings:outputSettings]; |
||||
|
||||
|
||||
int pixelBufferFormat = (self.grayscaleMode == YES) ? kCVPixelFormatType_420YpCbCr8BiPlanarFullRange : kCVPixelFormatType_32BGRA; |
||||
|
||||
self.recordPixelBufferAdaptor = |
||||
[[AVAssetWriterInputPixelBufferAdaptor alloc] |
||||
initWithAssetWriterInput:self.recordAssetWriterInput |
||||
sourcePixelBufferAttributes:[NSDictionary dictionaryWithObjectsAndKeys:[NSNumber numberWithInt:pixelBufferFormat], kCVPixelBufferPixelFormatTypeKey, nil]]; |
||||
|
||||
NSError* error = nil; |
||||
NSLog(@"Create AVAssetWriter with url: %@", [self videoFileURL]); |
||||
self.recordAssetWriter = [AVAssetWriter assetWriterWithURL:[self videoFileURL] |
||||
fileType:AVFileTypeMPEG4 |
||||
error:&error]; |
||||
if (error != nil) { |
||||
NSLog(@"[Camera] Unable to create AVAssetWriter: %@", error); |
||||
} |
||||
|
||||
[self.recordAssetWriter addInput:self.recordAssetWriterInput]; |
||||
self.recordAssetWriterInput.expectsMediaDataInRealTime = YES; |
||||
|
||||
NSLog(@"[Camera] created AVAssetWriter"); |
||||
} |
||||
|
||||
|
||||
- (void)createCaptureOutput; |
||||
{ |
||||
[self createVideoDataOutput]; |
||||
if (self.recordVideo == YES) { |
||||
[self createVideoFileOutput]; |
||||
} |
||||
} |
||||
|
||||
- (void)createCustomVideoPreview; |
||||
{ |
||||
[self.parentView.layer addSublayer:self.customPreviewLayer]; |
||||
} |
||||
|
||||
|
||||
#pragma mark - Protocol AVCaptureVideoDataOutputSampleBufferDelegate |
||||
|
||||
|
||||
- (void)captureOutput:(AVCaptureOutput *)captureOutput didOutputSampleBuffer:(CMSampleBufferRef)sampleBuffer fromConnection:(AVCaptureConnection *)connection |
||||
{ |
||||
if (self.delegate) { |
||||
|
||||
// convert from Core Media to Core Video |
||||
CVImageBufferRef imageBuffer = CMSampleBufferGetImageBuffer(sampleBuffer); |
||||
CVPixelBufferLockBaseAddress(imageBuffer, 0); |
||||
|
||||
void* bufferAddress; |
||||
size_t width; |
||||
size_t height; |
||||
size_t bytesPerRow; |
||||
|
||||
CGColorSpaceRef colorSpace; |
||||
CGContextRef context; |
||||
|
||||
int format_opencv; |
||||
|
||||
OSType format = CVPixelBufferGetPixelFormatType(imageBuffer); |
||||
if (format == kCVPixelFormatType_420YpCbCr8BiPlanarFullRange) { |
||||
|
||||
format_opencv = CV_8UC1; |
||||
|
||||
bufferAddress = CVPixelBufferGetBaseAddressOfPlane(imageBuffer, 0); |
||||
width = CVPixelBufferGetWidthOfPlane(imageBuffer, 0); |
||||
height = CVPixelBufferGetHeightOfPlane(imageBuffer, 0); |
||||
bytesPerRow = CVPixelBufferGetBytesPerRowOfPlane(imageBuffer, 0); |
||||
|
||||
} else { // expect kCVPixelFormatType_32BGRA |
||||
|
||||
format_opencv = CV_8UC4; |
||||
|
||||
bufferAddress = CVPixelBufferGetBaseAddress(imageBuffer); |
||||
width = CVPixelBufferGetWidth(imageBuffer); |
||||
height = CVPixelBufferGetHeight(imageBuffer); |
||||
bytesPerRow = CVPixelBufferGetBytesPerRow(imageBuffer); |
||||
|
||||
} |
||||
|
||||
// delegate image processing to the delegate |
||||
cv::Mat image(height, width, format_opencv, bufferAddress, bytesPerRow); |
||||
|
||||
cv::Mat* result = NULL; |
||||
CGImage* dstImage; |
||||
|
||||
if ([self.delegate respondsToSelector:@selector(processImage:)]) { |
||||
[self.delegate processImage:image]; |
||||
} |
||||
|
||||
// check if matrix data pointer or dimensions were changed by the delegate |
||||
bool iOSimage = false; |
||||
if (height == image.rows && width == image.cols && format_opencv == image.type() && bufferAddress == image.data && bytesPerRow == image.step) { |
||||
iOSimage = true; |
||||
} |
||||
|
||||
|
||||
// (create color space, create graphics context, render buffer) |
||||
CGBitmapInfo bitmapInfo; |
||||
|
||||
// basically we decide if it's a grayscale, rgb or rgba image |
||||
if (image.channels() == 1) { |
||||
colorSpace = CGColorSpaceCreateDeviceGray(); |
||||
bitmapInfo = kCGImageAlphaNone; |
||||
} else if (image.channels() == 3) { |
||||
colorSpace = CGColorSpaceCreateDeviceRGB(); |
||||
bitmapInfo = kCGImageAlphaNone; |
||||
if (iOSimage) { |
||||
bitmapInfo |= kCGBitmapByteOrder32Little; |
||||
} else { |
||||
bitmapInfo |= kCGBitmapByteOrder32Big; |
||||
} |
||||
} else { |
||||
colorSpace = CGColorSpaceCreateDeviceRGB(); |
||||
bitmapInfo = kCGImageAlphaPremultipliedFirst; |
||||
if (iOSimage) { |
||||
bitmapInfo |= kCGBitmapByteOrder32Little; |
||||
} else { |
||||
bitmapInfo |= kCGBitmapByteOrder32Big; |
||||
} |
||||
} |
||||
|
||||
if (iOSimage) { |
||||
context = CGBitmapContextCreate(bufferAddress, width, height, 8, bytesPerRow, colorSpace, bitmapInfo); |
||||
dstImage = CGBitmapContextCreateImage(context); |
||||
CGContextRelease(context); |
||||
} else { |
||||
|
||||
NSData *data = [NSData dataWithBytes:image.data length:image.elemSize()*image.total()]; |
||||
CGDataProviderRef provider = CGDataProviderCreateWithCFData((__bridge CFDataRef)data); |
||||
|
||||
// Creating CGImage from cv::Mat |
||||
dstImage = CGImageCreate(image.cols, // width |
||||
image.rows, // height |
||||
8, // bits per component |
||||
8 * image.elemSize(), // bits per pixel |
||||
image.step, // bytesPerRow |
||||
colorSpace, // colorspace |
||||
bitmapInfo, // bitmap info |
||||
provider, // CGDataProviderRef |
||||
NULL, // decode |
||||
false, // should interpolate |
||||
kCGRenderingIntentDefault // intent |
||||
); |
||||
|
||||
CGDataProviderRelease(provider); |
||||
} |
||||
|
||||
|
||||
// render buffer |
||||
dispatch_sync(dispatch_get_main_queue(), ^{ |
||||
self.customPreviewLayer.contents = (__bridge id)dstImage; |
||||
}); |
||||
|
||||
|
||||
if (self.recordVideo == YES) { |
||||
lastSampleTime = CMSampleBufferGetPresentationTimeStamp(sampleBuffer); |
||||
// CMTimeShow(lastSampleTime); |
||||
if (self.recordAssetWriter.status != AVAssetWriterStatusWriting) { |
||||
[self.recordAssetWriter startWriting]; |
||||
[self.recordAssetWriter startSessionAtSourceTime:lastSampleTime]; |
||||
if (self.recordAssetWriter.status != AVAssetWriterStatusWriting) { |
||||
NSLog(@"[Camera] Recording Error: asset writer status is not writing: %@", self.recordAssetWriter.error); |
||||
return; |
||||
} else { |
||||
NSLog(@"[Camera] Video recording started"); |
||||
} |
||||
} |
||||
|
||||
if (self.recordAssetWriterInput.readyForMoreMediaData) { |
||||
if (! [self.recordPixelBufferAdaptor appendPixelBuffer:imageBuffer |
||||
withPresentationTime:lastSampleTime] ) { |
||||
NSLog(@"Video Writing Error"); |
||||
} |
||||
} |
||||
|
||||
} |
||||
|
||||
|
||||
// cleanup |
||||
CGImageRelease(dstImage); |
||||
|
||||
CGColorSpaceRelease(colorSpace); |
||||
|
||||
CVPixelBufferUnlockBaseAddress(imageBuffer, 0); |
||||
} |
||||
} |
||||
|
||||
|
||||
- (void)updateOrientation; |
||||
{ |
||||
NSLog(@"rotate.."); |
||||
self.customPreviewLayer.bounds = CGRectMake(0, 0, self.parentView.frame.size.width, self.parentView.frame.size.height); |
||||
[self layoutPreviewLayer]; |
||||
} |
||||
|
||||
|
||||
- (void)saveVideo; |
||||
{ |
||||
if (self.recordVideo == NO) { |
||||
return; |
||||
} |
||||
|
||||
ALAssetsLibrary *library = [[ALAssetsLibrary alloc] init]; |
||||
if ([library videoAtPathIsCompatibleWithSavedPhotosAlbum:[self videoFileURL]]) { |
||||
[library writeVideoAtPathToSavedPhotosAlbum:[self videoFileURL] |
||||
completionBlock:^(NSURL *assetURL, NSError *error){}]; |
||||
} |
||||
} |
||||
|
||||
|
||||
- (NSURL *)videoFileURL; |
||||
{ |
||||
NSString *outputPath = [[NSString alloc] initWithFormat:@"%@%@", NSTemporaryDirectory(), @"output.mov"]; |
||||
NSURL *outputURL = [NSURL fileURLWithPath:outputPath]; |
||||
NSFileManager *fileManager = [NSFileManager defaultManager]; |
||||
if ([fileManager fileExistsAtPath:outputPath]) { |
||||
NSLog(@"file exists"); |
||||
} |
||||
return outputURL; |
||||
} |
||||
|
||||
|
||||
|
||||
- (NSString *)videoFileString; |
||||
{ |
||||
NSString *outputPath = [[NSString alloc] initWithFormat:@"%@%@", NSTemporaryDirectory(), @"output.mov"]; |
||||
return outputPath; |
||||
} |
||||
|
||||
@end |
||||
@ -0,0 +1,68 @@ |
||||
#include "perf_precomp.hpp" |
||||
|
||||
using namespace std; |
||||
using namespace cv; |
||||
using namespace perf; |
||||
using namespace testing; |
||||
using std::tr1::make_tuple; |
||||
using std::tr1::get; |
||||
|
||||
CV_ENUM(MatrixType, CV_16UC1, CV_16SC1, CV_32FC1) |
||||
CV_ENUM(MapType, CV_16SC2, CV_32FC1, CV_32FC2) |
||||
CV_ENUM(InterType, INTER_LINEAR, INTER_CUBIC, INTER_LANCZOS4, INTER_NEAREST) |
||||
|
||||
typedef TestBaseWithParam< tr1::tuple<Size, MatrixType, MapType, InterType> > TestRemap; |
||||
|
||||
PERF_TEST_P( TestRemap, Remap, |
||||
Combine( |
||||
Values( szVGA, sz1080p ),
|
||||
ValuesIn( MatrixType::all() ),
|
||||
ValuesIn( MapType::all() ),
|
||||
ValuesIn( InterType::all() )
|
||||
) |
||||
) |
||||
{ |
||||
Size sz; |
||||
int src_type, map1_type, inter_type; |
||||
|
||||
sz = get<0>(GetParam()); |
||||
src_type = get<1>(GetParam()); |
||||
map1_type = get<2>(GetParam()); |
||||
inter_type = get<3>(GetParam()); |
||||
|
||||
Mat src(sz, src_type); |
||||
Mat map1(sz, map1_type); |
||||
Mat dst(sz, src_type); |
||||
|
||||
Mat map2(map1_type == CV_32FC1 ? sz : Size(), CV_32FC1); |
||||
|
||||
RNG rng; |
||||
rng.fill(src, RNG::UNIFORM, 0, 256); |
||||
|
||||
for (int j = 0; j < map1.rows; ++j) |
||||
for (int i = 0; i < map1.cols; ++i) |
||||
switch (map1_type) |
||||
{ |
||||
case CV_32FC1: |
||||
map1.at<float>(j, i) = src.cols - i; |
||||
map2.at<float>(j, i) = j; |
||||
break; |
||||
case CV_32FC2: |
||||
map1.at<Vec2f>(j, i)[0] = src.cols - i; |
||||
map1.at<Vec2f>(j, i)[1] = j; |
||||
break; |
||||
case CV_16SC2: |
||||
map1.at<Vec2s>(j, i)[0] = src.cols - i; |
||||
map1.at<Vec2s>(j, i)[1] = j; |
||||
break; |
||||
default: |
||||
CV_Assert(0); |
||||
} |
||||
|
||||
|
||||
declare.in(src, WARMUP_RNG).out(dst).time(20); |
||||
|
||||
TEST_CYCLE() remap(src, dst, map1, map2, inter_type); |
||||
|
||||
SANITY_CHECK(dst); |
||||
} |
||||
File diff suppressed because it is too large
Load Diff
Some files were not shown because too many files have changed in this diff Show More
Loading…
Reference in new issue