OpenCL support improvement/refactoring
[x262.git] / common / opencl.c
index aa2bd04..5b1bb69 100644 (file)
@@ -4,6 +4,7 @@
  * Copyright (C) 2012-2013 x264 project
  *
  * Authors: Steve Borho <sborho@multicorewareinc.com>
+ *          Anton Mitrofanov <BugMaster@narod.ru>
  *
  * This program is free software; you can redistribute it and/or modify
  * it under the terms of the GNU General Public License as published by
  *****************************************************************************/
 
 #include "common.h"
-#if _WIN32
+
+#ifdef _WIN32
 #include <windows.h>
+#define ocl_open LoadLibrary( "OpenCL" )
+#define ocl_close FreeLibrary
+#define ocl_address GetProcAddress
 #else
 #include <dlfcn.h> //dlopen, dlsym, dlclose
+#if SYS_MACOSX
+#define ocl_open dlopen( "libOpenCL.dylib", RTLD_NOW )
+#else
+#define ocl_open dlopen( "libOpenCL.so", RTLD_NOW )
+#endif
+#define ocl_close dlclose
+#define ocl_address dlsym
 #endif
 
+#define LOAD_OCL_FUNC(name, continue_on_fail)\
+{\
+    ocl->name = (void*)ocl_address( ocl->library, #name );\
+    if( !continue_on_fail && !ocl->name )\
+        goto fail;\
+}
+
+/* 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;
+}
+
+void x264_opencl_close_library( x264_opencl_function_t *ocl )
+{
+    if( !ocl )
+        return;
+    ocl_close( ocl->library );
+    x264_free( ocl );
+}
+
 /* define from recent cl_ext.h, copied here in case headers are old */
 #define CL_DEVICE_SIMD_INSTRUCTION_WIDTH_AMD        0x4042
 
 /* Requires full include path in case of out-of-tree builds */
 #include "common/oclobj.h"
 
-static int x264_detect_switchable_graphics();
+static int x264_detect_switchable_graphics( void );
 
 /* 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, char *devname, char *devvendor, char *driverversion )
 {
-    cl_program program = NULL;
-    cl_int status;
-
     /* try to load cached program binary */
     FILE *fp = fopen( h->param.psz_clbin_file, "rb" );
     if( !fp )
         return NULL;
 
-    fseek( fp, 0L, SEEK_END );
+    x264_opencl_function_t *ocl = h->opencl.ocl;
+    cl_program program = NULL;
+    cl_int status;
+
+    fseek( fp, 0, SEEK_END );
     size_t size = ftell( fp );
     rewind( fp );
     uint8_t *binary;
@@ -75,7 +155,7 @@ static cl_program x264_opencl_cache_load( x264_t *h, char *devname, char *devven
     CHECK_STRING( x264_opencl_source_hash );
 #undef CHECK_STRING
 
-    program = clCreateProgramWithBinary( h->opencl.context, 1, &h->opencl.device, &size, &ptr, NULL, &status );
+    program = ocl->clCreateProgramWithBinary( h->opencl.context, 1, &h->opencl.device, &size, &ptr, NULL, &status );
     if( status != CL_SUCCESS )
         program = NULL;
 
@@ -96,13 +176,14 @@ static void x264_opencl_cache_save( x264_t *h, cl_program program, char *devname
         return;
     }
 
+    x264_opencl_function_t *ocl = h->opencl.ocl;
     size_t size;
-    cl_int status = clGetProgramInfo( program, CL_PROGRAM_BINARY_SIZES, sizeof(size_t), &size, NULL );
+    cl_int status = ocl->clGetProgramInfo( program, CL_PROGRAM_BINARY_SIZES, sizeof(size_t), &size, NULL );
     if( status == CL_SUCCESS )
     {
         uint8_t *binary;
         CHECKED_MALLOC( binary, size );
-        status = clGetProgramInfo( program, CL_PROGRAM_BINARIES, sizeof(uint8_t *), &binary, NULL );
+        status = ocl->clGetProgramInfo( program, CL_PROGRAM_BINARIES, sizeof(uint8_t *), &binary, NULL );
         if( status == CL_SUCCESS )
         {
             fputs( devname, fp );
@@ -133,15 +214,16 @@ fail:
  * 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;
     cl_int status;
 
     char devname[64];
     char devvendor[64];
     char driverversion[64];
-    status  = clGetDeviceInfo( h->opencl.device, CL_DEVICE_NAME,    sizeof(devname), devname, NULL );
-    status |= clGetDeviceInfo( h->opencl.device, CL_DEVICE_VENDOR,  sizeof(devvendor), devvendor, NULL );
-    status |= clGetDeviceInfo( h->opencl.device, CL_DRIVER_VERSION, sizeof(driverversion), driverversion, NULL );
+    status  = ocl->clGetDeviceInfo( h->opencl.device, CL_DEVICE_NAME,    sizeof(devname), devname, NULL );
+    status |= ocl->clGetDeviceInfo( h->opencl.device, CL_DEVICE_VENDOR,  sizeof(devvendor), devvendor, NULL );
+    status |= ocl->clGetDeviceInfo( h->opencl.device, CL_DRIVER_VERSION, sizeof(driverversion), driverversion, NULL );
     if( status != CL_SUCCESS )
         return NULL;
 
@@ -160,7 +242,7 @@ static cl_program x264_opencl_compile( x264_t *h )
 
         /* Detect AMD SouthernIsland or newer device (single-width registers) */
         cl_uint simdwidth = 4;
-        status = clGetDeviceInfo( h->opencl.device, CL_DEVICE_SIMD_INSTRUCTION_WIDTH_AMD, sizeof(cl_uint), &simdwidth, NULL );
+        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;
@@ -177,7 +259,7 @@ static cl_program x264_opencl_compile( x264_t *h )
         x264_log( h, X264_LOG_INFO, "Compiling OpenCL kernels...\n" );
         const char *strptr = (const char*)x264_opencl_source;
         size_t size = sizeof(x264_opencl_source);
-        program = clCreateProgramWithSource( h->opencl.context, 1, &strptr, &size, &status );
+        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" );
@@ -187,7 +269,7 @@ static cl_program x264_opencl_compile( x264_t *h )
 
     /* Build the program binary for the OpenCL device */
     const char *buildopts = vectorize ? "-DVECTORIZE=1" : "";
-    status = clBuildProgram( program, 1, &h->opencl.device, buildopts, NULL, NULL );
+    status = ocl->clBuildProgram( program, 1, &h->opencl.device, buildopts, NULL, NULL );
     if( status == CL_SUCCESS )
     {
         x264_opencl_cache_save( h, program, devname, devvendor, driverversion );
@@ -198,7 +280,7 @@ static cl_program x264_opencl_compile( x264_t *h )
 
     size_t build_log_len = 0;
 
-    status = clGetProgramBuildInfo( program, h->opencl.device, CL_PROGRAM_BUILD_LOG, build_log_len, NULL, &build_log_len );
+    status = ocl->clGetProgramBuildInfo( program, h->opencl.device, CL_PROGRAM_BUILD_LOG, build_log_len, NULL, &build_log_len );
     if( status != CL_SUCCESS )
     {
         x264_log( h, X264_LOG_WARNING, "OpenCL: Compilation failed, unable to query build log\n" );
@@ -213,7 +295,7 @@ static cl_program x264_opencl_compile( x264_t *h )
         return NULL;
     }
 
-    status = clGetProgramBuildInfo( program, h->opencl.device, CL_PROGRAM_BUILD_LOG, build_log_len, build_log, NULL );
+    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" );
@@ -234,42 +316,7 @@ fail:
     return NULL;
 }
 
-static void x264_opencl_free_lookahead( x264_t *h )
-{
-#define RELEASE( a, f ) if( a ) f( a );
-    RELEASE( h->opencl.intra_kernel, clReleaseKernel )
-    RELEASE( h->opencl.rowsum_intra_kernel, clReleaseKernel )
-    RELEASE( h->opencl.downscale_kernel1, clReleaseKernel )
-    RELEASE( h->opencl.downscale_kernel2, clReleaseKernel )
-    RELEASE( h->opencl.downscale_hpel_kernel, clReleaseKernel )
-    RELEASE( h->opencl.weightp_hpel_kernel, clReleaseKernel )
-    RELEASE( h->opencl.weightp_scaled_images_kernel, clReleaseKernel )
-    RELEASE( h->opencl.memset_kernel, clReleaseKernel )
-    RELEASE( h->opencl.hme_kernel, clReleaseKernel )
-    RELEASE( h->opencl.subpel_refine_kernel, clReleaseKernel )
-    RELEASE( h->opencl.mode_select_kernel, clReleaseKernel )
-    RELEASE( h->opencl.rowsum_inter_kernel, clReleaseKernel )
-    RELEASE( h->opencl.lookahead_program, clReleaseProgram )
-    RELEASE( h->opencl.row_satds[0], clReleaseMemObject )
-    RELEASE( h->opencl.row_satds[1], clReleaseMemObject )
-    RELEASE( h->opencl.frame_stats[0], clReleaseMemObject )
-    RELEASE( h->opencl.frame_stats[1], clReleaseMemObject )
-    RELEASE( h->opencl.mv_buffers[0], clReleaseMemObject )
-    RELEASE( h->opencl.mv_buffers[1], clReleaseMemObject )
-    RELEASE( h->opencl.mvp_buffer, clReleaseMemObject )
-    RELEASE( h->opencl.luma_16x16_image[0], clReleaseMemObject )
-    RELEASE( h->opencl.luma_16x16_image[1], clReleaseMemObject )
-    RELEASE( h->opencl.lowres_mv_costs, clReleaseMemObject )
-    RELEASE( h->opencl.lowres_costs[0], clReleaseMemObject )
-    RELEASE( h->opencl.lowres_costs[1], clReleaseMemObject )
-    RELEASE( h->opencl.page_locked_buffer, clReleaseMemObject )
-    RELEASE( h->opencl.weighted_luma_hpel, clReleaseMemObject )
-    for( int i = 0; i < NUM_IMAGE_SCALES; i++ )
-        RELEASE( h->opencl.weighted_scaled_images[i], clReleaseMemObject )
-#undef RELEASE
-}
-
-int x264_opencl_init_lookahead( x264_t *h )
+static int x264_opencl_lookahead_alloc( x264_t *h )
 {
     if( !h->param.rc.i_lookahead )
         return -1;
@@ -288,6 +335,7 @@ int x264_opencl_init_lookahead( x264_t *h )
         "mode_selection",
         "sum_inter_cost"
     };
+
     cl_kernel *kernels[] = {
         &h->opencl.intra_kernel,
         &h->opencl.rowsum_intra_kernel,
@@ -302,43 +350,42 @@ int x264_opencl_init_lookahead( x264_t *h )
         &h->opencl.mode_select_kernel,
         &h->opencl.rowsum_inter_kernel
     };
+
+    x264_opencl_function_t *ocl = h->opencl.ocl;
     cl_int status;
 
     h->opencl.lookahead_program = x264_opencl_compile( h );
     if( !h->opencl.lookahead_program )
-    {
-        x264_opencl_free_lookahead( h );
-        return -1;
-    }
+        goto fail;
 
     for( int i = 0; i < ARRAY_SIZE(kernelnames); i++ )
     {
-        *kernels[i] = clCreateKernel( h->opencl.lookahead_program, kernelnames[i], &status );
+        *kernels[i] = ocl->clCreateKernel( h->opencl.lookahead_program, kernelnames[i], &status );
         if( status != CL_SUCCESS )
         {
             x264_log( h, X264_LOG_WARNING, "OpenCL: Unable to compile kernel '%s' (%d)\n", kernelnames[i], status );
-            x264_opencl_free_lookahead( h );
-            return -1;
+            goto fail;
         }
     }
 
-    h->opencl.page_locked_buffer = clCreateBuffer( h->opencl.context, CL_MEM_WRITE_ONLY|CL_MEM_ALLOC_HOST_PTR, PAGE_LOCKED_BUF_SIZE, NULL, &status );
+    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 );
     if( status != CL_SUCCESS )
     {
         x264_log( h, X264_LOG_WARNING, "OpenCL: Unable to allocate page-locked buffer, error '%d'\n", status );
-        x264_opencl_free_lookahead( h );
-        return -1;
+        goto fail;
     }
-    h->opencl.page_locked_ptr = clEnqueueMapBuffer( h->opencl.queue, h->opencl.page_locked_buffer, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE,
-                                                    0, PAGE_LOCKED_BUF_SIZE, 0, NULL, NULL, &status );
+    h->opencl.page_locked_ptr = ocl->clEnqueueMapBuffer( h->opencl.queue, h->opencl.page_locked_buffer, CL_TRUE, CL_MAP_READ | CL_MAP_WRITE,
+                                                         0, PAGE_LOCKED_BUF_SIZE, 0, NULL, NULL, &status );
     if( status != CL_SUCCESS )
     {
         x264_log( h, X264_LOG_WARNING, "OpenCL: Unable to map page-locked buffer, error '%d'\n", status );
-        x264_opencl_free_lookahead( h );
-        return -1;
+        goto fail;
     }
 
     return 0;
+fail:
+    x264_opencl_lookahead_delete( h );
+    return -1;
 }
 
 static void x264_opencl_error_notify( const char *errinfo, const void *private_info, size_t cb, void *user_data )
@@ -352,13 +399,14 @@ static void x264_opencl_error_notify( const char *errinfo, const void *private_i
     x264_log( h, X264_LOG_ERROR, "OpenCL: fatal error, aborting encode\n" );
 }
 
-int x264_opencl_init( x264_t *h )
+int x264_opencl_lookahead_init( x264_t *h )
 {
+    x264_opencl_function_t *ocl = h->opencl.ocl;
     cl_int status;
     cl_uint numPlatforms;
     int ret = -1;
 
-    status = clGetPlatformIDs( 0, NULL, &numPlatforms );
+    status = ocl->clGetPlatformIDs( 0, NULL, &numPlatforms );
     if( status != CL_SUCCESS || numPlatforms == 0 )
     {
         x264_log( h, X264_LOG_WARNING, "OpenCL: Unable to query installed platforms\n");
@@ -366,7 +414,7 @@ int x264_opencl_init( x264_t *h )
     }
 
     cl_platform_id *platforms = (cl_platform_id*)x264_malloc( numPlatforms * sizeof(cl_platform_id) );
-    status = clGetPlatformIDs( numPlatforms, platforms, NULL );
+    status = ocl->clGetPlatformIDs( numPlatforms, platforms, NULL );
     if( status != CL_SUCCESS )
     {
         x264_log( h, X264_LOG_WARNING, "OpenCL: Unable to query installed platforms\n");
@@ -379,7 +427,7 @@ int x264_opencl_init( x264_t *h )
     for( cl_uint i = 0; i < numPlatforms; ++i )
     {
         cl_uint gpu_count = 0;
-        status = clGetDeviceIDs( platforms[i], CL_DEVICE_TYPE_GPU, 0, NULL, &gpu_count );
+        status = ocl->clGetDeviceIDs( platforms[i], CL_DEVICE_TYPE_GPU, 0, NULL, &gpu_count );
         if( status != CL_SUCCESS || !gpu_count )
             continue;
 
@@ -387,7 +435,7 @@ int x264_opencl_init( x264_t *h )
         if( !devices )
             continue;
 
-        status = clGetDeviceIDs( platforms[i], CL_DEVICE_TYPE_GPU, gpu_count, devices, NULL );
+        status = ocl->clGetDeviceIDs( platforms[i], CL_DEVICE_TYPE_GPU, gpu_count, devices, NULL );
         if( status != CL_SUCCESS )
         {
             x264_free( devices );
@@ -406,30 +454,30 @@ int x264_opencl_init( x264_t *h )
                 continue;
 
             cl_bool image_support;
-            clGetDeviceInfo( h->opencl.device, CL_DEVICE_IMAGE_SUPPORT, sizeof(cl_bool), &image_support, NULL );
+            ocl->clGetDeviceInfo( h->opencl.device, CL_DEVICE_IMAGE_SUPPORT, sizeof(cl_bool), &image_support, NULL );
             if( !image_support )
                 continue;
 
-            cl_context context = clCreateContext( NULL, 1, &h->opencl.device, (void*)x264_opencl_error_notify, (void*)h, &status );
+            cl_context context = ocl->clCreateContext( NULL, 1, &h->opencl.device, (void*)x264_opencl_error_notify, (void*)h, &status );
             if( status != CL_SUCCESS )
                 continue;
 
             cl_uint imagecount = 0;
-            clGetSupportedImageFormats( context, CL_MEM_READ_WRITE, CL_MEM_OBJECT_IMAGE2D, 0, NULL, &imagecount );
+            ocl->clGetSupportedImageFormats( context, CL_MEM_READ_WRITE, CL_MEM_OBJECT_IMAGE2D, 0, NULL, &imagecount );
             if( !imagecount )
             {
-                clReleaseContext( context );
+                ocl->clReleaseContext( context );
                 continue;
             }
 
             cl_image_format *imageType = x264_malloc( sizeof(cl_image_format) * imagecount );
             if( !imageType )
             {
-                clReleaseContext( context );
+                ocl->clReleaseContext( context );
                 continue;
             }
 
-            clGetSupportedImageFormats( context, CL_MEM_READ_WRITE, CL_MEM_OBJECT_IMAGE2D, imagecount, imageType, NULL );
+            ocl->clGetSupportedImageFormats( context, CL_MEM_READ_WRITE, CL_MEM_OBJECT_IMAGE2D, imagecount, imageType, NULL );
 
             int b_has_r = 0;
             int b_has_rgba = 0;
@@ -446,14 +494,14 @@ int x264_opencl_init( x264_t *h )
             if( !b_has_r || !b_has_rgba )
             {
                 char devname[64];
-                status = clGetDeviceInfo( h->opencl.device, CL_DEVICE_NAME, sizeof(devname), devname, NULL );
+                status = ocl->clGetDeviceInfo( h->opencl.device, CL_DEVICE_NAME, sizeof(devname), devname, 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", devname);
                 }
-                clReleaseContext( context );
+                ocl->clReleaseContext( context );
                 continue;
             }
 
@@ -461,14 +509,14 @@ int x264_opencl_init( x264_t *h )
             if( h->param.i_opencl_device )
             {
                 h->param.i_opencl_device--;
-                clReleaseContext( context );
+                ocl->clReleaseContext( context );
                 continue;
             }
 
-            h->opencl.queue = clCreateCommandQueue( context, h->opencl.device, 0, &status );
+            h->opencl.queue = ocl->clCreateCommandQueue( context, h->opencl.device, 0, &status );
             if( status != CL_SUCCESS )
             {
-                clReleaseContext( context );
+                ocl->clReleaseContext( context );
                 continue;
             }
 
@@ -492,14 +540,82 @@ int x264_opencl_init( x264_t *h )
     if( ret )
         x264_log( h, X264_LOG_WARNING, "OpenCL: Unable to find a compatible device\n");
     else
-        ret = x264_opencl_init_lookahead( h );
+        ret = x264_opencl_lookahead_alloc( h );
 
     return ret;
 }
 
+static void x264_opencl_lookahead_free( x264_t *h )
+{
+    x264_opencl_function_t *ocl = h->opencl.ocl;
+
+#define RELEASE( a, f ) do { if( a ) { ocl->f( a ); a = NULL; } } while( 0 )
+    RELEASE( h->opencl.downscale_hpel_kernel, clReleaseKernel );
+    RELEASE( h->opencl.downscale_kernel1, clReleaseKernel );
+    RELEASE( h->opencl.downscale_kernel2, clReleaseKernel );
+    RELEASE( h->opencl.weightp_hpel_kernel, clReleaseKernel );
+    RELEASE( h->opencl.weightp_scaled_images_kernel, clReleaseKernel );
+    RELEASE( h->opencl.memset_kernel, clReleaseKernel );
+    RELEASE( h->opencl.intra_kernel, clReleaseKernel );
+    RELEASE( h->opencl.rowsum_intra_kernel, clReleaseKernel );
+    RELEASE( h->opencl.hme_kernel, clReleaseKernel );
+    RELEASE( h->opencl.subpel_refine_kernel, clReleaseKernel );
+    RELEASE( h->opencl.mode_select_kernel, clReleaseKernel );
+    RELEASE( h->opencl.rowsum_inter_kernel, clReleaseKernel );
+
+    RELEASE( h->opencl.lookahead_program, clReleaseProgram );
+
+    RELEASE( h->opencl.page_locked_buffer, clReleaseMemObject );
+    RELEASE( h->opencl.luma_16x16_image[0], clReleaseMemObject );
+    RELEASE( h->opencl.luma_16x16_image[1], clReleaseMemObject );
+    for( int i = 0; i < NUM_IMAGE_SCALES; i++ )
+        RELEASE( h->opencl.weighted_scaled_images[i], clReleaseMemObject );
+    RELEASE( h->opencl.weighted_luma_hpel, clReleaseMemObject );
+    RELEASE( h->opencl.row_satds[0], clReleaseMemObject );
+    RELEASE( h->opencl.row_satds[1], clReleaseMemObject );
+    RELEASE( h->opencl.mv_buffers[0], clReleaseMemObject );
+    RELEASE( h->opencl.mv_buffers[1], clReleaseMemObject );
+    RELEASE( h->opencl.lowres_mv_costs, clReleaseMemObject );
+    RELEASE( h->opencl.mvp_buffer, clReleaseMemObject );
+    RELEASE( h->opencl.lowres_costs[0], clReleaseMemObject );
+    RELEASE( h->opencl.lowres_costs[1], clReleaseMemObject );
+    RELEASE( h->opencl.frame_stats[0], clReleaseMemObject );
+    RELEASE( h->opencl.frame_stats[1], clReleaseMemObject );
+#undef RELEASE
+}
+
+void x264_opencl_lookahead_delete( x264_t *h )
+{
+    x264_opencl_function_t *ocl = h->opencl.ocl;
+
+    if( !ocl )
+        return;
+
+    if( h->opencl.queue )
+        ocl->clFinish( h->opencl.queue );
+
+    x264_opencl_lookahead_free( h );
+
+    if( h->opencl.queue )
+    {
+        ocl->clReleaseCommandQueue( h->opencl.queue );
+        h->opencl.queue = NULL;
+    }
+    if( h->opencl.context )
+    {
+        ocl->clReleaseContext( h->opencl.context );
+        h->opencl.context = NULL;
+    }
+}
+
 void x264_opencl_frame_delete( x264_frame_t *frame )
 {
-#define RELEASEBUF(mem) if( mem ) clReleaseMemObject( mem );
+    x264_opencl_function_t *ocl = frame->opencl.ocl;
+
+    if( !ocl )
+        return;
+
+#define RELEASEBUF(mem) do { if( mem ) { ocl->clReleaseMemObject( mem ); mem = NULL; } } while( 0 )
     for( int j = 0; j < NUM_IMAGE_SCALES; j++ )
         RELEASEBUF( frame->opencl.scaled_image2Ds[j] );
     RELEASEBUF( frame->opencl.luma_hpel );
@@ -512,73 +628,68 @@ void x264_opencl_frame_delete( x264_frame_t *frame )
 #undef RELEASEBUF
 }
 
-void x264_opencl_free( x264_t *h )
-{
-    if( h->opencl.queue )
-        clFinish(h->opencl.queue );
-
-    x264_opencl_free_lookahead( h );
-
-    if( h->opencl.queue )
-        clReleaseCommandQueue( h->opencl.queue );
-    if( h->opencl.context )
-        clReleaseContext( h->opencl.context );
-}
-
 /* OpenCL misbehaves on hybrid laptops with Intel iGPU and AMD dGPU, so
  * we consult AMD's ADL interface to detect this situation and disable
  * OpenCL on these machines (Linux and Windows) */
-#ifndef _WIN32
-#define __stdcall
-#define HINSTANCE void *
+#ifdef _WIN32
+#define ADL_API_CALL
+#define ADL_CALLBACK __stdcall
+#define adl_close FreeLibrary
+#define adl_address GetProcAddress
+#else
+#define ADL_API_CALL
+#define ADL_CALLBACK
+#define adl_close dlclose
+#define adl_address dlsym
 #endif
-typedef void* ( __stdcall *ADL_MAIN_MALLOC_CALLBACK )( int );
-typedef int ( *ADL_MAIN_CONTROL_CREATE )(ADL_MAIN_MALLOC_CALLBACK, int );
-typedef int ( *ADL_ADAPTER_NUMBEROFADAPTERS_GET ) ( int* );
-typedef int ( *ADL_POWERXPRESS_SCHEME_GET ) ( int, int *, int *, int * );
-typedef int ( *ADL_MAIN_CONTROL_DESTROY )();
+
+typedef void* ( ADL_CALLBACK *ADL_MAIN_MALLOC_CALLBACK )( int );
+typedef int   ( ADL_API_CALL *ADL_MAIN_CONTROL_CREATE )( ADL_MAIN_MALLOC_CALLBACK, int );
+typedef int   ( ADL_API_CALL *ADL_ADAPTER_NUMBEROFADAPTERS_GET )( int * );
+typedef int   ( ADL_API_CALL *ADL_POWERXPRESS_SCHEME_GET )( int, int *, int *, int * );
+typedef int   ( ADL_API_CALL *ADL_MAIN_CONTROL_DESTROY )( void );
+
 #define ADL_OK 0
 #define ADL_PX_SCHEME_DYNAMIC 2
 
-void* __stdcall adl_malloc_wrapper( int iSize ) { return x264_malloc( iSize ); }
+static void* ADL_CALLBACK adl_malloc_wrapper( int iSize )
+{
+    return x264_malloc( iSize );
+}
 
-static int x264_detect_switchable_graphics()
+static int x264_detect_switchable_graphics( void )
 {
+    void *hDLL;
     ADL_MAIN_CONTROL_CREATE          ADL_Main_Control_Create;
     ADL_ADAPTER_NUMBEROFADAPTERS_GET ADL_Adapter_NumberOfAdapters_Get;
     ADL_POWERXPRESS_SCHEME_GET       ADL_PowerXpress_Scheme_Get;
     ADL_MAIN_CONTROL_DESTROY         ADL_Main_Control_Destroy;
-    HINSTANCE hDLL;
     int ret = 0;
 
-#if _WIN32
+#ifdef _WIN32
     hDLL = LoadLibrary( "atiadlxx.dll" );
     if( !hDLL )
         hDLL = LoadLibrary( "atiadlxy.dll" );
 #else
     hDLL = dlopen( "libatiadlxx.so", RTLD_LAZY|RTLD_GLOBAL );
-#define GetProcAddress dlsym
 #endif
     if( !hDLL )
-        return ret;
+        goto fail0;
 
-    ADL_Main_Control_Create = (ADL_MAIN_CONTROL_CREATE) GetProcAddress(hDLL, "ADL_Main_Control_Create");
-    ADL_Main_Control_Destroy = (ADL_MAIN_CONTROL_DESTROY) GetProcAddress(hDLL, "ADL_Main_Control_Destroy");
-    ADL_Adapter_NumberOfAdapters_Get = (ADL_ADAPTER_NUMBEROFADAPTERS_GET) GetProcAddress(hDLL, "ADL_Adapter_NumberOfAdapters_Get");
-    ADL_PowerXpress_Scheme_Get = (ADL_POWERXPRESS_SCHEME_GET) GetProcAddress(hDLL, "ADL_PowerXpress_Scheme_Get");
+    ADL_Main_Control_Create          = (ADL_MAIN_CONTROL_CREATE)adl_address(hDLL, "ADL_Main_Control_Create");
+    ADL_Main_Control_Destroy         = (ADL_MAIN_CONTROL_DESTROY)adl_address(hDLL, "ADL_Main_Control_Destroy");
+    ADL_Adapter_NumberOfAdapters_Get = (ADL_ADAPTER_NUMBEROFADAPTERS_GET)adl_address(hDLL, "ADL_Adapter_NumberOfAdapters_Get");
+    ADL_PowerXpress_Scheme_Get       = (ADL_POWERXPRESS_SCHEME_GET)adl_address(hDLL, "ADL_PowerXpress_Scheme_Get");
     if( !ADL_Main_Control_Destroy || !ADL_Main_Control_Destroy || !ADL_Adapter_NumberOfAdapters_Get ||
         !ADL_PowerXpress_Scheme_Get )
-        goto bail;
+        goto fail1;
 
-    if( ADL_OK != ADL_Main_Control_Create( adl_malloc_wrapper, 1) )
-        goto bail;
+    if( ADL_OK != ADL_Main_Control_Create( adl_malloc_wrapper, 1 ) )
+        goto fail1;
 
     int numAdapters = 0;
     if( ADL_OK != ADL_Adapter_NumberOfAdapters_Get( &numAdapters ) )
-    {
-        ADL_Main_Control_Destroy();
-        goto bail;
-    }
+        goto fail2;
 
     for( int i = 0; i < numAdapters; i++ )
     {
@@ -593,14 +704,10 @@ static int x264_detect_switchable_graphics()
         }
     }
 
+fail2:
     ADL_Main_Control_Destroy();
-
-bail:
-#if _WIN32
-    FreeLibrary( hDLL );
-#else
-    dlclose( hDLL );
-#endif
-
+fail1:
+    adl_close( hDLL );
+fail0:
     return ret;
 }