Commit 4881205b authored by marina.kolpakova's avatar marina.kolpakova

refactor logs

parent dca27b46
...@@ -47,6 +47,13 @@ ...@@ -47,6 +47,13 @@
//#define LOG_CUDA_CASCADE //#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 cv { namespace gpu { namespace device {
namespace icf { namespace icf {
...@@ -98,10 +105,8 @@ __global__ void detect(const cv::gpu::icf::Cascade cascade, const int* __restric ...@@ -98,10 +105,8 @@ __global__ void detect(const cv::gpu::icf::Cascade cascade, const int* __restric
float __device icf::Cascade::rescale(const icf::Level& level, uchar4& scaledRect, float __device icf::Cascade::rescale(const icf::Level& level, uchar4& scaledRect,
const int channel, const float threshold) const const int channel, const float threshold) const
{ {
#if defined LOG_CUDA_CASCADE dprintf("feature %d box %d %d %d %d\n", channel, scaledRect.x, scaledRect.y, scaledRect.z, scaledRect.w);
printf("feature %d box %d %d %d %d\n", channel, scaledRect.x, scaledRect.y, scaledRect.z, scaledRect.w); dprintf("rescale: %f [%f %f]\n",level.relScale, level.scaling[0], level.scaling[1]);
printf("rescale: %f [%f %f]\n",level.relScale, level.scaling[0], level.scaling[1]);
#endif
float relScale = level.relScale; float relScale = level.relScale;
float farea = (scaledRect.z - scaledRect.x) * (scaledRect.w - scaledRect.y); float farea = (scaledRect.z - scaledRect.x) * (scaledRect.w - scaledRect.y);
...@@ -122,19 +127,15 @@ float __device icf::Cascade::rescale(const icf::Level& level, uchar4& scaledRect ...@@ -122,19 +127,15 @@ float __device icf::Cascade::rescale(const icf::Level& level, uchar4& scaledRect
approx = expected_new_area / sarea; approx = expected_new_area / sarea;
} }
#if defined LOG_CUDA_CASCADE dprintf("new rect: %d box %d %d %d %d rel areas %f %f\n", channel,
printf("new rect: %d box %d %d %d %d rel areas %f %f\n", channel,
scaledRect.x, scaledRect.y, scaledRect.z, scaledRect.w, farea * relScale * relScale, sarea); scaledRect.x, scaledRect.y, scaledRect.z, scaledRect.w, farea * relScale * relScale, sarea);
#endif
// compensation areas rounding // compensation areas rounding
float rootThreshold = threshold / approx; float rootThreshold = threshold / approx;
// printf(" approx %f\n", rootThreshold); // printf(" approx %f\n", rootThreshold);
rootThreshold *= level.scaling[(int)(channel > 6)]; rootThreshold *= level.scaling[(int)(channel > 6)];
#if defined LOG_CUDA_CASCADE dprintf("approximation %f %f -> %f %f\n", approx, threshold, rootThreshold, level.scaling[(int)(channel > 6)]);
printf("approximation %f %f -> %f %f\n", approx, threshold, rootThreshold, level.scaling[(int)(channel > 6)]);
#endif
return rootThreshold; return rootThreshold;
} }
...@@ -143,14 +144,12 @@ typedef unsigned char uchar; ...@@ -143,14 +144,12 @@ typedef unsigned char uchar;
float __device get(const int* __restrict__ hogluv, const int pitch, float __device get(const int* __restrict__ hogluv, const int pitch,
const int x, const int y, int channel, uchar4 area) const int x, const int y, int channel, uchar4 area)
{ {
#if defined LOG_CUDA_CASCADE dprintf("feature box %d %d %d %d ", area.x, area.y, area.z, area.w);
printf("feature box %d %d %d %d ", area.x, area.y, area.z, area.w); dprintf("get for channel %d\n", channel);
printf("get for channel %d\n", channel); dprintf("extract feature for: [%d %d] [%d %d] [%d %d] [%d %d]\n",
printf("extract feature for: [%d %d] [%d %d] [%d %d] [%d %d]\n",
x + area.x, y + area.y, x + area.z, y + area.y, x + area.z,y + area.w, x + area.x, y + area.y, x + area.z, y + area.y, x + area.z,y + area.w,
x + area.x, y + area.w); x + area.x, y + area.w);
printf("at point %d %d with offset %d\n", x, y, 0); dprintf("at point %d %d with offset %d\n", x, y, 0);
#endif
const int* curr = hogluv + ((channel * 121) + y) * pitch; const int* curr = hogluv + ((channel * 121) + y) * pitch;
...@@ -159,9 +158,7 @@ float __device get(const int* __restrict__ hogluv, const int pitch, ...@@ -159,9 +158,7 @@ float __device get(const int* __restrict__ hogluv, const int pitch,
int c = curr[area.w * pitch + x + area.z]; int c = curr[area.w * pitch + x + area.z];
int d = curr[area.w * pitch + x + area.x]; int d = curr[area.w * pitch + x + area.x];
#if defined LOG_CUDA_CASCADE dprintf(" retruved integral values: %d %d %d %d\n", a, b, c, d);
printf(" retruved integral values: %d %d %d %d\n", a, b, c, d);
#endif
return (a - b + c - d); return (a - b + c - d);
} }
...@@ -176,13 +173,11 @@ void __device icf::Cascade::detectAt(const int* __restrict__ hogluv, const int p ...@@ -176,13 +173,11 @@ void __device icf::Cascade::detectAt(const int* __restrict__ hogluv, const int p
const int x = blockIdx.x * blockDim.x + threadIdx.x; const int x = blockIdx.x * blockDim.x + threadIdx.x;
// if (x > 0 || y > 0) return; // if (x > 0 || y > 0) return;
Level level = lls[0]; Level level = lls[blockIdx.z];
if (x >= level.workRect.x || y >= level.workRect.y) return; if (x >= level.workRect.x || y >= level.workRect.y) return;
#if defined LOG_CUDA_CASCADE dprintf("level: %d (%f %f) [%f %f] (%d %d) (%d %d)\n", level.octave, level.relScale, level.shrScale,
printf("level: %d (%f %f) [%f %f] (%d %d) (%d %d)\n", level.octave, level.relScale, level.shrScale,
level.scaling[0], level.scaling[1], level.workRect.x, level.workRect.y, level.objSize.x, level.objSize.y); level.scaling[0], level.scaling[1], level.workRect.x, level.workRect.y, level.objSize.x, level.objSize.y);
#endif
const Octave octave = ((const Octave*)octaves.ptr())[level.octave]; const Octave octave = ((const Octave*)octaves.ptr())[level.octave];
// printf("Octave: %d %d %d (%d %d) %f\n", octave.index, octave.stages, // printf("Octave: %d %d %d (%d %d) %f\n", octave.index, octave.stages,
...@@ -196,17 +191,15 @@ void __device icf::Cascade::detectAt(const int* __restrict__ hogluv, const int p ...@@ -196,17 +191,15 @@ void __device icf::Cascade::detectAt(const int* __restrict__ hogluv, const int p
for(; st < stEnd; ++st) for(; st < stEnd; ++st)
{ {
const float stage = stages(0, st); const float stage = stages(0, st);
#if defined LOG_CUDA_CASCADE dprintf("Stage: %f\n", stage);
printf("Stage: %f\n", stage);
#endif
{ {
const int nId = st * 3; const int nId = st * 3;
// work with root node // work with root node
const Node node = ((const Node*)nodes.ptr())[nId]; const Node node = ((const Node*)nodes.ptr())[nId];
#if defined LOG_CUDA_CASCADE
printf("Node: %d %f\n", node.feature, node.threshold); dprintf("Node: %d %f\n", node.feature, node.threshold);
#endif
const Feature feature = ((const Feature*)features.ptr())[node.feature]; const Feature feature = ((const Feature*)features.ptr())[node.feature];
uchar4 scaledRect = feature.rect; uchar4 scaledRect = feature.rect;
...@@ -214,14 +207,12 @@ void __device icf::Cascade::detectAt(const int* __restrict__ hogluv, const int p ...@@ -214,14 +207,12 @@ void __device icf::Cascade::detectAt(const int* __restrict__ hogluv, const int p
float sum = get(hogluv,pitch, x, y, feature.channel, scaledRect); float sum = get(hogluv,pitch, x, y, feature.channel, scaledRect);
#if defined LOG_CUDA_CASCADE dprintf("root feature %d %f\n",feature.channel, sum);
printf("root feature %d %f\n",feature.channel, sum);
#endif
int next = 1 + (int)(sum >= threshold); int next = 1 + (int)(sum >= threshold);
#if defined LOG_CUDA_CASCADE dprintf("go: %d (%f >= %f)\n\n" ,next, sum, threshold);
printf("go: %d (%f >= %f)\n\n" ,next, sum, threshold);
#endif
// leaves // leaves
const Node leaf = ((const Node*)nodes.ptr())[nId + next]; const Node leaf = ((const Node*)nodes.ptr())[nId + next];
const Feature fLeaf = ((const Feature*)features.ptr())[leaf.feature]; const Feature fLeaf = ((const Feature*)features.ptr())[leaf.feature];
...@@ -235,22 +226,17 @@ void __device icf::Cascade::detectAt(const int* __restrict__ hogluv, const int p ...@@ -235,22 +226,17 @@ void __device icf::Cascade::detectAt(const int* __restrict__ hogluv, const int p
detectionScore += impact; detectionScore += impact;
#if defined LOG_CUDA_CASCADE dprintf("decided: %d (%f >= %f) %d %f\n\n" ,next, sum, threshold, lShift, impact);
printf("decided: %d (%f >= %f) %d %f\n\n" ,next, sum, threshold, lShift, impact); dprintf("extracted stage:\n");
printf("extracted stage:\n"); dprintf("ct %f\n", stage);
printf("ct %f\n", stage); dprintf("computed score %f\n\n", detectionScore);
printf("computed score %f\n\n", detectionScore); dprintf("\n\n");
printf("\n\n");
#endif
} }
if (detectionScore <= stage) break; if (detectionScore <= stage || st - stBegin == 100) break;
} }
#if defined LOG_CUDA_CASCADE dprintf("x %d y %d: %d\n", x, y, st - stBegin);
// printf("x %d y %d: %d\n", x, y, st - stBegin);
#endif
if (st == stEnd) if (st == stEnd)
{ {
...@@ -264,7 +250,7 @@ void __device icf::Cascade::detectAt(const int* __restrict__ hogluv, const int p ...@@ -264,7 +250,7 @@ void __device icf::Cascade::detectAt(const int* __restrict__ hogluv, const int p
void icf::Cascade::detect(const cv::gpu::PtrStepSzi& hogluv, PtrStepSz<uchar4> objects, cudaStream_t stream) const void icf::Cascade::detect(const cv::gpu::PtrStepSzi& hogluv, PtrStepSz<uchar4> objects, cudaStream_t stream) const
{ {
dim3 block(32, 8, 1); dim3 block(32, 8, 1);
dim3 grid(ChannelStorage::FRAME_WIDTH / 32, ChannelStorage::FRAME_HEIGHT / 8, 1); dim3 grid(ChannelStorage::FRAME_WIDTH / 32, ChannelStorage::FRAME_HEIGHT / 8, 47);
device::detect<<<grid, block, 0, stream>>>(*this, hogluv, hogluv.step / sizeof(int), objects); device::detect<<<grid, block, 0, stream>>>(*this, hogluv, hogluv.step / sizeof(int), objects);
cudaSafeCall( cudaGetLastError() ); cudaSafeCall( cudaGetLastError() );
if (!stream) if (!stream)
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment