OpenCL: initial work on a new OpenCl wrapper, which will eventually phase out the existing one.

Based on an x264 patch by Anton Mitrofanov.



git-svn-id: svn://svn.handbrake.fr/HandBrake/trunk@5870 b64f7644-9d1e-0410-96f1-a4d463321fa5

6 files changed, 995 insertions, 70 deletions

diff --git a/libhb/hb.c b/libhb/hb.c
index 4186fa65a..21f38a186 100644
--- a/libhb/hb.c
+++ b/libhb/hb.c
@@ -8,6 +8,7 @@
  */
  
 #include "hb.h"
+#include "opencl.h"
 #include "hbffmpeg.h"
 #include <stdio.h>
 #include <unistd.h>
@@ -630,10 +631,8 @@ void hb_scan( hb_handle_t * h, const char * path, int title_index,
     }
     hb_log(" - logical processor count: %d", hb_get_cpu_count());
 
-#ifdef USE_OPENCL
     /* Print OpenCL info here so that it's in all scan and encode logs */
     hb_opencl_info_print();
-#endif
 
 #ifdef USE_QSV
     /* Print QSV info here so that it's in all scan and encode logs */
diff --git a/libhb/opencl.c b/libhb/opencl.c
new file mode 100644
index 000000000..cfc5e5747
--- /dev/null
+++ b/libhb/opencl.c
@@ -0,0 +1,324 @@
+/* opencl.c
+
+   Copyright (c) 2003-2013 HandBrake Team
+   This file is part of the HandBrake source code
+   Homepage: <http://handbrake.fr/>.
+   It may be used under the terms of the GNU General Public License v2.
+   For full terms see the file COPYING file or visit http://www.gnu.org/licenses/gpl-2.0.html
+ */
+
+#ifdef _WIN32
+#include <windows.h>
+#define HB_OCL_DLOPEN  LoadLibraryW(L"OpenCL")
+#define HB_OCL_DLSYM   GetProcAddress
+#define HB_OCL_DLCLOSE FreeLibrary
+#else
+#include <dlfcn.h>
+#ifdef __APPLE__
+#define HB_OCL_DLOPEN  dlopen("/System/Library/Frameworks/OpenCL.framework/OpenCL", RTLD_NOW)
+#else
+#define HB_OCL_DLOPEN  dlopen("libOpenCL.so", RTLD_NOW)
+#endif
+#define HB_OCL_DLSYM   dlsym
+#define HB_OCL_DLCLOSE dlclose
+#endif
+
+#include "common.h"
+#include "opencl.h"
+
+int hb_opencl_library_open(hb_opencl_library_t *opencl)
+{
+    if (opencl == NULL)
+    {
+        goto fail;
+    }
+
+    opencl->library = HB_OCL_DLOPEN;
+    if (opencl->library == NULL)
+    {
+        goto fail;
+    }
+
+#define HB_OCL_LOAD(func)                                                      \
+{                                                                              \
+    if ((opencl->func = (void*)HB_OCL_DLSYM(opencl->library, #func)) == NULL)  \
+    {                                                                          \
+        hb_log("hb_opencl_library_open: failed to load function '%s'", #func); \
+        goto fail;                                                             \
+    }                                                                          \
+}
+    HB_OCL_LOAD(clBuildProgram);
+    HB_OCL_LOAD(clCreateBuffer);
+    HB_OCL_LOAD(clCreateCommandQueue);
+    HB_OCL_LOAD(clCreateContextFromType);
+    HB_OCL_LOAD(clCreateKernel);
+    HB_OCL_LOAD(clCreateProgramWithBinary);
+    HB_OCL_LOAD(clCreateProgramWithSource);
+    HB_OCL_LOAD(clEnqueueCopyBuffer);
+    HB_OCL_LOAD(clEnqueueMapBuffer);
+    HB_OCL_LOAD(clEnqueueNDRangeKernel);
+    HB_OCL_LOAD(clEnqueueReadBuffer);
+    HB_OCL_LOAD(clEnqueueUnmapMemObject);
+    HB_OCL_LOAD(clEnqueueWriteBuffer);
+    HB_OCL_LOAD(clFlush);
+    HB_OCL_LOAD(clGetCommandQueueInfo);
+    HB_OCL_LOAD(clGetContextInfo);
+    HB_OCL_LOAD(clGetDeviceIDs);
+    HB_OCL_LOAD(clGetDeviceInfo);
+    HB_OCL_LOAD(clGetPlatformIDs);
+    HB_OCL_LOAD(clGetPlatformInfo);
+    HB_OCL_LOAD(clGetProgramBuildInfo);
+    HB_OCL_LOAD(clGetProgramInfo);
+    HB_OCL_LOAD(clReleaseCommandQueue);
+    HB_OCL_LOAD(clReleaseContext);
+    HB_OCL_LOAD(clReleaseEvent);
+    HB_OCL_LOAD(clReleaseKernel);
+    HB_OCL_LOAD(clReleaseProgram);
+    HB_OCL_LOAD(clSetKernelArg);
+    HB_OCL_LOAD(clWaitForEvents);
+    return 0;
+
+fail:
+    hb_opencl_library_close(opencl);
+    return -1;
+}
+
+void hb_opencl_library_close(hb_opencl_library_t *opencl)
+{
+    if (opencl != NULL)
+    {
+        if (opencl->library != NULL)
+        {
+            HB_OCL_DLCLOSE(opencl->library);
+        }
+        opencl->library = NULL;
+
+#define HB_OCL_UNLOAD(func) { opencl->func = NULL; }
+        HB_OCL_UNLOAD(clBuildProgram);
+        HB_OCL_UNLOAD(clCreateBuffer);
+        HB_OCL_UNLOAD(clCreateCommandQueue);
+        HB_OCL_UNLOAD(clCreateContextFromType);
+        HB_OCL_UNLOAD(clCreateKernel);
+        HB_OCL_UNLOAD(clCreateProgramWithBinary);
+        HB_OCL_UNLOAD(clCreateProgramWithSource);
+        HB_OCL_UNLOAD(clEnqueueCopyBuffer);
+        HB_OCL_UNLOAD(clEnqueueMapBuffer);
+        HB_OCL_UNLOAD(clEnqueueNDRangeKernel);
+        HB_OCL_UNLOAD(clEnqueueReadBuffer);
+        HB_OCL_UNLOAD(clEnqueueUnmapMemObject);
+        HB_OCL_UNLOAD(clEnqueueWriteBuffer);
+        HB_OCL_UNLOAD(clFlush);
+        HB_OCL_UNLOAD(clGetCommandQueueInfo);
+        HB_OCL_UNLOAD(clGetContextInfo);
+        HB_OCL_UNLOAD(clGetDeviceIDs);
+        HB_OCL_UNLOAD(clGetDeviceInfo);
+        HB_OCL_UNLOAD(clGetPlatformIDs);
+        HB_OCL_UNLOAD(clGetPlatformInfo);
+        HB_OCL_UNLOAD(clGetProgramBuildInfo);
+        HB_OCL_UNLOAD(clGetProgramInfo);
+        HB_OCL_UNLOAD(clReleaseCommandQueue);
+        HB_OCL_UNLOAD(clReleaseContext);
+        HB_OCL_UNLOAD(clReleaseEvent);
+        HB_OCL_UNLOAD(clReleaseKernel);
+        HB_OCL_UNLOAD(clReleaseProgram);
+        HB_OCL_UNLOAD(clSetKernelArg);
+        HB_OCL_UNLOAD(clWaitForEvents);
+    }
+}
+
+static int hb_opencl_device_is_supported(cl_device_type type,
+                                         const char *vendor,
+                                         const char *version)
+{
+    int major, minor;
+
+    // we only support OpenCL on GPUs
+    // disable on NVIDIA to to a bug (FIXME)
+    if (!(type & CL_DEVICE_TYPE_GPU) ||
+        !(strncmp(vendor, "NVIDIA", 6 /* strlen("NVIDIA") */)))
+    {
+        return 0;
+    }
+
+    // check OpenCL version; format:
+    // OpenCL<space><major_version.minor_version><space><vendor-specific information>
+    if (sscanf(version, "OpenCL %d.%d", &major, &minor) != 2)
+    {
+        return 0;
+    }
+
+    return (major > HB_OCL_MINVERSION_MAJOR) || (major == HB_OCL_MINVERSION_MAJOR &&
+                                                 minor >= HB_OCL_MINVERSION_MINOR);
+}
+
+int hb_opencl_available()
+{
+    static int opencl_available = -1;
+    if (opencl_available >= 0)
+    {
+        return opencl_available;
+    }
+    opencl_available = 0;
+
+    cl_device_type type;
+    char vendor[100], version[100];
+    cl_device_id *device_ids = NULL;
+    cl_platform_id *platform_ids = NULL;
+    hb_opencl_library_t lib, *opencl = &lib;
+    cl_uint i, j, num_platforms, num_devices;
+
+    /*
+     * Check whether we can load the OpenCL library, then check devices and make
+     * sure we support running OpenCL code on at least one of them.
+     */
+    if (hb_opencl_library_open(opencl) == 0)
+    {
+        if (opencl->clGetPlatformIDs(0, NULL, &num_platforms) != CL_SUCCESS || !num_platforms)
+        {
+            goto end;
+        }
+        if ((platform_ids = malloc(sizeof(cl_platform_id) * num_platforms)) == NULL)
+        {
+            goto end;
+        }
+        if (opencl->clGetPlatformIDs(num_platforms, platform_ids, NULL) != CL_SUCCESS)
+        {
+            goto end;
+        }
+        for (i = 0; i < num_platforms; i++)
+        {
+            if (opencl->clGetDeviceIDs(platform_ids[i], CL_DEVICE_TYPE_ALL, 0, NULL, &num_devices) != CL_SUCCESS || !num_devices)
+            {
+                goto end;
+            }
+            if ((device_ids = malloc(sizeof(cl_device_id) * num_devices)) == NULL)
+            {
+                goto end;
+            }
+            if (opencl->clGetDeviceIDs(platform_ids[i], CL_DEVICE_TYPE_ALL, num_devices, device_ids, NULL) != CL_SUCCESS)
+            {
+                goto end;
+            }
+            for (j = 0; j < num_devices; j++)
+            {
+                if (device_ids[j] != NULL)
+                {
+                    opencl->clGetDeviceInfo(device_ids[j], CL_DEVICE_VENDOR,  sizeof(vendor),
+                                            vendor,  NULL);
+                    opencl->clGetDeviceInfo(device_ids[j], CL_DEVICE_VERSION, sizeof(version),
+                                            version, NULL);
+                    opencl->clGetDeviceInfo(device_ids[j], CL_DEVICE_TYPE,    sizeof(type),
+                                            &type,  NULL);
+
+                    if (hb_opencl_device_is_supported(type,
+                                                      (const char*)vendor,
+                                                      (const char*)version))
+                    {
+                        opencl_available = 1;
+                        goto end;
+                    }
+                }
+            }
+            free(device_ids);
+            device_ids = NULL;
+        }
+    }
+
+end:
+    free(device_ids);
+    free(platform_ids);
+    hb_opencl_library_close(opencl);
+    return opencl_available;
+}
+
+void hb_opencl_info_print()
+{
+    /*
+     * Note: this function should not log any warnings or errors.
+     * Its only purpose is to list OpenCL-capable devices, so let's initialize
+     * only what we absolutely need here, rather than calling library_open().
+     */
+    hb_opencl_library_t lib, *opencl = &lib;
+    if ((opencl->library          = (void*)HB_OCL_DLOPEN)                                     == NULL ||
+        (opencl->clGetDeviceIDs   = (void*)HB_OCL_DLSYM(opencl->library, "clGetDeviceIDs"  )) == NULL ||
+        (opencl->clGetDeviceInfo  = (void*)HB_OCL_DLSYM(opencl->library, "clGetDeviceInfo" )) == NULL ||
+        (opencl->clGetPlatformIDs = (void*)HB_OCL_DLSYM(opencl->library, "clGetPlatformIDs")) == NULL)
+    {
+        // zero or insufficient OpenCL support
+        hb_log("OpenCL: library not available");
+        goto end;
+    }
+
+    cl_device_type type;
+    cl_device_id *device_ids;
+    cl_platform_id *platform_ids;
+    cl_uint i, j, k, num_platforms, num_devices;
+    char vendor[100], name[1024], version[100], driver[1024];
+
+    if (opencl->clGetPlatformIDs(0, NULL, &num_platforms) != CL_SUCCESS || !num_platforms)
+    {
+        goto end;
+    }
+    if ((platform_ids = malloc(sizeof(cl_platform_id) * num_platforms)) == NULL)
+    {
+        goto end;
+    }
+    if (opencl->clGetPlatformIDs(num_platforms, platform_ids, NULL) != CL_SUCCESS)
+    {
+        goto end;
+    }
+    for (i = 0, k = 1; i < num_platforms; i++)
+    {
+        if (opencl->clGetDeviceIDs(platform_ids[i], CL_DEVICE_TYPE_ALL, 0, NULL, &num_devices) != CL_SUCCESS || !num_devices)
+        {
+            goto end;
+        }
+        if ((device_ids = malloc(sizeof(cl_device_id) * num_devices)) == NULL)
+        {
+            goto end;
+        }
+        if (opencl->clGetDeviceIDs(platform_ids[i], CL_DEVICE_TYPE_ALL, num_devices, device_ids, NULL) != CL_SUCCESS)
+        {
+            goto end;
+        }
+        for (j = 0; j < num_devices; j++)
+        {
+            if (device_ids[j] != NULL)
+            {
+                opencl->clGetDeviceInfo(device_ids[j], CL_DEVICE_VENDOR,  sizeof(vendor),
+                                        vendor,  NULL);
+                opencl->clGetDeviceInfo(device_ids[j], CL_DEVICE_NAME,    sizeof(name),
+                                        name,    NULL);
+                opencl->clGetDeviceInfo(device_ids[j], CL_DEVICE_VERSION, sizeof(version),
+                                        version, NULL);
+                opencl->clGetDeviceInfo(device_ids[j], CL_DRIVER_VERSION, sizeof(driver),
+                                        driver, NULL);
+                opencl->clGetDeviceInfo(device_ids[j], CL_DEVICE_TYPE,    sizeof(type),
+                                        &type,  NULL);
+
+                // don't list unsupported devices
+                if (type & CL_DEVICE_TYPE_CPU)
+                {
+                    continue;
+                }
+                hb_log("OpenCL device #%d: %s %s", k++, vendor, name);
+                hb_log(" - OpenCL version: %s", version + 7 /* strlen("OpenCL ") */);
+                hb_log(" - driver version: %s", driver);
+                hb_log(" - device type: %s%s",
+                       type & CL_DEVICE_TYPE_CPU         ? "CPU"         :
+                       type & CL_DEVICE_TYPE_GPU         ? "GPU"         :
+                       type & CL_DEVICE_TYPE_CUSTOM      ? "Custom"      :
+                       type & CL_DEVICE_TYPE_ACCELERATOR ? "Accelerator" : "Unknown",
+                       type & CL_DEVICE_TYPE_DEFAULT     ? " (default)"  : "");
+                hb_log(" - supported: %s",
+                       hb_opencl_device_is_supported(type,
+                                                     (const char*)vendor,
+                                                     (const char*)version) ? "yes" : "no");
+            }
+        }
+        free(device_ids);
+    }
+
+end:
+    hb_opencl_library_close(opencl);
+}
diff --git a/libhb/opencl.h b/libhb/opencl.h
new file mode 100644
index 000000000..93171ebd0
--- /dev/null
+++ b/libhb/opencl.h
@@ -0,0 +1,668 @@
+/* opencl.h
+
+   Copyright (c) 2003-2013 HandBrake Team
+   This file is part of the HandBrake source code
+   Homepage: <http://handbrake.fr/>.
+   It may be used under the terms of the GNU General Public License v2.
+   For full terms see the file COPYING file or visit http://www.gnu.org/licenses/gpl-2.0.html
+ */
+
+#ifndef HB_OPENCL_H
+#define HB_OPENCL_H
+
+#include "extras/cl.h"
+
+// we only support OpenCL 1.1 or later
+#define HB_OCL_MINVERSION_MAJOR 1
+#define HB_OCL_MINVERSION_MINOR 1
+
+#define HB_OCL_FUNC_TYPE(name)      hb_opencl_##name##_func
+#define HB_OCL_FUNC_DECL(name)      HB_OCL_FUNC_TYPE(name) name
+#define HB_OCL_API(ret, attr, name) typedef ret (attr* HB_OCL_FUNC_TYPE(name))
+
+#pragma mark -
+#pragma mark OpenCL API
+
+/* Platform API */
+HB_OCL_API(cl_int, CL_API_CALL, clGetPlatformIDs)
+(cl_uint          /* num_entries */,
+ cl_platform_id * /* platforms */,
+ cl_uint *        /* num_platforms */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clGetPlatformInfo)
+(cl_platform_id   /* platform */,
+ cl_platform_info /* param_name */,
+ size_t           /* param_value_size */,
+ void *           /* param_value */,
+ size_t *         /* param_value_size_ret */);
+
+/* Device APIs */
+HB_OCL_API(cl_int, CL_API_CALL, clGetDeviceIDs)
+(cl_platform_id   /* platform */,
+ cl_device_type   /* device_type */,
+ cl_uint          /* num_entries */,
+ cl_device_id *   /* devices */,
+ cl_uint *        /* num_devices */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clGetDeviceInfo)
+(cl_device_id    /* device */,
+ cl_device_info  /* param_name */,
+ size_t          /* param_value_size */,
+ void *          /* param_value */,
+ size_t *        /* param_value_size_ret */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clCreateSubDevices)
+(cl_device_id                         /* in_device */,
+ const cl_device_partition_property * /* properties */,
+ cl_uint                              /* num_devices */,
+ cl_device_id *                       /* out_devices */,
+ cl_uint *                            /* num_devices_ret */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clRetainDevice)
+(cl_device_id /* device */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clReleaseDevice)
+(cl_device_id /* device */);
+
+/* Context APIs  */
+HB_OCL_API(cl_context, CL_API_CALL, clCreateContext)
+(const cl_context_properties * /* properties */,
+ cl_uint                 /* num_devices */,
+ const cl_device_id *    /* devices */,
+ void (CL_CALLBACK * /* pfn_notify */)(const char *, const void *, size_t, void *),
+ void *                  /* user_data */,
+ cl_int *                /* errcode_ret */);
+
+HB_OCL_API(cl_context, CL_API_CALL, clCreateContextFromType)
+(const cl_context_properties * /* properties */,
+ cl_device_type          /* device_type */,
+ void (CL_CALLBACK *     /* pfn_notify*/ )(const char *, const void *, size_t, void *),
+ void *                  /* user_data */,
+ cl_int *                /* errcode_ret */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clRetainContext)
+(cl_context /* context */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clReleaseContext)
+(cl_context /* context */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clGetContextInfo)
+(cl_context         /* context */,
+ cl_context_info    /* param_name */,
+ size_t             /* param_value_size */,
+ void *             /* param_value */,
+ size_t *           /* param_value_size_ret */);
+
+/* Command Queue APIs */
+HB_OCL_API(cl_command_queue, CL_API_CALL, clCreateCommandQueue)
+(cl_context                     /* context */,
+ cl_device_id                   /* device */,
+ cl_command_queue_properties    /* properties */,
+ cl_int *                       /* errcode_ret */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clRetainCommandQueue)
+(cl_command_queue /* command_queue */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clReleaseCommandQueue)
+(cl_command_queue /* command_queue */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clGetCommandQueueInfo)
+(cl_command_queue      /* command_queue */,
+ cl_command_queue_info /* param_name */,
+ size_t                /* param_value_size */,
+ void *                /* param_value */,
+ size_t *              /* param_value_size_ret */);
+
+/* Memory Object APIs */
+HB_OCL_API(cl_mem, CL_API_CALL, clCreateBuffer)
+(cl_context   /* context */,
+ cl_mem_flags /* flags */,
+ size_t       /* size */,
+ void *       /* host_ptr */,
+ cl_int *     /* errcode_ret */);
+
+HB_OCL_API(cl_mem, CL_API_CALL, clCreateSubBuffer)
+(cl_mem                   /* buffer */,
+ cl_mem_flags             /* flags */,
+ cl_buffer_create_type    /* buffer_create_type */,
+ const void *             /* buffer_create_info */,
+ cl_int *                 /* errcode_ret */);
+
+HB_OCL_API(cl_mem, CL_API_CALL, clCreateImage)
+(cl_context              /* context */,
+ cl_mem_flags            /* flags */,
+ const cl_image_format * /* image_format */,
+ const cl_image_desc *   /* image_desc */,
+ void *                  /* host_ptr */,
+ cl_int *                /* errcode_ret */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clRetainMemObject)
+(cl_mem /* memobj */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clReleaseMemObject)
+(cl_mem /* memobj */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clGetSupportedImageFormats)
+(cl_context           /* context */,
+ cl_mem_flags         /* flags */,
+ cl_mem_object_type   /* image_type */,
+ cl_uint              /* num_entries */,
+ cl_image_format *    /* image_formats */,
+ cl_uint *            /* num_image_formats */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clGetMemObjectInfo)
+(cl_mem           /* memobj */,
+ cl_mem_info      /* param_name */,
+ size_t           /* param_value_size */,
+ void *           /* param_value */,
+ size_t *         /* param_value_size_ret */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clGetImageInfo)
+(cl_mem           /* image */,
+ cl_image_info    /* param_name */,
+ size_t           /* param_value_size */,
+ void *           /* param_value */,
+ size_t *         /* param_value_size_ret */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clSetMemObjectDestructorCallback)
+(cl_mem /* memobj */,
+ void (CL_CALLBACK * /*pfn_notify*/)( cl_mem /* memobj */, void* /*user_data*/),
+ void * /*user_data */ );
+
+/* Sampler APIs */
+HB_OCL_API(cl_sampler, CL_API_CALL, clCreateSampler)
+(cl_context          /* context */,
+ cl_bool             /* normalized_coords */,
+ cl_addressing_mode  /* addressing_mode */,
+ cl_filter_mode      /* filter_mode */,
+ cl_int *            /* errcode_ret */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clRetainSampler)
+(cl_sampler /* sampler */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clReleaseSampler)
+(cl_sampler /* sampler */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clGetSamplerInfo)
+(cl_sampler         /* sampler */,
+ cl_sampler_info    /* param_name */,
+ size_t             /* param_value_size */,
+ void *             /* param_value */,
+ size_t *           /* param_value_size_ret */);
+
+/* Program Object APIs  */
+HB_OCL_API(cl_program, CL_API_CALL, clCreateProgramWithSource)
+(cl_context        /* context */,
+ cl_uint           /* count */,
+ const char **     /* strings */,
+ const size_t *    /* lengths */,
+ cl_int *          /* errcode_ret */);
+
+HB_OCL_API(cl_program, CL_API_CALL, clCreateProgramWithBinary)
+(cl_context                     /* context */,
+ cl_uint                        /* num_devices */,
+ const cl_device_id *           /* device_list */,
+ const size_t *                 /* lengths */,
+ const unsigned char **         /* binaries */,
+ cl_int *                       /* binary_status */,
+ cl_int *                       /* errcode_ret */);
+
+HB_OCL_API(cl_program, CL_API_CALL, clCreateProgramWithBuiltInKernels)
+(cl_context            /* context */,
+ cl_uint               /* num_devices */,
+ const cl_device_id *  /* device_list */,
+ const char *          /* kernel_names */,
+ cl_int *              /* errcode_ret */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clRetainProgram)
+(cl_program /* program */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clReleaseProgram)
+(cl_program /* program */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clBuildProgram)
+(cl_program           /* program */,
+ cl_uint              /* num_devices */,
+ const cl_device_id * /* device_list */,
+ const char *         /* options */,
+ void (CL_CALLBACK *  /* pfn_notify */)(cl_program /* program */, void * /* user_data */),
+ void *               /* user_data */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clCompileProgram)
+(cl_program           /* program */,
+ cl_uint              /* num_devices */,
+ const cl_device_id * /* device_list */,
+ const char *         /* options */,
+ cl_uint              /* num_input_headers */,
+ const cl_program *   /* input_headers */,
+ const char **        /* header_include_names */,
+ void (CL_CALLBACK *  /* pfn_notify */)(cl_program /* program */, void * /* user_data */),
+ void *               /* user_data */);
+
+HB_OCL_API(cl_program, CL_API_CALL, clLinkProgram)
+(cl_context           /* context */,
+ cl_uint              /* num_devices */,
+ const cl_device_id * /* device_list */,
+ const char *         /* options */,
+ cl_uint              /* num_input_programs */,
+ const cl_program *   /* input_programs */,
+ void (CL_CALLBACK *  /* pfn_notify */)(cl_program /* program */, void * /* user_data */),
+ void *               /* user_data */,
+ cl_int *             /* errcode_ret */ );
+
+
+HB_OCL_API(cl_int, CL_API_CALL, clUnloadPlatformCompiler)
+(cl_platform_id /* platform */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clGetProgramInfo)
+(cl_program         /* program */,
+ cl_program_info    /* param_name */,
+ size_t             /* param_value_size */,
+ void *             /* param_value */,
+ size_t *           /* param_value_size_ret */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clGetProgramBuildInfo)
+(cl_program            /* program */,
+ cl_device_id          /* device */,
+ cl_program_build_info /* param_name */,
+ size_t                /* param_value_size */,
+ void *                /* param_value */,
+ size_t *              /* param_value_size_ret */);
+
+/* Kernel Object APIs */
+HB_OCL_API(cl_kernel, CL_API_CALL, clCreateKernel)
+(cl_program      /* program */,
+ const char *    /* kernel_name */,
+ cl_int *        /* errcode_ret */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clCreateKernelsInProgram)
+(cl_program     /* program */,
+ cl_uint        /* num_kernels */,
+ cl_kernel *    /* kernels */,
+ cl_uint *      /* num_kernels_ret */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clRetainKernel)
+(cl_kernel    /* kernel */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clReleaseKernel)
+(cl_kernel   /* kernel */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clSetKernelArg)
+(cl_kernel    /* kernel */,
+ cl_uint      /* arg_index */,
+ size_t       /* arg_size */,
+ const void * /* arg_value */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clGetKernelInfo)
+(cl_kernel       /* kernel */,
+ cl_kernel_info  /* param_name */,
+ size_t          /* param_value_size */,
+ void *          /* param_value */,
+ size_t *        /* param_value_size_ret */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clGetKernelArgInfo)
+(cl_kernel       /* kernel */,
+ cl_uint         /* arg_indx */,
+ cl_kernel_arg_info  /* param_name */,
+ size_t          /* param_value_size */,
+ void *          /* param_value */,
+ size_t *        /* param_value_size_ret */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clGetKernelWorkGroupInfo)
+(cl_kernel                  /* kernel */,
+ cl_device_id               /* device */,
+ cl_kernel_work_group_info  /* param_name */,
+ size_t                     /* param_value_size */,
+ void *                     /* param_value */,
+ size_t *                   /* param_value_size_ret */);
+
+/* Event Object APIs */
+HB_OCL_API(cl_int, CL_API_CALL, clWaitForEvents)
+(cl_uint             /* num_events */,
+ const cl_event *    /* event_list */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clGetEventInfo)
+(cl_event         /* event */,
+ cl_event_info    /* param_name */,
+ size_t           /* param_value_size */,
+ void *           /* param_value */,
+ size_t *         /* param_value_size_ret */);
+
+HB_OCL_API(cl_event, CL_API_CALL, clCreateUserEvent)
+(cl_context    /* context */,
+ cl_int *      /* errcode_ret */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clRetainEvent)
+(cl_event /* event */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clReleaseEvent)
+(cl_event /* event */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clSetUserEventStatus)
+(cl_event   /* event */,
+ cl_int     /* execution_status */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clSetEventCallback)
+(cl_event    /* event */,
+ cl_int      /* command_exec_callback_type */,
+ void (CL_CALLBACK * /* pfn_notify */)(cl_event, cl_int, void *),
+ void *      /* user_data */);
+
+/* Profiling APIs */
+HB_OCL_API(cl_int, CL_API_CALL, clGetEventProfilingInfo)
+(cl_event            /* event */,
+ cl_profiling_info   /* param_name */,
+ size_t              /* param_value_size */,
+ void *              /* param_value */,
+ size_t *            /* param_value_size_ret */);
+
+/* Flush and Finish APIs */
+HB_OCL_API(cl_int, CL_API_CALL, clFlush)
+(cl_command_queue /* command_queue */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clFinish)
+(cl_command_queue /* command_queue */);
+
+/* Enqueued Commands APIs */
+HB_OCL_API(cl_int, CL_API_CALL, clEnqueueReadBuffer)
+(cl_command_queue    /* command_queue */,
+ cl_mem              /* buffer */,
+ cl_bool             /* blocking_read */,
+ size_t              /* offset */,
+ size_t              /* size */,
+ void *              /* ptr */,
+ cl_uint             /* num_events_in_wait_list */,
+ const cl_event *    /* event_wait_list */,
+ cl_event *          /* event */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clEnqueueReadBufferRect)
+(cl_command_queue    /* command_queue */,
+ cl_mem              /* buffer */,
+ cl_bool             /* blocking_read */,
+ const size_t *      /* buffer_offset */,
+ const size_t *      /* host_offset */,
+ const size_t *      /* region */,
+ size_t              /* buffer_row_pitch */,
+ size_t              /* buffer_slice_pitch */,
+ size_t              /* host_row_pitch */,
+ size_t              /* host_slice_pitch */,
+ void *              /* ptr */,
+ cl_uint             /* num_events_in_wait_list */,
+ const cl_event *    /* event_wait_list */,
+ cl_event *          /* event */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clEnqueueWriteBuffer)
+(cl_command_queue   /* command_queue */,
+ cl_mem             /* buffer */,
+ cl_bool            /* blocking_write */,
+ size_t             /* offset */,
+ size_t             /* size */,
+ const void *       /* ptr */,
+ cl_uint            /* num_events_in_wait_list */,
+ const cl_event *   /* event_wait_list */,
+ cl_event *         /* event */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clEnqueueWriteBufferRect)
+(cl_command_queue    /* command_queue */,
+ cl_mem              /* buffer */,
+ cl_bool             /* blocking_write */,
+ const size_t *      /* buffer_offset */,
+ const size_t *      /* host_offset */,
+ const size_t *      /* region */,
+ size_t              /* buffer_row_pitch */,
+ size_t              /* buffer_slice_pitch */,
+ size_t              /* host_row_pitch */,
+ size_t              /* host_slice_pitch */,
+ const void *        /* ptr */,
+ cl_uint             /* num_events_in_wait_list */,
+ const cl_event *    /* event_wait_list */,
+ cl_event *          /* event */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clEnqueueFillBuffer)
+(cl_command_queue   /* command_queue */,
+ cl_mem             /* buffer */,
+ const void *       /* pattern */,
+ size_t             /* pattern_size */,
+ size_t             /* offset */,
+ size_t             /* size */,
+ cl_uint            /* num_events_in_wait_list */,
+ const cl_event *   /* event_wait_list */,
+ cl_event *         /* event */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clEnqueueCopyBuffer)
+(cl_command_queue    /* command_queue */,
+ cl_mem              /* src_buffer */,
+ cl_mem              /* dst_buffer */,
+ size_t              /* src_offset */,
+ size_t              /* dst_offset */,
+ size_t              /* size */,
+ cl_uint             /* num_events_in_wait_list */,
+ const cl_event *    /* event_wait_list */,
+ cl_event *          /* event */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clEnqueueCopyBufferRect)
+(cl_command_queue    /* command_queue */,
+ cl_mem              /* src_buffer */,
+ cl_mem              /* dst_buffer */,
+ const size_t *      /* src_origin */,
+ const size_t *      /* dst_origin */,
+ const size_t *      /* region */,
+ size_t              /* src_row_pitch */,
+ size_t              /* src_slice_pitch */,
+ size_t              /* dst_row_pitch */,
+ size_t              /* dst_slice_pitch */,
+ cl_uint             /* num_events_in_wait_list */,
+ const cl_event *    /* event_wait_list */,
+ cl_event *          /* event */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clEnqueueReadImage)
+(cl_command_queue     /* command_queue */,
+ cl_mem               /* image */,
+ cl_bool              /* blocking_read */,
+ const size_t *       /* origin[3] */,
+ const size_t *       /* region[3] */,
+ size_t               /* row_pitch */,
+ size_t               /* slice_pitch */,
+ void *               /* ptr */,
+ cl_uint              /* num_events_in_wait_list */,
+ const cl_event *     /* event_wait_list */,
+ cl_event *           /* event */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clEnqueueWriteImage)
+(cl_command_queue    /* command_queue */,
+ cl_mem              /* image */,
+ cl_bool             /* blocking_write */,
+ const size_t *      /* origin[3] */,
+ const size_t *      /* region[3] */,
+ size_t              /* input_row_pitch */,
+ size_t              /* input_slice_pitch */,
+ const void *        /* ptr */,
+ cl_uint             /* num_events_in_wait_list */,
+ const cl_event *    /* event_wait_list */,
+ cl_event *          /* event */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clEnqueueFillImage)
+(cl_command_queue   /* command_queue */,
+ cl_mem             /* image */,
+ const void *       /* fill_color */,
+ const size_t *     /* origin[3] */,
+ const size_t *     /* region[3] */,
+ cl_uint            /* num_events_in_wait_list */,
+ const cl_event *   /* event_wait_list */,
+ cl_event *         /* event */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clEnqueueCopyImage)
+(cl_command_queue     /* command_queue */,
+ cl_mem               /* src_image */,
+ cl_mem               /* dst_image */,
+ const size_t *       /* src_origin[3] */,
+ const size_t *       /* dst_origin[3] */,
+ const size_t *       /* region[3] */,
+ cl_uint              /* num_events_in_wait_list */,
+ const cl_event *     /* event_wait_list */,
+ cl_event *           /* event */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clEnqueueCopyImageToBuffer)
+(cl_command_queue /* command_queue */,
+ cl_mem           /* src_image */,
+ cl_mem           /* dst_buffer */,
+ const size_t *   /* src_origin[3] */,
+ const size_t *   /* region[3] */,
+ size_t           /* dst_offset */,
+ cl_uint          /* num_events_in_wait_list */,
+ const cl_event * /* event_wait_list */,
+ cl_event *       /* event */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clEnqueueCopyBufferToImage)
+(cl_command_queue /* command_queue */,
+ cl_mem           /* src_buffer */,
+ cl_mem           /* dst_image */,
+ size_t           /* src_offset */,
+ const size_t *   /* dst_origin[3] */,
+ const size_t *   /* region[3] */,
+ cl_uint          /* num_events_in_wait_list */,
+ const cl_event * /* event_wait_list */,
+ cl_event *       /* event */);
+
+HB_OCL_API(void *, CL_API_CALL, clEnqueueMapBuffer)
+(cl_command_queue /* command_queue */,
+ cl_mem           /* buffer */,
+ cl_bool          /* blocking_map */,
+ cl_map_flags     /* map_flags */,
+ size_t           /* offset */,
+ size_t           /* size */,
+ cl_uint          /* num_events_in_wait_list */,
+ const cl_event * /* event_wait_list */,
+ cl_event *       /* event */,
+ cl_int *         /* errcode_ret */);
+
+HB_OCL_API(void *, CL_API_CALL, clEnqueueMapImage)
+(cl_command_queue  /* command_queue */,
+ cl_mem            /* image */,
+ cl_bool           /* blocking_map */,
+ cl_map_flags      /* map_flags */,
+ const size_t *    /* origin[3] */,
+ const size_t *    /* region[3] */,
+ size_t *          /* image_row_pitch */,
+ size_t *          /* image_slice_pitch */,
+ cl_uint           /* num_events_in_wait_list */,
+ const cl_event *  /* event_wait_list */,
+ cl_event *        /* event */,
+ cl_int *          /* errcode_ret */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clEnqueueUnmapMemObject)
+(cl_command_queue  /* command_queue */,
+ cl_mem            /* memobj */,
+ void *            /* mapped_ptr */,
+ cl_uint           /* num_events_in_wait_list */,
+ const cl_event *  /* event_wait_list */,
+ cl_event *        /* event */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clEnqueueMigrateMemObjects)
+(cl_command_queue       /* command_queue */,
+ cl_uint                /* num_mem_objects */,
+ const cl_mem *         /* mem_objects */,
+ cl_mem_migration_flags /* flags */,
+ cl_uint                /* num_events_in_wait_list */,
+ const cl_event *       /* event_wait_list */,
+ cl_event *             /* event */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clEnqueueNDRangeKernel)
+(cl_command_queue /* command_queue */,
+ cl_kernel        /* kernel */,
+ cl_uint          /* work_dim */,
+ const size_t *   /* global_work_offset */,
+ const size_t *   /* global_work_size */,
+ const size_t *   /* local_work_size */,
+ cl_uint          /* num_events_in_wait_list */,
+ const cl_event * /* event_wait_list */,
+ cl_event *       /* event */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clEnqueueTask)
+(cl_command_queue  /* command_queue */,
+ cl_kernel         /* kernel */,
+ cl_uint           /* num_events_in_wait_list */,
+ const cl_event *  /* event_wait_list */,
+ cl_event *        /* event */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clEnqueueNativeKernel)
+(cl_command_queue  /* command_queue */,
+ void (CL_CALLBACK * /*user_func*/)(void *),
+ void *            /* args */,
+ size_t            /* cb_args */,
+ cl_uint           /* num_mem_objects */,
+ const cl_mem *    /* mem_list */,
+ const void **     /* args_mem_loc */,
+ cl_uint           /* num_events_in_wait_list */,
+ const cl_event *  /* event_wait_list */,
+ cl_event *        /* event */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clEnqueueMarkerWithWaitList)
+(cl_command_queue  /* command_queue */,
+ cl_uint           /* num_events_in_wait_list */,
+ const cl_event *  /* event_wait_list */,
+ cl_event *        /* event */);
+
+HB_OCL_API(cl_int, CL_API_CALL, clEnqueueBarrierWithWaitList)
+(cl_command_queue  /* command_queue */,
+ cl_uint           /* num_events_in_wait_list */,
+ const cl_event *  /* event_wait_list */,
+ cl_event *        /* event */);
+
+
+/* Extension function access
+ *
+ * Returns the extension function address for the given function name,
+ * or NULL if a valid function can not be found.  The client must
+ * check to make sure the address is not NULL, before using or
+ * calling the returned function address.
+ */
+HB_OCL_API(void *, CL_API_CALL, clGetExtensionFunctionAddressForPlatform)
+(cl_platform_id /* platform */,
+ const char *   /* func_name */);
+
+#pragma mark -
+
+typedef struct hb_opencl_library_s
+{
+    void *library;
+
+    /* Pointers to select OpenCL API functions */
+    HB_OCL_FUNC_DECL(clBuildProgram);
+    HB_OCL_FUNC_DECL(clCreateBuffer);
+    HB_OCL_FUNC_DECL(clCreateCommandQueue);
+    HB_OCL_FUNC_DECL(clCreateContextFromType);
+    HB_OCL_FUNC_DECL(clCreateKernel);
+    HB_OCL_FUNC_DECL(clCreateProgramWithBinary);
+    HB_OCL_FUNC_DECL(clCreateProgramWithSource);
+    HB_OCL_FUNC_DECL(clEnqueueCopyBuffer);
+    HB_OCL_FUNC_DECL(clEnqueueMapBuffer);
+    HB_OCL_FUNC_DECL(clEnqueueNDRangeKernel);
+    HB_OCL_FUNC_DECL(clEnqueueReadBuffer);
+    HB_OCL_FUNC_DECL(clEnqueueUnmapMemObject);
+    HB_OCL_FUNC_DECL(clEnqueueWriteBuffer);
+    HB_OCL_FUNC_DECL(clFlush);
+    HB_OCL_FUNC_DECL(clGetCommandQueueInfo);
+    HB_OCL_FUNC_DECL(clGetContextInfo);
+    HB_OCL_FUNC_DECL(clGetDeviceIDs);
+    HB_OCL_FUNC_DECL(clGetDeviceInfo);
+    HB_OCL_FUNC_DECL(clGetPlatformIDs);
+    HB_OCL_FUNC_DECL(clGetPlatformInfo);
+    HB_OCL_FUNC_DECL(clGetProgramBuildInfo);
+    HB_OCL_FUNC_DECL(clGetProgramInfo);
+    HB_OCL_FUNC_DECL(clReleaseCommandQueue);
+    HB_OCL_FUNC_DECL(clReleaseContext);
+    HB_OCL_FUNC_DECL(clReleaseEvent);
+    HB_OCL_FUNC_DECL(clReleaseKernel);
+    HB_OCL_FUNC_DECL(clReleaseProgram);
+    HB_OCL_FUNC_DECL(clSetKernelArg);
+    HB_OCL_FUNC_DECL(clWaitForEvents);
+} hb_opencl_library_t;
+
+int  hb_opencl_library_open (hb_opencl_library_t *opencl);
+void hb_opencl_library_close(hb_opencl_library_t *opencl);
+
+int  hb_opencl_available();
+void hb_opencl_info_print();
+
+#endif//HB_OPENCL_H
diff --git a/libhb/openclwrapper.c b/libhb/openclwrapper.c
index 327eee372..b5faf7041 100644
--- a/libhb/openclwrapper.c
+++ b/libhb/openclwrapper.c
@@ -1144,71 +1144,6 @@ void hb_opencl_init()
     hb_get_opencl_env();
 }
 
-void hb_opencl_info_print()
-{
-    cl_uint i, numDevices;
-    cl_device_id *devices;
-    
-    if (hb_init_opencl_env(&gpu_env))
-    {
-        return;
-    }
-
-    if (clGetContextInfo(gpu_env.context, CL_CONTEXT_NUM_DEVICES,
-                         sizeof(numDevices), &numDevices, NULL) != CL_SUCCESS)
-    {
-        return;
-    }
-    
-    if ((devices = malloc(sizeof(cl_device_id) * numDevices)) == NULL)
-    {
-        return;
-    }
-    
-    if (clGetContextInfo(gpu_env.context, CL_CONTEXT_DEVICES,
-                         sizeof(cl_device_id) * numDevices, devices, NULL) != CL_SUCCESS)
-    {
-        return;
-    }
-
-    for (i = 0; i < numDevices; i++)
-    {
-        if (devices[i] != NULL)
-        {
-            char vendor[100], name[1024], version[1024];
-            cl_device_type device_type;
-            char *device_type_name = "Unknown";
-
-            clGetDeviceInfo(devices[i], CL_DEVICE_VENDOR,  sizeof(vendor),
-                            vendor,  NULL);
-            clGetDeviceInfo(devices[i], CL_DEVICE_NAME,    sizeof(name),
-                            name,    NULL);
-            clGetDeviceInfo(devices[i], CL_DRIVER_VERSION, sizeof(version),
-                            version, NULL);
-            clGetDeviceInfo(devices[i], CL_DEVICE_TYPE, sizeof(device_type),
-                            &device_type, NULL);
-
-            if (device_type & CL_DEVICE_TYPE_GPU)
-                 device_type_name = "GPU";
-            else
-            if (device_type & CL_DEVICE_TYPE_CPU)
-                 device_type_name = "CPU";
-            else
-            if (device_type & CL_DEVICE_TYPE_ACCELERATOR)
-                 device_type_name = "Accelerator";
-            else
-            if (device_type & CL_DEVICE_TYPE_CUSTOM)
-                 device_type_name = "Custom";
-
-            hb_log("GPU #%d: %s %s", i + 1, vendor, name);
-            hb_log(" - driver version: %s", version);
-            hb_log(" - OpenCL device type: %s%s",device_type_name,device_type & CL_DEVICE_TYPE_DEFAULT ? "/Default" : "");
-        }
-    }
-
-    free(devices);
-}
-
 int hb_use_buffers()
 {
     return useBuffers;
diff --git a/libhb/openclwrapper.h b/libhb/openclwrapper.h
index 1ee081da9..c7606afc0 100644
--- a/libhb/openclwrapper.h
+++ b/libhb/openclwrapper.h
@@ -71,7 +71,6 @@ int hb_create_kernel( char * kernelname, KernelEnv * env );
 int hb_release_kernel( KernelEnv * env );
 
 void hb_opencl_init();
-void hb_opencl_info_print();
 
 int hb_get_opencl_env();
 
diff --git a/libhb/stream.c b/libhb/stream.c
index 8a2e08dab..ac34cc550 100644
--- a/libhb/stream.c
+++ b/libhb/stream.c
@@ -1122,7 +1122,7 @@ hb_title_t * hb_stream_title_scan(hb_stream_t *stream, hb_title_t * title)
     title->hwd_support = 0;
 #endif
 #ifdef USE_OPENCL
-    if ( hb_confirm_gpu_type() == 0 )
+    if (hb_confirm_gpu_type() == 0 && hb_opencl_available() == 1)
         title->opencl_support = 1;
     else
         title->opencl_support = 0;
@@ -5688,7 +5688,7 @@ static hb_title_t *ffmpeg_title_scan( hb_stream_t *stream, hb_title_t *title )
     title->hwd_support = 0;
 #endif
 #ifdef USE_OPENCL
-    if (hb_confirm_gpu_type() == 0)
+    if (hb_confirm_gpu_type() == 0 && hb_opencl_available() == 1)
         title->opencl_support = 1;
     else
         title->opencl_support = 0;