fix the deinterlace issue

git-svn-id: svn://svn.handbrake.fr/HandBrake/branches/opencl@5275 b64f7644-9d1e-0410-96f1-a4d463321fa5

diff --git a/libhb/common.h b/libhb/common.h
index 5d6d899afff859c9c3471a287180f1764c2a98cc..f8662271f652fa23c518488f8e11c58b2d203f43 100644 (file)
--- a/libhb/common.h
+++ b/libhb/common.h
@@ -450,6 +450,8 @@ struct hb_job_s
                                         //  to non-I frames).
     int use_opencl;/* 0 is disable use of opencl. 1 is enable use of opencl */
     int use_hwd;
+    int use_decomb;
+    int use_detelecine;
 
 #ifdef __LIBHB__
     /* Internal data */

diff --git a/libhb/cropscale.c b/libhb/cropscale.c
index d548de2adae153517f049b3fcd8494f584945f5a..d2de647fc9d60a7002030854bdb14fd6e2584616 100644 (file)
--- a/libhb/cropscale.c
+++ b/libhb/cropscale.c
@@ -25,6 +25,8 @@ struct hb_filter_private_s
     
 #ifdef USE_OPENCL
     int                 use_dxva;
+    int                 use_decomb;
+    int                 use_detelecine;
     hb_oclscale_t       * os; //ocl scaler handler
 #endif 
     struct SwsContext * context;
