void x264_sei_version_write( x264_t *h, bs_t *s ) { int i; // random ID number generated according to ISO-11578 const uint8_t uuid[16] = { 0xdc, 0x45, 0xe9, 0xbd, 0xe6, 0xd9, 0x48, 0xb7, 0x96, 0x2c, 0xd8, 0x20, 0xd9, 0x23, 0xee, 0xef }; char *opts = x264_param2string( &h->param, 0 ); char *version = x264_malloc( 200 + strlen(opts) ); int length; sprintf( version, "x264 - core %d%s - H.264/MPEG-4 AVC codec - " "Copyleft 2003-2008 - http://www.videolan.org/x264.html - options: %s", X264_BUILD, X264_VERSION, opts ); length = strlen(version)+1+16; bs_write( s, 8, 0x5 ); // payload_type = user_data_unregistered // payload_size for( i = 0; i <= length-255; i += 255 ) bs_write( s, 8, 255 ); bs_write( s, 8, length-i ); for( i = 0; i < 16; i++ ) bs_write( s, 8, uuid[i] ); for( i = 0; i < length-16; i++ ) bs_write( s, 8, version[i] ); bs_rbsp_trailing( s ); x264_free( opts ); x264_free( version ); }
int x264_sei_version_write( x264_t *h, bs_t *s ) { // random ID number generated according to ISO-11578 static const uint8_t uuid[16] = { 0xdc, 0x45, 0xe9, 0xbd, 0xe6, 0xd9, 0x48, 0xb7, 0x96, 0x2c, 0xd8, 0x20, 0xd9, 0x23, 0xee, 0xef }; char *opts = x264_param2string( &h->param, 0 ); char *payload; int length; if( !opts ) return -1; CHECKED_MALLOC( payload, 200 + strlen( opts ) ); memcpy( payload, uuid, 16 ); sprintf( payload+16, "x264 - core %d%s - H.264/MPEG-4 AVC codec - " "Copy%s 2003-2011 - http://www.videolan.org/x264.html - options: %s", X264_BUILD, X264_VERSION, HAVE_GPL?"left":"right", opts ); length = strlen(payload)+1; x264_sei_write( s, (uint8_t *)payload, length, SEI_USER_DATA_UNREGISTERED ); x264_free( opts ); x264_free( payload ); return 0; fail: x264_free( opts ); return -1; }
float x264_pixel_ssim_wxh( x264_pixel_function_t *pf, uint8_t *pix1, int stride1, uint8_t *pix2, int stride2, int width, int height ) { int x, y, z; float ssim = 0.0; int (*sum0)[4] = x264_malloc(4 * (width/4+3) * sizeof(int)); int (*sum1)[4] = x264_malloc(4 * (width/4+3) * sizeof(int)); width >>= 2; height >>= 2; z = 0; for( y = 1; y < height; y++ ) { for( ; z <= y; z++ ) { XCHG( void*, sum0, sum1 ); for( x = 0; x < width; x+=2 ) pf->ssim_4x4x2_core( &pix1[4*(x+z*stride1)], stride1, &pix2[4*(x+z*stride2)], stride2, &sum0[x] ); } for( x = 0; x < width-1; x += 4 ) ssim += pf->ssim_end4( sum0+x, sum1+x, X264_MIN(4,width-x-1) ); } x264_free(sum0); x264_free(sum1); return ssim / ((width-1) * (height-1)); }
void x264_threadpool_delete( x264_threadpool_t *pool ) { x264_pthread_mutex_lock( &pool->run.mutex ); pool->exit = 1; x264_pthread_cond_broadcast( &pool->run.cv_fill ); x264_pthread_mutex_unlock( &pool->run.mutex ); for( int i = 0; i < pool->threads; i++ ) x264_pthread_join( pool->thread_handle[i], NULL ); x264_threadpool_list_delete( &pool->uninit ); x264_threadpool_list_delete( &pool->run ); x264_threadpool_list_delete( &pool->done ); x264_free( pool->thread_handle ); x264_free( pool ); }
static void free_filter( hnd_t handle ) { depth_hnd_t *h = handle; h->prev_filter.free( h->prev_hnd ); x264_cli_pic_clean( &h->buffer ); x264_free( h ); }
/**************************************************************************** * x264_slurp_file: ****************************************************************************/ char *x264_slurp_file( const char *filename ) { int b_error = 0; int i_size; char *buf; FILE *fh = fopen( filename, "rb" ); if( !fh ) return NULL; b_error |= fseek( fh, 0, SEEK_END ) < 0; b_error |= ( i_size = ftell( fh ) ) <= 0; b_error |= fseek( fh, 0, SEEK_SET ) < 0; if( b_error ) return NULL; buf = x264_malloc( i_size+2 ); if( buf == NULL ) return NULL; b_error |= fread( buf, 1, i_size, fh ) != i_size; if( buf[i_size-1] != '\n' ) buf[i_size++] = '\n'; buf[i_size] = 0; fclose( fh ); if( b_error ) { x264_free( buf ); return NULL; } return buf; }
void x264_opencl_close_library( x264_opencl_function_t *ocl ) { if( !ocl ) return; ocl_close( ocl->library ); x264_free( ocl ); }
void x264_frame_delete( x264_frame_t *frame ) { int i, j; for( i = 0; i < 12; i++ ) x264_free( frame->buffer[i] ); for( i = 0; i < X264_BFRAME_MAX+2; i++ ) for( j = 0; j < X264_BFRAME_MAX+2; j++ ) x264_free( frame->i_row_satds[i][j] ); x264_free( frame->i_row_bits ); x264_free( frame->i_row_qp ); x264_free( frame->mb_type ); x264_free( frame->mv[0] ); x264_free( frame->mv[1] ); x264_free( frame->ref[0] ); x264_free( frame->ref[1] ); x264_free( frame ); }
void x264_lookahead_delete( x264_t *h ) { if( h->param.i_sync_lookahead ) { h->lookahead->b_exit_thread = 1; x264_pthread_cond_broadcast( &h->lookahead->ifbuf.cv_fill ); x264_pthread_join( h->thread[h->param.i_threads]->thread_handle, NULL ); x264_macroblock_cache_end( h->thread[h->param.i_threads] ); x264_free( h->thread[h->param.i_threads] ); } x264_synch_frame_list_delete( &h->lookahead->ifbuf ); x264_synch_frame_list_delete( &h->lookahead->next ); if( h->lookahead->last_nonb ) x264_frame_push_unused( h, h->lookahead->last_nonb ); x264_synch_frame_list_delete( &h->lookahead->ofbuf ); x264_free( h->lookahead ); }
static int init( hnd_t *handle, cli_vid_filter_t *filter, video_info_t *info, x264_param_t *param, char *opt_string ) { int ret = 0; int change_fmt = (info->csp ^ param->i_csp) & X264_CSP_HIGH_DEPTH; int csp = ~(~info->csp ^ change_fmt); int bit_depth = 8*x264_cli_csp_depth_factor( csp ); if( opt_string ) { static const char * const optlist[] = { "bit_depth", NULL }; char **opts = x264_split_options( opt_string, optlist ); if( opts ) { char *str_bit_depth = x264_get_option( "bit_depth", opts ); bit_depth = x264_otoi( str_bit_depth, -1 ); ret = bit_depth < 8 || bit_depth > 16; csp = bit_depth > 8 ? csp | X264_CSP_HIGH_DEPTH : csp & ~X264_CSP_HIGH_DEPTH; change_fmt = (info->csp ^ csp) & X264_CSP_HIGH_DEPTH; free( opts ); } else ret = 1; } FAIL_IF_ERROR( bit_depth != BIT_DEPTH, "this filter supports only bit depth %d\n", BIT_DEPTH ); FAIL_IF_ERROR( ret, "unsupported bit depth conversion.\n" ); /* only add the filter to the chain if it's needed */ if( change_fmt || bit_depth != 8 * x264_cli_csp_depth_factor( csp ) ) { FAIL_IF_ERROR( !depth_filter_csp_is_supported(csp), "unsupported colorspace.\n" ); depth_hnd_t *h = x264_malloc( sizeof(depth_hnd_t) + (info->width+1)*sizeof(int16_t) ); if( !h ) return -1; h->error_buf = (int16_t*)(h + 1); h->dst_csp = csp; h->bit_depth = bit_depth; h->prev_hnd = *handle; h->prev_filter = *filter; if( x264_cli_pic_alloc( &h->buffer, h->dst_csp, info->width, info->height ) ) { x264_free( h ); return -1; } *handle = h; *filter = depth_filter; info->csp = h->dst_csp; } return 0; }
static void x264_threadpool_list_delete( x264_sync_frame_list_t *slist ) { for( int i = 0; slist->list[i]; i++ ) { x264_free( slist->list[i] ); slist->list[i] = NULL; } x264_sync_frame_list_delete( slist ); }
int x264_sei_version_write( x264_t *h, bs_t *s ) { int i; // random ID number generated according to ISO-11578 static const uint8_t uuid[16] = { 0xdc, 0x45, 0xe9, 0xbd, 0xe6, 0xd9, 0x48, 0xb7, 0x96, 0x2c, 0xd8, 0x20, 0xd9, 0x23, 0xee, 0xef }; char *opts = x264_param2string( &h->param, 0 ); char *version; int length; if( !opts ) return -1; CHECKED_MALLOC( version, 200 + strlen( opts ) ); sprintf( version, "x264 - core %d%s - H.264/MPEG-4 AVC codec - " "Copyleft 2003-2010 - http://www.videolan.org/x264.html - options: %s", X264_BUILD, X264_VERSION, opts ); length = strlen(version)+1+16; bs_realign( s ); bs_write( s, 8, SEI_USER_DATA_UNREGISTERED ); // payload_size for( i = 0; i <= length-255; i += 255 ) bs_write( s, 8, 255 ); bs_write( s, 8, length-i ); for( int j = 0; j < 16; j++ ) bs_write( s, 8, uuid[j] ); for( int j = 0; j < length-16; j++ ) bs_write( s, 8, version[j] ); bs_rbsp_trailing( s ); bs_flush( s ); x264_free( opts ); x264_free( version ); return 0; fail: x264_free( opts ); return -1; }
int main (int argc, char *argv[]) { GtkWidget *window; X264_Gtk *x264_gtk; x264_param_t *param; x264_param_t param_default; char *res; char *res_default; BIND_X264_TEXTDOMAIN(); gtk_init (&argc, &argv); window = x264_gtk_window_create (NULL); x264_gtk_shutdown (window); x264_gtk = x264_gtk_load (); param = x264_gtk_param_get (x264_gtk); /* do what you want with these data */ /* for example, displaying them and compare with default*/ res = x264_param2string (param, 0); printf ("%s\n", res); x264_param_default (¶m_default); res_default = x264_param2string (¶m_default, 0); printf ("\n%s\n", res_default); if (strcmp (res, res_default) == 0) printf (_("\nSame result !\n")); else printf (_("\nDifferent from default values\n")); x264_free (res); x264_free (res_default); x264_gtk_free (x264_gtk); g_free (param); return 1; }
void x264_speedcontrol_delete( x264_t *h ) { x264_speedcontrol_t *sc = h->sc; if( !sc ) return; x264_log( h, X264_LOG_INFO, "speedcontrol: avg preset=%.3f buffer min=%.3f max=%.3f\n", sc->stat.avg_preset / sc->stat.den, (float)sc->stat.min_buffer / sc->buffer_size, (float)sc->stat.max_buffer / sc->buffer_size ); // x264_log( h, X264_LOG_INFO, "speedcontrol: avg cplx=%.5f\n", sc->cplx_num / sc->cplx_den ); x264_free( sc ); }
/* Try to load the cached compiled program binary, verify the device context is * still valid before reuse */ static cl_program x264_opencl_cache_load( x264_t *h, const char *dev_name, const char *dev_vendor, const char *driver_version ) { /* try to load cached program binary */ FILE *fp = x264_fopen( h->param.psz_clbin_file, "rb" ); x264_opencl_function_t *ocl; cl_program program; uint8_t *binary; size_t size; const uint8_t *ptr; cl_int status; if( !fp ) return NULL; ocl = h->opencl.ocl; program = NULL; binary = NULL; fseek( fp, 0, SEEK_END ); size = ftell( fp ); rewind( fp ); CHECKED_MALLOC( binary, size ); if ( fread( binary, 1, size, fp ) != size ) goto fail; ptr = (const uint8_t*)binary; #define CHECK_STRING( STR )\ do {\ size_t len = strlen( STR );\ if( size <= len || strncmp( (char*)ptr, STR, len ) )\ goto fail;\ else {\ size -= (len+1); ptr += (len+1);\ }\ } while( 0 ) CHECK_STRING( dev_name ); CHECK_STRING( dev_vendor ); CHECK_STRING( driver_version ); CHECK_STRING( x264_opencl_source_hash ); #undef CHECK_STRING program = ocl->clCreateProgramWithBinary( h->opencl.context, 1, &h->opencl.device, &size, &ptr, NULL, &status ); if( status != CL_SUCCESS ) program = NULL; fail: fclose( fp ); x264_free( binary ); return program; }
HRESULT CX264Encoder::UnInit() { if ( m_px264Handle != NULL ) x264_encoder_close((x264_t*)m_px264Handle); m_px264Handle = NULL; if ( m_pPic != NULL ) x264_free( (x264_picture_t*)m_pPic ); m_pPic = NULL; return true; }
int x264_lookahead_init( x264_t *h ) { x264_lookahead_t *look; CHECKED_MALLOCZERO( look, sizeof(x264_lookahead_t) ); h->lookahead = look; look->i_last_keyframe = - h->param.i_keyint_max; return 0; fail: x264_free( look ); return -1; }
/* load the library and functions we require from it */ x264_opencl_function_t *x264_opencl_load_library( void ) { x264_opencl_function_t *ocl; #undef fail #define fail fail0 CHECKED_MALLOCZERO( ocl, sizeof(x264_opencl_function_t) ); #undef fail #define fail fail1 ocl->library = ocl_open; if( !ocl->library ) goto fail; #undef fail #define fail fail2 LOAD_OCL_FUNC( clBuildProgram, 0 ); LOAD_OCL_FUNC( clCreateBuffer, 0 ); LOAD_OCL_FUNC( clCreateCommandQueue, 0 ); LOAD_OCL_FUNC( clCreateContext, 0 ); LOAD_OCL_FUNC( clCreateImage2D, 0 ); LOAD_OCL_FUNC( clCreateKernel, 0 ); LOAD_OCL_FUNC( clCreateProgramWithBinary, 0 ); LOAD_OCL_FUNC( clCreateProgramWithSource, 0 ); LOAD_OCL_FUNC( clEnqueueCopyBuffer, 0 ); LOAD_OCL_FUNC( clEnqueueMapBuffer, 0 ); LOAD_OCL_FUNC( clEnqueueNDRangeKernel, 0 ); LOAD_OCL_FUNC( clEnqueueReadBuffer, 0 ); LOAD_OCL_FUNC( clEnqueueWriteBuffer, 0 ); LOAD_OCL_FUNC( clFinish, 0 ); LOAD_OCL_FUNC( clGetCommandQueueInfo, 0 ); LOAD_OCL_FUNC( clGetDeviceIDs, 0 ); LOAD_OCL_FUNC( clGetDeviceInfo, 0 ); LOAD_OCL_FUNC( clGetKernelWorkGroupInfo, 0 ); LOAD_OCL_FUNC( clGetPlatformIDs, 0 ); LOAD_OCL_FUNC( clGetProgramBuildInfo, 0 ); LOAD_OCL_FUNC( clGetProgramInfo, 0 ); LOAD_OCL_FUNC( clGetSupportedImageFormats, 0 ); LOAD_OCL_FUNC( clReleaseCommandQueue, 0 ); LOAD_OCL_FUNC( clReleaseContext, 0 ); LOAD_OCL_FUNC( clReleaseKernel, 0 ); LOAD_OCL_FUNC( clReleaseMemObject, 0 ); LOAD_OCL_FUNC( clReleaseProgram, 0 ); LOAD_OCL_FUNC( clSetKernelArg, 0 ); return ocl; #undef fail fail2: ocl_close( ocl->library ); fail1: x264_free( ocl ); fail0: return NULL; }
/* Save the compiled program binary to a file for later reuse. Device context * is also saved in the cache file so we do not reuse stale binaries */ static void x264_opencl_cache_save( x264_t *h, cl_program program, const char *dev_name, const char *dev_vendor, const char *driver_version ) { FILE *fp = x264_fopen( h->param.psz_clbin_file, "wb" ); x264_opencl_function_t *ocl; uint8_t *binary; size_t size; cl_int status; if( !fp ) { x264_log( h, X264_LOG_INFO, "OpenCL: unable to open clbin file for write\n" ); return; } ocl = h->opencl.ocl; binary = NULL; size = 0; status = ocl->clGetProgramInfo( program, CL_PROGRAM_BINARY_SIZES, sizeof(size_t), &size, NULL ); if( status != CL_SUCCESS || !size ) { x264_log( h, X264_LOG_INFO, "OpenCL: Unable to query program binary size, no cache file generated\n" ); goto fail; } CHECKED_MALLOC( binary, size ); status = ocl->clGetProgramInfo( program, CL_PROGRAM_BINARIES, sizeof(uint8_t *), &binary, NULL ); if( status != CL_SUCCESS ) { x264_log( h, X264_LOG_INFO, "OpenCL: Unable to query program binary, no cache file generated\n" ); goto fail; } fputs( dev_name, fp ); fputc( '\n', fp ); fputs( dev_vendor, fp ); fputc( '\n', fp ); fputs( driver_version, fp ); fputc( '\n', fp ); fputs( x264_opencl_source_hash, fp ); fputc( '\n', fp ); fwrite( binary, 1, size, fp ); fail: fclose( fp ); x264_free( binary ); return; }
void x264_frame_delete( x264_frame_t *frame ) { int i; for( i = 0; i < frame->i_plane; i++ ) { x264_free( frame->buffer[i] ); } for( i = 4; i < 11; i++ ) /* filtered planes */ { x264_free( frame->buffer[i] ); } x264_free( frame->mb_type ); x264_free( frame->mv[0] ); x264_free( frame->mv[1] ); x264_free( frame->ref[0] ); x264_free( frame->ref[1] ); x264_free( frame ); }
/**************************************************************************** * x264_realloc: ****************************************************************************/ void *x264_realloc( void *p, int i_size ) { #ifdef HAVE_MALLOC_H return realloc( p, i_size ); #else int i_old_size = 0; uint8_t * p_new; if( p ) { i_old_size = *( (int*) ( (uint8_t*) p ) - sizeof( void ** ) - sizeof( int ) ); } p_new = x264_malloc( i_size ); if( i_old_size > 0 && i_size > 0 ) { memcpy( p_new, p, ( i_old_size < i_size ) ? i_old_size : i_size ); } x264_free( p ); return p_new; #endif }
int x264_lookahead_init( x264_t *h, int i_slicetype_length ) { x264_lookahead_t *look; CHECKED_MALLOCZERO( look, sizeof(x264_lookahead_t) ); for( int i = 0; i < h->param.i_threads; i++ ) h->thread[i]->lookahead = look; look->i_last_keyframe = - h->param.i_keyint_max; look->b_analyse_keyframe = (h->param.rc.b_mb_tree || (h->param.rc.i_vbv_buffer_size && h->param.rc.i_lookahead)) && !h->param.rc.b_stat_read; look->i_slicetype_length = i_slicetype_length; /* init frame lists */ if( x264_sync_frame_list_init( &look->ifbuf, h->param.i_sync_lookahead+3 ) || x264_sync_frame_list_init( &look->next, h->frames.i_delay+3 ) || x264_sync_frame_list_init( &look->ofbuf, h->frames.i_delay+3 ) ) goto fail; if( !h->param.i_sync_lookahead ) return 0; x264_t *look_h = h->thread[h->param.i_threads]; *look_h = *h; if( x264_macroblock_cache_allocate( look_h ) ) goto fail; if( x264_macroblock_thread_allocate( look_h, 1 ) < 0 ) goto fail; if( x264_pthread_create( &look->thread_handle, NULL, (void*)x264_lookahead_thread, look_h ) ) goto fail; look->b_thread_active = 1; return 0; fail: x264_free( look ); return -1; }
sc_streamer sc_streamer_init_video(const char* stream_host, const char* room_name, sc_frame_rect capture_rect, sc_time start_time_stamp){ sc_streamer_setup_windows(); x264_param_t param; char *stream_uri = (char *) malloc (100); sprintf(stream_uri, "rtmp://%s/screenshare/%s", stream_host, room_name); sc_streamer streamer = {.start_time_stamp = start_time_stamp, .stream_uri = stream_uri, .room_name = room_name, .capture_rect = capture_rect, .frames = 0, .rtmpt = 0}; streamer.flv_out_handle = open_flv_buffer(); streamer.rtmp = open_RTMP_stream( stream_uri ); if(!RTMP_IsConnected(streamer.rtmp) || RTMP_IsTimedout(streamer.rtmp)) { printf("Using rtmpt \n"); streamer.rtmpt = 1; free(stream_uri); char *stream_uri = (char *) malloc (100); sprintf(stream_uri, "rtmpt://%s/screenshare/%s", stream_host, room_name); streamer.rtmp = open_RTMP_stream( stream_uri ); } write_RTMP_header(streamer.flv_out_handle, streamer.rtmp); x264_param_default_preset(¶m, "veryfast", "zerolatency"); param.i_log_level = X264_LOG_ERROR; //param.psz_dump_yuv = (char *)"/tmp/dump.y4m"; param.b_vfr_input = 1; param.i_keyint_max = 30; param.i_width = capture_rect.width; param.i_height = capture_rect.height; param.i_timebase_num = 1.0; param.i_timebase_den = SC_TimeBase; //Rate control/quality param.rc.i_rc_method = X264_RC_CRF; param.i_slice_max_size = 1024; param.rc.i_vbv_max_bitrate = 1024; param.rc.i_vbv_buffer_size = 2000; param.rc.f_rf_constant = 23; param.b_sliced_threads = 0; param.b_intra_refresh = 0; param.b_repeat_headers = 1; param.b_annexb = 0; x264_param_apply_profile(¶m, "main"); streamer.encoder = x264_encoder_open(¶m); set_RTMP_param( streamer.flv_out_handle, ¶m ); x264_nal_t *headers; int i_nal; x264_encoder_headers( streamer.encoder, &headers, &i_nal ); write_RTMP_headers( streamer.flv_out_handle, streamer.rtmp, headers ); headers = NULL; free(headers); streamer.rtmp_setup = 1; streamer.reconnect_tries = 0; return streamer; } sc_streamer sc_streamer_init_cursor(const char* stream_host, const char* room_name, sc_time start_time_stamp){ sc_streamer_setup_windows(); char *so_name = (char *) malloc (100); sprintf(so_name, "SC.SS.%s.Cursor", room_name); char *stream_uri = (char *) malloc (100); sprintf(stream_uri, "rtmp://%s/screenshare/%s-cursor", stream_host, room_name); sc_streamer streamer = {.start_time_stamp = start_time_stamp, .stream_uri = stream_uri, .room_name = room_name, .so_name = so_name, .rtmpt = 0, .have_inital_SO = 0}; streamer.flv_out_handle = open_flv_buffer(); streamer.rtmp = open_RTMP_stream( stream_uri ); if(!RTMP_IsConnected(streamer.rtmp) || RTMP_IsTimedout(streamer.rtmp)) { streamer.rtmpt = 1; free(stream_uri); char *stream_uri = (char *) malloc (100); sprintf(stream_uri, "rtmpt://%s/screenshare/%s-cursor", stream_host, room_name); streamer.rtmp = open_RTMP_stream( stream_uri ); } setup_shared_object(streamer.so_name, streamer.rtmp); streamer.rtmp_setup = 1; streamer.so_version = 0; streamer.reconnect_tries = 0; return streamer; } void sc_streamer_send_frame(sc_streamer *streamer, sc_frame *frame, sc_time frame_time_stamp) { x264_picture_t pic_in, pic_out; x264_picture_alloc(&pic_in, X264_CSP_I420, streamer->capture_rect.width, streamer->capture_rect.height); const size_t image_size = (streamer->capture_rect.width * streamer->capture_rect.height); x264_free(pic_in.img.plane[0]); pic_in.img.plane[0] = frame->framePtr; pic_in.img.plane[1] = frame->framePtr + image_size; pic_in.img.plane[2] = frame->framePtr + image_size + image_size / 4; pic_in.i_pts = floor((frame_time_stamp - streamer->start_time_stamp)); x264_nal_t* nals; int i_nals; int frame_size = x264_encoder_encode(streamer->encoder, &nals, &i_nals, &pic_in, &pic_out); if(frame_size > 0) { write_RTMP_frame( streamer->flv_out_handle, streamer->rtmp, nals[0].p_payload, frame_size, &pic_out ); streamer->frames++; } nals = NULL; free(nals); //x264_picture_clean(&pic_in); }
int x264_opencl_lookahead_init( x264_t *h ) { x264_opencl_function_t *ocl = h->opencl.ocl; cl_platform_id *platforms = NULL; cl_device_id *devices = NULL; cl_image_format *imageType = NULL; cl_context context = NULL; int ret = -1; cl_uint i; cl_uint numPlatforms = 0; cl_int status = ocl->clGetPlatformIDs( 0, NULL, &numPlatforms ); if( status != CL_SUCCESS || !numPlatforms ) { x264_log( h, X264_LOG_WARNING, "OpenCL: Unable to query installed platforms\n" ); goto fail; } platforms = (cl_platform_id*)x264_malloc( sizeof(cl_platform_id) * numPlatforms ); if( !platforms ) { x264_log( h, X264_LOG_WARNING, "OpenCL: malloc of installed platforms buffer failed\n" ); goto fail; } status = ocl->clGetPlatformIDs( numPlatforms, platforms, NULL ); if( status != CL_SUCCESS ) { x264_log( h, X264_LOG_WARNING, "OpenCL: Unable to query installed platforms\n" ); goto fail; } /* Select the first OpenCL platform with a GPU device that supports our * required image (texture) formats */ for( i = 0; i < numPlatforms; i++ ) { cl_uint gpu_count = 0; cl_uint gpu; status = ocl->clGetDeviceIDs( platforms[i], CL_DEVICE_TYPE_GPU, 0, NULL, &gpu_count ); if( status != CL_SUCCESS || !gpu_count ) continue; x264_free( devices ); devices = x264_malloc( sizeof(cl_device_id) * gpu_count ); if( !devices ) continue; status = ocl->clGetDeviceIDs( platforms[i], CL_DEVICE_TYPE_GPU, gpu_count, devices, NULL ); if( status != CL_SUCCESS ) continue; /* Find a GPU device that supports our image formats */ for( gpu = 0; gpu < gpu_count; gpu++ ) { cl_bool image_support; cl_uint imagecount; int b_has_r; int b_has_rgba; cl_uint j; h->opencl.device = devices[gpu]; /* if the user has specified an exact device ID, skip all other * GPUs. If this device matches, allow it to continue through the * checks for supported images, etc. */ if( h->param.opencl_device_id && devices[gpu] != (cl_device_id)h->param.opencl_device_id ) continue; image_support = 0; status = ocl->clGetDeviceInfo( h->opencl.device, CL_DEVICE_IMAGE_SUPPORT, sizeof(cl_bool), &image_support, NULL ); if( status != CL_SUCCESS || !image_support ) continue; if( context ) ocl->clReleaseContext( context ); context = ocl->clCreateContext( NULL, 1, &h->opencl.device, (void*)x264_opencl_error_notify, (void*)h, &status ); if( status != CL_SUCCESS || !context ) continue; imagecount = 0; status = ocl->clGetSupportedImageFormats( context, CL_MEM_READ_WRITE, CL_MEM_OBJECT_IMAGE2D, 0, NULL, &imagecount ); if( status != CL_SUCCESS || !imagecount ) continue; x264_free( imageType ); imageType = x264_malloc( sizeof(cl_image_format) * imagecount ); if( !imageType ) continue; status = ocl->clGetSupportedImageFormats( context, CL_MEM_READ_WRITE, CL_MEM_OBJECT_IMAGE2D, imagecount, imageType, NULL ); if( status != CL_SUCCESS ) continue; b_has_r = 0; b_has_rgba = 0; for( j = 0; j < imagecount; j++ ) { if( imageType[j].image_channel_order == CL_R && imageType[j].image_channel_data_type == CL_UNSIGNED_INT32 ) b_has_r = 1; else if( imageType[j].image_channel_order == CL_RGBA && imageType[j].image_channel_data_type == CL_UNSIGNED_INT8 ) b_has_rgba = 1; } if( !b_has_r || !b_has_rgba ) { char dev_name[64]; status = ocl->clGetDeviceInfo( h->opencl.device, CL_DEVICE_NAME, sizeof(dev_name), dev_name, NULL ); if( status == CL_SUCCESS ) { /* emit warning if we are discarding the user's explicit choice */ int level = h->param.opencl_device_id ? X264_LOG_WARNING : X264_LOG_DEBUG; x264_log( h, level, "OpenCL: %s does not support required image formats\n", dev_name ); } continue; } /* user selection of GPU device, skip N first matches */ if( h->param.i_opencl_device ) { h->param.i_opencl_device--; continue; } h->opencl.queue = ocl->clCreateCommandQueue( context, h->opencl.device, 0, &status ); if( status != CL_SUCCESS || !h->opencl.queue ) continue; h->opencl.context = context; context = NULL; ret = 0; break; } if( !ret ) break; } if( !h->param.psz_clbin_file ) h->param.psz_clbin_file = "x264_lookahead.clbin"; if( ret ) x264_log( h, X264_LOG_WARNING, "OpenCL: Unable to find a compatible device\n" ); else ret = x264_opencl_lookahead_alloc( h ); fail: if( context ) ocl->clReleaseContext( context ); x264_free( imageType ); x264_free( devices ); x264_free( platforms ); return ret; }
/* The OpenCL source under common/opencl will be merged into common/oclobj.h by * the Makefile. It defines a x264_opencl_source byte array which we will pass * to clCreateProgramWithSource(). We also attempt to use a cache file for the * compiled binary, stored in the current working folder. */ static cl_program x264_opencl_compile( x264_t *h ) { x264_opencl_function_t *ocl = h->opencl.ocl; cl_program program = NULL; char *build_log = NULL; char dev_name[64]; char dev_vendor[64]; char driver_version[64]; cl_int status; int vectorize; const char *buildopts; size_t build_log_len; FILE *log_file; status = ocl->clGetDeviceInfo( h->opencl.device, CL_DEVICE_NAME, sizeof(dev_name), dev_name, NULL ); status |= ocl->clGetDeviceInfo( h->opencl.device, CL_DEVICE_VENDOR, sizeof(dev_vendor), dev_vendor, NULL ); status |= ocl->clGetDeviceInfo( h->opencl.device, CL_DRIVER_VERSION, sizeof(driver_version), driver_version, NULL ); if( status != CL_SUCCESS ) return NULL; // Most AMD GPUs have vector registers vectorize = !strcmp( dev_vendor, "Advanced Micro Devices, Inc." ); h->opencl.b_device_AMD_SI = 0; if( vectorize ) { cl_uint simdwidth; /* Disable OpenCL on Intel/AMD switchable graphics devices */ if( x264_detect_switchable_graphics() ) { x264_log( h, X264_LOG_INFO, "OpenCL acceleration disabled, switchable graphics detected\n" ); return NULL; } /* Detect AMD SouthernIsland or newer device (single-width registers) */ simdwidth = 4; status = ocl->clGetDeviceInfo( h->opencl.device, CL_DEVICE_SIMD_INSTRUCTION_WIDTH_AMD, sizeof(cl_uint), &simdwidth, NULL ); if( status == CL_SUCCESS && simdwidth == 1 ) { vectorize = 0; h->opencl.b_device_AMD_SI = 1; } } x264_log( h, X264_LOG_INFO, "OpenCL acceleration enabled with %s %s %s\n", dev_vendor, dev_name, h->opencl.b_device_AMD_SI ? "(SI)" : "" ); program = x264_opencl_cache_load( h, dev_name, dev_vendor, driver_version ); if( !program ) { const char *strptr; size_t size; /* clCreateProgramWithSource() requires a pointer variable, you cannot just use &x264_opencl_source */ x264_log( h, X264_LOG_INFO, "Compiling OpenCL kernels...\n" ); strptr = (const char*)x264_opencl_source; size = sizeof(x264_opencl_source); program = ocl->clCreateProgramWithSource( h->opencl.context, 1, &strptr, &size, &status ); if( status != CL_SUCCESS || !program ) { x264_log( h, X264_LOG_WARNING, "OpenCL: unable to create program\n" ); return NULL; } } /* Build the program binary for the OpenCL device */ buildopts = vectorize ? "-DVECTORIZE=1" : ""; status = ocl->clBuildProgram( program, 1, &h->opencl.device, buildopts, NULL, NULL ); if( status == CL_SUCCESS ) { x264_opencl_cache_save( h, program, dev_name, dev_vendor, driver_version ); return program; } /* Compile failure, should not happen with production code. */ build_log_len = 0; status = ocl->clGetProgramBuildInfo( program, h->opencl.device, CL_PROGRAM_BUILD_LOG, 0, NULL, &build_log_len ); if( status != CL_SUCCESS || !build_log_len ) { x264_log( h, X264_LOG_WARNING, "OpenCL: Compilation failed, unable to query build log\n" ); goto fail; } build_log = x264_malloc( build_log_len ); if( !build_log ) { x264_log( h, X264_LOG_WARNING, "OpenCL: Compilation failed, unable to alloc build log\n" ); goto fail; } status = ocl->clGetProgramBuildInfo( program, h->opencl.device, CL_PROGRAM_BUILD_LOG, build_log_len, build_log, NULL ); if( status != CL_SUCCESS ) { x264_log( h, X264_LOG_WARNING, "OpenCL: Compilation failed, unable to get build log\n" ); goto fail; } log_file = x264_fopen( "x264_kernel_build_log.txt", "w" ); if( !log_file ) { x264_log( h, X264_LOG_WARNING, "OpenCL: Compilation failed, unable to create file x264_kernel_build_log.txt\n" ); goto fail; } fwrite( build_log, 1, build_log_len, log_file ); fclose( log_file ); x264_log( h, X264_LOG_WARNING, "OpenCL: kernel build errors written to x264_kernel_build_log.txt\n" ); fail: x264_free( build_log ); if( program ) ocl->clReleaseProgram( program ); return NULL; }
/* {{{ [fold] void x264_visualize_close( x264_t *h ) */ void x264_visualize_close( x264_t *h ) { x264_free(h->visualize); }
void x264_cli_pic_clean( cli_pic_t *pic ) { for( int i = 0; i < pic->img.planes; i++ ) x264_free( pic->img.plane[i] ); memset( pic, 0, sizeof(cli_pic_t) ); }