x86: faster AVX2 iDCT, AVX deblock_luma_h, deblock_luma_h_intra
[x262.git] / common / opencl.c
1 /*****************************************************************************
2  * opencl.c: OpenCL initialization and kernel compilation
3  *****************************************************************************
4  * Copyright (C) 2012-2013 x264 project
5  *
6  * Authors: Steve Borho <sborho@multicorewareinc.com>
7  *          Anton Mitrofanov <BugMaster@narod.ru>
8  *
9  * This program is free software; you can redistribute it and/or modify
10  * it under the terms of the GNU General Public License as published by
11  * the Free Software Foundation; either version 2 of the License, or
12  * (at your option) any later version.
13  *
14  * This program is distributed in the hope that it will be useful,
15  * but WITHOUT ANY WARRANTY; without even the implied warranty of
16  * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
17  * GNU General Public License for more details.
18  *
19  * You should have received a copy of the GNU General Public License
20  * along with this program; if not, write to the Free Software
21  * Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA  02111, USA.
22  *
23  * This program is also available under a commercial proprietary license.
24  * For more information, contact us at licensing@x264.com.
25  *****************************************************************************/
26
27 #include "common.h"
28
29 #ifdef _WIN32
30 #include <windows.h>
31 #define ocl_open LoadLibrary( "OpenCL" )
32 #define ocl_close FreeLibrary
33 #define ocl_address GetProcAddress
34 #else
35 #include <dlfcn.h> //dlopen, dlsym, dlclose
36 #if SYS_MACOSX
37 #define ocl_open dlopen( "/System/Library/Frameworks/OpenCL.framework/OpenCL", RTLD_NOW )
38 #else
39 #define ocl_open dlopen( "libOpenCL.so", RTLD_NOW )
40 #endif
41 #define ocl_close dlclose
42 #define ocl_address dlsym
43 #endif
44
45 #define LOAD_OCL_FUNC(name, continue_on_fail)\
46 {\
47     ocl->name = (void*)ocl_address( ocl->library, #name );\
48     if( !continue_on_fail && !ocl->name )\
49         goto fail;\
50 }
51
52 /* load the library and functions we require from it */
53 x264_opencl_function_t *x264_opencl_load_library( void )
54 {
55     x264_opencl_function_t *ocl;
56 #undef fail
57 #define fail fail0
58     CHECKED_MALLOCZERO( ocl, sizeof(x264_opencl_function_t) );
59 #undef fail
60 #define fail fail1
61     ocl->library = ocl_open;
62     if( !ocl->library )
63         goto fail;
64 #undef fail
65 #define fail fail2
66     LOAD_OCL_FUNC( clBuildProgram, 0 );
67     LOAD_OCL_FUNC( clCreateBuffer, 0 );
68     LOAD_OCL_FUNC( clCreateCommandQueue, 0 );
69     LOAD_OCL_FUNC( clCreateContext, 0 );
70     LOAD_OCL_FUNC( clCreateImage2D, 0 );
71     LOAD_OCL_FUNC( clCreateKernel, 0 );
72     LOAD_OCL_FUNC( clCreateProgramWithBinary, 0 );
73     LOAD_OCL_FUNC( clCreateProgramWithSource, 0 );
74     LOAD_OCL_FUNC( clEnqueueCopyBuffer, 0 );
75     LOAD_OCL_FUNC( clEnqueueMapBuffer, 0 );
76     LOAD_OCL_FUNC( clEnqueueNDRangeKernel, 0 );
77     LOAD_OCL_FUNC( clEnqueueReadBuffer, 0 );
78     LOAD_OCL_FUNC( clEnqueueWriteBuffer, 0 );
79     LOAD_OCL_FUNC( clFinish, 0 );
80     LOAD_OCL_FUNC( clGetCommandQueueInfo, 0 );
81     LOAD_OCL_FUNC( clGetDeviceIDs, 0 );
82     LOAD_OCL_FUNC( clGetDeviceInfo, 0 );
83     LOAD_OCL_FUNC( clGetKernelWorkGroupInfo, 0 );
84     LOAD_OCL_FUNC( clGetPlatformIDs, 0 );
85     LOAD_OCL_FUNC( clGetProgramBuildInfo, 0 );
86     LOAD_OCL_FUNC( clGetProgramInfo, 0 );
87     LOAD_OCL_FUNC( clGetSupportedImageFormats, 0 );
88     LOAD_OCL_FUNC( clReleaseCommandQueue, 0 );
89     LOAD_OCL_FUNC( clReleaseContext, 0 );
90     LOAD_OCL_FUNC( clReleaseKernel, 0 );
91     LOAD_OCL_FUNC( clReleaseMemObject, 0 );
92     LOAD_OCL_FUNC( clReleaseProgram, 0 );
93     LOAD_OCL_FUNC( clSetKernelArg, 0 );
94     return ocl;
95 #undef fail
96 fail2:
97     ocl_close( ocl->library );
98 fail1:
99     x264_free( ocl );
100 fail0:
101     return NULL;
102 }
103
104 void x264_opencl_close_library( x264_opencl_function_t *ocl )
105 {
106     if( !ocl )
107         return;
108     ocl_close( ocl->library );
109     x264_free( ocl );
110 }
111
112 /* define from recent cl_ext.h, copied here in case headers are old */
113 #define CL_DEVICE_SIMD_INSTRUCTION_WIDTH_AMD        0x4042
114
115 /* Requires full include path in case of out-of-tree builds */
116 #include "common/oclobj.h"
117
118 static int x264_detect_switchable_graphics( void );
119
120 /* Try to load the cached compiled program binary, verify the device context is
121  * still valid before reuse */
122 static cl_program x264_opencl_cache_load( x264_t *h, char *dev_name, char *dev_vendor, char *driver_version )
123 {
124     /* try to load cached program binary */
125     FILE *fp = fopen( h->param.psz_clbin_file, "rb" );
126     if( !fp )
127         return NULL;
128
129     x264_opencl_function_t *ocl = h->opencl.ocl;
130     cl_program program = NULL;
131     uint8_t *binary = NULL;
132
133     fseek( fp, 0, SEEK_END );
134     size_t size = ftell( fp );
135     rewind( fp );
136     CHECKED_MALLOC( binary, size );
137
138     fread( binary, 1, size, fp );
139     const uint8_t *ptr = (const uint8_t*)binary;
140
141 #define CHECK_STRING( STR )\
142     do {\
143         size_t len = strlen( STR );\
144         if( size <= len || strncmp( (char*)ptr, STR, len ) )\
145             goto fail;\
146         else {\
147             size -= (len+1); ptr += (len+1);\
148         }\
149     } while( 0 )
150
151     CHECK_STRING( dev_name );
152     CHECK_STRING( dev_vendor );
153     CHECK_STRING( driver_version );
154     CHECK_STRING( x264_opencl_source_hash );
155 #undef CHECK_STRING
156
157     cl_int status;
158     program = ocl->clCreateProgramWithBinary( h->opencl.context, 1, &h->opencl.device, &size, &ptr, NULL, &status );
159     if( status != CL_SUCCESS )
160         program = NULL;
161
162 fail:
163     fclose( fp );
164     x264_free( binary );
165     return program;
166 }
167
168 /* Save the compiled program binary to a file for later reuse.  Device context
169  * is also saved in the cache file so we do not reuse stale binaries */
170 static void x264_opencl_cache_save( x264_t *h, cl_program program, char *dev_name, char *dev_vendor, char *driver_version )
171 {
172     FILE *fp = fopen( h->param.psz_clbin_file, "wb" );
173     if( !fp )
174     {
175         x264_log( h, X264_LOG_INFO, "OpenCL: unable to open clbin file for write\n" );
176         return;
177     }
178
179     x264_opencl_function_t *ocl = h->opencl.ocl;
180     uint8_t *binary = NULL;
181
182     size_t size = 0;
183     cl_int status = ocl->clGetProgramInfo( program, CL_PROGRAM_BINARY_SIZES, sizeof(size_t), &size, NULL );
184     if( status != CL_SUCCESS || !size )
185     {
186         x264_log( h, X264_LOG_INFO, "OpenCL: Unable to query program binary size, no cache file generated\n" );
187         goto fail;
188     }
189
190     CHECKED_MALLOC( binary, size );
191     status = ocl->clGetProgramInfo( program, CL_PROGRAM_BINARIES, sizeof(uint8_t *), &binary, NULL );
192     if( status != CL_SUCCESS )
193     {
194         x264_log( h, X264_LOG_INFO, "OpenCL: Unable to query program binary, no cache file generated\n" );
195         goto fail;
196     }
197
198     fputs( dev_name, fp );
199     fputc( '\n', fp );
200     fputs( dev_vendor, fp );
201     fputc( '\n', fp );
202     fputs( driver_version, fp );
203     fputc( '\n', fp );
204     fputs( x264_opencl_source_hash, fp );
205     fputc( '\n', fp );
206     fwrite( binary, 1, size, fp );
207
208 fail:
209     fclose( fp );
210     x264_free( binary );
211     return;
212 }
213
214 /* The OpenCL source under common/opencl will be merged into common/oclobj.h by
215  * the Makefile. It defines a x264_opencl_source byte array which we will pass
216  * to clCreateProgramWithSource().  We also attempt to use a cache file for the
217  * compiled binary, stored in the current working folder. */
218 static cl_program x264_opencl_compile( x264_t *h )
219 {
220     x264_opencl_function_t *ocl = h->opencl.ocl;
221     cl_program program = NULL;
222     char *build_log = NULL;
223
224     char dev_name[64];
225     char dev_vendor[64];
226     char driver_version[64];
227     cl_int status;
228     status  = ocl->clGetDeviceInfo( h->opencl.device, CL_DEVICE_NAME,    sizeof(dev_name), dev_name, NULL );
229     status |= ocl->clGetDeviceInfo( h->opencl.device, CL_DEVICE_VENDOR,  sizeof(dev_vendor), dev_vendor, NULL );
230     status |= ocl->clGetDeviceInfo( h->opencl.device, CL_DRIVER_VERSION, sizeof(driver_version), driver_version, NULL );
231     if( status != CL_SUCCESS )
232         return NULL;
233
234     // Most AMD GPUs have vector registers
235     int vectorize = !strcmp( dev_vendor, "Advanced Micro Devices, Inc." );
236     h->opencl.b_device_AMD_SI = 0;
237
238     if( vectorize )
239     {
240         /* Disable OpenCL on Intel/AMD switchable graphics devices */
241         if( x264_detect_switchable_graphics() )
242         {
243             x264_log( h, X264_LOG_INFO, "OpenCL acceleration disabled, switchable graphics detected\n" );
244             return NULL;
245         }
246
247         /* Detect AMD SouthernIsland or newer device (single-width registers) */
248         cl_uint simdwidth = 4;
249         status = ocl->clGetDeviceInfo( h->opencl.device, CL_DEVICE_SIMD_INSTRUCTION_WIDTH_AMD, sizeof(cl_uint), &simdwidth, NULL );
250         if( status == CL_SUCCESS && simdwidth == 1 )
251         {
252             vectorize = 0;
253             h->opencl.b_device_AMD_SI = 1;
254         }
255     }
256
257     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)" : "" );
258
259     program = x264_opencl_cache_load( h, dev_name, dev_vendor, driver_version );
260     if( !program )
261     {
262         /* clCreateProgramWithSource() requires a pointer variable, you cannot just use &x264_opencl_source */
263         x264_log( h, X264_LOG_INFO, "Compiling OpenCL kernels...\n" );
264         const char *strptr = (const char*)x264_opencl_source;
265         size_t size = sizeof(x264_opencl_source);
266         program = ocl->clCreateProgramWithSource( h->opencl.context, 1, &strptr, &size, &status );
267         if( status != CL_SUCCESS || !program )
268         {
269             x264_log( h, X264_LOG_WARNING, "OpenCL: unable to create program\n" );
270             return NULL;
271         }
272     }
273
274     /* Build the program binary for the OpenCL device */
275     const char *buildopts = vectorize ? "-DVECTORIZE=1" : "";
276     status = ocl->clBuildProgram( program, 1, &h->opencl.device, buildopts, NULL, NULL );
277     if( status == CL_SUCCESS )
278     {
279         x264_opencl_cache_save( h, program, dev_name, dev_vendor, driver_version );
280         return program;
281     }
282
283     /* Compile failure, should not happen with production code. */
284
285     size_t build_log_len = 0;
286     status = ocl->clGetProgramBuildInfo( program, h->opencl.device, CL_PROGRAM_BUILD_LOG, 0, NULL, &build_log_len );
287     if( status != CL_SUCCESS || !build_log_len )
288     {
289         x264_log( h, X264_LOG_WARNING, "OpenCL: Compilation failed, unable to query build log\n" );
290         goto fail;
291     }
292
293     build_log = x264_malloc( build_log_len );
294     if( !build_log )
295     {
296         x264_log( h, X264_LOG_WARNING, "OpenCL: Compilation failed, unable to alloc build log\n" );
297         goto fail;
298     }
299
300     status = ocl->clGetProgramBuildInfo( program, h->opencl.device, CL_PROGRAM_BUILD_LOG, build_log_len, build_log, NULL );
301     if( status != CL_SUCCESS )
302     {
303         x264_log( h, X264_LOG_WARNING, "OpenCL: Compilation failed, unable to get build log\n" );
304         goto fail;
305     }
306
307     FILE *log_file = fopen( "x264_kernel_build_log.txt", "w" );
308     if( !log_file )
309     {
310         x264_log( h, X264_LOG_WARNING, "OpenCL: Compilation failed, unable to create file x264_kernel_build_log.txt\n" );
311         goto fail;
312     }
313     fwrite( build_log, 1, build_log_len, log_file );
314     fclose( log_file );
315     x264_log( h, X264_LOG_WARNING, "OpenCL: kernel build errors written to x264_kernel_build_log.txt\n" );
316
317 fail:
318     x264_free( build_log );
319     if( program )
320         ocl->clReleaseProgram( program );
321     return NULL;
322 }
323
324 static int x264_opencl_lookahead_alloc( x264_t *h )
325 {
326     if( !h->param.rc.i_lookahead )
327         return -1;
328
329     static const char *kernelnames[] = {
330         "mb_intra_cost_satd_8x8",
331         "sum_intra_cost",
332         "downscale_hpel",
333         "downscale1",
334         "downscale2",
335         "memset_int16",
336         "weightp_scaled_images",
337         "weightp_hpel",
338         "hierarchical_motion",
339         "subpel_refine",
340         "mode_selection",
341         "sum_inter_cost"
342     };
343
344     cl_kernel *kernels[] = {
345         &h->opencl.intra_kernel,
346         &h->opencl.rowsum_intra_kernel,
347         &h->opencl.downscale_hpel_kernel,
348         &h->opencl.downscale_kernel1,
349         &h->opencl.downscale_kernel2,
350         &h->opencl.memset_kernel,
351         &h->opencl.weightp_scaled_images_kernel,
352         &h->opencl.weightp_hpel_kernel,
353         &h->opencl.hme_kernel,
354         &h->opencl.subpel_refine_kernel,
355         &h->opencl.mode_select_kernel,
356         &h->opencl.rowsum_inter_kernel
357     };
358
359     x264_opencl_function_t *ocl = h->opencl.ocl;
360     cl_int status;
361
362     h->opencl.lookahead_program = x264_opencl_compile( h );
363     if( !h->opencl.lookahead_program )
364         goto fail;
365
366     for( int i = 0; i < ARRAY_SIZE(kernelnames); i++ )
367     {
368         *kernels[i] = ocl->clCreateKernel( h->opencl.lookahead_program, kernelnames[i], &status );
369         if( status != CL_SUCCESS )
370         {
371             x264_log( h, X264_LOG_WARNING, "OpenCL: Unable to compile kernel '%s' (%d)\n", kernelnames[i], status );
372             goto fail;
373         }
374     }
375
376     h->opencl.page_locked_buffer = ocl->clCreateBuffer( h->opencl.context, CL_MEM_WRITE_ONLY|CL_MEM_ALLOC_HOST_PTR, PAGE_LOCKED_BUF_SIZE, NULL, &status );
377     if( status != CL_SUCCESS )
378     {
379         x264_log( h, X264_LOG_WARNING, "OpenCL: Unable to allocate page-locked buffer, error '%d'\n", status );
380         goto fail;
381     }
382     h->opencl.page_locked_ptr = ocl->clEnqueueMapBuffer( h->opencl.queue, h->opencl.page_locked_buffer, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE,
383                                                          0, PAGE_LOCKED_BUF_SIZE, 0, NULL, NULL, &status );
384     if( status != CL_SUCCESS )
385     {
386         x264_log( h, X264_LOG_WARNING, "OpenCL: Unable to map page-locked buffer, error '%d'\n", status );
387         goto fail;
388     }
389
390     return 0;
391 fail:
392     x264_opencl_lookahead_delete( h );
393     return -1;
394 }
395
396 static void CL_CALLBACK x264_opencl_error_notify( const char *errinfo, const void *private_info, size_t cb, void *user_data )
397 {
398     /* Any error notification can be assumed to be fatal to the OpenCL context.
399      * We need to stop using it immediately to prevent further damage. */
400     x264_t *h = (x264_t*)user_data;
401     h->param.b_opencl = 0;
402     h->opencl.b_fatal_error = 1;
403     x264_log( h, X264_LOG_ERROR, "OpenCL: %s\n", errinfo );
404     x264_log( h, X264_LOG_ERROR, "OpenCL: fatal error, aborting encode\n" );
405 }
406
407 int x264_opencl_lookahead_init( x264_t *h )
408 {
409     x264_opencl_function_t *ocl = h->opencl.ocl;
410     cl_platform_id *platforms = NULL;
411     cl_device_id *devices = NULL;
412     cl_image_format *imageType = NULL;
413     cl_context context = NULL;
414     int ret = -1;
415
416     cl_uint numPlatforms = 0;
417     cl_int status = ocl->clGetPlatformIDs( 0, NULL, &numPlatforms );
418     if( status != CL_SUCCESS || !numPlatforms )
419     {
420         x264_log( h, X264_LOG_WARNING, "OpenCL: Unable to query installed platforms\n" );
421         goto fail;
422     }
423     platforms = (cl_platform_id*)x264_malloc( sizeof(cl_platform_id) * numPlatforms );
424     if( !platforms )
425     {
426         x264_log( h, X264_LOG_WARNING, "OpenCL: malloc of installed platforms buffer failed\n" );
427         goto fail;
428     }
429     status = ocl->clGetPlatformIDs( numPlatforms, platforms, NULL );
430     if( status != CL_SUCCESS )
431     {
432         x264_log( h, X264_LOG_WARNING, "OpenCL: Unable to query installed platforms\n" );
433         goto fail;
434     }
435
436     /* Select the first OpenCL platform with a GPU device that supports our
437      * required image (texture) formats */
438     for( cl_uint i = 0; i < numPlatforms; i++ )
439     {
440         cl_uint gpu_count = 0;
441         status = ocl->clGetDeviceIDs( platforms[i], CL_DEVICE_TYPE_GPU, 0, NULL, &gpu_count );
442         if( status != CL_SUCCESS || !gpu_count )
443             continue;
444
445         x264_free( devices );
446         devices = x264_malloc( sizeof(cl_device_id) * gpu_count );
447         if( !devices )
448             continue;
449
450         status = ocl->clGetDeviceIDs( platforms[i], CL_DEVICE_TYPE_GPU, gpu_count, devices, NULL );
451         if( status != CL_SUCCESS )
452             continue;
453
454         /* Find a GPU device that supports our image formats */
455         for( cl_uint gpu = 0; gpu < gpu_count; gpu++ )
456         {
457             h->opencl.device = devices[gpu];
458
459             /* if the user has specified an exact device ID, skip all other
460              * GPUs.  If this device matches, allow it to continue through the
461              * checks for supported images, etc.  */
462             if( h->param.opencl_device_id && devices[gpu] != (cl_device_id)h->param.opencl_device_id )
463                 continue;
464
465             cl_bool image_support = 0;
466             status = ocl->clGetDeviceInfo( h->opencl.device, CL_DEVICE_IMAGE_SUPPORT, sizeof(cl_bool), &image_support, NULL );
467             if( status != CL_SUCCESS || !image_support )
468                 continue;
469
470             if( context )
471                 ocl->clReleaseContext( context );
472             context = ocl->clCreateContext( NULL, 1, &h->opencl.device, (void*)x264_opencl_error_notify, (void*)h, &status );
473             if( status != CL_SUCCESS || !context )
474                 continue;
475
476             cl_uint imagecount = 0;
477             status = ocl->clGetSupportedImageFormats( context, CL_MEM_READ_WRITE, CL_MEM_OBJECT_IMAGE2D, 0, NULL, &imagecount );
478             if( status != CL_SUCCESS || !imagecount )
479                 continue;
480
481             x264_free( imageType );
482             imageType = x264_malloc( sizeof(cl_image_format) * imagecount );
483             if( !imageType )
484                 continue;
485
486             status = ocl->clGetSupportedImageFormats( context, CL_MEM_READ_WRITE, CL_MEM_OBJECT_IMAGE2D, imagecount, imageType, NULL );
487             if( status != CL_SUCCESS )
488                 continue;
489
490             int b_has_r = 0;
491             int b_has_rgba = 0;
492             for( cl_uint j = 0; j < imagecount; j++ )
493             {
494                 if( imageType[j].image_channel_order == CL_R &&
495                     imageType[j].image_channel_data_type == CL_UNSIGNED_INT32 )
496                     b_has_r = 1;
497                 else if( imageType[j].image_channel_order == CL_RGBA &&
498                          imageType[j].image_channel_data_type == CL_UNSIGNED_INT8 )
499                     b_has_rgba = 1;
500             }
501             if( !b_has_r || !b_has_rgba )
502             {
503                 char dev_name[64];
504                 status = ocl->clGetDeviceInfo( h->opencl.device, CL_DEVICE_NAME, sizeof(dev_name), dev_name, NULL );
505                 if( status == CL_SUCCESS )
506                 {
507                     /* emit warning if we are discarding the user's explicit choice */
508                     int level = h->param.opencl_device_id ? X264_LOG_WARNING : X264_LOG_DEBUG;
509                     x264_log( h, level, "OpenCL: %s does not support required image formats\n", dev_name );
510                 }
511                 continue;
512             }
513
514             /* user selection of GPU device, skip N first matches */
515             if( h->param.i_opencl_device )
516             {
517                 h->param.i_opencl_device--;
518                 continue;
519             }
520
521             h->opencl.queue = ocl->clCreateCommandQueue( context, h->opencl.device, 0, &status );
522             if( status != CL_SUCCESS || !h->opencl.queue )
523                 continue;
524
525             h->opencl.context = context;
526             context = NULL;
527
528             ret = 0;
529             break;
530         }
531
532         if( !ret )
533             break;
534     }
535
536     if( !h->param.psz_clbin_file )
537         h->param.psz_clbin_file = "x264_lookahead.clbin";
538
539     if( ret )
540         x264_log( h, X264_LOG_WARNING, "OpenCL: Unable to find a compatible device\n" );
541     else
542         ret = x264_opencl_lookahead_alloc( h );
543
544 fail:
545     if( context )
546         ocl->clReleaseContext( context );
547     x264_free( imageType );
548     x264_free( devices );
549     x264_free( platforms );
550     return ret;
551 }
552
553 static void x264_opencl_lookahead_free( x264_t *h )
554 {
555     x264_opencl_function_t *ocl = h->opencl.ocl;
556
557 #define RELEASE( a, f ) do { if( a ) { ocl->f( a ); a = NULL; } } while( 0 )
558     RELEASE( h->opencl.downscale_hpel_kernel, clReleaseKernel );
559     RELEASE( h->opencl.downscale_kernel1, clReleaseKernel );
560     RELEASE( h->opencl.downscale_kernel2, clReleaseKernel );
561     RELEASE( h->opencl.weightp_hpel_kernel, clReleaseKernel );
562     RELEASE( h->opencl.weightp_scaled_images_kernel, clReleaseKernel );
563     RELEASE( h->opencl.memset_kernel, clReleaseKernel );
564     RELEASE( h->opencl.intra_kernel, clReleaseKernel );
565     RELEASE( h->opencl.rowsum_intra_kernel, clReleaseKernel );
566     RELEASE( h->opencl.hme_kernel, clReleaseKernel );
567     RELEASE( h->opencl.subpel_refine_kernel, clReleaseKernel );
568     RELEASE( h->opencl.mode_select_kernel, clReleaseKernel );
569     RELEASE( h->opencl.rowsum_inter_kernel, clReleaseKernel );
570
571     RELEASE( h->opencl.lookahead_program, clReleaseProgram );
572
573     RELEASE( h->opencl.page_locked_buffer, clReleaseMemObject );
574     RELEASE( h->opencl.luma_16x16_image[0], clReleaseMemObject );
575     RELEASE( h->opencl.luma_16x16_image[1], clReleaseMemObject );
576     for( int i = 0; i < NUM_IMAGE_SCALES; i++ )
577         RELEASE( h->opencl.weighted_scaled_images[i], clReleaseMemObject );
578     RELEASE( h->opencl.weighted_luma_hpel, clReleaseMemObject );
579     RELEASE( h->opencl.row_satds[0], clReleaseMemObject );
580     RELEASE( h->opencl.row_satds[1], clReleaseMemObject );
581     RELEASE( h->opencl.mv_buffers[0], clReleaseMemObject );
582     RELEASE( h->opencl.mv_buffers[1], clReleaseMemObject );
583     RELEASE( h->opencl.lowres_mv_costs, clReleaseMemObject );
584     RELEASE( h->opencl.mvp_buffer, clReleaseMemObject );
585     RELEASE( h->opencl.lowres_costs[0], clReleaseMemObject );
586     RELEASE( h->opencl.lowres_costs[1], clReleaseMemObject );
587     RELEASE( h->opencl.frame_stats[0], clReleaseMemObject );
588     RELEASE( h->opencl.frame_stats[1], clReleaseMemObject );
589 #undef RELEASE
590 }
591
592 void x264_opencl_lookahead_delete( x264_t *h )
593 {
594     x264_opencl_function_t *ocl = h->opencl.ocl;
595
596     if( !ocl )
597         return;
598
599     if( h->opencl.queue )
600         ocl->clFinish( h->opencl.queue );
601
602     x264_opencl_lookahead_free( h );
603
604     if( h->opencl.queue )
605     {
606         ocl->clReleaseCommandQueue( h->opencl.queue );
607         h->opencl.queue = NULL;
608     }
609     if( h->opencl.context )
610     {
611         ocl->clReleaseContext( h->opencl.context );
612         h->opencl.context = NULL;
613     }
614 }
615
616 void x264_opencl_frame_delete( x264_frame_t *frame )
617 {
618     x264_opencl_function_t *ocl = frame->opencl.ocl;
619
620     if( !ocl )
621         return;
622
623 #define RELEASEBUF(mem) do { if( mem ) { ocl->clReleaseMemObject( mem ); mem = NULL; } } while( 0 )
624     for( int j = 0; j < NUM_IMAGE_SCALES; j++ )
625         RELEASEBUF( frame->opencl.scaled_image2Ds[j] );
626     RELEASEBUF( frame->opencl.luma_hpel );
627     RELEASEBUF( frame->opencl.inv_qscale_factor );
628     RELEASEBUF( frame->opencl.intra_cost );
629     RELEASEBUF( frame->opencl.lowres_mvs0 );
630     RELEASEBUF( frame->opencl.lowres_mvs1 );
631     RELEASEBUF( frame->opencl.lowres_mv_costs0 );
632     RELEASEBUF( frame->opencl.lowres_mv_costs1 );
633 #undef RELEASEBUF
634 }
635
636 /* OpenCL misbehaves on hybrid laptops with Intel iGPU and AMD dGPU, so
637  * we consult AMD's ADL interface to detect this situation and disable
638  * OpenCL on these machines (Linux and Windows) */
639 #ifdef _WIN32
640 #define ADL_API_CALL
641 #define ADL_CALLBACK __stdcall
642 #define adl_close FreeLibrary
643 #define adl_address GetProcAddress
644 #else
645 #define ADL_API_CALL
646 #define ADL_CALLBACK
647 #define adl_close dlclose
648 #define adl_address dlsym
649 #endif
650
651 typedef void* ( ADL_CALLBACK *ADL_MAIN_MALLOC_CALLBACK )( int );
652 typedef int   ( ADL_API_CALL *ADL_MAIN_CONTROL_CREATE )( ADL_MAIN_MALLOC_CALLBACK, int );
653 typedef int   ( ADL_API_CALL *ADL_ADAPTER_NUMBEROFADAPTERS_GET )( int * );
654 typedef int   ( ADL_API_CALL *ADL_POWERXPRESS_SCHEME_GET )( int, int *, int *, int * );
655 typedef int   ( ADL_API_CALL *ADL_MAIN_CONTROL_DESTROY )( void );
656
657 #define ADL_OK 0
658 #define ADL_PX_SCHEME_DYNAMIC 2
659
660 static void* ADL_CALLBACK adl_malloc_wrapper( int iSize )
661 {
662     return x264_malloc( iSize );
663 }
664
665 static int x264_detect_switchable_graphics( void )
666 {
667     void *hDLL;
668     ADL_MAIN_CONTROL_CREATE          ADL_Main_Control_Create;
669     ADL_ADAPTER_NUMBEROFADAPTERS_GET ADL_Adapter_NumberOfAdapters_Get;
670     ADL_POWERXPRESS_SCHEME_GET       ADL_PowerXpress_Scheme_Get;
671     ADL_MAIN_CONTROL_DESTROY         ADL_Main_Control_Destroy;
672     int ret = 0;
673
674 #ifdef _WIN32
675     hDLL = LoadLibrary( "atiadlxx.dll" );
676     if( !hDLL )
677         hDLL = LoadLibrary( "atiadlxy.dll" );
678 #else
679     hDLL = dlopen( "libatiadlxx.so", RTLD_LAZY|RTLD_GLOBAL );
680 #endif
681     if( !hDLL )
682         goto fail0;
683
684     ADL_Main_Control_Create          = (ADL_MAIN_CONTROL_CREATE)adl_address(hDLL, "ADL_Main_Control_Create");
685     ADL_Main_Control_Destroy         = (ADL_MAIN_CONTROL_DESTROY)adl_address(hDLL, "ADL_Main_Control_Destroy");
686     ADL_Adapter_NumberOfAdapters_Get = (ADL_ADAPTER_NUMBEROFADAPTERS_GET)adl_address(hDLL, "ADL_Adapter_NumberOfAdapters_Get");
687     ADL_PowerXpress_Scheme_Get       = (ADL_POWERXPRESS_SCHEME_GET)adl_address(hDLL, "ADL_PowerXpress_Scheme_Get");
688     if( !ADL_Main_Control_Destroy || !ADL_Main_Control_Destroy || !ADL_Adapter_NumberOfAdapters_Get ||
689         !ADL_PowerXpress_Scheme_Get )
690         goto fail1;
691
692     if( ADL_OK != ADL_Main_Control_Create( adl_malloc_wrapper, 1 ) )
693         goto fail1;
694
695     int numAdapters = 0;
696     if( ADL_OK != ADL_Adapter_NumberOfAdapters_Get( &numAdapters ) )
697         goto fail2;
698
699     for( int i = 0; i < numAdapters; i++ )
700     {
701         int PXSchemeRange, PXSchemeCurrentState, PXSchemeDefaultState;
702         if( ADL_OK != ADL_PowerXpress_Scheme_Get( i, &PXSchemeRange, &PXSchemeCurrentState, &PXSchemeDefaultState) )
703             break;
704
705         if( PXSchemeRange >= ADL_PX_SCHEME_DYNAMIC )
706         {
707             ret = 1;
708             break;
709         }
710     }
711
712 fail2:
713     ADL_Main_Control_Destroy();
714 fail1:
715     adl_close( hDLL );
716 fail0:
717     return ret;
718 }