@@ -69,6 +71,8 @@ static int hb_crop_scale_init( hb_filter_object_t * filter,
     pv->height_out = init->height;
 #ifdef USE_OPENCL
     pv->use_dxva = init->use_dxva;
+    pv->use_decomb = init->job->use_decomb;
+    pv->use_detelecine = init->job->use_detelecine;
 
        if ( pv->job->use_opencl )
     {
@@ -327,7 +331,8 @@ static int hb_crop_scale_work( hb_filter_object_t * filter,
 #ifdef USE_OPENCL
     if ( (in->f.fmt == pv->pix_fmt_out &&
          !pv->crop[0] && !pv->crop[1] && !pv->crop[2] && !pv->crop[3] &&
-         in->f.width == pv->width_out && in->f.height == pv->height_out) ||
+         in->f.width == pv->width_out && in->f.height == pv->height_out) && 
+         (pv->use_decomb == 0) && (pv->use_detelecine == 0) ||
          (pv->use_dxva && in->f.width == pv->width_out && in->f.height == pv->height_out) )
     {
         *buf_out = in;

diff --git a/libhb/decavcodec.c b/libhb/decavcodec.c
index 4e3f02f6ea855ef478c33640ab296cc3cce11c9b..3484a11d7e0eb4ded26e2cc9b1894ad07e54102b 100644 (file)
--- a/libhb/decavcodec.c
+++ b/libhb/decavcodec.c
@@ -587,7 +587,7 @@ static hb_buffer_t *copy_frame( hb_work_private_t *pv, AVFrame *frame )
         hb_buffer_t *buf;
         int ww, hh;
 
-        if( (w > pv->job->width || h > pv->job->height) && (pv->job->use_opencl) )
+        if( (w > pv->job->width || h > pv->job->height) && (pv->job->use_opencl) && (pv->job->use_decomb == 0) &&  (pv->job->use_detelecine == 0) )
         {
             buf = hb_video_buffer_init( pv->job->width, pv->job->height );
             ww = pv->job->width;
@@ -603,7 +603,7 @@ static hb_buffer_t *copy_frame( hb_work_private_t *pv, AVFrame *frame )
         {
             pv->dst_frame = malloc( ww * hh * 3 / 2 );
         }
-        if( hb_va_extract( pv->dxva2, pv->dst_frame, frame, pv->job->width, pv->job->height, pv->job->title->crop, pv->os, pv->job->use_opencl ) == HB_WORK_ERROR )
+        if( hb_va_extract( pv->dxva2, pv->dst_frame, frame, pv->job->width, pv->job->height, pv->job->title->crop, pv->os, pv->job->use_opencl, pv->job->use_decomb, pv->job->use_detelecine ) == HB_WORK_ERROR )
         {
             hb_log( "hb_va_Extract failed!!!!!!" );
         }

diff --git a/libhb/oclnv12toyuv.c b/libhb/oclnv12toyuv.c
index 138a345b5e192587488fd2d6efab567f112fe2a3..26b725bfa6a0809997d760f047f6c94f94908b65 100644 (file)
--- a/libhb/oclnv12toyuv.c
+++ b/libhb/oclnv12toyuv.c
@@ -114,6 +114,8 @@ static int hb_nv12toyuv( void **userdata, KernelEnv *kenv )
 
     uint8_t *bufi2 = userdata[5];
     int p = (int)userdata[6];
+    int decomb = (int)userdata[7];
+    int detelecine = (int)userdata[8];
     int i;
     if( hb_init_nv12toyuv_ocl( kenv, w, h, dxva2 ) )
         return -1;
@@ -150,7 +152,7 @@ static int hb_nv12toyuv( void **userdata, KernelEnv *kenv )
     size_t gdim[2] = {w>>1, h>>1};
     OCLCHECK( clEnqueueNDRangeKernel, kenv->command_queue, kenv->kernel, 2, NULL, gdim, NULL, 0, NULL, NULL );
 
-    if( crop[0] || crop[1] || crop[2] || crop[3] )
+    if( (crop[0] || crop[1] || crop[2] || crop[3]) && (decomb == 0) && (detelecine == 0) )
     {
         AVPicture           pic_in;
         AVPicture           pic_crop;
@@ -200,9 +202,9 @@ static int hb_nv12toyuv_reg_kernel( void )
  * nv12 to yuv interface
  * bufi is input frame of nv12, w is input frame width, h is input frame height
  */
-int hb_ocl_nv12toyuv( uint8_t *bufi[], int p, int w, int h, int *crop, hb_va_dxva2_t *dxva2 )
+int hb_ocl_nv12toyuv( uint8_t *bufi[], int p, int w, int h, int *crop, hb_va_dxva2_t *dxva2, int decomb, int detelecine )
 {
-    void *userdata[7];
+    void *userdata[9];
     userdata[0] = (void*)w;
     userdata[1] = (void*)h;
     userdata[2] = bufi[0];
@@ -210,6 +212,8 @@ int hb_ocl_nv12toyuv( uint8_t *bufi[], int p, int w, int h, int *crop, hb_va_dxv
     userdata[4] = dxva2;
     userdata[5] = bufi[1];
     userdata[6] = (void*)p;
+    userdata[7] = decomb;
+    userdata[8] = detelecine;
     if( hb_nv12toyuv_reg_kernel() )
         return -1;
     if( hb_run_kernel( "nv12toyuv", userdata ) )

diff --git a/libhb/oclnv12toyuv.h b/libhb/oclnv12toyuv.h
index 5098d805e800f851786555e6fc5e59edad48d580..28e09a8a55ecc1aea2f3f2be6a6063b51b3ad691 100644 (file)
--- a/libhb/oclnv12toyuv.h
+++ b/libhb/oclnv12toyuv.h
@@ -23,7 +23,7 @@
  * bufi is input frame of nv12, w is input frame width, h is input frame height
  */
 #ifdef USE_HWD
-int hb_ocl_nv12toyuv( uint8_t *bufi[], int p, int w, int h, int *crop, hb_va_dxva2_t *dxva2 );
+int hb_ocl_nv12toyuv( uint8_t *bufi[], int p, int w, int h, int *crop, hb_va_dxva2_t *dxva2, int decomb, int detelecine );
 #endif
 #endif
 #endif

diff --git a/libhb/vadxva2.c b/libhb/vadxva2.c
index d67a5b3779a54444177f1a01245c0277c381ef25..0ba8dfa4b60431dbb5b569a5a49c4401abef92a8 100644 (file)
--- a/libhb/vadxva2.c
+++ b/libhb/vadxva2.c
@@ -608,7 +608,7 @@ void hb_init_filter( cl_mem src, int srcwidth, int srcheight, uint8_t* dst, int
  *  nv12 to yuv with opencl and with C reference
  *  scale with opencl
  */
-int hb_va_extract( hb_va_dxva2_t *dxva2, uint8_t *dst, AVFrame *frame, int job_w, int job_h, int *crop, hb_oclscale_t *os, int use_opencl )
+int hb_va_extract( hb_va_dxva2_t *dxva2, uint8_t *dst, AVFrame *frame, int job_w, int job_h, int *crop, hb_oclscale_t *os, int use_opencl, int use_decomb, int use_detelecine )
 
 {
     LPDIRECT3DSURFACE9 d3d = (LPDIRECT3DSURFACE9)(uintptr_t)frame->data[3];
@@ -634,15 +634,20 @@ int hb_va_extract( hb_va_dxva2_t *dxva2, uint8_t *dst, AVFrame *frame, int job_w
 #ifdef USE_OPENCL
         if( ( dxva2->width > job_w || dxva2->height > job_h ) && (use_opencl))
         {
-            hb_ocl_nv12toyuv( plane, lock.Pitch,  dxva2->width, dxva2->height, crop, dxva2 );
-
-            static int init_flag = 0;
-            if(init_flag == 0){
-                scale_init( dxva2->width - crop[2] - crop[3], dxva2->height - crop[0] - crop[1], job_w, job_h );
-                init_flag = 1;
-            }
-
-            hb_init_filter( dxva2->cl_mem_yuv, dxva2->width - crop[2] - crop[3], dxva2->height - crop[0] - crop[1], dst, job_w, job_h, crop ); 
+            hb_ocl_nv12toyuv( plane, lock.Pitch,  dxva2->width, dxva2->height, crop, dxva2, use_decomb, use_detelecine );
+            if ( use_decomb || use_detelecine )
+                hb_read_opencl_buffer( dxva2->cl_mem_yuv, dst, dxva2->width*dxva2->height*3/2 );
+            else
+            {                                                                          
+                static int init_flag = 0;
+                if(init_flag == 0)
+                {
+                    scale_init( dxva2->width - crop[2] - crop[3], dxva2->height - crop[0] - crop[1], job_w, job_h );
+                    init_flag = 1;
+                }
+
+                hb_init_filter( dxva2->cl_mem_yuv, dxva2->width - crop[2] - crop[3], dxva2->height - crop[0] - crop[1], dst, job_w, job_h, crop ); 
+            }                          
         }
         else
 #endif

diff --git a/libhb/vadxva2.h b/libhb/vadxva2.h
index bb1a3778a9c5cb7fe2c185b88eace0dc0d04a0e4..f039db7953fe973bbc0481fb9cf5ca89ed9a8567 100644 (file)
--- a/libhb/vadxva2.h
+++ b/libhb/vadxva2.h
@@ -204,7 +204,7 @@ static const hb_dx_mode_t dxva2_modes[] =
 };
 
 int hb_va_get_frame_buf( hb_va_dxva2_t *dxva2, AVCodecContext *p_context, AVFrame *frame );
-int hb_va_extract( hb_va_dxva2_t *dxva2, uint8_t *dst, AVFrame *frame, int job_w, int job_h, int *crop, hb_oclscale_t *os, int use_opencl );
+int hb_va_extract( hb_va_dxva2_t *dxva2, uint8_t *dst, AVFrame *frame, int job_w, int job_h, int *crop, hb_oclscale_t *os, int use_opencl, int use_decomb, int use_detelecine );
 enum PixelFormat hb_ffmpeg_get_format( AVCodecContext *, const enum PixelFormat * );
 hb_va_dxva2_t *hb_va_create_dxva2( hb_va_dxva2_t *dxva2, int codec_id );
 void hb_va_new_dxva2( hb_va_dxva2_t *dxva2, AVCodecContext *p_context );

diff --git a/test/test.c b/test/test.c
index 3094eea8815fd1dfde9156136c2f96e51ccd6d59..d13ec7b2f667179e5f94a830e1cb2a2baae8c3ff 100644 (file)
--- a/test/test.c
+++ b/test/test.c
@@ -1379,6 +1379,9 @@ static int HandleEvents( hb_handle_t * h )
             
             hb_filter_object_t * filter;
 
+            job->use_detelecine = detelecine;
+            job->use_decomb = decomb;
+
             /* Add selected filters */
             if( detelecine )
             {