From 06e7a9ace3439f08350a75a6f9a491f1bc4161a3 Mon Sep 17 00:00:00 2001 From: handbrake Date: Thu, 28 Feb 2013 09:20:50 +0000 Subject: [PATCH] fix the deinterlace issue git-svn-id: svn://svn.handbrake.fr/HandBrake/branches/opencl@5275 b64f7644-9d1e-0410-96f1-a4d463321fa5 --- libhb/common.h | 2 ++ libhb/cropscale.c | 7 ++++++- libhb/decavcodec.c | 4 ++-- libhb/oclnv12toyuv.c | 10 +++++++--- libhb/oclnv12toyuv.h | 2 +- libhb/vadxva2.c | 25 +++++++++++++++---------- libhb/vadxva2.h | 2 +- test/test.c | 3 +++ 8 files changed, 37 insertions(+), 18 deletions(-) diff --git a/libhb/common.h b/libhb/common.h index 5d6d899af..f8662271f 100644 --- 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 d548de2ad..d2de647fc 100644 --- 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 4e3f02f6e..3484a11d7 100644 --- 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 138a345b5..26b725bfa 100644 --- 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 5098d805e..28e09a8a5 100644 --- 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 d67a5b377..0ba8dfa4b 100644 --- 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 bb1a3778a..f039db795 100644 --- 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 3094eea88..d13ec7b2f 100644 --- 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 ) { -- 2.49.0