virtual ~SCascade();
virtual bool load(const FileNode& fn);
virtual void detect(InputArray image, InputArray rois, OutputArray objects, Stream& stream = Stream::Null()) const;
- virtual void detect(InputArray image, InputArray rois, OutputArray objects, const int level, Stream& stream = Stream::Null()) const;
void genRoi(InputArray roi, OutputArray mask, Stream& stream = Stream::Null()) const;
};
Apply cascade to an input frame and return the vector of Decection objcts.
.. ocv:function:: void detect(InputArray image, InputArray rois, OutputArray objects, Stream& stream = Stream::Null()) const
-.. ocv:function:: void detect(InputArray image, InputArray rois, OutputArray objects, const int level, Stream& stream = Stream::Null()) const
:param image: a frame on which detector will be applied.
:param stream: a high-level CUDA stream abstraction used for asynchronous execution.
- :param level: used for execution cascade on specific scales pyramid level.
-
SCascade::genRoi
--------------------------
// Param objects is an output array of Detections represented as GpuMat of detections (SCascade::Detection)
// The first element of the matrix is actually a count of detections.
// Param stream is stream is a high-level CUDA stream abstraction used for asynchronous execution
- // Param level used for execution cascade on specific scales pyramid level.
virtual void detect(InputArray image, InputArray rois, OutputArray objects, Stream& stream = Stream::Null()) const;
- virtual void detect(InputArray image, InputArray rois, OutputArray objects, const int level, Stream& stream = Stream::Null()) const;
// Convert ROI matrix into the suitable for detect method.
// Param roi is an input matrix of the same size as the image.
#include <stdio.h>
#include <float.h>
-// #define LOG_CUDA_CASCADE
-
-#if defined LOG_CUDA_CASCADE
-# define dprintf(format, ...) \
- do { printf(format, __VA_ARGS__); } while (0)
-#else
-# define dprintf(format, ...)
-#endif
-
namespace cv { namespace gpu { namespace device {
namespace icf {
template<typename Policy>
void CascadeInvoker<Policy>::operator()(const PtrStepSzb& roi, const PtrStepSzi& hogluv,
- PtrStepSz<uchar4> objects, PtrStepSzi counter, const int downscales, const int scale, const cudaStream_t& stream) const
+ PtrStepSz<uchar4> objects, PtrStepSzi counter, const int downscales, const cudaStream_t& stream) const
{
int fw = roi.rows;
int fh = roi.cols;
- dim3 grid(fw, fh / Policy::STA_Y, (scale == -1) ? downscales : 1);
+ dim3 grid(fw, fh / Policy::STA_Y, downscales);
uint* ctr = (uint*)(counter.ptr(0));
Detection* det = (Detection*)objects.ptr();
cudaChannelFormatDesc desc = cudaCreateChannelDesc<int>();
cudaSafeCall( cudaBindTexture2D(0, thogluv, hogluv.data, desc, hogluv.cols, hogluv.rows, hogluv.step));
- cudaChannelFormatDesc desc_roi = cudaCreateChannelDesc<float2>();
- cudaSafeCall( cudaBindTexture2D(0, troi, roi.data, desc_roi, roi.cols / 8, roi.rows, roi.step));
+ cudaChannelFormatDesc desc_roi = cudaCreateChannelDesc<typename Policy::roi_type>();
+ cudaSafeCall( cudaBindTexture2D(0, troi, roi.data, desc_roi, roi.cols / Policy::STA_Y, roi.rows, roi.step));
const CascadeInvoker<Policy> inv = *this;
- if (scale == -1)
- {
- soft_cascade<Policy, false><<<grid, Policy::block(), 0, stream>>>(inv, det, max_det, ctr, 0);
- cudaSafeCall( cudaGetLastError());
+ soft_cascade<Policy, false><<<grid, Policy::block(), 0, stream>>>(inv, det, max_det, ctr, 0);
+ cudaSafeCall( cudaGetLastError());
- grid = dim3(fw, fh / Policy::STA_Y, scales - downscales);
- soft_cascade<Policy, true><<<grid, Policy::block(), 0, stream>>>(inv, det, max_det, ctr, downscales);
- }
- else
- {
- if (scale >= downscales)
- soft_cascade<Policy, true><<<grid, Policy::block(), 0, stream>>>(inv, det, max_det, ctr, scale);
- else
- soft_cascade<Policy, false><<<grid, Policy::block(), 0, stream>>>(inv, det, max_det, ctr, scale);
- }
+ grid = dim3(fw, fh / Policy::STA_Y, scales - downscales);
+ soft_cascade<Policy, true><<<grid, Policy::block(), 0, stream>>>(inv, det, max_det, ctr, downscales);
if (!stream)
{
}
template void CascadeInvoker<GK107PolicyX4>::operator()(const PtrStepSzb& roi, const PtrStepSzi& hogluv,
- PtrStepSz<uchar4> objects, PtrStepSzi counter, const int downscales, const int scale, const cudaStream_t& stream) const;
+ PtrStepSz<uchar4> objects, PtrStepSzi counter, const int downscales, const cudaStream_t& stream) const;
}
}}}
\ No newline at end of file
struct GK107PolicyX4
{
enum {WARP = 32, STA_X = WARP, STA_Y = 8, SHRINKAGE = 4};
+ typedef float2 roi_type;
static const dim3 block()
{
- return dim3(GK107PolicyX4::STA_X, GK107PolicyX4::STA_Y);
+ return dim3(STA_X, STA_Y);
}
};
int scales;
void operator()(const PtrStepSzb& roi, const PtrStepSzi& hogluv, PtrStepSz<uchar4> objects,
- PtrStepSzi counter, const int downscales, const int csale = -1, const cudaStream_t& stream = 0) const;
+ PtrStepSzi counter, const int downscales, const cudaStream_t& stream = 0) const;
template<bool isUp>
__device void detect(Detection* objects, const uint ndetections, uint* ctr, const int downscales) const;
leaves.upload(hleaves);
}
- void detect(int scale, const cv::gpu::GpuMat& roi, const cv::gpu::GpuMat& count, cv::gpu::GpuMat& objects, const cudaStream_t& stream) const
+ void detect(const cv::gpu::GpuMat& roi, const cv::gpu::GpuMat& count, cv::gpu::GpuMat& objects, const cudaStream_t& stream) const
{
cudaMemset(count.data, 0, sizeof(Detection));
cudaSafeCall( cudaGetLastError());
device::icf::CascadeInvoker<device::icf::GK107PolicyX4> invoker
= device::icf::CascadeInvoker<device::icf::GK107PolicyX4>(levels, octaves, stages, nodes, leaves);
- invoker(roi, hogluv, objects, count, downscales, scale, stream);
+ invoker(roi, hogluv, objects, count, downscales, stream);
}
void preprocess(const cv::gpu::GpuMat& colored, Stream& s)
objects = GpuMat(objects, cv::Rect( sizeof(Detection), 0, objects.cols - sizeof(Detection), 1));
cudaStream_t stream = StreamAccessor::getStream(s);
- flds.detect(-1, rois, tmp, objects, stream);
-}
-
-void cv::gpu::SCascade::detect(InputArray image, InputArray _rois, OutputArray _objects, const int level, Stream& s) const
-{
- CV_Assert(fields);
-
- const GpuMat colored = image.getGpuMat();
- // only color images are supperted
- CV_Assert(colored.type() == CV_8UC3 || colored.type() == CV_32SC1);
-
- Fields& flds = *fields;
- if (colored.type() == CV_8UC3)
- {
- // only this window size allowed
- // CV_Assert(colored.cols == Fields::FRAME_WIDTH && colored.rows == Fields::FRAME_HEIGHT);
- flds.preprocess(colored, s);
- }
- else
- {
- colored.copyTo(flds.hogluv);
- }
-
- GpuMat rois = _rois.getGpuMat(), objects = _objects.getGpuMat();
-
- GpuMat tmp = GpuMat(objects, cv::Rect(0, 0, sizeof(Detection), 1));
- objects = GpuMat(objects, cv::Rect( sizeof(Detection), 0, objects.cols - sizeof(Detection), 1));
- cudaStream_t stream = StreamAccessor::getStream(s);
-
- flds.detect(level, rois, tmp, objects, stream);
+ flds.detect(rois, tmp, objects, stream);
}
void cv::gpu::SCascade::genRoi(InputArray _roi, OutputArray _mask, Stream& stream) const
}
-typedef ::testing::TestWithParam<std::tr1::tuple<cv::gpu::DeviceInfo, std::string, std::string, int> > SCascadeTestLevel;
-GPU_TEST_P(SCascadeTestLevel, detect,
- testing::Combine(
- ALL_DEVICES,
- testing::Values(std::string("cv/cascadeandhog/sc_cvpr_2012_to_opencv.xml")),
- testing::Values(std::string("../cv/cascadeandhog/bahnhof/image_00000000_0.png")),
- testing::Range(0, 47)
- ))
-{
- cv::gpu::setDevice(GET_PARAM(0).deviceID());
-
- cv::gpu::SCascade cascade;
-
- cv::FileStorage fs(perf::TestBase::getDataPath(GET_PARAM(1)), cv::FileStorage::READ);
- ASSERT_TRUE(fs.isOpened());
-
- ASSERT_TRUE(cascade.load(fs.getFirstTopLevelNode()));
-
- cv::Mat coloredCpu = cv::imread(cvtest::TS::ptr()->get_data_path() + GET_PARAM(2));
- ASSERT_FALSE(coloredCpu.empty());
-
- typedef cv::gpu::SCascade::Detection Detection;
- GpuMat colored(coloredCpu), objectBoxes(1, 100 * sizeof(Detection), CV_8UC1), rois(colored.size(), CV_8UC1);
- rois.setTo(1);
-
- cv::gpu::GpuMat trois;
- cascade.genRoi(rois, trois);
- objectBoxes.setTo(0);
- int level = GET_PARAM(3);
- cascade.detect(colored, trois, objectBoxes, level);
-
- cv::Mat dt(objectBoxes);
-
- Detection* dts = ((Detection*)dt.data) + 1;
- int* count = dt.ptr<int>(0);
-
- cv::Mat result(coloredCpu);
-
- printTotal(std::cout, *count);
- for (int i = 0; i < *count; ++i)
- {
- Detection d = dts[i];
- print(std::cout, d);
- cv::rectangle(result, cv::Rect(d.x, d.y, d.w, d.h), cv::Scalar(255, 0, 0, 255), 1);
- }
-
- writeResult(result, level);
- SHOW(result);
-}
+// typedef ::testing::TestWithParam<std::tr1::tuple<cv::gpu::DeviceInfo, std::string, std::string, int> > SCascadeTestLevel;
+// GPU_TEST_P(SCascadeTestLevel, detect,
+// testing::Combine(
+// ALL_DEVICES,
+// testing::Values(std::string("cv/cascadeandhog/sc_cvpr_2012_to_opencv.xml")),
+// testing::Values(std::string("../cv/cascadeandhog/bahnhof/image_00000000_0.png")),
+// testing::Range(0, 47)
+// ))
+// {
+// cv::gpu::setDevice(GET_PARAM(0).deviceID());
+
+// cv::gpu::SCascade cascade;
+
+// cv::FileStorage fs(perf::TestBase::getDataPath(GET_PARAM(1)), cv::FileStorage::READ);
+// ASSERT_TRUE(fs.isOpened());
+
+// ASSERT_TRUE(cascade.load(fs.getFirstTopLevelNode()));
+
+// cv::Mat coloredCpu = cv::imread(cvtest::TS::ptr()->get_data_path() + GET_PARAM(2));
+// ASSERT_FALSE(coloredCpu.empty());
+
+// typedef cv::gpu::SCascade::Detection Detection;
+// GpuMat colored(coloredCpu), objectBoxes(1, 100 * sizeof(Detection), CV_8UC1), rois(colored.size(), CV_8UC1);
+// rois.setTo(1);
+
+// cv::gpu::GpuMat trois;
+// cascade.genRoi(rois, trois);
+// objectBoxes.setTo(0);
+// int level = GET_PARAM(3);
+// cascade.detect(colored, trois, objectBoxes, level);
+
+// cv::Mat dt(objectBoxes);
+
+// Detection* dts = ((Detection*)dt.data) + 1;
+// int* count = dt.ptr<int>(0);
+
+// cv::Mat result(coloredCpu);
+
+// printTotal(std::cout, *count);
+// for (int i = 0; i < *count; ++i)
+// {
+// Detection d = dts[i];
+// print(std::cout, d);
+// cv::rectangle(result, cv::Rect(d.x, d.y, d.w, d.h), cv::Scalar(255, 0, 0, 255), 1);
+// }
+
+// writeResult(result, level);
+// SHOW(result);
+// }
TEST(SCascadeTest, readCascade)
{