summary | shortlog | log | commit | commitdiff | tree
raw | patch | inline | side by side (parent: 726df4b)
raw | patch | inline | side by side (parent: 726df4b)
author | Djordje Senicic <d-senicic1@ti.com> | |
Thu, 1 Sep 2016 14:52:54 +0000 (10:52 -0400) | ||
committer | Djordje Senicic <d-senicic1@ti.com> | |
Thu, 1 Sep 2016 14:52:54 +0000 (10:52 -0400) |
modules/video/src/bgfg_gaussmix2.cpp | patch | blob | history | |
modules/video/src/opencl/bgfg_mog2.cl | patch | blob | history |
index 136360a626e54d1d02631fbc29866aef1126cbbc..9b5ca4bd210500c6d75c09f54a13f06254312939 100644 (file)
{
create_ocl_apply_kernel();
#ifdef CV_TIOPENCL
- //kernel_getBg.create("getBackgroundImage2_kernel", ocl::video::bgfg_mog2_oclsrc, format( "-v -D CN=%d -D NMIXTURES=%d -DHAVE_TIMOG2LIB %s", nchannels, nmixtures,
- // bShadowDetection ? "-DSHADOW_DETECT" : "" ));
- kernel_getBg.create("getBackgroundImage2_kernel", ocl::video::bgfg_mog2_oclsrc, format( "-D CN=%d -D NMIXTURES=%d %s -DTIDSP_MOG2 ", nchannels, nmixtures,
- bShadowDetection ? "-DSHADOW_DETECT" : "" ));
-#else
- kernel_getBg.create("getBackgroundImage2_kernel", ocl::video::bgfg_mog2_oclsrc, format( "-D CN=%d -D NMIXTURES=%d %s", nchannels, nmixtures,
- bShadowDetection ? "-DSHADOW_DETECT" : "" ));
+ int do_tidsp = ((u_bgmodelUsedModes.cols % 8) == 0);
+ int subline_cache = 8;
+ if(do_tidsp) {
+ if((u_bgmodelUsedModes.cols % 128) == 0) subline_cache = 128;
+ else if((u_bgmodelUsedModes.cols % 64) == 0) subline_cache = 64;
+ else if((u_bgmodelUsedModes.cols % 32) == 0) subline_cache = 32;
+ else if((u_bgmodelUsedModes.cols % 16) == 0) subline_cache = 16;
+ //kernel_getBg.create("getBackgroundImage2_kernel", ocl::video::bgfg_mog2_oclsrc, format( "-v -D CN=%d -D NMIXTURES=%d -DHAVE_TIMOG2LIB %s", nchannels, nmixtures,
+ // bShadowDetection ? "-DSHADOW_DETECT" : "" ));
+ kernel_getBg.create("getBackgroundImage2_kernel", ocl::video::bgfg_mog2_oclsrc, format( "-D CN=%d -D NMIXTURES=%d %s -DTIDSP_MOG2 -D SUBLINE_CACHE=%d ", nchannels, nmixtures,
+ bShadowDetection ? "-DSHADOW_DETECT" : "", subline_cache));
+ } else
#endif
+ {
+ kernel_getBg.create("getBackgroundImage2_kernel", ocl::video::bgfg_mog2_oclsrc, format( "-D CN=%d -D NMIXTURES=%d %s", nchannels, nmixtures,
+ bShadowDetection ? "-DSHADOW_DETECT" : "" ));
+ }
if (kernel_apply.empty() || kernel_getBg.empty())
opencl_ON = false;
}
{
int nchannels = CV_MAT_CN(frameType);
#ifdef CV_TIOPENCL
+ int do_tidsp = ((u_bgmodelUsedModes.cols % 8) == 0);
+ int subline_cache = 8;
+ if(do_tidsp) {
+ if((u_bgmodelUsedModes.cols % 128) == 0) subline_cache = 128;
+ else if((u_bgmodelUsedModes.cols % 64) == 0) subline_cache = 64;
+ else if((u_bgmodelUsedModes.cols % 32) == 0) subline_cache = 32;
+ else if((u_bgmodelUsedModes.cols % 16) == 0) subline_cache = 16;
// Macro HAVE_TIMOG2LIB would trigger linking with external DSP library (as specified in core/src/ocl.cpp)
// String opts = format("-v -D CN=%d -D NMIXTURES=%d%s -DHAVE_TIMOG2LIB -DTIDSP_MOG2", nchannels, nmixtures, bShadowDetection ? " -DSHADOW_DETECT" : "");
- String opts = format("-D CN=%d -D NMIXTURES=%d%s -DTIDSP_MOG2 ", nchannels, nmixtures, bShadowDetection ? " -DSHADOW_DETECT" : "");
- kernel_apply.create("mog2_kernel", ocl::video::bgfg_mog2_oclsrc, opts);
-#else
- String opts = format("-D CN=%d -D NMIXTURES=%d%s", nchannels, nmixtures, bShadowDetection ? " -D SHADOW_DETECT" : "");
- kernel_apply.create("mog2_kernel", ocl::video::bgfg_mog2_oclsrc, opts);
+ String opts = format("-D CN=%d -D NMIXTURES=%d%s -DTIDSP_MOG2 -D SUBLINE_CACHE=%d", nchannels, nmixtures, bShadowDetection ? " -DSHADOW_DETECT" : "", subline_cache);
+ kernel_apply.create("mog2_kernel", ocl::video::bgfg_mog2_oclsrc, opts);
+ } else
#endif
+ {
+ String opts = format("-D CN=%d -D NMIXTURES=%d%s", nchannels, nmixtures, bShadowDetection ? " -D SHADOW_DETECT" : "");
+ kernel_apply.create("mog2_kernel", ocl::video::bgfg_mog2_oclsrc, opts);
+ }
}
#endif
index b8e8f67e40d20a0e0cbed2277be481457437ac41..4e1f210d7231c0ef287b7370e367387ca98163a9 100644 (file)
#include <dsp_c.h>
#include <edmamgr.h>
-#define SUBLINE_CACHE 64
__kernel void mog2_kernel(__global const uchar* frame, int frame_step, int frame_offset, int frame_row, int frame_col, //uchar || uchar3
__global uchar* modesUsed, //uchar
__global uchar* weight, //float
EdmaMgr_wait(evOUT);
EdmaMgr_free(evIN);
EdmaMgr_free(evOUT);
- printf ("TIDSP Modified MOG2 clk=%d frame_row=%d frame_col=%d \n", clk_tot, frame_row, frame_col);
+ printf ("TIDSP Modified MOG2 clk=%d frame_row=%d frame_col=%d (%p %p %p) prune=%f\n", clk_tot, frame_row, frame_col, line_weight, line_variance, line_mean, prune);
}
__kernel void getBackgroundImage2_kernel(__global const uchar* modesUsed,