Merge pull request #367 from PCSX2/gabest11-master

GSdx-opencl

3rdparty/opencl/CL/cl_d3d10.h

@@ -0,0 +1,126 @@
+/**********************************************************************************
+ * Copyright (c) 2008-2013 The Khronos Group Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a
+ * copy of this software and/or associated documentation files (the
+ * "Materials"), to deal in the Materials without restriction, including
+ * without limitation the rights to use, copy, modify, merge, publish,
+ * distribute, sublicense, and/or sell copies of the Materials, and to
+ * permit persons to whom the Materials are furnished to do so, subject to
+ * the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included
+ * in all copies or substantial portions of the Materials.
+ *
+ * THE MATERIALS ARE PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+ * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+ * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
+ * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
+ * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
+ * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
+ * MATERIALS OR THE USE OR OTHER DEALINGS IN THE MATERIALS.
+ **********************************************************************************/
+
+/* $Revision: 11708 $ on $Date: 2010-06-13 23:36:24 -0700 (Sun, 13 Jun 2010) $ */
+
+#ifndef __OPENCL_CL_D3D10_H
+#define __OPENCL_CL_D3D10_H
+
+#include <d3d10.h>
+#include <CL/cl.h>
+#include <CL/cl_platform.h>
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+/******************************************************************************
+ * cl_khr_d3d10_sharing                                                       */
+#define cl_khr_d3d10_sharing 1
+
+typedef cl_uint cl_d3d10_device_source_khr;
+typedef cl_uint cl_d3d10_device_set_khr;
+
+/******************************************************************************/
+
+/* Error Codes */
+#define CL_INVALID_D3D10_DEVICE_KHR                  -1002
+#define CL_INVALID_D3D10_RESOURCE_KHR                -1003
+#define CL_D3D10_RESOURCE_ALREADY_ACQUIRED_KHR       -1004
+#define CL_D3D10_RESOURCE_NOT_ACQUIRED_KHR           -1005
+
+/* cl_d3d10_device_source_nv */
+#define CL_D3D10_DEVICE_KHR                          0x4010
+#define CL_D3D10_DXGI_ADAPTER_KHR                    0x4011
+
+/* cl_d3d10_device_set_nv */
+#define CL_PREFERRED_DEVICES_FOR_D3D10_KHR           0x4012
+#define CL_ALL_DEVICES_FOR_D3D10_KHR                 0x4013
+
+/* cl_context_info */
+#define CL_CONTEXT_D3D10_DEVICE_KHR                  0x4014
+#define CL_CONTEXT_D3D10_PREFER_SHARED_RESOURCES_KHR 0x402C
+
+/* cl_mem_info */
+#define CL_MEM_D3D10_RESOURCE_KHR                    0x4015
+
+/* cl_image_info */
+#define CL_IMAGE_D3D10_SUBRESOURCE_KHR               0x4016
+
+/* cl_command_type */
+#define CL_COMMAND_ACQUIRE_D3D10_OBJECTS_KHR         0x4017
+#define CL_COMMAND_RELEASE_D3D10_OBJECTS_KHR         0x4018
+
+/******************************************************************************/
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *clGetDeviceIDsFromD3D10KHR_fn)(
+    cl_platform_id             platform,
+    cl_d3d10_device_source_khr d3d_device_source,
+    void *                     d3d_object,
+    cl_d3d10_device_set_khr    d3d_device_set,
+    cl_uint                    num_entries,
+    cl_device_id *             devices,
+    cl_uint *                  num_devices) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_mem (CL_API_CALL *clCreateFromD3D10BufferKHR_fn)(
+    cl_context     context,
+    cl_mem_flags   flags,
+    ID3D10Buffer * resource,
+    cl_int *       errcode_ret) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_mem (CL_API_CALL *clCreateFromD3D10Texture2DKHR_fn)(
+    cl_context        context,
+    cl_mem_flags      flags,
+    ID3D10Texture2D * resource,
+    UINT              subresource,
+    cl_int *          errcode_ret) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_mem (CL_API_CALL *clCreateFromD3D10Texture3DKHR_fn)(
+    cl_context        context,
+    cl_mem_flags      flags,
+    ID3D10Texture3D * resource,
+    UINT              subresource,
+    cl_int *          errcode_ret) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *clEnqueueAcquireD3D10ObjectsKHR_fn)(
+    cl_command_queue command_queue,
+    cl_uint          num_objects,
+    const cl_mem *   mem_objects,
+    cl_uint          num_events_in_wait_list,
+    const cl_event * event_wait_list,
+    cl_event *       event) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *clEnqueueReleaseD3D10ObjectsKHR_fn)(
+    cl_command_queue command_queue,
+    cl_uint          num_objects,
+    const cl_mem *   mem_objects,
+    cl_uint          num_events_in_wait_list,
+    const cl_event * event_wait_list,
+    cl_event *       event) CL_API_SUFFIX__VERSION_1_0;
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif  /* __OPENCL_CL_D3D10_H */
+

3rdparty/opencl/CL/cl_d3d11.h

@@ -0,0 +1,132 @@
+/**********************************************************************************
+ * Copyright (c) 2008-2013 The Khronos Group Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a
+ * copy of this software and/or associated documentation files (the
+ * "Materials"), to deal in the Materials without restriction, including
+ * without limitation the rights to use, copy, modify, merge, publish,
+ * distribute, sublicense, and/or sell copies of the Materials, and to
+ * permit persons to whom the Materials are furnished to do so, subject to
+ * the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included
+ * in all copies or substantial portions of the Materials.
+ *
+ * THE MATERIALS ARE PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+ * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+ * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
+ * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
+ * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
+ * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
+ * MATERIALS OR THE USE OR OTHER DEALINGS IN THE MATERIALS.
+ **********************************************************************************/
+
+/* $Revision: 11708 $ on $Date: 2010-06-13 23:36:24 -0700 (Sun, 13 Jun 2010) $ */
+
+#ifndef __OPENCL_CL_D3D11_H
+#define __OPENCL_CL_D3D11_H
+
+#include <d3d11.h>
+#include <CL/cl.h>
+#include <CL/cl_platform.h>
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+/******************************************************************************
+ * cl_khr_d3d11_sharing                                                       */
+#define cl_khr_d3d11_sharing 1
+
+typedef cl_uint cl_d3d11_device_source_khr;
+typedef cl_uint cl_d3d11_device_set_khr;
+
+/******************************************************************************/
+
+/* Error Codes */
+#define CL_INVALID_D3D11_DEVICE_KHR                  -1006
+#define CL_INVALID_D3D11_RESOURCE_KHR                -1007
+#define CL_D3D11_RESOURCE_ALREADY_ACQUIRED_KHR       -1008
+#define CL_D3D11_RESOURCE_NOT_ACQUIRED_KHR           -1009
+
+/* cl_d3d11_device_source */
+#define CL_D3D11_DEVICE_KHR                          0x4019
+#define CL_D3D11_DXGI_ADAPTER_KHR                    0x401A
+
+/* cl_d3d11_device_set */
+#define CL_PREFERRED_DEVICES_FOR_D3D11_KHR           0x401B
+#define CL_ALL_DEVICES_FOR_D3D11_KHR                 0x401C
+
+/* cl_context_info */
+#define CL_CONTEXT_D3D11_DEVICE_KHR                  0x401D
+#define CL_CONTEXT_D3D11_PREFER_SHARED_RESOURCES_KHR 0x402D
+
+/* cl_mem_info */
+#define CL_MEM_D3D11_RESOURCE_KHR                    0x401E
+
+/* cl_image_info */
+#define CL_IMAGE_D3D11_SUBRESOURCE_KHR               0x401F
+
+/* cl_command_type */
+#define CL_COMMAND_ACQUIRE_D3D11_OBJECTS_KHR         0x4020
+#define CL_COMMAND_RELEASE_D3D11_OBJECTS_KHR         0x4021
+
+// object types
+
+#define CL_D3D11_OBJECT_BUFFER                0x3000
+#define CL_D3D11_OBJECT_TEXTURE2D             0x3001
+#define CL_D3D11_OBJECT_TEXTURE3D             0x3003
+
+/******************************************************************************/
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *clGetDeviceIDsFromD3D11KHR_fn)(
+    cl_platform_id             platform,
+    cl_d3d11_device_source_khr d3d_device_source,
+    void *                     d3d_object,
+    cl_d3d11_device_set_khr    d3d_device_set,
+    cl_uint                    num_entries,
+    cl_device_id *             devices,
+    cl_uint *                  num_devices) CL_API_SUFFIX__VERSION_1_2;
+
+typedef CL_API_ENTRY cl_mem (CL_API_CALL *clCreateFromD3D11BufferKHR_fn)(
+    cl_context     context,
+    cl_mem_flags   flags,
+    ID3D11Buffer * resource,
+    cl_int *       errcode_ret) CL_API_SUFFIX__VERSION_1_2;
+
+typedef CL_API_ENTRY cl_mem (CL_API_CALL *clCreateFromD3D11Texture2DKHR_fn)(
+    cl_context        context,
+    cl_mem_flags      flags,
+    ID3D11Texture2D * resource,
+    UINT              subresource,
+    cl_int *          errcode_ret) CL_API_SUFFIX__VERSION_1_2;
+
+typedef CL_API_ENTRY cl_mem (CL_API_CALL *clCreateFromD3D11Texture3DKHR_fn)(
+    cl_context        context,
+    cl_mem_flags      flags,
+    ID3D11Texture3D * resource,
+    UINT              subresource,
+    cl_int *          errcode_ret) CL_API_SUFFIX__VERSION_1_2;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *clEnqueueAcquireD3D11ObjectsKHR_fn)(
+    cl_command_queue command_queue,
+    cl_uint          num_objects,
+    const cl_mem *   mem_objects,
+    cl_uint          num_events_in_wait_list,
+    const cl_event * event_wait_list,
+    cl_event *       event) CL_API_SUFFIX__VERSION_1_2;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *clEnqueueReleaseD3D11ObjectsKHR_fn)(
+    cl_command_queue command_queue,
+    cl_uint          num_objects,
+    const cl_mem *   mem_objects,
+    cl_uint          num_events_in_wait_list,
+    const cl_event * event_wait_list,
+    cl_event *       event) CL_API_SUFFIX__VERSION_1_2;
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif  /* __OPENCL_CL_D3D11_H */
+

3rdparty/opencl/CL/cl_d3d9.h

@@ -0,0 +1,331 @@
+/**********************************************************************************
+ * Copyright (c) 2008-2013 The Khronos Group Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a
+ * copy of this software and/or associated documentation files (the
+ * "Materials"), to deal in the Materials without restriction, including
+ * without limitation the rights to use, copy, modify, merge, publish,
+ * distribute, sublicense, and/or sell copies of the Materials, and to
+ * permit persons to whom the Materials are furnished to do so, subject to
+ * the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included
+ * in all copies or substantial portions of the Materials.
+ *
+ * THE MATERIALS ARE PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+ * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+ * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
+ * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
+ * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
+ * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
+ * MATERIALS OR THE USE OR OTHER DEALINGS IN THE MATERIALS.
+ **********************************************************************************/
+
+/* $Revision: 11708 $ on $Date: 2010-06-13 23:36:24 -0700 (Sun, 13 Jun 2010) $ */
+
+#ifndef __OPENCL_CL_D3D9_H
+#define __OPENCL_CL_D3D9_H
+
+#include <CL/cl.h>
+#include <CL/cl_platform.h>
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+/******************************************************************************/
+/* cl_khr_dx9_media_sharing                                                   */
+#define cl_khr_dx9_media_sharing 1
+
+typedef cl_uint             cl_dx9_media_adapter_type_khr;
+typedef cl_uint             cl_dx9_media_adapter_set_khr;
+    
+#if defined(_WIN32)
+#include <d3d9.h>
+typedef struct _cl_dx9_surface_info_khr
+{
+    IDirect3DSurface9 *resource;
+    HANDLE shared_handle;
+} cl_dx9_surface_info_khr;
+#endif
+
+
+/******************************************************************************/
+
+/* Error Codes */
+#define CL_INVALID_DX9_MEDIA_ADAPTER_KHR                -1010
+#define CL_INVALID_DX9_MEDIA_SURFACE_KHR                -1011
+#define CL_DX9_MEDIA_SURFACE_ALREADY_ACQUIRED_KHR       -1012
+#define CL_DX9_MEDIA_SURFACE_NOT_ACQUIRED_KHR           -1013
+
+/* cl_media_adapter_type_khr */
+#define CL_ADAPTER_D3D9_KHR                              0x2020
+#define CL_ADAPTER_D3D9EX_KHR                            0x2021
+#define CL_ADAPTER_DXVA_KHR                              0x2022
+
+/* cl_media_adapter_set_khr */
+#define CL_PREFERRED_DEVICES_FOR_DX9_MEDIA_ADAPTER_KHR   0x2023
+#define CL_ALL_DEVICES_FOR_DX9_MEDIA_ADAPTER_KHR         0x2024
+
+/* cl_context_info */
+#define CL_CONTEXT_ADAPTER_D3D9_KHR                      0x2025
+#define CL_CONTEXT_ADAPTER_D3D9EX_KHR                    0x2026
+#define CL_CONTEXT_ADAPTER_DXVA_KHR                      0x2027
+
+/* cl_mem_info */
+#define CL_MEM_DX9_MEDIA_ADAPTER_TYPE_KHR                0x2028
+#define CL_MEM_DX9_MEDIA_SURFACE_INFO_KHR                0x2029
+
+/* cl_image_info */
+#define CL_IMAGE_DX9_MEDIA_PLANE_KHR                     0x202A
+
+/* cl_command_type */
+#define CL_COMMAND_ACQUIRE_DX9_MEDIA_SURFACES_KHR        0x202B
+#define CL_COMMAND_RELEASE_DX9_MEDIA_SURFACES_KHR        0x202C
+
+/******************************************************************************/
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *clGetDeviceIDsFromDX9MediaAdapterKHR_fn)(
+    cl_platform_id                   platform,
+    cl_uint                          num_media_adapters,
+    cl_dx9_media_adapter_type_khr *  media_adapter_type,
+    void *                           media_adapters[],
+    cl_dx9_media_adapter_set_khr     media_adapter_set,
+    cl_uint                          num_entries,
+    cl_device_id *                   devices,
+    cl_uint *                        num_devices) CL_API_SUFFIX__VERSION_1_2;
+
+typedef CL_API_ENTRY cl_mem (CL_API_CALL *clCreateFromDX9MediaSurfaceKHR_fn)(
+    cl_context                    context,
+    cl_mem_flags                  flags,
+    cl_dx9_media_adapter_type_khr adapter_type,
+    void *                        surface_info,
+    cl_uint                       plane,                                                                          
+    cl_int *                      errcode_ret) CL_API_SUFFIX__VERSION_1_2;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *clEnqueueAcquireDX9MediaSurfacesKHR_fn)(
+    cl_command_queue command_queue,
+    cl_uint          num_objects,
+    const cl_mem *   mem_objects,
+    cl_uint          num_events_in_wait_list,
+    const cl_event * event_wait_list,
+    cl_event *       event) CL_API_SUFFIX__VERSION_1_2;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *clEnqueueReleaseDX9MediaSurfacesKHR_fn)(
+    cl_command_queue command_queue,
+    cl_uint          num_objects,
+    const cl_mem *   mem_objects,
+    cl_uint          num_events_in_wait_list,
+    const cl_event * event_wait_list,
+    cl_event *       event) CL_API_SUFFIX__VERSION_1_2;
+
+#if defined CL_DX9_MEDIA_SHARING_INTEL_EXT
+
+#ifndef _WIN32
+#include <d3d9.h>
+#endif
+#include <d3d9.h>
+#include <dxvahd.h>
+#include <wtypes.h>
+#include <d3d9types.h>
+ 
+/******************************************************************************/    
+/* cl_intel_dx9_media_sharing extension                                       */
+#define cl_intel_dx9_media_sharing 1
+
+typedef cl_uint cl_dx9_device_source_intel;
+typedef cl_uint cl_dx9_device_set_intel;
+
+/******************************************************************************/
+
+// Error Codes
+#define CL_INVALID_DX9_DEVICE_INTEL                   -1010
+#define CL_INVALID_DX9_RESOURCE_INTEL                 -1011
+#define CL_DX9_RESOURCE_ALREADY_ACQUIRED_INTEL        -1012
+#define CL_DX9_RESOURCE_NOT_ACQUIRED_INTEL            -1013
+
+// cl_dx9_device_source_intel
+#define CL_D3D9_DEVICE_INTEL                          0x4022
+#define CL_D3D9EX_DEVICE_INTEL                        0x4070
+#define CL_DXVA_DEVICE_INTEL                          0x4071
+
+// cl_dx9_device_set_intel
+#define CL_PREFERRED_DEVICES_FOR_DX9_INTEL            0x4024
+#define CL_ALL_DEVICES_FOR_DX9_INTEL                  0x4025
+
+// cl_context_info
+#define CL_CONTEXT_D3D9_DEVICE_INTEL                  0x4026
+#define CL_CONTEXT_D3D9EX_DEVICE_INTEL                0x4072
+#define CL_CONTEXT_DXVA_DEVICE_INTEL                  0x4073
+
+// cl_mem_info
+#define CL_MEM_DX9_RESOURCE_INTEL                     0x4027
+#define CL_MEM_DX9_SHARED_HANDLE_INTEL                0x4074
+
+// cl_image_info
+#define CL_IMAGE_DX9_PLANE_INTEL                      0x4075
+
+// cl_command_type
+#define CL_COMMAND_ACQUIRE_DX9_OBJECTS_INTEL          0x402A
+#define CL_COMMAND_RELEASE_DX9_OBJECTS_INTEL          0x402B
+
+//packed YUV channel order
+#define CL_YUYV_INTEL                                 0x4076
+#define CL_UYVY_INTEL                                 0x4077
+#define CL_YVYU_INTEL                                 0x4078
+#define CL_VYUY_INTEL                                 0x4079
+
+/******************************************************************************/
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL* clGetDeviceIDsFromDX9INTEL_fn)(
+    cl_platform_id              /*platform*/,
+    cl_dx9_device_source_intel  /*dx9_device_source*/,
+    void*                       /*dx9_object*/,
+    cl_dx9_device_set_intel     /*dx9_device_set*/,
+    cl_uint                     /*num_entries*/, 
+    cl_device_id*               /*devices*/, 
+    cl_uint*                    /*num_devices*/);
+    
+typedef CL_API_ENTRY cl_mem (CL_API_CALL *clCreateFromDX9MediaSurfaceINTEL_fn)(
+    cl_context                  /*context*/,
+    cl_mem_flags                /*flags*/,
+    IDirect3DSurface9 *         /*resource*/,
+    HANDLE                      /*sharedHandle*/,
+    UINT                        /*plane*/,
+    cl_int *                    /*errcode_ret*/);
+    
+typedef CL_API_ENTRY cl_int (CL_API_CALL *clEnqueueAcquireDX9ObjectsINTEL_fn)(
+    cl_command_queue            /*command_queue*/,
+    cl_uint                     /*num_objects*/,
+    const cl_mem *              /*mem_objects*/,
+    cl_uint                     /*num_events_in_wait_list*/,
+    const cl_event *            /*event_wait_list*/,
+    cl_event *                  /*event*/);
+    
+typedef CL_API_ENTRY cl_int (CL_API_CALL *clEnqueueReleaseDX9ObjectsINTEL_fn)(
+    cl_command_queue            /*command_queue*/,
+    cl_uint                     /*num_objects*/,
+    cl_mem *                    /*mem_objects*/,
+    cl_uint                     /*num_events_in_wait_list*/,
+    const cl_event *            /*event_wait_list*/,
+    cl_event *                  /*event*/);
+
+#endif // CL_DX9_MEDIA_SHARING_INTEL_EXT
+
+#if defined CL_DX9_MEDIA_SHARING_NV_EXT
+
+#ifndef _WIN32
+#include <d3d9.h>
+#endif
+    
+/******************************************************************************
+ * cl_nv_d3d9_sharing                                                         */
+
+typedef cl_uint cl_d3d9_device_source_nv;
+typedef cl_uint cl_d3d9_device_set_nv;
+
+/******************************************************************************/
+
+// Error Codes
+#define CL_INVALID_D3D9_DEVICE_NV              -1010
+#define CL_INVALID_D3D9_RESOURCE_NV            -1011
+#define CL_D3D9_RESOURCE_ALREADY_ACQUIRED_NV   -1012
+#define CL_D3D9_RESOURCE_NOT_ACQUIRED_NV       -1013
+
+// cl_d3d9_device_source_nv
+#define CL_D3D9_DEVICE_NV                      0x4022
+#define CL_D3D9_ADAPTER_NAME_NV                0x4023
+
+// cl_d3d9_device_set_nv
+#define CL_PREFERRED_DEVICES_FOR_D3D9_NV       0x4024
+#define CL_ALL_DEVICES_FOR_D3D9_NV             0x4025
+
+// cl_context_info
+#define CL_CONTEXT_D3D9_DEVICE_NV              0x4026
+
+// cl_mem_info
+#define CL_MEM_D3D9_RESOURCE_NV                0x4027
+
+// cl_image_info
+#define CL_IMAGE_D3D9_FACE_NV                  0x4028
+#define CL_IMAGE_D3D9_LEVEL_NV                 0x4029
+
+// cl_command_type
+#define CL_COMMAND_ACQUIRE_D3D9_OBJECTS_NV     0x402A
+#define CL_COMMAND_RELEASE_D3D9_OBJECTS_NV     0x402B
+
+/******************************************************************************/
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *clGetDeviceIDsFromD3D9NV_fn)(
+    cl_platform_id            platform,
+    cl_d3d9_device_source_nv  d3d_device_source,
+    void *                    d3d_object,
+    cl_d3d9_device_set_nv     d3d_device_set,
+    cl_uint                   num_entries, 
+    cl_device_id *            devices, 
+    cl_uint *                 num_devices) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_mem (CL_API_CALL *clCreateFromD3D9VertexBufferNV_fn)(
+    cl_context               context,
+    cl_mem_flags             flags,
+    IDirect3DVertexBuffer9 * resource,
+    cl_int *                 errcode_ret) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_mem (CL_API_CALL *clCreateFromD3D9IndexBufferNV_fn)(
+    cl_context              context,
+    cl_mem_flags            flags,
+    IDirect3DIndexBuffer9 * resource,
+    cl_int *                errcode_ret) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_mem (CL_API_CALL *clCreateFromD3D9SurfaceNV_fn)(
+    cl_context          context,
+    cl_mem_flags        flags,
+    IDirect3DSurface9 * resource,
+    cl_int *            errcode_ret) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_mem (CL_API_CALL *clCreateFromD3D9TextureNV_fn)(
+    cl_context         context,
+    cl_mem_flags       flags,
+    IDirect3DTexture9 *resource,
+    UINT               miplevel,
+    cl_int *           errcode_ret) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_mem (CL_API_CALL *clCreateFromD3D9CubeTextureNV_fn)(
+    cl_context              context,
+    cl_mem_flags            flags,
+    IDirect3DCubeTexture9 * resource,
+    D3DCUBEMAP_FACES        facetype,
+    UINT                    miplevel,
+    cl_int *                errcode_ret) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_mem (CL_API_CALL *clCreateFromD3D9VolumeTextureNV_fn)(
+    cl_context                context,
+    cl_mem_flags              flags,
+    IDirect3DVolumeTexture9 * resource,
+    UINT                      miplevel,
+    cl_int *                  errcode_ret) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *clEnqueueAcquireD3D9ObjectsNV_fn)(
+    cl_command_queue command_queue,
+    cl_uint num_objects,
+    const cl_mem *mem_objects,
+    cl_uint num_events_in_wait_list,
+    const cl_event *event_wait_list,
+    cl_event *event) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *clEnqueueReleaseD3D9ObjectsNV_fn)(
+    cl_command_queue command_queue,
+    cl_uint num_objects,
+    cl_mem *mem_objects,
+    cl_uint num_events_in_wait_list,
+    const cl_event *event_wait_list,
+    cl_event *event) CL_API_SUFFIX__VERSION_1_0;
+
+#endif // CL_DX9_MEDIA_SHARING_NV_EXT
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif  /* __OPENCL_CL_DX9_MEDIA_SHARING_H */
+

3rdparty/opencl/CL/cl_egl.h

@@ -0,0 +1,133 @@
+/*******************************************************************************
+ * Copyright (c) 2008-2010 The Khronos Group Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a
+ * copy of this software and/or associated documentation files (the
+ * "Materials"), to deal in the Materials without restriction, including
+ * without limitation the rights to use, copy, modify, merge, publish,
+ * distribute, sublicense, and/or sell copies of the Materials, and to
+ * permit persons to whom the Materials are furnished to do so, subject to
+ * the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included
+ * in all copies or substantial portions of the Materials.
+ *
+ * THE MATERIALS ARE PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+ * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+ * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
+ * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
+ * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
+ * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
+ * MATERIALS OR THE USE OR OTHER DEALINGS IN THE MATERIALS.
+ ******************************************************************************/
+
+#ifndef __OPENCL_CL_EGL_H
+#define __OPENCL_CL_EGL_H
+
+#ifdef __APPLE__
+
+#else
+#include <CL/cl.h>
+#include <EGL/egl.h>
+#include <EGL/eglext.h>
+#endif  
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+
+/* Command type for events created with clEnqueueAcquireEGLObjectsKHR */
+#define CL_COMMAND_EGL_FENCE_SYNC_OBJECT_KHR  0x202F
+#define CL_COMMAND_ACQUIRE_EGL_OBJECTS_KHR    0x202D
+#define CL_COMMAND_RELEASE_EGL_OBJECTS_KHR    0x202E
+
+/* Error type for clCreateFromEGLImageKHR */
+#define CL_INVALID_EGL_OBJECT_KHR             -1093
+#define CL_EGL_RESOURCE_NOT_ACQUIRED_KHR      -1092
+
+/* CLeglImageKHR is an opaque handle to an EGLImage */
+typedef void* CLeglImageKHR;
+
+/* CLeglDisplayKHR is an opaque handle to an EGLDisplay */
+typedef void* CLeglDisplayKHR;
+
+/* CLeglSyncKHR is an opaque handle to an EGLSync object */
+typedef void* CLeglSyncKHR;
+
+/* properties passed to clCreateFromEGLImageKHR */
+typedef intptr_t cl_egl_image_properties_khr;
+
+
+#define cl_khr_egl_image 1
+
+extern CL_API_ENTRY cl_mem CL_API_CALL
+clCreateFromEGLImageKHR(cl_context                  /* context */,
+                        CLeglDisplayKHR             /* egldisplay */,
+                        CLeglImageKHR               /* eglimage */,
+                        cl_mem_flags                /* flags */,
+                        const cl_egl_image_properties_khr * /* properties */,
+                        cl_int *                    /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_mem (CL_API_CALL *clCreateFromEGLImageKHR_fn)(
+	cl_context                  context,
+	CLeglDisplayKHR             egldisplay,
+	CLeglImageKHR               eglimage,
+	cl_mem_flags                flags,
+	const cl_egl_image_properties_khr * properties,
+	cl_int *                    errcode_ret);
+
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+clEnqueueAcquireEGLObjectsKHR(cl_command_queue /* command_queue */,
+                              cl_uint          /* num_objects */,
+                              const cl_mem *   /* mem_objects */,
+                              cl_uint          /* num_events_in_wait_list */,
+                              const cl_event * /* event_wait_list */,
+                              cl_event *       /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *clEnqueueAcquireEGLObjectsKHR_fn)(
+	cl_command_queue command_queue,
+	cl_uint          num_objects,
+	const cl_mem *   mem_objects,
+	cl_uint          num_events_in_wait_list,
+	const cl_event * event_wait_list,
+	cl_event *       event);
+
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+clEnqueueReleaseEGLObjectsKHR(cl_command_queue /* command_queue */,
+                              cl_uint          /* num_objects */,
+                              const cl_mem *   /* mem_objects */,
+                              cl_uint          /* num_events_in_wait_list */,
+                              const cl_event * /* event_wait_list */,
+                              cl_event *       /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *clEnqueueReleaseEGLObjectsKHR_fn)(
+	cl_command_queue command_queue,
+	cl_uint          num_objects,
+	const cl_mem *   mem_objects,
+	cl_uint          num_events_in_wait_list,
+	const cl_event * event_wait_list,
+	cl_event *       event);
+
+
+#define cl_khr_egl_event 1
+
+extern CL_API_ENTRY cl_event CL_API_CALL
+clCreateEventFromEGLSyncKHR(cl_context      /* context */,
+                            CLeglSyncKHR    /* sync */,
+                            CLeglDisplayKHR /* display */,
+                            cl_int *        /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_event (CL_API_CALL *clCreateEventFromEGLSyncKHR_fn)(
+	cl_context      context,
+	CLeglSyncKHR    sync,
+	CLeglDisplayKHR display,
+	cl_int *        errcode_ret);
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif /* __OPENCL_CL_EGL_H */

3rdparty/opencl/CL/cl_ext.h

@@ -0,0 +1,458 @@
+/*******************************************************************************
+ * Copyright (c) 2008-2013 The Khronos Group Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a
+ * copy of this software and/or associated documentation files (the
+ * "Materials"), to deal in the Materials without restriction, including
+ * without limitation the rights to use, copy, modify, merge, publish,
+ * distribute, sublicense, and/or sell copies of the Materials, and to
+ * permit persons to whom the Materials are furnished to do so, subject to
+ * the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included
+ * in all copies or substantial portions of the Materials.
+ *
+ * THE MATERIALS ARE PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+ * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+ * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
+ * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
+ * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
+ * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
+ * MATERIALS OR THE USE OR OTHER DEALINGS IN THE MATERIALS.
+ ******************************************************************************/
+
+/* $Revision: 11928 $ on $Date: 2010-07-13 09:04:56 -0700 (Tue, 13 Jul 2010) $ */
+
+/* cl_ext.h contains OpenCL extensions which don't have external */
+/* (OpenGL, D3D) dependencies.                                   */
+
+#ifndef __CL_EXT_H
+#define __CL_EXT_H
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+#ifdef __APPLE__
+        #include <OpenCL/cl.h>
+    #include <AvailabilityMacros.h>
+#else
+        #include <CL/cl.h>
+#endif
+
+/* cl_khr_fp16 extension - no extension #define since it has no functions  */
+#define CL_DEVICE_HALF_FP_CONFIG                    0x1033
+
+/* Memory object destruction
+ *
+ * Apple extension for use to manage externally allocated buffers used with cl_mem objects with CL_MEM_USE_HOST_PTR
+ *
+ * Registers a user callback function that will be called when the memory object is deleted and its resources 
+ * freed. Each call to clSetMemObjectCallbackFn registers the specified user callback function on a callback 
+ * stack associated with memobj. The registered user callback functions are called in the reverse order in 
+ * which they were registered. The user callback functions are called and then the memory object is deleted 
+ * and its resources freed. This provides a mechanism for the application (and libraries) using memobj to be 
+ * notified when the memory referenced by host_ptr, specified when the memory object is created and used as 
+ * the storage bits for the memory object, can be reused or freed.
+ *
+ * The application may not call CL api's with the cl_mem object passed to the pfn_notify.
+ *
+ * Please check for the "cl_APPLE_SetMemObjectDestructor" extension using clGetDeviceInfo(CL_DEVICE_EXTENSIONS)
+ * before using.
+ */
+#define cl_APPLE_SetMemObjectDestructor 1
+cl_int  CL_API_ENTRY clSetMemObjectDestructorAPPLE(  cl_mem /* memobj */, 
+                                        void (* /*pfn_notify*/)( cl_mem /* memobj */, void* /*user_data*/), 
+                                        void * /*user_data */ )             CL_EXT_SUFFIX__VERSION_1_0;  
+
+
+/* Context Logging Functions
+ *
+ * The next three convenience functions are intended to be used as the pfn_notify parameter to clCreateContext().
+ * Please check for the "cl_APPLE_ContextLoggingFunctions" extension using clGetDeviceInfo(CL_DEVICE_EXTENSIONS)
+ * before using.
+ *
+ * clLogMessagesToSystemLog fowards on all log messages to the Apple System Logger 
+ */
+#define cl_APPLE_ContextLoggingFunctions 1
+extern void CL_API_ENTRY clLogMessagesToSystemLogAPPLE(  const char * /* errstr */, 
+                                            const void * /* private_info */, 
+                                            size_t       /* cb */, 
+                                            void *       /* user_data */ )  CL_EXT_SUFFIX__VERSION_1_0;
+
+/* clLogMessagesToStdout sends all log messages to the file descriptor stdout */
+extern void CL_API_ENTRY clLogMessagesToStdoutAPPLE(   const char * /* errstr */, 
+                                          const void * /* private_info */, 
+                                          size_t       /* cb */, 
+                                          void *       /* user_data */ )    CL_EXT_SUFFIX__VERSION_1_0;
+
+/* clLogMessagesToStderr sends all log messages to the file descriptor stderr */
+extern void CL_API_ENTRY clLogMessagesToStderrAPPLE(   const char * /* errstr */, 
+                                          const void * /* private_info */, 
+                                          size_t       /* cb */, 
+                                          void *       /* user_data */ )    CL_EXT_SUFFIX__VERSION_1_0;
+
+
+/************************ 
+* cl_khr_icd extension *                                                  
+************************/
+#define cl_khr_icd 1
+
+/* cl_platform_info                                                        */
+#define CL_PLATFORM_ICD_SUFFIX_KHR                  0x0920
+
+/* Additional Error Codes                                                  */
+#define CL_PLATFORM_NOT_FOUND_KHR                   -1001
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+clIcdGetPlatformIDsKHR(cl_uint          /* num_entries */,
+                       cl_platform_id * /* platforms */,
+                       cl_uint *        /* num_platforms */);
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *clIcdGetPlatformIDsKHR_fn)(
+    cl_uint          /* num_entries */,
+    cl_platform_id * /* platforms */,
+    cl_uint *        /* num_platforms */);
+
+
+/* Extension: cl_khr_image2D_buffer
+ *
+ * This extension allows a 2D image to be created from a cl_mem buffer without a copy.
+ * The type associated with a 2D image created from a buffer in an OpenCL program is image2d_t.
+ * Both the sampler and sampler-less read_image built-in functions are supported for 2D images
+ * and 2D images created from a buffer.  Similarly, the write_image built-ins are also supported
+ * for 2D images created from a buffer.
+ *
+ * When the 2D image from buffer is created, the client must specify the width,
+ * height, image format (i.e. channel order and channel data type) and optionally the row pitch
+ *
+ * The pitch specified must be a multiple of CL_DEVICE_IMAGE_PITCH_ALIGNMENT pixels.
+ * The base address of the buffer must be aligned to CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT pixels.
+ */
+    
+/*************************************
+ * cl_khr_initalize_memory extension *
+ *************************************/
+    
+#define CL_CONTEXT_MEMORY_INITIALIZE_KHR            0x2030
+    
+    
+/**************************************
+ * cl_khr_terminate_context extension *
+ **************************************/
+    
+#define CL_DEVICE_TERMINATE_CAPABILITY_KHR          0x2031
+#define CL_CONTEXT_TERMINATE_KHR                    0x2032
+
+#define cl_khr_terminate_context 1
+extern CL_API_ENTRY cl_int CL_API_CALL clTerminateContextKHR(cl_context /* context */) CL_EXT_SUFFIX__VERSION_1_2;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *clTerminateContextKHR_fn)(cl_context /* context */) CL_EXT_SUFFIX__VERSION_1_2;
+    
+    
+/*
+ * Extension: cl_khr_spir
+ *
+ * This extension adds support to create an OpenCL program object from a 
+ * Standard Portable Intermediate Representation (SPIR) instance
+ */
+
+#define CL_DEVICE_SPIR_VERSIONS                     0x40E0
+#define CL_PROGRAM_BINARY_TYPE_INTERMEDIATE         0x40E1
+
+
+/******************************************
+* cl_nv_device_attribute_query extension *
+******************************************/
+/* cl_nv_device_attribute_query extension - no extension #define since it has no functions */
+#define CL_DEVICE_COMPUTE_CAPABILITY_MAJOR_NV       0x4000
+#define CL_DEVICE_COMPUTE_CAPABILITY_MINOR_NV       0x4001
+#define CL_DEVICE_REGISTERS_PER_BLOCK_NV            0x4002
+#define CL_DEVICE_WARP_SIZE_NV                      0x4003
+#define CL_DEVICE_GPU_OVERLAP_NV                    0x4004
+#define CL_DEVICE_KERNEL_EXEC_TIMEOUT_NV            0x4005
+#define CL_DEVICE_INTEGRATED_MEMORY_NV              0x4006
+
+/*********************************
+* cl_amd_device_attribute_query *
+*********************************/
+#define CL_DEVICE_PROFILING_TIMER_OFFSET_AMD        0x4036
+
+/*********************************
+* cl_arm_printf extension
+*********************************/
+#define CL_PRINTF_CALLBACK_ARM                      0x40B0
+#define CL_PRINTF_BUFFERSIZE_ARM                    0x40B1
+
+/*********************************
+* cl_intel_accelerator extension *
+*********************************/
+#define cl_intel_accelerator 1
+#define cl_intel_motion_estimation 1
+
+typedef struct _cl_accelerator_intel*     cl_accelerator_intel;
+typedef cl_uint                           cl_accelerator_type_intel;
+typedef cl_uint                           cl_accelerator_info_intel;
+
+typedef struct _cl_motion_estimation_desc_intel {
+    cl_uint mb_block_type;
+    cl_uint subpixel_mode;
+    cl_uint sad_adjust_mode;
+    cl_uint search_path_type;
+} cl_motion_estimation_desc_intel;
+
+/* Error Codes */
+#define CL_INVALID_ACCELERATOR_INTEL            -1094
+#define CL_INVALID_ACCELERATOR_TYPE_INTEL       -1095
+#define CL_INVALID_ACCELERATOR_DESCRIPTOR_INTEL -1096
+#define CL_ACCELERATOR_TYPE_NOT_SUPPORTED_INTEL -1097
+
+/* Deprecated Error Codes */
+#define CL_INVALID_ACCELERATOR_INTEL_DEPRECATED            -6000
+#define CL_INVALID_ACCELERATOR_TYPE_INTEL_DEPRECATED       -6001
+#define CL_INVALID_ACCELERATOR_DESCRIPTOR_INTEL_DEPRECATED -6002
+#define CL_ACCELERATOR_TYPE_NOT_SUPPORTED_INTEL_DEPRECATED -6003
+
+/* cl_accelerator_type_intel */
+#define CL_ACCELERATOR_TYPE_MOTION_ESTIMATION_INTEL     0x0
+
+/* cl_accelerator_info_intel */
+#define CL_ACCELERATOR_DESCRIPTOR_INTEL                 0x4090
+#define CL_ACCELERATOR_REFERENCE_COUNT_INTEL            0x4091
+#define CL_ACCELERATOR_CONTEXT_INTEL                    0x4092
+#define CL_ACCELERATOR_TYPE_INTEL                       0x4093
+
+/*cl_motion_detect_desc_intel flags */
+#define CL_ME_MB_TYPE_16x16_INTEL                       0x0
+#define CL_ME_MB_TYPE_8x8_INTEL                         0x1
+#define CL_ME_MB_TYPE_4x4_INTEL                         0x2
+
+#define CL_ME_SUBPIXEL_MODE_INTEGER_INTEL               0x0
+#define CL_ME_SUBPIXEL_MODE_HPEL_INTEL                  0x1
+#define CL_ME_SUBPIXEL_MODE_QPEL_INTEL                  0x2
+
+#define CL_ME_SAD_ADJUST_MODE_NONE_INTEL                0x0
+#define CL_ME_SAD_ADJUST_MODE_HAAR_INTEL                0x1
+
+#define CL_ME_SEARCH_PATH_RADIUS_2_2_INTEL              0x0
+#define CL_ME_SEARCH_PATH_RADIUS_4_4_INTEL              0x1
+#define CL_ME_SEARCH_PATH_RADIUS_16_12_INTEL            0x5
+
+extern CL_API_ENTRY cl_accelerator_intel CL_API_CALL
+clCreateAcceleratorINTEL(
+    cl_context                  /* context */,
+    cl_accelerator_type_intel   /* accelerator_type */,
+    size_t                      /* descriptor_size */,
+    const void*                 /* descriptor */,
+    cl_int*                     /* errcode_ret */ ) CL_EXT_SUFFIX__VERSION_1_2;
+
+typedef CL_API_ENTRY cl_accelerator_intel 
+    (CL_API_CALL *clCreateAcceleratorINTEL_fn)(
+    cl_context                  /* context */,
+    cl_accelerator_type_intel   /* accelerator_type */,
+    size_t                      /* descriptor_size */,
+    const void*                 /* descriptor */,
+    cl_int*                     /* errcode_ret */ ) CL_EXT_SUFFIX__VERSION_1_2;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+clGetAcceleratorInfoINTEL
+(
+    cl_accelerator_intel        /* accelerator */,
+    cl_accelerator_info_intel   /* param_name */,
+    size_t                      /* param_value_size */,
+    void*                       /* param_value */,
+    size_t*                     /* param_value_size_ret */ ) CL_EXT_SUFFIX__VERSION_1_2;
+
+typedef CL_API_ENTRY cl_int 
+    (CL_API_CALL *clGetAcceleratorInfoINTEL_fn)(
+    cl_accelerator_intel        /* accelerator */,
+    cl_accelerator_info_intel   /* param_name */,
+    size_t                      /* param_value_size */,
+    void*                       /* param_value */,
+    size_t*                     /* param_value_size_ret */ ) CL_EXT_SUFFIX__VERSION_1_2;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+clRetainAcceleratorINTEL(
+    cl_accelerator_intel        /* accelerator */ ) CL_EXT_SUFFIX__VERSION_1_2;
+
+typedef CL_API_ENTRY cl_int
+    (CL_API_CALL *clRetainAcceleratorINTEL_fn)(
+    cl_accelerator_intel        /* accelerator */ ) CL_EXT_SUFFIX__VERSION_1_2;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+clReleaseAcceleratorINTEL(
+    cl_accelerator_intel        /* accelerator */ ) CL_EXT_SUFFIX__VERSION_1_2;
+
+typedef CL_API_ENTRY cl_int
+    (CL_API_CALL *clReleaseAcceleratorINTEL_fn)(
+    cl_accelerator_intel        /* accelerator */ ) CL_EXT_SUFFIX__VERSION_1_2;
+
+#ifdef CL_VERSION_1_1
+   /***********************************
+    * cl_ext_device_fission extension *
+    ***********************************/
+    #define cl_ext_device_fission   1
+    
+    extern CL_API_ENTRY cl_int CL_API_CALL
+    clReleaseDeviceEXT( cl_device_id /*device*/ ) CL_EXT_SUFFIX__VERSION_1_1; 
+    
+    typedef CL_API_ENTRY cl_int 
+    (CL_API_CALL *clReleaseDeviceEXT_fn)( cl_device_id /*device*/ ) CL_EXT_SUFFIX__VERSION_1_1;
+
+    extern CL_API_ENTRY cl_int CL_API_CALL
+    clRetainDeviceEXT( cl_device_id /*device*/ ) CL_EXT_SUFFIX__VERSION_1_1; 
+    
+    typedef CL_API_ENTRY cl_int 
+    (CL_API_CALL *clRetainDeviceEXT_fn)( cl_device_id /*device*/ ) CL_EXT_SUFFIX__VERSION_1_1;
+
+    typedef cl_ulong  cl_device_partition_property_ext;
+    extern CL_API_ENTRY cl_int CL_API_CALL
+    clCreateSubDevicesEXT(  cl_device_id /*in_device*/,
+                            const cl_device_partition_property_ext * /* properties */,
+                            cl_uint /*num_entries*/,
+                            cl_device_id * /*out_devices*/,
+                            cl_uint * /*num_devices*/ ) CL_EXT_SUFFIX__VERSION_1_1;
+
+    typedef CL_API_ENTRY cl_int 
+    ( CL_API_CALL * clCreateSubDevicesEXT_fn)(  cl_device_id /*in_device*/,
+                                                const cl_device_partition_property_ext * /* properties */,
+                                                cl_uint /*num_entries*/,
+                                                cl_device_id * /*out_devices*/,
+                                                cl_uint * /*num_devices*/ ) CL_EXT_SUFFIX__VERSION_1_1;
+
+    /* cl_device_partition_property_ext */
+    #define CL_DEVICE_PARTITION_EQUALLY_EXT             0x4050
+    #define CL_DEVICE_PARTITION_BY_COUNTS_EXT           0x4051
+    #define CL_DEVICE_PARTITION_BY_NAMES_EXT            0x4052
+    #define CL_DEVICE_PARTITION_BY_NAMES_INTEL          0x4052
+    #define CL_DEVICE_PARTITION_BY_AFFINITY_DOMAIN_EXT  0x4053
+    
+    /* clDeviceGetInfo selectors */
+    #define CL_DEVICE_PARENT_DEVICE_EXT                 0x4054
+    #define CL_DEVICE_PARTITION_TYPES_EXT               0x4055
+    #define CL_DEVICE_AFFINITY_DOMAINS_EXT              0x4056
+    #define CL_DEVICE_REFERENCE_COUNT_EXT               0x4057
+    #define CL_DEVICE_PARTITION_STYLE_EXT               0x4058
+    
+    /* error codes */
+    #define CL_DEVICE_PARTITION_FAILED_EXT              -1057
+    #define CL_INVALID_PARTITION_COUNT_EXT              -1058
+    #define CL_INVALID_PARTITION_NAME_EXT               -1059
+    
+    /* CL_AFFINITY_DOMAINs */
+    #define CL_AFFINITY_DOMAIN_L1_CACHE_EXT             0x1
+    #define CL_AFFINITY_DOMAIN_L2_CACHE_EXT             0x2
+    #define CL_AFFINITY_DOMAIN_L3_CACHE_EXT             0x3
+    #define CL_AFFINITY_DOMAIN_L4_CACHE_EXT             0x4
+    #define CL_AFFINITY_DOMAIN_NUMA_EXT                 0x10
+    #define CL_AFFINITY_DOMAIN_NEXT_FISSIONABLE_EXT     0x100
+    
+    /* cl_device_partition_property_ext list terminators */
+    #define CL_PROPERTIES_LIST_END_EXT                  ((cl_device_partition_property_ext) 0)
+    #define CL_PARTITION_BY_COUNTS_LIST_END_EXT         ((cl_device_partition_property_ext) 0)
+    #define CL_PARTITION_BY_NAMES_LIST_END_EXT          ((cl_device_partition_property_ext) 0 - 1)
+    #define CL_PARTITION_BY_NAMES_LIST_END_INTEL        ((cl_device_partition_property_ext) 0 - 1)
+
+    #define CL_QUEUE_THREAD_LOCAL_EXEC_ENABLE_INTEL      (1 << 31)
+
+
+/*********************************
+* cl_qcom_ext_host_ptr extension
+*********************************/
+
+#define CL_MEM_EXT_HOST_PTR_QCOM                  (1 << 29)
+
+#define CL_DEVICE_EXT_MEM_PADDING_IN_BYTES_QCOM   0x40A0      
+#define CL_DEVICE_PAGE_SIZE_QCOM                  0x40A1
+#define CL_IMAGE_ROW_ALIGNMENT_QCOM               0x40A2
+#define CL_IMAGE_SLICE_ALIGNMENT_QCOM             0x40A3
+#define CL_MEM_HOST_UNCACHED_QCOM                 0x40A4
+#define CL_MEM_HOST_WRITEBACK_QCOM                0x40A5
+#define CL_MEM_HOST_WRITETHROUGH_QCOM             0x40A6
+#define CL_MEM_HOST_WRITE_COMBINING_QCOM          0x40A7
+
+typedef cl_uint                                   cl_image_pitch_info_qcom;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+clGetDeviceImageInfoQCOM(cl_device_id             device,
+                         size_t                   image_width,
+                         size_t                   image_height,
+                         const cl_image_format   *image_format,
+                         cl_image_pitch_info_qcom param_name,
+                         size_t                   param_value_size,
+                         void                    *param_value,
+                         size_t                  *param_value_size_ret);
+
+typedef struct _cl_mem_ext_host_ptr
+{
+    /* Type of external memory allocation. */
+    /* Legal values will be defined in layered extensions. */
+    cl_uint  allocation_type;
+            
+    /* Host cache policy for this external memory allocation. */
+    cl_uint  host_cache_policy;
+
+} cl_mem_ext_host_ptr;
+
+/*********************************
+* cl_qcom_ion_host_ptr extension
+*********************************/
+
+#define CL_MEM_ION_HOST_PTR_QCOM                  0x40A8
+
+typedef struct _cl_mem_ion_host_ptr
+{
+    /* Type of external memory allocation. */
+    /* Must be CL_MEM_ION_HOST_PTR_QCOM for ION allocations. */
+    cl_mem_ext_host_ptr  ext_host_ptr;
+
+    /* ION file descriptor */
+    int                  ion_filedesc;
+            
+    /* Host pointer to the ION allocated memory */
+    void*                ion_hostptr;
+
+} cl_mem_ion_host_ptr;
+
+#endif /* CL_VERSION_1_1 */
+
+
+#ifdef CL_VERSION_2_0
+/*********************************
+* cl_khr_sub_groups extension
+*********************************/
+#define cl_khr_sub_groups 1
+
+typedef cl_uint  cl_kernel_sub_group_info;
+
+/* cl_khr_sub_group_info */
+#define CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE_KHR	0x2033
+#define CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE_KHR		0x2034
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+clGetKernelSubGroupInfoKHR(cl_kernel /* in_kernel */,
+						   cl_device_id /*in_device*/,
+						   const cl_kernel_sub_group_info /* param_name */,
+						   size_t /*input_value_size*/,
+						   const void * /*input_value*/,
+						   size_t /*param_value_size*/,
+						   void* /*param_value*/,
+						   size_t* /*param_value_size_ret*/ ) CL_EXT_SUFFIX__VERSION_2_0;
+						   
+typedef CL_API_ENTRY cl_int
+     ( CL_API_CALL * clGetKernelSubGroupInfoKHR_fn)(cl_kernel /* in_kernel */,
+						      cl_device_id /*in_device*/,
+						      const cl_kernel_sub_group_info /* param_name */,
+						      size_t /*input_value_size*/,
+						      const void * /*input_value*/,
+						      size_t /*param_value_size*/,
+						      void* /*param_value*/,
+						      size_t* /*param_value_size_ret*/ ) CL_EXT_SUFFIX__VERSION_2_0;
+#endif /* CL_VERSION_2_0 */
+
+#ifdef __cplusplus
+}
+#endif
+
+
+#endif /* __CL_EXT_H */

3rdparty/opencl/CL/cl_gl.h

@@ -0,0 +1,162 @@
+/**********************************************************************************
+ * Copyright (c) 2008 - 2013 The Khronos Group Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a
+ * copy of this software and/or associated documentation files (the
+ * "Materials"), to deal in the Materials without restriction, including
+ * without limitation the rights to use, copy, modify, merge, publish,
+ * distribute, sublicense, and/or sell copies of the Materials, and to
+ * permit persons to whom the Materials are furnished to do so, subject to
+ * the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included
+ * in all copies or substantial portions of the Materials.
+ *
+ * THE MATERIALS ARE PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+ * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+ * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
+ * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
+ * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
+ * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
+ * MATERIALS OR THE USE OR OTHER DEALINGS IN THE MATERIALS.
+ **********************************************************************************/
+
+#ifndef __OPENCL_CL_GL_H
+#define __OPENCL_CL_GL_H
+
+#ifdef __APPLE__
+#include <OpenCL/cl.h>
+#else
+#include <CL/cl.h>
+#endif	
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+typedef cl_uint     cl_gl_object_type;
+typedef cl_uint     cl_gl_texture_info;
+typedef cl_uint     cl_gl_platform_info;
+typedef struct __GLsync *cl_GLsync;
+
+/* cl_gl_object_type = 0x2000 - 0x200F enum values are currently taken           */
+#define CL_GL_OBJECT_BUFFER                     0x2000
+#define CL_GL_OBJECT_TEXTURE2D                  0x2001
+#define CL_GL_OBJECT_TEXTURE3D                  0x2002
+#define CL_GL_OBJECT_RENDERBUFFER               0x2003
+#define CL_GL_OBJECT_TEXTURE2D_ARRAY            0x200E
+#define CL_GL_OBJECT_TEXTURE1D                  0x200F
+#define CL_GL_OBJECT_TEXTURE1D_ARRAY            0x2010
+#define CL_GL_OBJECT_TEXTURE_BUFFER             0x2011
+
+/* cl_gl_texture_info           */
+#define CL_GL_TEXTURE_TARGET                    0x2004
+#define CL_GL_MIPMAP_LEVEL                      0x2005
+#define CL_GL_NUM_SAMPLES                       0x2012
+
+
+extern CL_API_ENTRY cl_mem CL_API_CALL
+clCreateFromGLBuffer(cl_context     /* context */,
+                     cl_mem_flags   /* flags */,
+                     cl_GLuint      /* bufobj */,
+                     int *          /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_mem CL_API_CALL
+clCreateFromGLTexture(cl_context      /* context */,
+                      cl_mem_flags    /* flags */,
+                      cl_GLenum       /* target */,
+                      cl_GLint        /* miplevel */,
+                      cl_GLuint       /* texture */,
+                      cl_int *        /* errcode_ret */) CL_API_SUFFIX__VERSION_1_2;
+    
+extern CL_API_ENTRY cl_mem CL_API_CALL
+clCreateFromGLRenderbuffer(cl_context   /* context */,
+                           cl_mem_flags /* flags */,
+                           cl_GLuint    /* renderbuffer */,
+                           cl_int *     /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+clGetGLObjectInfo(cl_mem                /* memobj */,
+                  cl_gl_object_type *   /* gl_object_type */,
+                  cl_GLuint *           /* gl_object_name */) CL_API_SUFFIX__VERSION_1_0;
+                  
+extern CL_API_ENTRY cl_int CL_API_CALL
+clGetGLTextureInfo(cl_mem               /* memobj */,
+                   cl_gl_texture_info   /* param_name */,
+                   size_t               /* param_value_size */,
+                   void *               /* param_value */,
+                   size_t *             /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+clEnqueueAcquireGLObjects(cl_command_queue      /* command_queue */,
+                          cl_uint               /* num_objects */,
+                          const cl_mem *        /* mem_objects */,
+                          cl_uint               /* num_events_in_wait_list */,
+                          const cl_event *      /* event_wait_list */,
+                          cl_event *            /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+extern CL_API_ENTRY cl_int CL_API_CALL
+clEnqueueReleaseGLObjects(cl_command_queue      /* command_queue */,
+                          cl_uint               /* num_objects */,
+                          const cl_mem *        /* mem_objects */,
+                          cl_uint               /* num_events_in_wait_list */,
+                          const cl_event *      /* event_wait_list */,
+                          cl_event *            /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+
+/* Deprecated OpenCL 1.1 APIs */
+extern CL_API_ENTRY CL_EXT_PREFIX__VERSION_1_1_DEPRECATED cl_mem CL_API_CALL
+clCreateFromGLTexture2D(cl_context      /* context */,
+                        cl_mem_flags    /* flags */,
+                        cl_GLenum       /* target */,
+                        cl_GLint        /* miplevel */,
+                        cl_GLuint       /* texture */,
+                        cl_int *        /* errcode_ret */) CL_EXT_SUFFIX__VERSION_1_1_DEPRECATED;
+    
+extern CL_API_ENTRY CL_EXT_PREFIX__VERSION_1_1_DEPRECATED cl_mem CL_API_CALL
+clCreateFromGLTexture3D(cl_context      /* context */,
+                        cl_mem_flags    /* flags */,
+                        cl_GLenum       /* target */,
+                        cl_GLint        /* miplevel */,
+                        cl_GLuint       /* texture */,
+                        cl_int *        /* errcode_ret */) CL_EXT_SUFFIX__VERSION_1_1_DEPRECATED;
+    
+/* cl_khr_gl_sharing extension  */
+    
+#define cl_khr_gl_sharing 1
+    
+typedef cl_uint     cl_gl_context_info;
+    
+/* Additional Error Codes  */
+#define CL_INVALID_GL_SHAREGROUP_REFERENCE_KHR  -1000
+    
+/* cl_gl_context_info  */
+#define CL_CURRENT_DEVICE_FOR_GL_CONTEXT_KHR    0x2006
+#define CL_DEVICES_FOR_GL_CONTEXT_KHR           0x2007
+    
+/* Additional cl_context_properties  */
+#define CL_GL_CONTEXT_KHR                       0x2008
+#define CL_EGL_DISPLAY_KHR                      0x2009
+#define CL_GLX_DISPLAY_KHR                      0x200A
+#define CL_WGL_HDC_KHR                          0x200B
+#define CL_CGL_SHAREGROUP_KHR                   0x200C
+    
+extern CL_API_ENTRY cl_int CL_API_CALL
+clGetGLContextInfoKHR(const cl_context_properties * /* properties */,
+                      cl_gl_context_info            /* param_name */,
+                      size_t                        /* param_value_size */,
+                      void *                        /* param_value */,
+                      size_t *                      /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+    
+typedef CL_API_ENTRY cl_int (CL_API_CALL *clGetGLContextInfoKHR_fn)(
+    const cl_context_properties * properties,
+    cl_gl_context_info            param_name,
+    size_t                        param_value_size,
+    void *                        param_value,
+    size_t *                      param_value_size_ret);
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif  /* __OPENCL_CL_GL_H */

3rdparty/opencl/CL/cl_gl_ext.h

@@ -0,0 +1,69 @@
+/**********************************************************************************
+ * Copyright (c) 2008-2013 The Khronos Group Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a
+ * copy of this software and/or associated documentation files (the
+ * "Materials"), to deal in the Materials without restriction, including
+ * without limitation the rights to use, copy, modify, merge, publish,
+ * distribute, sublicense, and/or sell copies of the Materials, and to
+ * permit persons to whom the Materials are furnished to do so, subject to
+ * the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included
+ * in all copies or substantial portions of the Materials.
+ *
+ * THE MATERIALS ARE PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+ * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+ * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
+ * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
+ * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
+ * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
+ * MATERIALS OR THE USE OR OTHER DEALINGS IN THE MATERIALS.
+ **********************************************************************************/
+
+/* $Revision: 11708 $ on $Date: 2010-06-13 23:36:24 -0700 (Sun, 13 Jun 2010) $ */
+
+/* cl_gl_ext.h contains vendor (non-KHR) OpenCL extensions which have           */
+/* OpenGL dependencies.                                                         */
+
+#ifndef __OPENCL_CL_GL_EXT_H
+#define __OPENCL_CL_GL_EXT_H
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+#ifdef __APPLE__
+    #include <OpenCL/cl_gl.h>
+#else
+    #include <CL/cl_gl.h>
+#endif
+
+/*
+ * For each extension, follow this template
+ *  cl_VEN_extname extension  */
+/* #define cl_VEN_extname 1
+ * ... define new types, if any
+ * ... define new tokens, if any
+ * ... define new APIs, if any
+ *
+ *  If you need GLtypes here, mirror them with a cl_GLtype, rather than including a GL header
+ *  This allows us to avoid having to decide whether to include GL headers or GLES here.
+ */
+
+/* 
+ *  cl_khr_gl_event  extension
+ *  See section 9.9 in the OpenCL 1.1 spec for more information
+ */
+#define CL_COMMAND_GL_FENCE_SYNC_OBJECT_KHR     0x200D
+
+extern CL_API_ENTRY cl_event CL_API_CALL
+clCreateEventFromGLsyncKHR(cl_context           /* context */,
+                           cl_GLsync            /* cl_GLsync */,
+                           cl_int *             /* errcode_ret */) CL_EXT_SUFFIX__VERSION_1_1;
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif	/* __OPENCL_CL_GL_EXT_H  */

3rdparty/opencl/CL/opencl.h

@@ -0,0 +1,54 @@
+/*******************************************************************************
+ * Copyright (c) 2008-2013 The Khronos Group Inc.
+ *
+ * Permission is hereby granted, free of charge, to any person obtaining a
+ * copy of this software and/or associated documentation files (the
+ * "Materials"), to deal in the Materials without restriction, including
+ * without limitation the rights to use, copy, modify, merge, publish,
+ * distribute, sublicense, and/or sell copies of the Materials, and to
+ * permit persons to whom the Materials are furnished to do so, subject to
+ * the following conditions:
+ *
+ * The above copyright notice and this permission notice shall be included
+ * in all copies or substantial portions of the Materials.
+ *
+ * THE MATERIALS ARE PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+ * EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+ * MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
+ * IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
+ * CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
+ * TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
+ * MATERIALS OR THE USE OR OTHER DEALINGS IN THE MATERIALS.
+ ******************************************************************************/
+
+/* $Revision: 11708 $ on $Date: 2010-06-13 23:36:24 -0700 (Sun, 13 Jun 2010) $ */
+
+#ifndef __OPENCL_H
+#define __OPENCL_H
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+#ifdef __APPLE__
+
+#include <OpenCL/cl.h>
+#include <OpenCL/cl_gl.h>
+#include <OpenCL/cl_gl_ext.h>
+#include <OpenCL/cl_ext.h>
+
+#else
+
+#include <CL/cl.h>
+#include <CL/cl_gl.h>
+#include <CL/cl_gl_ext.h>
+#include <CL/cl_ext.h>
+
+#endif
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif  /* __OPENCL_H   */
+

3rdparty/opencl/opencl.cpp

@@ -0,0 +1,824 @@
+#include "CL/cl.h"
+
+typedef cl_int (CL_API_CALL * clGetPlatformIDsPtr)(cl_uint num_entries, cl_platform_id* platforms, cl_uint* num_platforms) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clGetPlatformInfoPtr)(cl_platform_id platform, cl_platform_info param_name, size_t param_value_size, void* param_value, size_t* param_value_size_ret) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clGetDeviceIDsPtr)(cl_platform_id platform, cl_device_type device_type, cl_uint num_entries, cl_device_id* devices, cl_uint* num_devices) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clGetDeviceInfoPtr)(cl_device_id device, cl_device_info param_name, size_t param_value_size, void* param_value, size_t* param_value_size_ret) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clCreateSubDevicesPtr)(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) CL_API_SUFFIX__VERSION_1_2;
+typedef cl_int (CL_API_CALL * clRetainDevicePtr)(cl_device_id device) CL_API_SUFFIX__VERSION_1_2;
+typedef cl_int (CL_API_CALL * clReleaseDevicePtr)(cl_device_id device) CL_API_SUFFIX__VERSION_1_2;
+typedef cl_context (CL_API_CALL * clCreateContextPtr)(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) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_context (CL_API_CALL * clCreateContextFromTypePtr)(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) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clRetainContextPtr)(cl_context context) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clReleaseContextPtr)(cl_context context) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clGetContextInfoPtr)(cl_context context, cl_context_info param_name, size_t param_value_size, void* param_value, size_t* param_value_size_ret) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_command_queue (CL_API_CALL * clCreateCommandQueueWithPropertiesPtr)(cl_context context, cl_device_id device, const cl_queue_properties* properties, cl_int* errcode_ret) CL_API_SUFFIX__VERSION_2_0;
+typedef cl_int (CL_API_CALL * clRetainCommandQueuePtr)(cl_command_queue command_queue) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clReleaseCommandQueuePtr)(cl_command_queue command_queue) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clGetCommandQueueInfoPtr)(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) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_mem (CL_API_CALL * clCreateBufferPtr)(cl_context context, cl_mem_flags flags, size_t size, void* host_ptr, cl_int* errcode_ret) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_mem (CL_API_CALL * clCreateSubBufferPtr)(cl_mem buffer, cl_mem_flags flags, cl_buffer_create_type buffer_create_type, const void* buffer_create_info, cl_int* errcode_ret) CL_API_SUFFIX__VERSION_1_1;
+typedef cl_mem (CL_API_CALL * clCreateImagePtr)(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) CL_API_SUFFIX__VERSION_1_2;
+typedef cl_mem (CL_API_CALL * clCreatePipePtr)(cl_context context, cl_mem_flags flags, cl_uint pipe_packet_size, cl_uint pipe_max_packets, const cl_pipe_properties* properties, cl_int* errcode_ret) CL_API_SUFFIX__VERSION_2_0;
+typedef cl_int (CL_API_CALL * clRetainMemObjectPtr)(cl_mem memobj) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clReleaseMemObjectPtr)(cl_mem memobj) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clGetSupportedImageFormatsPtr)(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) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clGetMemObjectInfoPtr)(cl_mem memobj, cl_mem_info param_name, size_t param_value_size, void* param_value, size_t* param_value_size_ret) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clGetImageInfoPtr)(cl_mem image, cl_image_info param_name, size_t param_value_size, void* param_value, size_t* param_value_size_ret) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clGetPipeInfoPtr)(cl_mem pipe, cl_pipe_info param_name, size_t param_value_size, void* param_value, size_t* param_value_size_ret) CL_API_SUFFIX__VERSION_2_0;
+typedef cl_int (CL_API_CALL * clSetMemObjectDestructorCallbackPtr)(cl_mem memobj, void (CL_CALLBACK* pfn_notify)(cl_mem memobj, void* user_data), void* user_data) CL_API_SUFFIX__VERSION_1_1;
+typedef void* (CL_API_CALL * clSVMAllocPtr)(cl_context context, cl_svm_mem_flags flags, size_t size, cl_uint alignment) CL_API_SUFFIX__VERSION_2_0;
+typedef void (CL_API_CALL * clSVMFreePtr)(cl_context context, void* svm_pointer) CL_API_SUFFIX__VERSION_2_0;
+typedef cl_sampler (CL_API_CALL * clCreateSamplerWithPropertiesPtr)(cl_context context, const cl_sampler_properties* normalized_coords, cl_int* errcode_ret) CL_API_SUFFIX__VERSION_2_0;
+typedef cl_int (CL_API_CALL * clRetainSamplerPtr)(cl_sampler sampler) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clReleaseSamplerPtr)(cl_sampler sampler) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clGetSamplerInfoPtr)(cl_sampler sampler, cl_sampler_info param_name, size_t param_value_size, void* param_value, size_t* param_value_size_ret) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_program (CL_API_CALL * clCreateProgramWithSourcePtr)(cl_context context, cl_uint count, const char** strings, const size_t* lengths, cl_int* errcode_ret) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_program (CL_API_CALL * clCreateProgramWithBinaryPtr)(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) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_program (CL_API_CALL * clCreateProgramWithBuiltInKernelsPtr)(cl_context context, cl_uint num_devices, const cl_device_id* device_list, const char* kernel_names, cl_int* errcode_ret) CL_API_SUFFIX__VERSION_1_2;
+typedef cl_int (CL_API_CALL * clRetainProgramPtr)(cl_program program) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clReleaseProgramPtr)(cl_program program) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clBuildProgramPtr)(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) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clCompileProgramPtr)(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) CL_API_SUFFIX__VERSION_1_2;
+typedef cl_program (CL_API_CALL * clLinkProgramPtr)(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) CL_API_SUFFIX__VERSION_1_2;
+typedef cl_int (CL_API_CALL * clUnloadPlatformCompilerPtr)(cl_platform_id platform) CL_API_SUFFIX__VERSION_1_2;
+typedef cl_int (CL_API_CALL * clGetProgramInfoPtr)(cl_program program, cl_program_info param_name, size_t param_value_size, void* param_value, size_t* param_value_size_ret) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clGetProgramBuildInfoPtr)(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) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_kernel (CL_API_CALL * clCreateKernelPtr)(cl_program program, const char* kernel_name, cl_int* errcode_ret) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clCreateKernelsInProgramPtr)(cl_program program, cl_uint num_kernels, cl_kernel* kernels, cl_uint* num_kernels_ret) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clRetainKernelPtr)(cl_kernel kernel) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clReleaseKernelPtr)(cl_kernel kernel) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clSetKernelArgPtr)(cl_kernel kernel, cl_uint arg_index, size_t arg_size, const void* arg_value) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clSetKernelArgSVMPointerPtr)(cl_kernel kernel, cl_uint arg_index, const void* arg_value) CL_API_SUFFIX__VERSION_2_0;
+typedef cl_int (CL_API_CALL * clSetKernelExecInfoPtr)(cl_kernel kernel, cl_kernel_exec_info param_name, size_t param_value_size, const void* param_value) CL_API_SUFFIX__VERSION_2_0;
+typedef cl_int (CL_API_CALL * clGetKernelInfoPtr)(cl_kernel kernel, cl_kernel_info param_name, size_t param_value_size, void* param_value, size_t* param_value_size_ret) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clGetKernelArgInfoPtr)(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) CL_API_SUFFIX__VERSION_1_2;
+typedef cl_int (CL_API_CALL * clGetKernelWorkGroupInfoPtr)(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) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clWaitForEventsPtr)(cl_uint num_events, const cl_event* event_list) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clGetEventInfoPtr)(cl_event event, cl_event_info param_name, size_t param_value_size, void* param_value, size_t* param_value_size_ret) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_event (CL_API_CALL * clCreateUserEventPtr)(cl_context context, cl_int* errcode_ret) CL_API_SUFFIX__VERSION_1_1;
+typedef cl_int (CL_API_CALL * clRetainEventPtr)(cl_event event) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clReleaseEventPtr)(cl_event event) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clSetUserEventStatusPtr)(cl_event event, cl_int execution_status) CL_API_SUFFIX__VERSION_1_1;
+typedef cl_int (CL_API_CALL * clSetEventCallbackPtr)(cl_event event, cl_int command_exec_callback_type, void (CL_CALLBACK* pfn_notify)(cl_event, cl_int, void*), void* user_data) CL_API_SUFFIX__VERSION_1_1;
+typedef cl_int (CL_API_CALL * clGetEventProfilingInfoPtr)(cl_event event, cl_profiling_info param_name, size_t param_value_size, void* param_value, size_t* param_value_size_ret) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clFlushPtr)(cl_command_queue command_queue) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clFinishPtr)(cl_command_queue command_queue) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clEnqueueReadBufferPtr)(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) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clEnqueueReadBufferRectPtr)(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) CL_API_SUFFIX__VERSION_1_1;
+typedef cl_int (CL_API_CALL * clEnqueueWriteBufferPtr)(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) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clEnqueueWriteBufferRectPtr)(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) CL_API_SUFFIX__VERSION_1_1;
+typedef cl_int (CL_API_CALL * clEnqueueFillBufferPtr)(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) CL_API_SUFFIX__VERSION_1_2;
+typedef cl_int (CL_API_CALL * clEnqueueCopyBufferPtr)(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) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clEnqueueCopyBufferRectPtr)(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) CL_API_SUFFIX__VERSION_1_1;
+typedef cl_int (CL_API_CALL * clEnqueueReadImagePtr)(cl_command_queue command_queue, cl_mem image, cl_bool blocking_read, const size_t* origin, const size_t* region, 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) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clEnqueueWriteImagePtr)(cl_command_queue command_queue, cl_mem image, cl_bool blocking_write, const size_t* origin, const size_t* region, 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) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clEnqueueFillImagePtr)(cl_command_queue command_queue, cl_mem image, const void* fill_color, const size_t* origin, const size_t* region, cl_uint num_events_in_wait_list, const cl_event* event_wait_list, cl_event* event) CL_API_SUFFIX__VERSION_1_2;
+typedef cl_int (CL_API_CALL * clEnqueueCopyImagePtr)(cl_command_queue command_queue, cl_mem src_image, cl_mem dst_image, const size_t* src_origin, const size_t* dst_origin, const size_t* region, cl_uint num_events_in_wait_list, const cl_event* event_wait_list, cl_event* event) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clEnqueueCopyImageToBufferPtr)(cl_command_queue command_queue, cl_mem src_image, cl_mem dst_buffer, const size_t* src_origin, const size_t* region, size_t dst_offset, cl_uint num_events_in_wait_list, const cl_event* event_wait_list, cl_event* event) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clEnqueueCopyBufferToImagePtr)(cl_command_queue command_queue, cl_mem src_buffer, cl_mem dst_image, size_t src_offset, const size_t* dst_origin, const size_t* region, cl_uint num_events_in_wait_list, const cl_event* event_wait_list, cl_event* event) CL_API_SUFFIX__VERSION_1_0;
+typedef void* (CL_API_CALL * clEnqueueMapBufferPtr)(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) CL_API_SUFFIX__VERSION_1_0;
+typedef void* (CL_API_CALL * clEnqueueMapImagePtr)(cl_command_queue command_queue, cl_mem image, cl_bool blocking_map, cl_map_flags map_flags, const size_t* origin, const size_t* region, 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) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clEnqueueUnmapMemObjectPtr)(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) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clEnqueueMigrateMemObjectsPtr)(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) CL_API_SUFFIX__VERSION_1_2;
+typedef cl_int (CL_API_CALL * clEnqueueNDRangeKernelPtr)(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) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clEnqueueNativeKernelPtr)(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) CL_API_SUFFIX__VERSION_1_0;
+typedef cl_int (CL_API_CALL * clEnqueueMarkerWithWaitListPtr)(cl_command_queue command_queue, cl_uint num_events_in_wait_list, const cl_event* event_wait_list, cl_event* event) CL_API_SUFFIX__VERSION_1_2;
+typedef cl_int (CL_API_CALL * clEnqueueBarrierWithWaitListPtr)(cl_command_queue command_queue, cl_uint num_events_in_wait_list, const cl_event* event_wait_list, cl_event* event) CL_API_SUFFIX__VERSION_1_2;
+typedef cl_int (CL_API_CALL * clEnqueueSVMFreePtr)(cl_command_queue command_queue, cl_uint num_svm_pointers, void* svm_pointers[], void (CL_CALLBACK* /*pfn_free_func*/)(cl_command_queue queue, cl_uint num_svm_pointers, void* svm_pointers[], void* user_data), void* user_data, cl_uint num_events_in_wait_list, const cl_event* event_wait_list, cl_event* event) CL_API_SUFFIX__VERSION_2_0;
+typedef cl_int (CL_API_CALL * clEnqueueSVMMemcpyPtr)(cl_command_queue command_queue, cl_bool blocking_copy, void* dst_ptr, const void* src_ptr, size_t size, cl_uint num_events_in_wait_list, const cl_event* event_wait_list, cl_event* event) CL_API_SUFFIX__VERSION_2_0;
+typedef cl_int (CL_API_CALL * clEnqueueSVMMemFillPtr)(cl_command_queue command_queue, void* svm_ptr, const void* pattern, size_t pattern_size, size_t size, cl_uint num_events_in_wait_list, const cl_event* event_wait_list, cl_event* event) CL_API_SUFFIX__VERSION_2_0;
+typedef cl_int (CL_API_CALL * clEnqueueSVMMapPtr)(cl_command_queue command_queue, cl_bool blocking_map, cl_map_flags flags, void* svm_ptr, size_t size, cl_uint num_events_in_wait_list, const cl_event* event_wait_list, cl_event* event) CL_API_SUFFIX__VERSION_2_0;
+typedef cl_int (CL_API_CALL * clEnqueueSVMUnmapPtr)(cl_command_queue command_queue, void* svm_ptr, cl_uint num_events_in_wait_list, const cl_event* event_wait_list, cl_event* event) CL_API_SUFFIX__VERSION_2_0;
+typedef void* (CL_API_CALL * clGetExtensionFunctionAddressForPlatformPtr)(cl_platform_id platform, const char* func_name) CL_API_SUFFIX__VERSION_1_2;
+typedef cl_mem (CL_API_CALL * clCreateImage2DPtr)(cl_context context, cl_mem_flags flags, const cl_image_format* image_format, size_t image_width, size_t image_height, size_t image_row_pitch, void* host_ptr, cl_int* errcode_ret);
+typedef cl_mem (CL_API_CALL * clCreateImage3DPtr)(cl_context context, cl_mem_flags flags, const cl_image_format* image_format, size_t image_width, size_t image_height, size_t image_depth, size_t image_row_pitch, size_t image_slice_pitch, void* host_ptr, cl_int* errcode_ret);
+typedef cl_int (CL_API_CALL * clEnqueueMarkerPtr)(cl_command_queue command_queue, cl_event* event);
+typedef cl_int (CL_API_CALL * clEnqueueWaitForEventsPtr)(cl_command_queue command_queue, cl_uint num_events, const cl_event* event_list);
+typedef cl_int (CL_API_CALL * clEnqueueBarrierPtr)(cl_command_queue command_queue);
+typedef cl_int (CL_API_CALL * clUnloadCompilerPtr)(void);
+typedef void* (CL_API_CALL * clGetExtensionFunctionAddressPtr)(const char* func_name);
+typedef cl_command_queue (CL_API_CALL * clCreateCommandQueuePtr)(cl_context context, cl_device_id device, cl_command_queue_properties properties, cl_int* errcode_ret);
+typedef cl_sampler (CL_API_CALL * clCreateSamplerPtr)(cl_context context, cl_bool normalized_coords, cl_addressing_mode addressing_mode, cl_filter_mode filter_mode, cl_int* errcode_ret);
+typedef cl_int (CL_API_CALL * clEnqueueTaskPtr)(cl_command_queue command_queue, cl_kernel kernel, cl_uint num_events_in_wait_list, const cl_event* event_wait_list, cl_event* event);
+
+static clGetPlatformIDsPtr cl_GetPlatformIDs = NULL;
+static clGetPlatformInfoPtr cl_GetPlatformInfo = NULL;
+static clGetDeviceIDsPtr cl_GetDeviceIDs = NULL;
+static clGetDeviceInfoPtr cl_GetDeviceInfo = NULL;
+static clCreateSubDevicesPtr cl_CreateSubDevices = NULL;
+static clRetainDevicePtr cl_RetainDevice = NULL;
+static clReleaseDevicePtr cl_ReleaseDevice = NULL;
+static clCreateContextPtr cl_CreateContext = NULL;
+static clCreateContextFromTypePtr cl_CreateContextFromType = NULL;
+static clRetainContextPtr cl_RetainContext = NULL;
+static clReleaseContextPtr cl_ReleaseContext = NULL;
+static clGetContextInfoPtr cl_GetContextInfo = NULL;
+static clCreateCommandQueueWithPropertiesPtr cl_CreateCommandQueueWithProperties = NULL;
+static clRetainCommandQueuePtr cl_RetainCommandQueue = NULL;
+static clReleaseCommandQueuePtr cl_ReleaseCommandQueue = NULL;
+static clGetCommandQueueInfoPtr cl_GetCommandQueueInfo = NULL;
+static clCreateBufferPtr cl_CreateBuffer = NULL;
+static clCreateSubBufferPtr cl_CreateSubBuffer = NULL;
+static clCreateImagePtr cl_CreateImage = NULL;
+static clCreatePipePtr cl_CreatePipe = NULL;
+static clRetainMemObjectPtr cl_RetainMemObject = NULL;
+static clReleaseMemObjectPtr cl_ReleaseMemObject = NULL;
+static clGetSupportedImageFormatsPtr cl_GetSupportedImageFormats = NULL;
+static clGetMemObjectInfoPtr cl_GetMemObjectInfo = NULL;
+static clGetImageInfoPtr cl_GetImageInfo = NULL;
+static clGetPipeInfoPtr cl_GetPipeInfo = NULL;
+static clSetMemObjectDestructorCallbackPtr cl_SetMemObjectDestructorCallback = NULL;
+static clSVMAllocPtr cl_SVMAlloc = NULL;
+static clSVMFreePtr cl_SVMFree = NULL;
+static clCreateSamplerWithPropertiesPtr cl_CreateSamplerWithProperties = NULL;
+static clRetainSamplerPtr cl_RetainSampler = NULL;
+static clReleaseSamplerPtr cl_ReleaseSampler = NULL;
+static clGetSamplerInfoPtr cl_GetSamplerInfo = NULL;
+static clCreateProgramWithSourcePtr cl_CreateProgramWithSource = NULL;
+static clCreateProgramWithBinaryPtr cl_CreateProgramWithBinary = NULL;
+static clCreateProgramWithBuiltInKernelsPtr cl_CreateProgramWithBuiltInKernels = NULL;
+static clRetainProgramPtr cl_RetainProgram = NULL;
+static clReleaseProgramPtr cl_ReleaseProgram = NULL;
+static clBuildProgramPtr cl_BuildProgram = NULL;
+static clCompileProgramPtr cl_CompileProgram = NULL;
+static clLinkProgramPtr cl_LinkProgram = NULL;
+static clUnloadPlatformCompilerPtr cl_UnloadPlatformCompiler = NULL;
+static clGetProgramInfoPtr cl_GetProgramInfo = NULL;
+static clGetProgramBuildInfoPtr cl_GetProgramBuildInfo = NULL;
+static clCreateKernelPtr cl_CreateKernel = NULL;
+static clCreateKernelsInProgramPtr cl_CreateKernelsInProgram = NULL;
+static clRetainKernelPtr cl_RetainKernel = NULL;
+static clReleaseKernelPtr cl_ReleaseKernel = NULL;
+static clSetKernelArgPtr cl_SetKernelArg = NULL;
+static clSetKernelArgSVMPointerPtr cl_SetKernelArgSVMPointer = NULL;
+static clSetKernelExecInfoPtr cl_SetKernelExecInfo = NULL;
+static clGetKernelInfoPtr cl_GetKernelInfo = NULL;
+static clGetKernelArgInfoPtr cl_GetKernelArgInfo = NULL;
+static clGetKernelWorkGroupInfoPtr cl_GetKernelWorkGroupInfo = NULL;
+static clWaitForEventsPtr cl_WaitForEvents = NULL;
+static clGetEventInfoPtr cl_GetEventInfo = NULL;
+static clCreateUserEventPtr cl_CreateUserEvent = NULL;
+static clRetainEventPtr cl_RetainEvent = NULL;
+static clReleaseEventPtr cl_ReleaseEvent = NULL;
+static clSetUserEventStatusPtr cl_SetUserEventStatus = NULL;
+static clSetEventCallbackPtr cl_SetEventCallback = NULL;
+static clGetEventProfilingInfoPtr cl_GetEventProfilingInfo = NULL;
+static clFlushPtr cl_Flush = NULL;
+static clFinishPtr cl_Finish = NULL;
+static clEnqueueReadBufferPtr cl_EnqueueReadBuffer = NULL;
+static clEnqueueReadBufferRectPtr cl_EnqueueReadBufferRect = NULL;
+static clEnqueueWriteBufferPtr cl_EnqueueWriteBuffer = NULL;
+static clEnqueueWriteBufferRectPtr cl_EnqueueWriteBufferRect = NULL;
+static clEnqueueFillBufferPtr cl_EnqueueFillBuffer = NULL;
+static clEnqueueCopyBufferPtr cl_EnqueueCopyBuffer = NULL;
+static clEnqueueCopyBufferRectPtr cl_EnqueueCopyBufferRect = NULL;
+static clEnqueueReadImagePtr cl_EnqueueReadImage = NULL;
+static clEnqueueWriteImagePtr cl_EnqueueWriteImage = NULL;
+static clEnqueueFillImagePtr cl_EnqueueFillImage = NULL;
+static clEnqueueCopyImagePtr cl_EnqueueCopyImage = NULL;
+static clEnqueueCopyImageToBufferPtr cl_EnqueueCopyImageToBuffer = NULL;
+static clEnqueueCopyBufferToImagePtr cl_EnqueueCopyBufferToImage = NULL;
+static clEnqueueMapBufferPtr cl_EnqueueMapBuffer = NULL;
+static clEnqueueMapImagePtr cl_EnqueueMapImage = NULL;
+static clEnqueueUnmapMemObjectPtr cl_EnqueueUnmapMemObject = NULL;
+static clEnqueueMigrateMemObjectsPtr cl_EnqueueMigrateMemObjects = NULL;
+static clEnqueueNDRangeKernelPtr cl_EnqueueNDRangeKernel = NULL;
+static clEnqueueNativeKernelPtr cl_EnqueueNativeKernel = NULL;
+static clEnqueueMarkerWithWaitListPtr cl_EnqueueMarkerWithWaitList = NULL;
+static clEnqueueBarrierWithWaitListPtr cl_EnqueueBarrierWithWaitList = NULL;
+static clEnqueueSVMFreePtr cl_EnqueueSVMFree = NULL;
+static clEnqueueSVMMemcpyPtr cl_EnqueueSVMMemcpy = NULL;
+static clEnqueueSVMMemFillPtr cl_EnqueueSVMMemFill = NULL;
+static clEnqueueSVMMapPtr cl_EnqueueSVMMap = NULL;
+static clEnqueueSVMUnmapPtr cl_EnqueueSVMUnmap = NULL;
+static clGetExtensionFunctionAddressForPlatformPtr cl_GetExtensionFunctionAddressForPlatform = NULL;
+static clCreateImage2DPtr cl_CreateImage2D = NULL;
+static clCreateImage3DPtr cl_CreateImage3D = NULL;
+static clEnqueueMarkerPtr cl_EnqueueMarker = NULL;
+static clEnqueueWaitForEventsPtr cl_EnqueueWaitForEvents = NULL;
+static clEnqueueBarrierPtr cl_EnqueueBarrier = NULL;
+static clUnloadCompilerPtr cl_UnloadCompiler = NULL;
+static clGetExtensionFunctionAddressPtr cl_GetExtensionFunctionAddress = NULL;
+static clCreateCommandQueuePtr cl_CreateCommandQueue = NULL;
+static clCreateSamplerPtr cl_CreateSampler = NULL;
+static clEnqueueTaskPtr cl_EnqueueTask = NULL;
+
+#include <Windows.h>
+
+static struct Loader
+{
+	struct Loader()
+	{
+		HMODULE hModule = LoadLibrary("OpenCL.dll");
+
+		if(hModule == NULL) return;
+
+		*(void**)&cl_GetPlatformIDs = GetProcAddress(hModule, "clGetPlatformIDs");
+		*(void**)&cl_GetPlatformInfo = GetProcAddress(hModule, "clGetPlatformInfo");
+		*(void**)&cl_GetDeviceIDs = GetProcAddress(hModule, "clGetDeviceIDs");
+		*(void**)&cl_GetDeviceInfo = GetProcAddress(hModule, "clGetDeviceInfo");
+		*(void**)&cl_CreateSubDevices = GetProcAddress(hModule, "clCreateSubDevices");
+		*(void**)&cl_RetainDevice = GetProcAddress(hModule, "clRetainDevice");
+		*(void**)&cl_ReleaseDevice = GetProcAddress(hModule, "clReleaseDevice");
+		*(void**)&cl_CreateContext = GetProcAddress(hModule, "clCreateContext");
+		*(void**)&cl_CreateContextFromType = GetProcAddress(hModule, "clCreateContextFromType");
+		*(void**)&cl_RetainContext = GetProcAddress(hModule, "clRetainContext");
+		*(void**)&cl_ReleaseContext = GetProcAddress(hModule, "clReleaseContext");
+		*(void**)&cl_GetContextInfo = GetProcAddress(hModule, "clGetContextInfo");
+		*(void**)&cl_CreateCommandQueueWithProperties = GetProcAddress(hModule, "clCreateCommandQueueWithProperties");
+		*(void**)&cl_RetainCommandQueue = GetProcAddress(hModule, "clRetainCommandQueue");
+		*(void**)&cl_ReleaseCommandQueue = GetProcAddress(hModule, "clReleaseCommandQueue");
+		*(void**)&cl_GetCommandQueueInfo = GetProcAddress(hModule, "clGetCommandQueueInfo");
+		*(void**)&cl_CreateBuffer = GetProcAddress(hModule, "clCreateBuffer");
+		*(void**)&cl_CreateSubBuffer = GetProcAddress(hModule, "clCreateSubBuffer");
+		*(void**)&cl_CreateImage = GetProcAddress(hModule, "clCreateImage");
+		*(void**)&cl_CreatePipe = GetProcAddress(hModule, "clCreatePipe");
+		*(void**)&cl_RetainMemObject = GetProcAddress(hModule, "clRetainMemObject");
+		*(void**)&cl_ReleaseMemObject = GetProcAddress(hModule, "clReleaseMemObject");
+		*(void**)&cl_GetSupportedImageFormats = GetProcAddress(hModule, "clGetSupportedImageFormats");
+		*(void**)&cl_GetMemObjectInfo = GetProcAddress(hModule, "clGetMemObjectInfo");
+		*(void**)&cl_GetImageInfo = GetProcAddress(hModule, "clGetImageInfo");
+		*(void**)&cl_GetPipeInfo = GetProcAddress(hModule, "clGetPipeInfo");
+		*(void**)&cl_SetMemObjectDestructorCallback = GetProcAddress(hModule, "clSetMemObjectDestructorCallback");
+		*(void**)&cl_SVMAlloc = GetProcAddress(hModule, "clSVMAlloc");
+		*(void**)&cl_SVMFree = GetProcAddress(hModule, "clSVMFree");
+		*(void**)&cl_CreateSamplerWithProperties = GetProcAddress(hModule, "clCreateSamplerWithProperties");
+		*(void**)&cl_RetainSampler = GetProcAddress(hModule, "clRetainSampler");
+		*(void**)&cl_ReleaseSampler = GetProcAddress(hModule, "clReleaseSampler");
+		*(void**)&cl_GetSamplerInfo = GetProcAddress(hModule, "clGetSamplerInfo");
+		*(void**)&cl_CreateProgramWithSource = GetProcAddress(hModule, "clCreateProgramWithSource");
+		*(void**)&cl_CreateProgramWithBinary = GetProcAddress(hModule, "clCreateProgramWithBinary");
+		*(void**)&cl_CreateProgramWithBuiltInKernels = GetProcAddress(hModule, "clCreateProgramWithBuiltInKernels");
+		*(void**)&cl_RetainProgram = GetProcAddress(hModule, "clRetainProgram");
+		*(void**)&cl_ReleaseProgram = GetProcAddress(hModule, "clReleaseProgram");
+		*(void**)&cl_BuildProgram = GetProcAddress(hModule, "clBuildProgram");
+		*(void**)&cl_CompileProgram = GetProcAddress(hModule, "clCompileProgram");
+		*(void**)&cl_LinkProgram = GetProcAddress(hModule, "clLinkProgram");
+		*(void**)&cl_UnloadPlatformCompiler = GetProcAddress(hModule, "clUnloadPlatformCompiler");
+		*(void**)&cl_GetProgramInfo = GetProcAddress(hModule, "clGetProgramInfo");
+		*(void**)&cl_GetProgramBuildInfo = GetProcAddress(hModule, "clGetProgramBuildInfo");
+		*(void**)&cl_CreateKernel = GetProcAddress(hModule, "clCreateKernel");
+		*(void**)&cl_CreateKernelsInProgram = GetProcAddress(hModule, "clCreateKernelsInProgram");
+		*(void**)&cl_RetainKernel = GetProcAddress(hModule, "clRetainKernel");
+		*(void**)&cl_ReleaseKernel = GetProcAddress(hModule, "clReleaseKernel");
+		*(void**)&cl_SetKernelArg = GetProcAddress(hModule, "clSetKernelArg");
+		*(void**)&cl_SetKernelArgSVMPointer = GetProcAddress(hModule, "clSetKernelArgSVMPointer");
+		*(void**)&cl_SetKernelExecInfo = GetProcAddress(hModule, "clSetKernelExecInfo");
+		*(void**)&cl_GetKernelInfo = GetProcAddress(hModule, "clGetKernelInfo");
+		*(void**)&cl_GetKernelArgInfo = GetProcAddress(hModule, "clGetKernelArgInfo");
+		*(void**)&cl_GetKernelWorkGroupInfo = GetProcAddress(hModule, "clGetKernelWorkGroupInfo");
+		*(void**)&cl_WaitForEvents = GetProcAddress(hModule, "clWaitForEvents");
+		*(void**)&cl_GetEventInfo = GetProcAddress(hModule, "clGetEventInfo");
+		*(void**)&cl_CreateUserEvent = GetProcAddress(hModule, "clCreateUserEvent");
+		*(void**)&cl_RetainEvent = GetProcAddress(hModule, "clRetainEvent");
+		*(void**)&cl_ReleaseEvent = GetProcAddress(hModule, "clReleaseEvent");
+		*(void**)&cl_SetUserEventStatus = GetProcAddress(hModule, "clSetUserEventStatus");
+		*(void**)&cl_SetEventCallback = GetProcAddress(hModule, "clSetEventCallback");
+		*(void**)&cl_GetEventProfilingInfo = GetProcAddress(hModule, "clGetEventProfilingInfo");
+		*(void**)&cl_Flush = GetProcAddress(hModule, "clFlush");
+		*(void**)&cl_Finish = GetProcAddress(hModule, "clFinish");
+		*(void**)&cl_EnqueueReadBuffer = GetProcAddress(hModule, "clEnqueueReadBuffer");
+		*(void**)&cl_EnqueueReadBufferRect = GetProcAddress(hModule, "clEnqueueReadBufferRect");
+		*(void**)&cl_EnqueueWriteBuffer = GetProcAddress(hModule, "clEnqueueWriteBuffer");
+		*(void**)&cl_EnqueueWriteBufferRect = GetProcAddress(hModule, "clEnqueueWriteBufferRect");
+		*(void**)&cl_EnqueueFillBuffer = GetProcAddress(hModule, "clEnqueueFillBuffer");
+		*(void**)&cl_EnqueueCopyBuffer = GetProcAddress(hModule, "clEnqueueCopyBuffer");
+		*(void**)&cl_EnqueueCopyBufferRect = GetProcAddress(hModule, "clEnqueueCopyBufferRect");
+		*(void**)&cl_EnqueueReadImage = GetProcAddress(hModule, "clEnqueueReadImage");
+		*(void**)&cl_EnqueueWriteImage = GetProcAddress(hModule, "clEnqueueWriteImage");
+		*(void**)&cl_EnqueueFillImage = GetProcAddress(hModule, "clEnqueueFillImage");
+		*(void**)&cl_EnqueueCopyImage = GetProcAddress(hModule, "clEnqueueCopyImage");
+		*(void**)&cl_EnqueueCopyImageToBuffer = GetProcAddress(hModule, "clEnqueueCopyImageToBuffer");
+		*(void**)&cl_EnqueueCopyBufferToImage = GetProcAddress(hModule, "clEnqueueCopyBufferToImage");
+		*(void**)&cl_EnqueueMapBuffer = GetProcAddress(hModule, "clEnqueueMapBuffer");
+		*(void**)&cl_EnqueueMapImage = GetProcAddress(hModule, "clEnqueueMapImage");
+		*(void**)&cl_EnqueueUnmapMemObject = GetProcAddress(hModule, "clEnqueueUnmapMemObject");
+		*(void**)&cl_EnqueueMigrateMemObjects = GetProcAddress(hModule, "clEnqueueMigrateMemObjects");
+		*(void**)&cl_EnqueueNDRangeKernel = GetProcAddress(hModule, "clEnqueueNDRangeKernel");
+		*(void**)&cl_EnqueueNativeKernel = GetProcAddress(hModule, "clEnqueueNativeKernel");
+		*(void**)&cl_EnqueueMarkerWithWaitList = GetProcAddress(hModule, "clEnqueueMarkerWithWaitList");
+		*(void**)&cl_EnqueueBarrierWithWaitList = GetProcAddress(hModule, "clEnqueueBarrierWithWaitList");
+		*(void**)&cl_EnqueueSVMFree = GetProcAddress(hModule, "clEnqueueSVMFree");
+		*(void**)&cl_EnqueueSVMMemcpy = GetProcAddress(hModule, "clEnqueueSVMMemcpy");
+		*(void**)&cl_EnqueueSVMMemFill = GetProcAddress(hModule, "clEnqueueSVMMemFill");
+		*(void**)&cl_EnqueueSVMMap = GetProcAddress(hModule, "clEnqueueSVMMap");
+		*(void**)&cl_EnqueueSVMUnmap = GetProcAddress(hModule, "clEnqueueSVMUnmap");
+		*(void**)&cl_GetExtensionFunctionAddressForPlatform = GetProcAddress(hModule, "clGetExtensionFunctionAddressForPlatform");
+		*(void**)&cl_CreateImage2D = GetProcAddress(hModule, "clCreateImage2D");
+		*(void**)&cl_CreateImage3D = GetProcAddress(hModule, "clCreateImage3D");
+		*(void**)&cl_EnqueueMarker = GetProcAddress(hModule, "clEnqueueMarker");
+		*(void**)&cl_EnqueueWaitForEvents = GetProcAddress(hModule, "clEnqueueWaitForEvents");
+		*(void**)&cl_EnqueueBarrier = GetProcAddress(hModule, "clEnqueueBarrier");
+		*(void**)&cl_UnloadCompiler = GetProcAddress(hModule, "clUnloadCompiler");
+		*(void**)&cl_GetExtensionFunctionAddress = GetProcAddress(hModule, "clGetExtensionFunctionAddress");
+		*(void**)&cl_CreateCommandQueue = GetProcAddress(hModule, "clCreateCommandQueue");
+		*(void**)&cl_CreateSampler = GetProcAddress(hModule, "clCreateSampler");
+		*(void**)&cl_EnqueueTask = GetProcAddress(hModule, "clEnqueueTask");
+	}
+} s_loader;
+
+cl_int CL_API_CALL clGetPlatformIDs(cl_uint num_entries, cl_platform_id* platforms, cl_uint* num_platforms) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_GetPlatformIDs(num_entries, platforms, num_platforms);
+}
+
+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) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_GetPlatformInfo(platform, param_name, param_value_size, param_value, param_value_size_ret);
+}
+
+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) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_GetDeviceIDs(platform, device_type, num_entries, devices, num_devices);
+}
+
+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) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_GetDeviceInfo(device, param_name, param_value_size, param_value, param_value_size_ret);
+}
+
+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) CL_API_SUFFIX__VERSION_1_2
+{
+	return cl_CreateSubDevices(in_device, properties, num_devices, out_devices, num_devices_ret);
+}
+
+cl_int CL_API_CALL clRetainDevice(cl_device_id device) CL_API_SUFFIX__VERSION_1_2
+{
+	return cl_RetainDevice(device);
+}
+
+cl_int CL_API_CALL clReleaseDevice(cl_device_id device) CL_API_SUFFIX__VERSION_1_2
+{
+	return cl_ReleaseDevice(device);
+}
+
+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) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_CreateContext(properties, num_devices, devices, pfn_notify, user_data, errcode_ret);
+}
+
+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) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_CreateContextFromType(properties, device_type, pfn_notify, user_data, errcode_ret);
+}
+
+cl_int CL_API_CALL clRetainContext(cl_context context) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_RetainContext(context);
+}
+
+cl_int CL_API_CALL clReleaseContext(cl_context context) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_ReleaseContext(context);
+}
+
+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) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_GetContextInfo(context, param_name, param_value_size, param_value, param_value_size_ret);
+}
+
+cl_command_queue CL_API_CALL clCreateCommandQueueWithProperties(cl_context context, cl_device_id device, const cl_queue_properties* properties, cl_int* errcode_ret) CL_API_SUFFIX__VERSION_2_0
+{
+	return cl_CreateCommandQueueWithProperties(context, device, properties, errcode_ret);
+}
+
+cl_int CL_API_CALL clRetainCommandQueue(cl_command_queue command_queue) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_RetainCommandQueue(command_queue);
+}
+
+cl_int CL_API_CALL clReleaseCommandQueue(cl_command_queue command_queue) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_ReleaseCommandQueue(command_queue);
+}
+
+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) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_GetCommandQueueInfo(command_queue, param_name, param_value_size, param_value, param_value_size_ret);
+}
+
+cl_mem CL_API_CALL clCreateBuffer(cl_context context, cl_mem_flags flags, size_t size, void* host_ptr, cl_int* errcode_ret) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_CreateBuffer(context, flags, size, host_ptr, errcode_ret);
+}
+
+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) CL_API_SUFFIX__VERSION_1_1
+{
+	return cl_CreateSubBuffer(buffer, flags, buffer_create_type, buffer_create_info, errcode_ret);
+}
+
+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) CL_API_SUFFIX__VERSION_1_2
+{
+	return cl_CreateImage(context, flags, image_format, image_desc, host_ptr, errcode_ret);
+}
+
+cl_mem CL_API_CALL clCreatePipe(cl_context context, cl_mem_flags flags, cl_uint pipe_packet_size, cl_uint pipe_max_packets, const cl_pipe_properties* properties, cl_int* errcode_ret) CL_API_SUFFIX__VERSION_2_0
+{
+	return cl_CreatePipe(context, flags, pipe_packet_size, pipe_max_packets, properties, errcode_ret);
+}
+
+cl_int CL_API_CALL clRetainMemObject(cl_mem memobj) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_RetainMemObject(memobj);
+}
+
+cl_int CL_API_CALL clReleaseMemObject(cl_mem memobj) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_ReleaseMemObject(memobj);
+}
+
+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) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_GetSupportedImageFormats(context, flags, image_type, num_entries, image_formats, num_image_formats);
+}
+
+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) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_GetMemObjectInfo(memobj, param_name, param_value_size, param_value, param_value_size_ret);
+}
+
+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) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_GetImageInfo(image, param_name, param_value_size, param_value, param_value_size_ret);
+}
+
+cl_int CL_API_CALL clGetPipeInfo(cl_mem pipe, cl_pipe_info param_name, size_t param_value_size, void* param_value, size_t* param_value_size_ret) CL_API_SUFFIX__VERSION_2_0
+{
+	return cl_GetPipeInfo(pipe, param_name, param_value_size, param_value, param_value_size_ret);
+}
+
+cl_int CL_API_CALL clSetMemObjectDestructorCallback(cl_mem memobj, void (CL_CALLBACK* pfn_notify)(cl_mem memobj, void* user_data), void* user_data) CL_API_SUFFIX__VERSION_1_1
+{
+	return cl_SetMemObjectDestructorCallback(memobj, pfn_notify, user_data);
+}
+
+void* CL_API_CALL clSVMAlloc(cl_context context, cl_svm_mem_flags flags, size_t size, cl_uint alignment) CL_API_SUFFIX__VERSION_2_0
+{
+	return cl_SVMAlloc(context, flags, size, alignment);
+}
+
+void CL_API_CALL clSVMFree(cl_context context, void* svm_pointer) CL_API_SUFFIX__VERSION_2_0
+{
+	cl_SVMFree(context, svm_pointer);
+}
+
+cl_sampler CL_API_CALL clCreateSamplerWithProperties(cl_context context, const cl_sampler_properties* normalized_coords, cl_int* errcode_ret) CL_API_SUFFIX__VERSION_2_0
+{
+	return cl_CreateSamplerWithProperties(context, normalized_coords, errcode_ret);
+}
+
+cl_int CL_API_CALL clRetainSampler(cl_sampler sampler) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_RetainSampler(sampler);
+}
+
+cl_int CL_API_CALL clReleaseSampler(cl_sampler sampler) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_ReleaseSampler(sampler);
+}
+
+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) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_GetSamplerInfo(sampler, param_name, param_value_size, param_value, param_value_size_ret);
+}
+
+cl_program CL_API_CALL clCreateProgramWithSource(cl_context context, cl_uint count, const char** strings, const size_t* lengths, cl_int* errcode_ret) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_CreateProgramWithSource(context, count, strings, lengths, errcode_ret);
+}
+
+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) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_CreateProgramWithBinary(context, num_devices, device_list, lengths, binaries, binary_status, errcode_ret);
+}
+
+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) CL_API_SUFFIX__VERSION_1_2
+{
+	return cl_CreateProgramWithBuiltInKernels(context, num_devices, device_list, kernel_names, errcode_ret);
+}
+
+cl_int CL_API_CALL clRetainProgram(cl_program program) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_RetainProgram(program);
+}
+
+cl_int CL_API_CALL clReleaseProgram(cl_program program) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_ReleaseProgram(program);
+}
+
+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) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_BuildProgram(program, num_devices, device_list, options, pfn_notify, user_data);
+}
+
+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) CL_API_SUFFIX__VERSION_1_2
+{
+	return cl_CompileProgram(program, num_devices, device_list, options, num_input_headers, input_headers, header_include_names, pfn_notify, user_data);
+}
+
+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) CL_API_SUFFIX__VERSION_1_2
+{
+	return cl_LinkProgram(context, num_devices, device_list, options, num_input_programs, input_programs, pfn_notify, user_data, errcode_ret);
+}
+
+cl_int CL_API_CALL clUnloadPlatformCompiler(cl_platform_id platform) CL_API_SUFFIX__VERSION_1_2
+{
+	return cl_UnloadPlatformCompiler(platform);
+}
+
+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) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_GetProgramInfo(program, param_name, param_value_size, param_value, param_value_size_ret);
+}
+
+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) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_GetProgramBuildInfo(program, device, param_name, param_value_size, param_value, param_value_size_ret);
+}
+
+cl_kernel CL_API_CALL clCreateKernel(cl_program program, const char* kernel_name, cl_int* errcode_ret) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_CreateKernel(program, kernel_name, errcode_ret);
+}
+
+cl_int CL_API_CALL clCreateKernelsInProgram(cl_program program, cl_uint num_kernels, cl_kernel* kernels, cl_uint* num_kernels_ret) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_CreateKernelsInProgram(program, num_kernels, kernels, num_kernels_ret);
+}
+
+cl_int CL_API_CALL clRetainKernel(cl_kernel kernel) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_RetainKernel(kernel);
+}
+
+cl_int CL_API_CALL clReleaseKernel(cl_kernel kernel) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_ReleaseKernel(kernel);
+}
+
+cl_int CL_API_CALL clSetKernelArg(cl_kernel kernel, cl_uint arg_index, size_t arg_size, const void* arg_value) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_SetKernelArg(kernel, arg_index, arg_size, arg_value);
+}
+
+cl_int CL_API_CALL clSetKernelArgSVMPointer(cl_kernel kernel, cl_uint arg_index, const void* arg_value) CL_API_SUFFIX__VERSION_2_0
+{
+	return cl_SetKernelArgSVMPointer(kernel, arg_index, arg_value);
+}
+
+cl_int CL_API_CALL clSetKernelExecInfo(cl_kernel kernel, cl_kernel_exec_info param_name, size_t param_value_size, const void* param_value) CL_API_SUFFIX__VERSION_2_0
+{
+	return cl_SetKernelExecInfo(kernel, param_name, param_value_size, param_value);
+}
+
+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) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_GetKernelInfo(kernel, param_name, param_value_size, param_value, param_value_size_ret);
+}
+
+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) CL_API_SUFFIX__VERSION_1_2
+{
+	return cl_GetKernelArgInfo(kernel, arg_indx, param_name, param_value_size, param_value, param_value_size_ret);
+}
+
+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) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_GetKernelWorkGroupInfo(kernel, device, param_name, param_value_size, param_value, param_value_size_ret);
+}
+
+cl_int CL_API_CALL clWaitForEvents(cl_uint num_events, const cl_event* event_list) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_WaitForEvents(num_events, event_list);
+}
+
+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) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_GetEventInfo(event, param_name, param_value_size, param_value, param_value_size_ret);
+}
+
+cl_event CL_API_CALL clCreateUserEvent(cl_context context, cl_int* errcode_ret) CL_API_SUFFIX__VERSION_1_1
+{
+	return cl_CreateUserEvent(context, errcode_ret);
+}
+
+cl_int CL_API_CALL clRetainEvent(cl_event event) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_RetainEvent(event);
+}
+
+cl_int CL_API_CALL clReleaseEvent(cl_event event) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_ReleaseEvent(event);
+}
+
+cl_int CL_API_CALL clSetUserEventStatus(cl_event event, cl_int execution_status) CL_API_SUFFIX__VERSION_1_1
+{
+	return cl_SetUserEventStatus(event, execution_status);
+}
+
+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) CL_API_SUFFIX__VERSION_1_1
+{
+	return cl_SetEventCallback(event, command_exec_callback_type, pfn_notify, user_data);
+}
+
+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) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_GetEventProfilingInfo(event, param_name, param_value_size, param_value, param_value_size_ret);
+}
+
+cl_int CL_API_CALL clFlush(cl_command_queue command_queue) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_Flush(command_queue);
+}
+
+cl_int CL_API_CALL clFinish(cl_command_queue command_queue) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_Finish(command_queue);
+}
+
+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) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_EnqueueReadBuffer(command_queue, buffer, blocking_read, offset, size, ptr, num_events_in_wait_list, event_wait_list, event);
+}
+
+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) CL_API_SUFFIX__VERSION_1_1
+{
+	return cl_EnqueueReadBufferRect(command_queue, buffer, blocking_read, buffer_offset, host_offset, region, buffer_row_pitch, buffer_slice_pitch, host_row_pitch, host_slice_pitch, ptr, num_events_in_wait_list, event_wait_list, event);
+}
+
+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) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_EnqueueWriteBuffer(command_queue, buffer, blocking_write, offset, size, ptr, num_events_in_wait_list, event_wait_list, event);
+}
+
+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) CL_API_SUFFIX__VERSION_1_1
+{
+	return cl_EnqueueWriteBufferRect(command_queue, buffer, blocking_write, buffer_offset, host_offset, region, buffer_row_pitch, buffer_slice_pitch, host_row_pitch, host_slice_pitch, ptr, num_events_in_wait_list, event_wait_list, event);
+}
+
+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) CL_API_SUFFIX__VERSION_1_2
+{
+	return cl_EnqueueFillBuffer(command_queue, buffer, pattern, pattern_size, offset, size, num_events_in_wait_list, event_wait_list, event);
+}
+
+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) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_EnqueueCopyBuffer(command_queue, src_buffer, dst_buffer, src_offset, dst_offset, size, num_events_in_wait_list, event_wait_list, event);
+}
+
+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) CL_API_SUFFIX__VERSION_1_1
+{
+	return cl_EnqueueCopyBufferRect(command_queue, src_buffer, dst_buffer, src_origin, dst_origin, region, src_row_pitch, src_slice_pitch, dst_row_pitch, dst_slice_pitch, num_events_in_wait_list, event_wait_list, event);
+}
+
+cl_int CL_API_CALL clEnqueueReadImage(cl_command_queue command_queue, cl_mem image, cl_bool blocking_read, const size_t* origin, const size_t* region, 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) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_EnqueueReadImage(command_queue, image, blocking_read, origin, region, row_pitch, slice_pitch, ptr, num_events_in_wait_list, event_wait_list, event);
+}
+
+cl_int CL_API_CALL clEnqueueWriteImage(cl_command_queue command_queue, cl_mem image, cl_bool blocking_write, const size_t* origin, const size_t* region, 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) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_EnqueueWriteImage(command_queue, image, blocking_write, origin, region, input_row_pitch, input_slice_pitch, ptr, num_events_in_wait_list, event_wait_list, event);
+}
+
+cl_int CL_API_CALL clEnqueueFillImage(cl_command_queue command_queue, cl_mem image, const void* fill_color, const size_t* origin, const size_t* region, cl_uint num_events_in_wait_list, const cl_event* event_wait_list, cl_event* event) CL_API_SUFFIX__VERSION_1_2
+{
+	return cl_EnqueueFillImage(command_queue, image, fill_color, origin, region, num_events_in_wait_list, event_wait_list, event);
+}
+
+cl_int CL_API_CALL clEnqueueCopyImage(cl_command_queue command_queue, cl_mem src_image, cl_mem dst_image, const size_t* src_origin, const size_t* dst_origin, const size_t* region, cl_uint num_events_in_wait_list, const cl_event* event_wait_list, cl_event* event) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_EnqueueCopyImage(command_queue, src_image, dst_image, src_origin, dst_origin, region, num_events_in_wait_list, event_wait_list, event);
+}
+
+cl_int CL_API_CALL clEnqueueCopyImageToBuffer(cl_command_queue command_queue, cl_mem src_image, cl_mem dst_buffer, const size_t* src_origin, const size_t* region, size_t dst_offset, cl_uint num_events_in_wait_list, const cl_event* event_wait_list, cl_event* event) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_EnqueueCopyImageToBuffer(command_queue, src_image, dst_buffer, src_origin, region, dst_offset, num_events_in_wait_list, event_wait_list, event);
+}
+
+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, const size_t* region, cl_uint num_events_in_wait_list, const cl_event* event_wait_list, cl_event* event) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_EnqueueCopyBufferToImage(command_queue, src_buffer, dst_image, src_offset, dst_origin, region, num_events_in_wait_list, event_wait_list, event);
+}
+
+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) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_EnqueueMapBuffer(command_queue, buffer, blocking_map, map_flags, offset, size, num_events_in_wait_list, event_wait_list, event, errcode_ret);
+}
+
+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, const size_t* region, 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) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_EnqueueMapImage(command_queue, image, blocking_map, map_flags, origin, region, image_row_pitch, image_slice_pitch, num_events_in_wait_list, event_wait_list, event, errcode_ret);
+}
+
+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) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_EnqueueUnmapMemObject(command_queue, memobj, mapped_ptr, num_events_in_wait_list, event_wait_list, event);
+}
+
+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) CL_API_SUFFIX__VERSION_1_2
+{
+	return cl_EnqueueMigrateMemObjects(command_queue, num_mem_objects, mem_objects, flags, num_events_in_wait_list, event_wait_list, event);
+}
+
+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) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_EnqueueNDRangeKernel(command_queue, kernel, work_dim, global_work_offset, global_work_size, local_work_size, num_events_in_wait_list, event_wait_list, event);
+}
+
+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) CL_API_SUFFIX__VERSION_1_0
+{
+	return cl_EnqueueNativeKernel(command_queue, user_func, args, cb_args, num_mem_objects, mem_list, args_mem_loc, num_events_in_wait_list, event_wait_list, event);
+}
+
+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) CL_API_SUFFIX__VERSION_1_2
+{
+	return cl_EnqueueMarkerWithWaitList(command_queue, num_events_in_wait_list, event_wait_list, event);
+}
+
+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) CL_API_SUFFIX__VERSION_1_2
+{
+	return cl_EnqueueBarrierWithWaitList(command_queue, num_events_in_wait_list, event_wait_list, event);
+}
+
+cl_int CL_API_CALL clEnqueueSVMFree(cl_command_queue command_queue, cl_uint num_svm_pointers, void* svm_pointers[], void (CL_CALLBACK* pfn_free_func)(cl_command_queue queue, cl_uint num_svm_pointers, void* svm_pointers[], void* user_data), void* user_data, cl_uint num_events_in_wait_list, const cl_event* event_wait_list, cl_event* event) CL_API_SUFFIX__VERSION_2_0
+{
+	return cl_EnqueueSVMFree(command_queue, num_svm_pointers, svm_pointers, pfn_free_func, user_data, num_events_in_wait_list, event_wait_list, event);
+}
+
+cl_int CL_API_CALL clEnqueueSVMMemcpy(cl_command_queue command_queue, cl_bool blocking_copy, void* dst_ptr, const void* src_ptr, size_t size, cl_uint num_events_in_wait_list, const cl_event* event_wait_list, cl_event* event) CL_API_SUFFIX__VERSION_2_0
+{
+	return cl_EnqueueSVMMemcpy(command_queue, blocking_copy, dst_ptr, src_ptr, size, num_events_in_wait_list, event_wait_list, event);
+}
+
+cl_int CL_API_CALL clEnqueueSVMMemFill(cl_command_queue command_queue, void* svm_ptr, const void* pattern, size_t pattern_size, size_t size, cl_uint num_events_in_wait_list, const cl_event* event_wait_list, cl_event* event) CL_API_SUFFIX__VERSION_2_0
+{
+	return cl_EnqueueSVMMemFill(command_queue, svm_ptr, pattern, pattern_size, size, num_events_in_wait_list, event_wait_list, event);
+}
+
+cl_int CL_API_CALL clEnqueueSVMMap(cl_command_queue command_queue, cl_bool blocking_map, cl_map_flags flags, void* svm_ptr, size_t size, cl_uint num_events_in_wait_list, const cl_event* event_wait_list, cl_event* event) CL_API_SUFFIX__VERSION_2_0
+{
+	return cl_EnqueueSVMMap(command_queue, blocking_map, flags, svm_ptr, size, num_events_in_wait_list, event_wait_list, event);
+}
+
+cl_int CL_API_CALL clEnqueueSVMUnmap(cl_command_queue command_queue, void* svm_ptr, cl_uint num_events_in_wait_list, const cl_event* event_wait_list, cl_event* event) CL_API_SUFFIX__VERSION_2_0
+{
+	return cl_EnqueueSVMUnmap(command_queue, svm_ptr, num_events_in_wait_list, event_wait_list, event);
+}
+
+void* CL_API_CALL clGetExtensionFunctionAddressForPlatform(cl_platform_id platform, const char* func_name) CL_API_SUFFIX__VERSION_1_2
+{
+	return cl_GetExtensionFunctionAddressForPlatform(platform, func_name);
+}
+
+cl_mem CL_API_CALL clCreateImage2D(cl_context context, cl_mem_flags flags, const cl_image_format* image_format, size_t image_width, size_t image_height, size_t image_row_pitch, void* host_ptr, cl_int* errcode_ret)
+{
+	return cl_CreateImage2D(context, flags, image_format, image_width, image_height, image_row_pitch, host_ptr, errcode_ret);
+}
+
+cl_mem CL_API_CALL clCreateImage3D(cl_context context, cl_mem_flags flags, const cl_image_format* image_format, size_t image_width, size_t image_height, size_t image_depth, size_t image_row_pitch, size_t image_slice_pitch, void* host_ptr, cl_int* errcode_ret)
+{
+	return cl_CreateImage3D(context, flags, image_format, image_width, image_height, image_depth, image_row_pitch, image_slice_pitch, host_ptr, errcode_ret);
+}
+
+cl_int CL_API_CALL clEnqueueMarker(cl_command_queue command_queue, cl_event* event)
+{
+	return cl_EnqueueMarker(command_queue, event);
+}
+
+cl_int CL_API_CALL clEnqueueWaitForEvents(cl_command_queue command_queue, cl_uint num_events, const cl_event* event_list)
+{
+	return cl_EnqueueWaitForEvents(command_queue, num_events, event_list);
+}
+
+cl_int CL_API_CALL clEnqueueBarrier(cl_command_queue command_queue)
+{
+	return cl_EnqueueBarrier(command_queue);
+}
+
+cl_int CL_API_CALL clUnloadCompiler(void)
+{
+	return cl_UnloadCompiler();
+}
+
+void* CL_API_CALL clGetExtensionFunctionAddress(const char* func_name)
+{
+	return cl_GetExtensionFunctionAddress(func_name);
+}
+
+cl_command_queue CL_API_CALL clCreateCommandQueue(cl_context context, cl_device_id device, cl_command_queue_properties properties, cl_int* errcode_ret)
+{
+	return cl_CreateCommandQueue(context, device, properties, errcode_ret);
+}
+
+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)
+{
+	return cl_CreateSampler(context, normalized_coords, addressing_mode, filter_mode, errcode_ret);
+}
+
+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)
+{
+	return cl_EnqueueTask(command_queue, kernel, num_events_in_wait_list, event_wait_list, event);
+}

3rdparty/opencl/opencl.vcxproj

@@ -0,0 +1,160 @@
+<?xml version="1.0" encoding="utf-8"?>
+<Project DefaultTargets="Build" ToolsVersion="12.0" xmlns="http://schemas.microsoft.com/developer/msbuild/2003">
+  <ItemGroup Label="ProjectConfigurations">
+    <ProjectConfiguration Include="Debug|Win32">
+      <Configuration>Debug</Configuration>
+      <Platform>Win32</Platform>
+    </ProjectConfiguration>
+    <ProjectConfiguration Include="Debug|x64">
+      <Configuration>Debug</Configuration>
+      <Platform>x64</Platform>
+    </ProjectConfiguration>
+    <ProjectConfiguration Include="Release|Win32">
+      <Configuration>Release</Configuration>
+      <Platform>Win32</Platform>
+    </ProjectConfiguration>
+    <ProjectConfiguration Include="Release|x64">
+      <Configuration>Release</Configuration>
+      <Platform>x64</Platform>
+    </ProjectConfiguration>
+  </ItemGroup>
+  <PropertyGroup Label="Globals">
+    <ProjectGuid>{D80D4A75-C385-41BD-AE62-83D2E2B595A7}</ProjectGuid>
+    <Keyword>Win32Proj</Keyword>
+    <RootNamespace>opencl</RootNamespace>
+  </PropertyGroup>
+  <Import Project="$(VCTargetsPath)\Microsoft.Cpp.Default.props" />
+  <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'" Label="Configuration">
+    <ConfigurationType>StaticLibrary</ConfigurationType>
+    <UseDebugLibraries>true</UseDebugLibraries>
+    <PlatformToolset>v120</PlatformToolset>
+    <CharacterSet>MultiByte</CharacterSet>
+  </PropertyGroup>
+  <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Debug|x64'" Label="Configuration">
+    <ConfigurationType>StaticLibrary</ConfigurationType>
+    <UseDebugLibraries>true</UseDebugLibraries>
+    <PlatformToolset>v120</PlatformToolset>
+    <CharacterSet>MultiByte</CharacterSet>
+  </PropertyGroup>
+  <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Release|Win32'" Label="Configuration">
+    <ConfigurationType>StaticLibrary</ConfigurationType>
+    <UseDebugLibraries>false</UseDebugLibraries>
+    <PlatformToolset>v120</PlatformToolset>
+    <WholeProgramOptimization>true</WholeProgramOptimization>
+    <CharacterSet>MultiByte</CharacterSet>
+  </PropertyGroup>
+  <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Release|x64'" Label="Configuration">
+    <ConfigurationType>StaticLibrary</ConfigurationType>
+    <UseDebugLibraries>false</UseDebugLibraries>
+    <PlatformToolset>v120</PlatformToolset>
+    <WholeProgramOptimization>true</WholeProgramOptimization>
+    <CharacterSet>MultiByte</CharacterSet>
+  </PropertyGroup>
+  <Import Project="$(VCTargetsPath)\Microsoft.Cpp.props" />
+  <ImportGroup Label="ExtensionSettings">
+  </ImportGroup>
+  <ImportGroup Label="PropertySheets" Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'">
+    <Import Project="$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props" Condition="exists('$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props')" Label="LocalAppDataPlatform" />
+    <Import Project="..\DefaultProjectRootDir.props" />
+    <Import Project="..\3rdparty.props" />
+  </ImportGroup>
+  <ImportGroup Condition="'$(Configuration)|$(Platform)'=='Debug|x64'" Label="PropertySheets">
+    <Import Project="$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props" Condition="exists('$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props')" Label="LocalAppDataPlatform" />
+    <Import Project="..\DefaultProjectRootDir.props" />
+    <Import Project="..\3rdparty.props" />
+  </ImportGroup>
+  <ImportGroup Label="PropertySheets" Condition="'$(Configuration)|$(Platform)'=='Release|Win32'">
+    <Import Project="$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props" Condition="exists('$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props')" Label="LocalAppDataPlatform" />
+    <Import Project="..\DefaultProjectRootDir.props" />
+    <Import Project="..\3rdparty.props" />
+  </ImportGroup>
+  <ImportGroup Condition="'$(Configuration)|$(Platform)'=='Release|x64'" Label="PropertySheets">
+    <Import Project="$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props" Condition="exists('$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props')" Label="LocalAppDataPlatform" />
+    <Import Project="..\DefaultProjectRootDir.props" />
+    <Import Project="..\3rdparty.props" />
+  </ImportGroup>
+  <PropertyGroup Label="UserMacros" />
+  <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Debug|x64'" />
+  <PropertyGroup Condition="'$(Configuration)|$(Platform)'=='Release|x64'" />
+  <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Debug|Win32'">
+    <ClCompile>
+      <PrecompiledHeader>
+      </PrecompiledHeader>
+      <WarningLevel>Level3</WarningLevel>
+      <Optimization>Disabled</Optimization>
+      <PreprocessorDefinitions>WIN32;_DEBUG;_LIB;%(PreprocessorDefinitions)</PreprocessorDefinitions>
+    </ClCompile>
+    <Link>
+      <SubSystem>Windows</SubSystem>
+      <GenerateDebugInformation>true</GenerateDebugInformation>
+    </Link>
+    <Lib />
+  </ItemDefinitionGroup>
+  <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Debug|x64'">
+    <ClCompile>
+      <PrecompiledHeader>
+      </PrecompiledHeader>
+      <WarningLevel>Level3</WarningLevel>
+      <Optimization>Disabled</Optimization>
+      <PreprocessorDefinitions>WIN32;_DEBUG;_LIB;%(PreprocessorDefinitions)</PreprocessorDefinitions>
+    </ClCompile>
+    <Link>
+      <SubSystem>Windows</SubSystem>
+      <GenerateDebugInformation>true</GenerateDebugInformation>
+    </Link>
+  </ItemDefinitionGroup>
+  <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Release|Win32'">
+    <ClCompile>
+      <WarningLevel>Level3</WarningLevel>
+      <PrecompiledHeader>
+      </PrecompiledHeader>
+      <Optimization>MaxSpeed</Optimization>
+      <FunctionLevelLinking>true</FunctionLevelLinking>
+      <IntrinsicFunctions>true</IntrinsicFunctions>
+      <PreprocessorDefinitions>WIN32;NDEBUG;_LIB;%(PreprocessorDefinitions)</PreprocessorDefinitions>
+    </ClCompile>
+    <Link>
+      <SubSystem>Windows</SubSystem>
+      <GenerateDebugInformation>true</GenerateDebugInformation>
+      <EnableCOMDATFolding>true</EnableCOMDATFolding>
+      <OptimizeReferences>true</OptimizeReferences>
+    </Link>
+    <Lib />
+  </ItemDefinitionGroup>
+  <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Release|x64'">
+    <ClCompile>
+      <WarningLevel>Level3</WarningLevel>
+      <PrecompiledHeader>
+      </PrecompiledHeader>
+      <Optimization>MaxSpeed</Optimization>
+      <FunctionLevelLinking>true</FunctionLevelLinking>
+      <IntrinsicFunctions>true</IntrinsicFunctions>
+      <PreprocessorDefinitions>WIN32;NDEBUG;_LIB;%(PreprocessorDefinitions)</PreprocessorDefinitions>
+    </ClCompile>
+    <Link>
+      <SubSystem>Windows</SubSystem>
+      <GenerateDebugInformation>true</GenerateDebugInformation>
+      <EnableCOMDATFolding>true</EnableCOMDATFolding>
+      <OptimizeReferences>true</OptimizeReferences>
+    </Link>
+  </ItemDefinitionGroup>
+  <ItemGroup>
+    <ClInclude Include="CL\cl.h" />
+    <ClInclude Include="CL\cl.hpp" />
+    <ClInclude Include="CL\cl_d3d10.h" />
+    <ClInclude Include="CL\cl_d3d11.h" />
+    <ClInclude Include="CL\cl_d3d9.h" />
+    <ClInclude Include="CL\cl_egl.h" />
+    <ClInclude Include="CL\cl_ext.h" />
+    <ClInclude Include="CL\cl_gl.h" />
+    <ClInclude Include="CL\cl_gl_ext.h" />
+    <ClInclude Include="CL\cl_platform.h" />
+    <ClInclude Include="CL\opencl.h" />
+  </ItemGroup>
+  <ItemGroup>
+    <ClCompile Include="opencl.cpp" />
+  </ItemGroup>
+  <Import Project="$(VCTargetsPath)\Microsoft.Cpp.targets" />
+  <ImportGroup Label="ExtensionTargets">
+  </ImportGroup>
+</Project>

3rdparty/opencl/opencl.vcxproj.filters

@@ -0,0 +1,57 @@
+<?xml version="1.0" encoding="utf-8"?>
+<Project ToolsVersion="4.0" xmlns="http://schemas.microsoft.com/developer/msbuild/2003">
+  <ItemGroup>
+    <Filter Include="Source Files">
+      <UniqueIdentifier>{4FC737F1-C7A5-4376-A066-2A32D752A2FF}</UniqueIdentifier>
+      <Extensions>cpp;c;cc;cxx;def;odl;idl;hpj;bat;asm;asmx</Extensions>
+    </Filter>
+    <Filter Include="Header Files">
+      <UniqueIdentifier>{93995380-89BD-4b04-88EB-625FBE52EBFB}</UniqueIdentifier>
+      <Extensions>h;hh;hpp;hxx;hm;inl;inc;xsd</Extensions>
+    </Filter>
+    <Filter Include="Resource Files">
+      <UniqueIdentifier>{67DA6AB6-F800-4c08-8B7A-83BB121AAD01}</UniqueIdentifier>
+      <Extensions>rc;ico;cur;bmp;dlg;rc2;rct;bin;rgs;gif;jpg;jpeg;jpe;resx;tiff;tif;png;wav;mfcribbon-ms</Extensions>
+    </Filter>
+  </ItemGroup>
+  <ItemGroup>
+    <ClInclude Include="CL\cl.h">
+      <Filter>Header Files</Filter>
+    </ClInclude>
+    <ClInclude Include="CL\cl.hpp">
+      <Filter>Header Files</Filter>
+    </ClInclude>
+    <ClInclude Include="CL\cl_d3d10.h">
+      <Filter>Header Files</Filter>
+    </ClInclude>
+    <ClInclude Include="CL\cl_d3d11.h">
+      <Filter>Header Files</Filter>
+    </ClInclude>
+    <ClInclude Include="CL\cl_d3d9.h">
+      <Filter>Header Files</Filter>
+    </ClInclude>
+    <ClInclude Include="CL\cl_egl.h">
+      <Filter>Header Files</Filter>
+    </ClInclude>
+    <ClInclude Include="CL\cl_ext.h">
+      <Filter>Header Files</Filter>
+    </ClInclude>
+    <ClInclude Include="CL\cl_gl.h">
+      <Filter>Header Files</Filter>
+    </ClInclude>
+    <ClInclude Include="CL\cl_gl_ext.h">
+      <Filter>Header Files</Filter>
+    </ClInclude>
+    <ClInclude Include="CL\cl_platform.h">
+      <Filter>Header Files</Filter>
+    </ClInclude>
+    <ClInclude Include="CL\opencl.h">
+      <Filter>Header Files</Filter>
+    </ClInclude>
+  </ItemGroup>
+  <ItemGroup>
+    <ClCompile Include="opencl.cpp">
+      <Filter>Source Files</Filter>
+    </ClCompile>
+  </ItemGroup>
+</Project>

pcsx2_suite_2013.sln

@@ -1,7 +1,7 @@
 
 Microsoft Visual Studio Solution File, Format Version 12.00
 # Visual Studio 2013
-VisualStudioVersion = 12.0.30723.0
+VisualStudioVersion = 12.0.30825.0
 MinimumVisualStudioVersion = 10.0.40219.1
 Project("{2150E333-8FDC-42A3-9474-1A3956D46DE8}") = "Nulls", "Nulls", "{E1828E40-2FBB-48FE-AE7F-5587755DCE0E}"
 EndProject
@@ -40,6 +40,9 @@ EndProject
 Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "SPU2-X", "plugins\spu2-x\src\Windows\Spu2-X_vs2013.vcxproj", "{5307BBB7-EBB9-4AA4-8CB6-A94EC473C8C4}"
 EndProject
 Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "GSdx", "plugins\GSdx\GSdx_vs2013.vcxproj", "{18E42F6F-3A62-41EE-B42F-79366C4F1E95}"
+	ProjectSection(ProjectDependencies) = postProject
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7} = {D80D4A75-C385-41BD-AE62-83D2E2B595A7}
+	EndProjectSection
 EndProject
 Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "SoundTouch", "3rdparty\soundtouch\SoundTouch_vs2013.vcxproj", "{E9B51944-7E6D-4BCD-83F2-7BBD5A46182D}"
 EndProject
@@ -91,6 +94,8 @@ Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "DEV9ghzdrk", "plugins\dev9g
 EndProject
 Project("{2150E333-8FDC-42A3-9474-1A3956D46DE8}") = "Experimental", "Experimental", "{7A407562-D70F-4F0A-9D3E-B32506416003}"
 EndProject
+Project("{8BC9CEB8-8B4A-11D0-8D11-00A0C91BC942}") = "opencl", "3rdparty\opencl\opencl.vcxproj", "{D80D4A75-C385-41BD-AE62-83D2E2B595A7}"
+EndProject
 Global
 	GlobalSection(SolutionConfigurationPlatforms) = preSolution
 		Debug AVX|Win32 = Debug AVX|Win32
@@ -1154,6 +1159,58 @@ Global
 		{BBE4E5FB-530A-4D18-A633-35AF0577B7F3}.Release|Win32.ActiveCfg = Release|Win32
 		{BBE4E5FB-530A-4D18-A633-35AF0577B7F3}.Release|Win32.Build.0 = Release|Win32
 		{BBE4E5FB-530A-4D18-A633-35AF0577B7F3}.Release|x64.ActiveCfg = Release|Win32
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Debug AVX|Win32.ActiveCfg = Debug|Win32
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Debug AVX|Win32.Build.0 = Debug|Win32
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Debug AVX|x64.ActiveCfg = Debug|x64
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Debug AVX|x64.Build.0 = Debug|x64
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Debug AVX2|Win32.ActiveCfg = Debug|Win32
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Debug AVX2|Win32.Build.0 = Debug|Win32
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Debug AVX2|x64.ActiveCfg = Debug|x64
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Debug AVX2|x64.Build.0 = Debug|x64
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Debug SSE2|Win32.ActiveCfg = Debug|Win32
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Debug SSE2|Win32.Build.0 = Debug|Win32
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Debug SSE2|x64.ActiveCfg = Debug|x64
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Debug SSE2|x64.Build.0 = Debug|x64
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Debug SSE4|Win32.ActiveCfg = Debug|Win32
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Debug SSE4|Win32.Build.0 = Debug|Win32
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Debug SSE4|x64.ActiveCfg = Debug|x64
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Debug SSE4|x64.Build.0 = Debug|x64
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Debug SSSE3|Win32.ActiveCfg = Debug|Win32
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Debug SSSE3|Win32.Build.0 = Debug|Win32
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Debug SSSE3|x64.ActiveCfg = Debug|x64
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Debug SSSE3|x64.Build.0 = Debug|x64
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Debug|Win32.ActiveCfg = Debug|Win32
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Debug|Win32.Build.0 = Debug|Win32
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Debug|x64.ActiveCfg = Debug|x64
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Debug|x64.Build.0 = Debug|x64
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Devel|Win32.ActiveCfg = Debug|Win32
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Devel|Win32.Build.0 = Debug|Win32
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Devel|x64.ActiveCfg = Debug|x64
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Devel|x64.Build.0 = Debug|x64
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Release AVX|Win32.ActiveCfg = Release|Win32
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Release AVX|Win32.Build.0 = Release|Win32
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Release AVX|x64.ActiveCfg = Release|x64
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Release AVX|x64.Build.0 = Release|x64
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Release AVX2|Win32.ActiveCfg = Release|Win32
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Release AVX2|Win32.Build.0 = Release|Win32
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Release AVX2|x64.ActiveCfg = Release|x64
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Release AVX2|x64.Build.0 = Release|x64
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Release SSE2|Win32.ActiveCfg = Release|Win32
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Release SSE2|Win32.Build.0 = Release|Win32
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Release SSE2|x64.ActiveCfg = Release|x64
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Release SSE2|x64.Build.0 = Release|x64
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Release SSE4|Win32.ActiveCfg = Release|Win32
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Release SSE4|Win32.Build.0 = Release|Win32
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Release SSE4|x64.ActiveCfg = Release|x64
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Release SSE4|x64.Build.0 = Release|x64
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Release SSSE3|Win32.ActiveCfg = Release|Win32
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Release SSSE3|Win32.Build.0 = Release|Win32
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Release SSSE3|x64.ActiveCfg = Release|x64
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Release SSSE3|x64.Build.0 = Release|x64
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Release|Win32.ActiveCfg = Release|Win32
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Release|Win32.Build.0 = Release|Win32
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Release|x64.ActiveCfg = Release|x64
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7}.Release|x64.Build.0 = Release|x64
 	EndGlobalSection
 	GlobalSection(SolutionProperties) = preSolution
 		HideSolutionNode = FALSE
@@ -1189,5 +1246,6 @@ Global
 		{E613DA9F-41B4-4613-9911-E418EF5533BC} = {7A407562-D70F-4F0A-9D3E-B32506416003}
 		{BBE4E5FB-530A-4D18-A633-35AF0577B7F3} = {7A407562-D70F-4F0A-9D3E-B32506416003}
 		{7A407562-D70F-4F0A-9D3E-B32506416003} = {703FD00B-D7A0-41E3-BD03-CEC86B385DAF}
+		{D80D4A75-C385-41BD-AE62-83D2E2B595A7} = {78EBE642-7A4D-4EA7-86BE-5639C6646C38}
 	EndGlobalSection
 EndGlobal

plugins/GSdx/CMakeLists.txt

@@ -104,6 +104,7 @@ set(GSdxSources
     GSPerfMon.cpp
     GSRasterizer.cpp
     GSRenderer.cpp
+    GSRendererCL.cpp
     GSRendererHW.cpp
     GSRendererNull.cpp
     GSRendererOGL.cpp

plugins/GSdx/GPURenderer.h

@@ -123,13 +123,13 @@ protected:
 		int maxcount = std::max<int>(m_maxcount * 3 / 2, 10000);
 		Vertex* vertices = (Vertex*)_aligned_malloc(sizeof(Vertex) * maxcount, 32);
 
-		if (!vertices)
+		if(vertices == NULL)
 		{
 			printf("GSdx: failed to allocate %d bytes for verticles.\n", sizeof(Vertex) * maxcount);
 			throw GSDXError();
 		}
 
-		if (m_vertices != NULL)
+		if(m_vertices != NULL)
 		{
 			memcpy(vertices, m_vertices, sizeof(Vertex) * m_maxcount);
 			_aligned_free(m_vertices);

plugins/GSdx/GS.cpp

@@ -27,6 +27,7 @@
 #include "GSDeviceNull.h"
 #include "GSDeviceOGL.h"
 #include "GSRendererOGL.h"
+#include "GSRendererCL.h"
 
 #ifdef _WINDOWS
 
@@ -203,6 +204,7 @@ static int _GSopen(void** dsp, char* title, int renderer, int threads = -1)
 	}
 
 	GSWnd* wnd[2];
+
 	try
 	{
 		if(s_renderer != renderer)
@@ -216,36 +218,23 @@ static int _GSopen(void** dsp, char* title, int renderer, int threads = -1)
 			s_gs = NULL;
 		}
 
-		if(renderer == 15)
+		switch(renderer)
-		{
-			#ifdef _WINDOWS
-	
-			dev = new GSDevice11();
-
-			if(dev == NULL)
-			{
-				return -1;
-			}
-
-			delete s_gs;
-
-			s_gs = new GSRendererCS();
-
-			s_renderer = renderer;
-
-			#endif
-		}
-		else
-		{
-			switch(renderer / 3)
 		{
 		default:
-				#ifdef _WINDOWS
+#ifdef _WINDOWS
-				case 0: dev = new GSDevice9(); break;
+		case 0: case 1: case 2: case 14:
-				case 1: dev = new GSDevice11(); break;
+			dev = new GSDevice9(); 
-				#endif
+			break;
-				case 3: dev = new GSDeviceNull(); break;
+		case 3: case 4: case 5: case 15:
-				case 4: dev = new GSDeviceOGL(); break;
+			dev = new GSDevice11(); 
+			break;
+#endif
+		case 9: case 10: case 11: case 16:
+			dev = new GSDeviceNull(); 
+			break;
+		case 12: case 13: case 17:
+			dev = new GSDeviceOGL(); 
+			break;
 		}
 
 		if(dev == NULL)
@@ -255,39 +244,48 @@ static int _GSopen(void** dsp, char* title, int renderer, int threads = -1)
 
 		if(s_gs == NULL)
 		{
-				switch(renderer % 3)
-				{
-					default:
-					case 0:
 			switch(renderer)
 			{
 			default:
 #ifdef _WINDOWS
-							case 0: s_gs = (GSRenderer*)new GSRendererDX9(); break;
+			case 0:
-							case 3: s_gs = (GSRenderer*)new GSRendererDX11(); break;
+				s_gs = (GSRenderer*)new GSRendererDX9();
-#endif
-							case 12: s_gs = (GSRenderer*)new GSRendererOGL(); break;
-						}
 				break;
-					case 1:
+			case 3: 
+				s_gs = (GSRenderer*)new GSRendererDX11(); 
+				break;
+#endif
+			case 12: 
+				s_gs = (GSRenderer*)new GSRendererOGL(); 
+				break;
+			case 1: case 4: case 10: case 13:
 				s_gs = new GSRendererSW(threads);
 				break;
-					case 2:
+			case 2: case 5: case 11:
 				s_gs = new GSRendererNull();
 				break;
+			case 14: case 15: case 16: case 17:
+#ifdef ENABLE_OPENCL
+				s_gs = new GSRendererCL();
+#endif
+				break;
 			}
 
 			s_renderer = renderer;
 		}
-		}
 
 		if (s_gs->m_wnd == NULL)
 		{
 #ifdef _WINDOWS
-			if (renderer / 3 == 4)
+			switch(renderer)
+			{
+			case 12: case 13: case 17:
 				s_gs->m_wnd = new GSWndWGL();
-			else
+				break;
+			default:
 				s_gs->m_wnd = new GSWndDX();
+				break;
+			}
 #else
 #ifdef ENABLE_GLES
 			wnd[0] = NULL;
@@ -681,9 +679,11 @@ EXPORT_C GSkeyEvent(GSKeyEventData* e)
 {
 	try
 	{
-		if (gsopen_done)
+		if(gsopen_done)
+		{
 			s_gs->KeyEvent(e);
 		}
+	}
 	catch (GSDXRecoverableError)
 	{
 	}
@@ -1218,15 +1218,11 @@ EXPORT_C GSBenchmark(HWND hwnd, HINSTANCE hinst, LPSTR lpszCmdLine, int nCmdShow
 {
 	::SetPriorityClass(::GetCurrentProcess(), HIGH_PRIORITY_CLASS);
 
-	FILE* file = fopen("c:\\temp1\\log.txt", "a");
+	Console console("GSdx", true);
-
-	fprintf(file, "-------------------------\n\n");
 
 	if(1)
 	{
-		GSLocalMemory * pMem = new GSLocalMemory();
+		GSLocalMemory* mem = new GSLocalMemory();		
-		GSLocalMemory& mem(*pMem);
-		
 
 		static struct {int psm; const char* name;} s_format[] =
 		{
@@ -1258,7 +1254,7 @@ EXPORT_C GSBenchmark(HWND hwnd, HINSTANCE hinst, LPSTR lpszCmdLine, int nCmdShow
 			int w = 1 << tbw;
 			int h = 1 << tbw;
 
-			fprintf(file, "%d x %d\n\n", w, h);
+			printf("%d x %d\n\n", w, h);
 
 			for(size_t i = 0; i < countof(s_format); i++)
 			{
@@ -1308,7 +1304,7 @@ EXPORT_C GSBenchmark(HWND hwnd, HINSTANCE hinst, LPSTR lpszCmdLine, int nCmdShow
 
 				clock_t start, end;
 
-				_ftprintf(file, _T("[%4s] "), s_format[i].name);
+				printf("[%4s] ", s_format[i].name);
 
 				start = clock();
 
@@ -1317,12 +1313,12 @@ EXPORT_C GSBenchmark(HWND hwnd, HINSTANCE hinst, LPSTR lpszCmdLine, int nCmdShow
 					int x = 0;
 					int y = 0;
 
-					(mem.*wi)(x, y, ptr, trlen, BITBLTBUF, TRXPOS, TRXREG);
+					(mem->*wi)(x, y, ptr, trlen, BITBLTBUF, TRXPOS, TRXREG);
 				}
 
 				end = clock();
 
-				fprintf(file, "%6d %6d | ", (int)((float)trlen * n / (end - start) / 1000), (int)((float)(w * h) * n / (end - start) / 1000));
+				printf("%6d %6d | ", (int)((float)trlen * n / (end - start) / 1000), (int)((float)(w * h) * n / (end - start) / 1000));
 
 				start = clock();
 
@@ -1331,25 +1327,25 @@ EXPORT_C GSBenchmark(HWND hwnd, HINSTANCE hinst, LPSTR lpszCmdLine, int nCmdShow
 					int x = 0;
 					int y = 0;
 
-					(mem.*ri)(x, y, ptr, trlen, BITBLTBUF, TRXPOS, TRXREG);
+					(mem->*ri)(x, y, ptr, trlen, BITBLTBUF, TRXPOS, TRXREG);
 				}
 
 				end = clock();
 
-				fprintf(file, "%6d %6d | ", (int)((float)trlen * n / (end - start) / 1000), (int)((float)(w * h) * n / (end - start) / 1000));
+				printf("%6d %6d | ", (int)((float)trlen * n / (end - start) / 1000), (int)((float)(w * h) * n / (end - start) / 1000));
 
-				const GSOffset* o = mem.GetOffset(TEX0.TBP0, TEX0.TBW, TEX0.PSM);
+				const GSOffset* o = mem->GetOffset(TEX0.TBP0, TEX0.TBW, TEX0.PSM);
 
 				start = clock();
 
 				for(int j = 0; j < n; j++)
 				{
-					(mem.*rtx)(o, r, ptr, w * 4, TEXA);
+					(mem->*rtx)(o, r, ptr, w * 4, TEXA);
 				}
 
 				end = clock();
 
-				fprintf(file, "%6d %6d ", (int)((float)len * n / (end - start) / 1000), (int)((float)(w * h) * n / (end - start) / 1000));
+				printf("%6d %6d ", (int)((float)len * n / (end - start) / 1000), (int)((float)(w * h) * n / (end - start) / 1000));
 
 				if(psm.pal > 0)
 				{
@@ -1357,32 +1353,30 @@ EXPORT_C GSBenchmark(HWND hwnd, HINSTANCE hinst, LPSTR lpszCmdLine, int nCmdShow
 
 					for(int j = 0; j < n; j++)
 					{
-						(mem.*rtxP)(o, r, ptr, w, TEXA);
+						(mem->*rtxP)(o, r, ptr, w, TEXA);
 					}
 
 					end = clock();
 
-					fprintf(file, "| %6d %6d ", (int)((float)len * n / (end - start) / 1000), (int)((float)(w * h) * n / (end - start) / 1000));
+					printf("| %6d %6d ", (int)((float)len * n / (end - start) / 1000), (int)((float)(w * h) * n / (end - start) / 1000));
 				}
 
-				fprintf(file, "\n");
+				printf("\n");
-
-				fflush(file);
 			}
 
-			fprintf(file, "\n");
+			printf("\n");
 		}
 
 		_aligned_free(ptr);
-		delete pMem;
+
+		delete mem;
 	}
 
 	//
 
 	if(0)
 	{
-		GSLocalMemory * pMem2 = new GSLocalMemory();
+		GSLocalMemory* mem = new GSLocalMemory();
-		GSLocalMemory& mem2(*pMem2);
 
 		uint8* ptr = (uint8*)_aligned_malloc(1024 * 1024 * 4, 32);
 
@@ -1413,13 +1407,13 @@ EXPORT_C GSBenchmark(HWND hwnd, HINSTANCE hinst, LPSTR lpszCmdLine, int nCmdShow
 		int x = 0;
 		int y = 0;
 
-		(mem2.*wi)(x, y, ptr, trlen, BITBLTBUF, TRXPOS, TRXREG);
+		(mem->*wi)(x, y, ptr, trlen, BITBLTBUF, TRXPOS, TRXREG);
-		delete pMem2;
+
+		delete mem;
 	}
 
 	//
 
-	fclose(file);
 	PostQuitMessage(0);
 }
 

plugins/GSdx/GSDevice9.h

@@ -173,7 +173,6 @@ public: // TODO
 	// Shaders...
 
 	hash_map<uint32, GSVertexShader9 > m_vs;
-	D3DXHANDLE m_vs_params;
 	hash_map<uint32, CComPtr<IDirect3DPixelShader9> > m_ps;
 	hash_map<uint32, Direct3DSamplerState9* > m_ps_ss;
 	hash_map<uint32, Direct3DDepthStencilState9* > m_om_dss;

plugins/GSdx/GSDialog.cpp

@@ -145,6 +145,8 @@ void GSDialog::ComboBoxInit(UINT id, const vector<GSSetting>& settings, uint32 s
 			ComboBoxAppend(id, str.c_str(), (LPARAM)s.id, s.id == selid);
 		}
 	}
+
+	ComboBoxFixDroppedWidth(id);
 }
 
 int GSDialog::ComboBoxAppend(UINT id, const char* str, LPARAM data, bool select)
@@ -178,3 +180,49 @@ bool GSDialog::ComboBoxGetSelData(UINT id, INT_PTR& data)
 
 	return false;
 }
+
+void GSDialog::ComboBoxFixDroppedWidth(UINT id)
+{
+	HWND hWnd = GetDlgItem(m_hWnd, id);
+
+	int count = (int)SendMessage(hWnd, CB_GETCOUNT, 0, 0);
+
+	if(count > 0)
+	{
+		HDC hDC = GetDC(hWnd);
+
+		SelectObject(hDC, (HFONT)SendMessage(hWnd, WM_GETFONT, 0, 0));
+
+		int width = (int)SendMessage(hWnd, CB_GETDROPPEDWIDTH, 0, 0);
+
+		for(int i = 0; i < count; i++)
+		{
+			int len = (int)SendMessage(hWnd, CB_GETLBTEXTLEN, i, 0);
+
+			if(len > 0)
+			{
+				char* buff = new char[len + 1];
+
+				SendMessage(hWnd, CB_GETLBTEXT, i, (LPARAM)buff);
+
+				SIZE size;
+				
+				if(GetTextExtentPoint32(hDC, buff, strlen(buff), &size))
+				{
+					size.cx += 10;
+
+					if(size.cx > width) width = size.cx;
+				}
+
+				delete [] buff;
+			}
+		}
+
+		ReleaseDC(hWnd, hDC);
+
+		if(width > 0)
+		{
+			SendMessage(hWnd, CB_SETDROPPEDWIDTH, width, 0);
+		}
+	}
+}

plugins/GSdx/GSDialog.h

@@ -53,4 +53,5 @@ public:
 	void ComboBoxInit(UINT id, const vector<GSSetting>& settings, uint32 selid, uint32 maxid = ~0);
 	int ComboBoxAppend(UINT id, const char* str, LPARAM data = 0, bool select = false);
 	bool ComboBoxGetSelData(UINT id, INT_PTR& data);
+	void ComboBoxFixDroppedWidth(UINT id);
 };

plugins/GSdx/GSDrawScanline.cpp

@@ -2711,7 +2711,7 @@ bool GSDrawScanline::TestAlpha(T& test, T& fm, T& zm, const T& ga)
 
 	case AFAIL_RGB_ONLY:
 		zm |= t;
-		fm |= t & T::xff000000();
+		fm |= t & T::xff000000(); // fpsm 16 bit => & 0xffff8000?
 		break;
 
 	default:

plugins/GSdx/GSLinuxDialog.cpp

@@ -31,30 +31,29 @@ GtkWidget* CreateRenderComboBox()
 
 	render_combo_box = gtk_combo_box_new_text ();
 
-	for(size_t i = 6; i < theApp.m_gs_renderers.size(); i++)
+	for(auto s = theApp.m_gs_renderers.begin(); s != theApp.m_gs_renderers.end(); s++)
 	{
-		const GSSetting& s = theApp.m_gs_renderers[i];
+		string label = s->name;
 
-		string label = s.name;
+		if(!s->note.empty()) label += format(" (%s)", s->note.c_str());
-
-		if(!s.note.empty()) label += format(" (%s)", s.note.c_str());
 
 		// Add some tags to ease users selection
-		switch (i) {
+		switch (s->id) {
-			// better use opengl instead of SDL
+			// Supported opengl
-			case 6:
+			case 12:
-			case 7:
+			case 13:
-				label += " (removed)";
+			case 17:
 				break;
 
 			// (dev only) for any NULL stuff
-			case 8:
+			case 10:
-			case 9:
+			case 11:
+			case 16:
 				label += " (debug only)";
 				break;
 
 			default:
-				break;
+				continue;
 		}
 
 		gtk_combo_box_append_text(GTK_COMBO_BOX(render_combo_box), label.c_str());
@@ -62,13 +61,15 @@ GtkWidget* CreateRenderComboBox()
 
 	switch (theApp.GetConfig("renderer", 0)) {
 		// Note the value are based on m_gs_renderers vector on GSdx.cpp
-		case 10: renderer_box_position = 2; break;
+		case 10: renderer_box_position = 0; break;
-		case 11: renderer_box_position = 3; break;
+		case 16: renderer_box_position = 1; break;
-		case 12: renderer_box_position = 4; break;
+		case 11: renderer_box_position = 2; break;
-		case 13: renderer_box_position = 5; break;
+		case 12: renderer_box_position = 3; break;
+		case 13: renderer_box_position = 4; break;
+		case 17: renderer_box_position = 5; break;
 
 		// Fallback to openGL SW
-		default: renderer_box_position = 5; break;
+		default: renderer_box_position = 4; break;
 	}
 	gtk_combo_box_set_active(GTK_COMBO_BOX(render_combo_box), renderer_box_position);
 	return render_combo_box;
@@ -483,10 +484,12 @@ override_GL_ARB_shading_language_420pack = -1
 		if (gtk_combo_box_get_active(GTK_COMBO_BOX(render_combo_box)) != -1) {
 			// Note the value are based on m_gs_renderers vector on GSdx.cpp
 			switch (gtk_combo_box_get_active(GTK_COMBO_BOX(render_combo_box))) {
-				case 2: theApp.SetConfig("renderer", 10); break;
+				case 0: theApp.SetConfig("renderer", 10); break;
-				case 3: theApp.SetConfig("renderer", 11); break;
+				case 1: theApp.SetConfig("renderer", 16); break;
-				case 4: theApp.SetConfig("renderer", 12); break;
+				case 2: theApp.SetConfig("renderer", 11); break;
-				case 5: theApp.SetConfig("renderer", 13); break;
+				case 3: theApp.SetConfig("renderer", 12); break;
+				case 4: theApp.SetConfig("renderer", 13); break;
+				case 5: theApp.SetConfig("renderer", 17); break;
 
 				// Fallback to SW opengl
 				default: theApp.SetConfig("renderer", 13); break;

plugins/GSdx/GSLocalMemory.cpp

@@ -692,14 +692,14 @@ void GSLocalMemory::WriteImageColumn(int l, int r, int y, int h, const uint8* sr
 		{
 			switch(psm)
 			{
-			case PSM_PSMCT32: WriteColumn32<alignment, 0xffffffff>(y, BlockPtr32(x, y, bp, bw), &src[x * 4], srcpitch); break;
+			case PSM_PSMCT32: GSBlock::WriteColumn32<alignment, 0xffffffff>(y, BlockPtr32(x, y, bp, bw), &src[x * 4], srcpitch); break;
-			case PSM_PSMCT16: WriteColumn16<alignment>(y, BlockPtr16(x, y, bp, bw), &src[x * 2], srcpitch); break;
+			case PSM_PSMCT16: GSBlock::WriteColumn16<alignment>(y, BlockPtr16(x, y, bp, bw), &src[x * 2], srcpitch); break;
-			case PSM_PSMCT16S: WriteColumn16<alignment>(y, BlockPtr16S(x, y, bp, bw), &src[x * 2], srcpitch); break;
+			case PSM_PSMCT16S: GSBlock::WriteColumn16<alignment>(y, BlockPtr16S(x, y, bp, bw), &src[x * 2], srcpitch); break;
-			case PSM_PSMT8: WriteColumn8<alignment>(y, BlockPtr8(x, y, bp, bw), &src[x], srcpitch); break;
+			case PSM_PSMT8: GSBlock::WriteColumn8<alignment>(y, BlockPtr8(x, y, bp, bw), &src[x], srcpitch); break;
-			case PSM_PSMT4: WriteColumn4<alignment>(y, BlockPtr4(x, y, bp, bw), &src[x >> 1], srcpitch); break;
+			case PSM_PSMT4: GSBlock::WriteColumn4<alignment>(y, BlockPtr4(x, y, bp, bw), &src[x >> 1], srcpitch); break;
-			case PSM_PSMZ32: WriteColumn32<alignment, 0xffffffff>(y, BlockPtr32Z(x, y, bp, bw), &src[x * 4], srcpitch); break;
+			case PSM_PSMZ32: GSBlock::WriteColumn32<alignment, 0xffffffff>(y, BlockPtr32Z(x, y, bp, bw), &src[x * 4], srcpitch); break;
-			case PSM_PSMZ16: WriteColumn16<alignment>(y, BlockPtr16Z(x, y, bp, bw), &src[x * 2], srcpitch); break;
+			case PSM_PSMZ16: GSBlock::WriteColumn16<alignment>(y, BlockPtr16Z(x, y, bp, bw), &src[x * 2], srcpitch); break;
-			case PSM_PSMZ16S: WriteColumn16<alignment>(y, BlockPtr16SZ(x, y, bp, bw), &src[x * 2], srcpitch); break;
+			case PSM_PSMZ16S: GSBlock::WriteColumn16<alignment>(y, BlockPtr16SZ(x, y, bp, bw), &src[x * 2], srcpitch); break;
 			// TODO
 			default: __assume(0);
 			}
@@ -719,14 +719,14 @@ void GSLocalMemory::WriteImageBlock(int l, int r, int y, int h, const uint8* src
 		{
 			switch(psm)
 			{
-			case PSM_PSMCT32: WriteBlock32<alignment, 0xffffffff>(BlockPtr32(x, y, bp, bw), &src[x * 4], srcpitch); break;
+			case PSM_PSMCT32: GSBlock::WriteBlock32<alignment, 0xffffffff>(BlockPtr32(x, y, bp, bw), &src[x * 4], srcpitch); break;
-			case PSM_PSMCT16: WriteBlock16<alignment>(BlockPtr16(x, y, bp, bw), &src[x * 2], srcpitch); break;
+			case PSM_PSMCT16: GSBlock::WriteBlock16<alignment>(BlockPtr16(x, y, bp, bw), &src[x * 2], srcpitch); break;
-			case PSM_PSMCT16S: WriteBlock16<alignment>(BlockPtr16S(x, y, bp, bw), &src[x * 2], srcpitch); break;
+			case PSM_PSMCT16S: GSBlock::WriteBlock16<alignment>(BlockPtr16S(x, y, bp, bw), &src[x * 2], srcpitch); break;
-			case PSM_PSMT8: WriteBlock8<alignment>(BlockPtr8(x, y, bp, bw), &src[x], srcpitch); break;
+			case PSM_PSMT8: GSBlock::WriteBlock8<alignment>(BlockPtr8(x, y, bp, bw), &src[x], srcpitch); break;
-			case PSM_PSMT4: WriteBlock4<alignment>(BlockPtr4(x, y, bp, bw), &src[x >> 1], srcpitch); break;
+			case PSM_PSMT4: GSBlock::WriteBlock4<alignment>(BlockPtr4(x, y, bp, bw), &src[x >> 1], srcpitch); break;
-			case PSM_PSMZ32: WriteBlock32<alignment, 0xffffffff>(BlockPtr32Z(x, y, bp, bw), &src[x * 4], srcpitch); break;
+			case PSM_PSMZ32: GSBlock::WriteBlock32<alignment, 0xffffffff>(BlockPtr32Z(x, y, bp, bw), &src[x * 4], srcpitch); break;
-			case PSM_PSMZ16: WriteBlock16<alignment>(BlockPtr16Z(x, y, bp, bw), &src[x * 2], srcpitch); break;
+			case PSM_PSMZ16: GSBlock::WriteBlock16<alignment>(BlockPtr16Z(x, y, bp, bw), &src[x * 2], srcpitch); break;
-			case PSM_PSMZ16S: WriteBlock16<alignment>(BlockPtr16SZ(x, y, bp, bw), &src[x * 2], srcpitch); break;
+			case PSM_PSMZ16S: GSBlock::WriteBlock16<alignment>(BlockPtr16SZ(x, y, bp, bw), &src[x * 2], srcpitch); break;
 			// TODO
 			default: __assume(0);
 			}
@@ -801,27 +801,27 @@ void GSLocalMemory::WriteImageTopBottom(int l, int r, int y, int h, const uint8*
 			{
 			case PSM_PSMCT32:
 			case PSM_PSMZ32:
-				ReadColumn32(y, dst, buff, 32);
+				GSBlock::ReadColumn32(y, dst, buff, 32);
 				memcpy(&buff[32], &src[x * 4], 32);
-				WriteColumn32<32, 0xffffffff>(y, dst, buff, 32);
+				GSBlock::WriteColumn32<32, 0xffffffff>(y, dst, buff, 32);
 				break;
 			case PSM_PSMCT16:
 			case PSM_PSMCT16S:
 			case PSM_PSMZ16:
 			case PSM_PSMZ16S:
-				ReadColumn16(y, dst, buff, 32);
+				GSBlock::ReadColumn16(y, dst, buff, 32);
 				memcpy(&buff[32], &src[x * 2], 32);
-				WriteColumn16<32>(y, dst, buff, 32);
+				GSBlock::WriteColumn16<32>(y, dst, buff, 32);
 				break;
 			case PSM_PSMT8:
-				ReadColumn8(y, dst, buff, 16);
+				GSBlock::ReadColumn8(y, dst, buff, 16);
 				for(int i = 0, j = y2; i < h2; i++, j++) memcpy(&buff[j * 16], &src[i * srcpitch + x], 16);
-				WriteColumn8<32>(y, dst, buff, 16);
+				GSBlock::WriteColumn8<32>(y, dst, buff, 16);
 				break;
 			case PSM_PSMT4:
-				ReadColumn4(y, dst, buff, 16);
+				GSBlock::ReadColumn4(y, dst, buff, 16);
 				for(int i = 0, j = y2; i < h2; i++, j++) memcpy(&buff[j * 16], &src[i * srcpitch + (x >> 1)], 16);
-				WriteColumn4<32>(y, dst, buff, 16);
+				GSBlock::WriteColumn4<32>(y, dst, buff, 16);
 				break;
 			// TODO
 			default:
@@ -888,27 +888,27 @@ void GSLocalMemory::WriteImageTopBottom(int l, int r, int y, int h, const uint8*
 			{
 			case PSM_PSMCT32:
 			case PSM_PSMZ32:
-				ReadColumn32(y, dst, buff, 32);
+				GSBlock::ReadColumn32(y, dst, buff, 32);
 				memcpy(&buff[0], &src[x * 4], 32);
-				WriteColumn32<32, 0xffffffff>(y, dst, buff, 32);
+				GSBlock::WriteColumn32<32, 0xffffffff>(y, dst, buff, 32);
 				break;
 			case PSM_PSMCT16:
 			case PSM_PSMCT16S:
 			case PSM_PSMZ16:
 			case PSM_PSMZ16S:
-				ReadColumn16(y, dst, buff, 32);
+				GSBlock::ReadColumn16(y, dst, buff, 32);
 				memcpy(&buff[0], &src[x * 2], 32);
-				WriteColumn16<32>(y, dst, buff, 32);
+				GSBlock::WriteColumn16<32>(y, dst, buff, 32);
 				break;
 			case PSM_PSMT8:
-				ReadColumn8(y, dst, buff, 16);
+				GSBlock::ReadColumn8(y, dst, buff, 16);
 				for(int i = 0; i < h; i++) memcpy(&buff[i * 16], &src[i * srcpitch + x], 16);
-				WriteColumn8<32>(y, dst, buff, 16);
+				GSBlock::WriteColumn8<32>(y, dst, buff, 16);
 				break;
 			case PSM_PSMT4:
-				ReadColumn4(y, dst, buff, 16);
+				GSBlock::ReadColumn4(y, dst, buff, 16);
 				for(int i = 0; i < h; i++) memcpy(&buff[i * 16], &src[i * srcpitch + (x >> 1)], 16);
-				WriteColumn4<32>(y, dst, buff, 16);
+				GSBlock::WriteColumn4<32>(y, dst, buff, 16);
 				break;
 			// TODO
 			default:
@@ -1060,7 +1060,7 @@ void GSLocalMemory::WriteImage24(int& tx, int& ty, const uint8* src, int len, GI
 		{
 			for(int x = tx; x < tw; x += 8)
 			{
-				UnpackAndWriteBlock24(src + (x - tx) * 3, srcpitch, BlockPtr32(x, y, bp, bw));
+				GSBlock::UnpackAndWriteBlock24(src + (x - tx) * 3, srcpitch, BlockPtr32(x, y, bp, bw));
 			}
 		}
 
@@ -1094,7 +1094,7 @@ void GSLocalMemory::WriteImage8H(int& tx, int& ty, const uint8* src, int len, GI
 		{
 			for(int x = tx; x < tw; x += 8)
 			{
-				UnpackAndWriteBlock8H(src + (x - tx), srcpitch, BlockPtr32(x, y, bp, bw));
+				GSBlock::UnpackAndWriteBlock8H(src + (x - tx), srcpitch, BlockPtr32(x, y, bp, bw));
 			}
 		}
 
@@ -1128,7 +1128,7 @@ void GSLocalMemory::WriteImage4HL(int& tx, int& ty, const uint8* src, int len, G
 		{
 			for(int x = tx; x < tw; x += 8)
 			{
-				UnpackAndWriteBlock4HL(src + (x - tx) / 2, srcpitch, BlockPtr32(x, y, bp, bw));
+				GSBlock::UnpackAndWriteBlock4HL(src + (x - tx) / 2, srcpitch, BlockPtr32(x, y, bp, bw));
 			}
 		}
 
@@ -1162,7 +1162,7 @@ void GSLocalMemory::WriteImage4HH(int& tx, int& ty, const uint8* src, int len, G
 		{
 			for(int x = tx; x < tw; x += 8)
 			{
-				UnpackAndWriteBlock4HH(src + (x - tx) / 2, srcpitch, BlockPtr32(x, y, bp, bw));
+				GSBlock::UnpackAndWriteBlock4HH(src + (x - tx) / 2, srcpitch, BlockPtr32(x, y, bp, bw));
 			}
 		}
 
@@ -1196,7 +1196,7 @@ void GSLocalMemory::WriteImage24Z(int& tx, int& ty, const uint8* src, int len, G
 		{
 			for(int x = tx; x < tw; x += 8)
 			{
-				UnpackAndWriteBlock24(src + (x - tx) * 3, srcpitch, BlockPtr32Z(x, y, bp, bw));
+				GSBlock::UnpackAndWriteBlock24(src + (x - tx) * 3, srcpitch, BlockPtr32Z(x, y, bp, bw));
 			}
 		}
 
@@ -1612,7 +1612,7 @@ void GSLocalMemory::ReadTexture32(const GSOffset* RESTRICT o, const GSVector4i&
 {
 	FOREACH_BLOCK_START(r, 8, 8, 32)
 	{
-		ReadBlock32(src, dst, dstpitch);
+		GSBlock::ReadBlock32(src, dst, dstpitch);
 	}
 	FOREACH_BLOCK_END
 }
@@ -1623,7 +1623,7 @@ void GSLocalMemory::ReadTexture24(const GSOffset* RESTRICT o, const GSVector4i&
 	{
 		FOREACH_BLOCK_START(r, 8, 8, 32)
 		{
-			ReadAndExpandBlock24<true>(src, dst, dstpitch, TEXA);
+			GSBlock::ReadAndExpandBlock24<true>(src, dst, dstpitch, TEXA);
 		}
 		FOREACH_BLOCK_END
 	}
@@ -1631,7 +1631,7 @@ void GSLocalMemory::ReadTexture24(const GSOffset* RESTRICT o, const GSVector4i&
 	{
 		FOREACH_BLOCK_START(r, 8, 8, 32)
 		{
-			ReadAndExpandBlock24<false>(src, dst, dstpitch, TEXA);
+			GSBlock::ReadAndExpandBlock24<false>(src, dst, dstpitch, TEXA);
 		}
 		FOREACH_BLOCK_END
 	}
@@ -1643,7 +1643,7 @@ void GSLocalMemory::ReadTexture16(const GSOffset* RESTRICT o, const GSVector4i&
 	{
 		FOREACH_BLOCK_START(r, 16, 8, 32)
 		{
-			ReadAndExpandBlock16<true>(src, dst, dstpitch, TEXA);
+			GSBlock::ReadAndExpandBlock16<true>(src, dst, dstpitch, TEXA);
 		}
 		FOREACH_BLOCK_END
 	}
@@ -1651,7 +1651,7 @@ void GSLocalMemory::ReadTexture16(const GSOffset* RESTRICT o, const GSVector4i&
 	{
 		FOREACH_BLOCK_START(r, 16, 8, 32)
 		{
-			ReadAndExpandBlock16<false>(src, dst, dstpitch, TEXA);
+			GSBlock::ReadAndExpandBlock16<false>(src, dst, dstpitch, TEXA);
 		}
 		FOREACH_BLOCK_END
 	}
@@ -1663,7 +1663,7 @@ void GSLocalMemory::ReadTexture8(const GSOffset* RESTRICT o, const GSVector4i& r
 
 	FOREACH_BLOCK_START(r, 16, 16, 32)
 	{
-		ReadAndExpandBlock8_32(src, dst, dstpitch, pal);
+		GSBlock::ReadAndExpandBlock8_32(src, dst, dstpitch, pal);
 	}
 	FOREACH_BLOCK_END
 }
@@ -1674,7 +1674,7 @@ void GSLocalMemory::ReadTexture4(const GSOffset* RESTRICT o, const GSVector4i& r
 
 	FOREACH_BLOCK_START(r, 32, 16, 32)
 	{
-		ReadAndExpandBlock4_32(src, dst, dstpitch, pal);
+		GSBlock::ReadAndExpandBlock4_32(src, dst, dstpitch, pal);
 	}
 	FOREACH_BLOCK_END
 }
@@ -1685,7 +1685,7 @@ void GSLocalMemory::ReadTexture8H(const GSOffset* RESTRICT o, const GSVector4i&
 
 	FOREACH_BLOCK_START(r, 8, 8, 32)
 	{
-		ReadAndExpandBlock8H_32(src, dst, dstpitch, pal);
+		GSBlock::ReadAndExpandBlock8H_32(src, dst, dstpitch, pal);
 	}
 	FOREACH_BLOCK_END
 }
@@ -1696,7 +1696,7 @@ void GSLocalMemory::ReadTexture4HL(const GSOffset* RESTRICT o, const GSVector4i&
 
 	FOREACH_BLOCK_START(r, 8, 8, 32)
 	{
-		ReadAndExpandBlock4HL_32(src, dst, dstpitch, pal);
+		GSBlock::ReadAndExpandBlock4HL_32(src, dst, dstpitch, pal);
 	}
 	FOREACH_BLOCK_END
 }
@@ -1707,7 +1707,7 @@ void GSLocalMemory::ReadTexture4HH(const GSOffset* RESTRICT o, const GSVector4i&
 
 	FOREACH_BLOCK_START(r, 8, 8, 32)
 	{
-		ReadAndExpandBlock4HH_32(src, dst, dstpitch, pal);
+		GSBlock::ReadAndExpandBlock4HH_32(src, dst, dstpitch, pal);
 	}
 	FOREACH_BLOCK_END
 }
@@ -1718,7 +1718,7 @@ void GSLocalMemory::ReadTextureBlock32(uint32 bp, uint8* dst, int dstpitch, cons
 {
 	ALIGN_STACK(32);
 
-	ReadBlock32(BlockPtr(bp), dst, dstpitch);
+	GSBlock::ReadBlock32(BlockPtr(bp), dst, dstpitch);
 }
 
 void GSLocalMemory::ReadTextureBlock24(uint32 bp, uint8* dst, int dstpitch, const GIFRegTEXA& TEXA) const
@@ -1727,11 +1727,11 @@ void GSLocalMemory::ReadTextureBlock24(uint32 bp, uint8* dst, int dstpitch, cons
 
 	if(TEXA.AEM)
 	{
-		ReadAndExpandBlock24<true>(BlockPtr(bp), dst, dstpitch, TEXA);
+		GSBlock::ReadAndExpandBlock24<true>(BlockPtr(bp), dst, dstpitch, TEXA);
 	}
 	else
 	{
-		ReadAndExpandBlock24<false>(BlockPtr(bp), dst, dstpitch, TEXA);
+		GSBlock::ReadAndExpandBlock24<false>(BlockPtr(bp), dst, dstpitch, TEXA);
 	}
 }
 
@@ -1741,11 +1741,11 @@ void GSLocalMemory::ReadTextureBlock16(uint32 bp, uint8* dst, int dstpitch, cons
 
 	if(TEXA.AEM)
 	{
-		ReadAndExpandBlock16<true>(BlockPtr(bp), dst, dstpitch, TEXA);
+		GSBlock::ReadAndExpandBlock16<true>(BlockPtr(bp), dst, dstpitch, TEXA);
 	}
 	else
 	{
-		ReadAndExpandBlock16<false>(BlockPtr(bp), dst, dstpitch, TEXA);
+		GSBlock::ReadAndExpandBlock16<false>(BlockPtr(bp), dst, dstpitch, TEXA);
 	}
 }
 
@@ -1753,35 +1753,35 @@ void GSLocalMemory::ReadTextureBlock8(uint32 bp, uint8* dst, int dstpitch, const
 {
 	ALIGN_STACK(32);
 
-	ReadAndExpandBlock8_32(BlockPtr(bp), dst, dstpitch, m_clut);
+	GSBlock::ReadAndExpandBlock8_32(BlockPtr(bp), dst, dstpitch, m_clut);
 }
 
 void GSLocalMemory::ReadTextureBlock4(uint32 bp, uint8* dst, int dstpitch, const GIFRegTEXA& TEXA) const
 {
 	ALIGN_STACK(32);
 
-	ReadAndExpandBlock4_32(BlockPtr(bp), dst, dstpitch, m_clut);
+	GSBlock::ReadAndExpandBlock4_32(BlockPtr(bp), dst, dstpitch, m_clut);
 }
 
 void GSLocalMemory::ReadTextureBlock8H(uint32 bp, uint8* dst, int dstpitch, const GIFRegTEXA& TEXA) const
 {
 	ALIGN_STACK(32);
 
-	ReadAndExpandBlock8H_32(BlockPtr(bp), dst, dstpitch, m_clut);
+	GSBlock::ReadAndExpandBlock8H_32(BlockPtr(bp), dst, dstpitch, m_clut);
 }
 
 void GSLocalMemory::ReadTextureBlock4HL(uint32 bp, uint8* dst, int dstpitch, const GIFRegTEXA& TEXA) const
 {
 	ALIGN_STACK(32);
 
-	ReadAndExpandBlock4HL_32(BlockPtr(bp), dst, dstpitch, m_clut);
+	GSBlock::ReadAndExpandBlock4HL_32(BlockPtr(bp), dst, dstpitch, m_clut);
 }
 
 void GSLocalMemory::ReadTextureBlock4HH(uint32 bp, uint8* dst, int dstpitch, const GIFRegTEXA& TEXA) const
 {
 	ALIGN_STACK(32);
 
-	ReadAndExpandBlock4HH_32(BlockPtr(bp), dst, dstpitch, m_clut);
+	GSBlock::ReadAndExpandBlock4HH_32(BlockPtr(bp), dst, dstpitch, m_clut);
 }
 
 ///////////////////
@@ -1870,7 +1870,7 @@ void GSLocalMemory::ReadTexture8P(const GSOffset* RESTRICT o, const GSVector4i&
 {
 	FOREACH_BLOCK_START(r, 16, 16, 8)
 	{
-		ReadBlock8(src, dst, dstpitch);
+		GSBlock::ReadBlock8(src, dst, dstpitch);
 	}
 	FOREACH_BLOCK_END
 }
@@ -1879,7 +1879,7 @@ void GSLocalMemory::ReadTexture4P(const GSOffset* RESTRICT o, const GSVector4i&
 {
 	FOREACH_BLOCK_START(r, 32, 16, 8)
 	{
-		ReadBlock4P(src, dst, dstpitch);
+		GSBlock::ReadBlock4P(src, dst, dstpitch);
 	}
 	FOREACH_BLOCK_END
 }
@@ -1888,7 +1888,7 @@ void GSLocalMemory::ReadTexture8HP(const GSOffset* RESTRICT o, const GSVector4i&
 {
 	FOREACH_BLOCK_START(r, 8, 8, 8)
 	{
-		ReadBlock8HP(src, dst, dstpitch);
+		GSBlock::ReadBlock8HP(src, dst, dstpitch);
 	}
 	FOREACH_BLOCK_END
 }
@@ -1897,7 +1897,7 @@ void GSLocalMemory::ReadTexture4HLP(const GSOffset* RESTRICT o, const GSVector4i
 {
 	FOREACH_BLOCK_START(r, 8, 8, 8)
 	{
-		ReadBlock4HLP(src, dst, dstpitch);
+		GSBlock::ReadBlock4HLP(src, dst, dstpitch);
 	}
 	FOREACH_BLOCK_END
 }
@@ -1906,7 +1906,7 @@ void GSLocalMemory::ReadTexture4HHP(const GSOffset* RESTRICT o, const GSVector4i
 {
 	FOREACH_BLOCK_START(r, 8, 8, 8)
 	{
-		ReadBlock4HHP(src, dst, dstpitch);
+		GSBlock::ReadBlock4HHP(src, dst, dstpitch);
 	}
 	FOREACH_BLOCK_END
 }
@@ -1915,35 +1915,35 @@ void GSLocalMemory::ReadTexture4HHP(const GSOffset* RESTRICT o, const GSVector4i
 
 void GSLocalMemory::ReadTextureBlock8P(uint32 bp, uint8* dst, int dstpitch, const GIFRegTEXA& TEXA) const
 {
-	ReadBlock8(BlockPtr(bp), dst, dstpitch);
+	GSBlock::ReadBlock8(BlockPtr(bp), dst, dstpitch);
 }
 
 void GSLocalMemory::ReadTextureBlock4P(uint32 bp, uint8* dst, int dstpitch, const GIFRegTEXA& TEXA) const
 {
 	ALIGN_STACK(32);
 
-	ReadBlock4P(BlockPtr(bp), dst, dstpitch);
+	GSBlock::ReadBlock4P(BlockPtr(bp), dst, dstpitch);
 }
 
 void GSLocalMemory::ReadTextureBlock8HP(uint32 bp, uint8* dst, int dstpitch, const GIFRegTEXA& TEXA) const
 {
 	ALIGN_STACK(32);
 
-	ReadBlock8HP(BlockPtr(bp), dst, dstpitch);
+	GSBlock::ReadBlock8HP(BlockPtr(bp), dst, dstpitch);
 }
 
 void GSLocalMemory::ReadTextureBlock4HLP(uint32 bp, uint8* dst, int dstpitch, const GIFRegTEXA& TEXA) const
 {
 	ALIGN_STACK(32);
 
-	ReadBlock4HLP(BlockPtr(bp), dst, dstpitch);
+	GSBlock::ReadBlock4HLP(BlockPtr(bp), dst, dstpitch);
 }
 
 void GSLocalMemory::ReadTextureBlock4HHP(uint32 bp, uint8* dst, int dstpitch, const GIFRegTEXA& TEXA) const
 {
 	ALIGN_STACK(32);
 
-	ReadBlock4HHP(BlockPtr(bp), dst, dstpitch);
+	GSBlock::ReadBlock4HHP(BlockPtr(bp), dst, dstpitch);
 }
 
 //
@@ -2082,3 +2082,45 @@ uint32* GSOffset::GetPages(const GSVector4i& rect, uint32* pages, GSVector4i* bb
 
 	return pages;
 }
+
+GSVector4i* GSOffset::GetPagesAsBits(const GSVector4i& rect, GSVector4i* pages, GSVector4i* bbox)
+{
+	if(pages == NULL)
+	{
+		pages = (GSVector4i*)_aligned_malloc(sizeof(GSVector4i) * 4, 16);
+	}
+
+	pages[0] = GSVector4i::zero();
+	pages[1] = GSVector4i::zero();
+	pages[2] = GSVector4i::zero();
+	pages[3] = GSVector4i::zero();
+
+	GSVector2i bs = (bp & 31) == 0 ? GSLocalMemory::m_psm[psm].pgs : GSLocalMemory::m_psm[psm].bs;
+
+	GSVector4i r = rect.ralign<Align_Outside>(bs);
+
+	if(bbox != NULL) *bbox = r;
+
+	r = r.sra32(3);
+
+	bs.x >>= 3;
+	bs.y >>= 3;
+
+	for(int y = r.top; y < r.bottom; y += bs.y)
+	{
+		uint32 base = block.row[y];
+
+		for(int x = r.left; x < r.right; x += bs.x)
+		{
+			uint32 n = (base + block.col[x]) >> 5;
+
+			if(n < MAX_PAGES)
+			{
+				((uint32*)pages)[n >> 5] |= 1 << (n & 31);
+			}
+		}
+	}
+
+	return pages;
+
+}

plugins/GSdx/GSLocalMemory.h

@@ -54,6 +54,7 @@ public:
 	enum {EOP = 0xffffffff};
 
 	uint32* GetPages(const GSVector4i& rect, uint32* pages = NULL, GSVector4i* bbox = NULL);
+	GSVector4i* GetPagesAsBits(const GSVector4i& rect, GSVector4i* pages = NULL, GSVector4i* bbox = NULL); // free returned value with _aligned_free
 };
 
 struct GSPixelOffset
@@ -76,7 +77,7 @@ struct GSPixelOffset4
 	uint32 fbp, zbp, fpsm, zpsm, bw;
 };
 
-class GSLocalMemory : public GSBlock
+class GSLocalMemory : public GSAlignedClass<32>
 {
 public:
 	typedef uint32 (*pixelAddress)(int x, int y, uint32 bp, uint32 bw);

plugins/GSdx/GSRendererCL.h

@@ -0,0 +1,258 @@
+/*
+ *	Copyright (C) 2007-2009 Gabest
+ *	http://www.gabest.org
+ *
+ *  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
+ *  the Free Software Foundation; either version 2, or (at your option)
+ *  any later version.
+ *
+ *  This Program is distributed in the hope that it will be useful,
+ *  but WITHOUT ANY WARRANTY; without even the implied warranty of
+ *  MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the
+ *  GNU General Public License for more details.
+ *
+ *  You should have received a copy of the GNU General Public License
+ *  along with GNU Make; see the file COPYING.  If not, write to
+ *  the Free Software Foundation, Inc., 51 Franklin Street, Fifth Floor, Boston, MA  02110-1301, USA USA.
+ *  http://www.gnu.org/copyleft/gpl.html
+ *
+ */
+
+#pragma once
+
+#include "GSRenderer.h"
+
+#ifdef ENABLE_OPENCL
+
+__aligned(struct, 32) GSVertexCL
+{
+	GSVector4 p, t;
+};
+
+class GSRendererCL : public GSRenderer
+{
+	typedef void (GSRendererCL::*ConvertVertexBufferPtr)(GSVertexCL* RESTRICT dst, const GSVertex* RESTRICT src, size_t count);
+
+	ConvertVertexBufferPtr m_cvb[4][2][2];
+
+	template<uint32 primclass, uint32 tme, uint32 fst>
+	void ConvertVertexBuffer(GSVertexCL* RESTRICT dst, const GSVertex* RESTRICT src, size_t count);
+
+	union PrimSelector
+	{
+		struct
+		{
+			uint32 prim:2; // 0
+		};
+
+		uint32 key;
+
+		operator uint32() const { return key; }
+	};
+
+	union TileSelector
+	{
+		struct
+		{
+			uint32 prim:2; // 0
+			uint32 mode:2; // 2
+			uint32 clear:1; // 4
+		};
+
+		uint32 key;
+
+		operator uint32() const { return key; }
+	};
+
+	union TFXSelector
+	{
+		struct
+		{
+			uint32 fpsm:3; // 0
+			uint32 zpsm:3; // 3
+			uint32 ztst:2; // 6 (0: off, 1: write, 2: test (ge), 3: test (g))
+			uint32 atst:3; // 8
+			uint32 afail:2; // 11
+			uint32 iip:1; // 13
+			uint32 tfx:3; // 14
+			uint32 tcc:1; // 17
+			uint32 fst:1; // 18
+			uint32 ltf:1; // 19
+			uint32 tlu:1; // 20
+			uint32 fge:1; // 21
+			uint32 date:1; // 22
+			uint32 abe:1; // 23
+			uint32 aba:2; // 24
+			uint32 abb:2; // 26
+			uint32 abc:2; // 28
+			uint32 abd:2; // 30
+
+			uint32 pabe:1; // 32
+			uint32 aa1:1; // 33
+			uint32 fwrite:1; // 34
+			uint32 ftest:1; // 35
+			uint32 rfb:1; // 36
+			uint32 zwrite:1; // 37
+			uint32 ztest:1; // 38
+			uint32 rzb:1; // 39
+			uint32 wms:2; // 40
+			uint32 wmt:2; // 42
+			uint32 datm:1; // 44
+			uint32 colclamp:1; // 45
+			uint32 fba:1; // 46
+			uint32 dthe:1; // 47
+			uint32 prim:2; // 48
+			uint32 lcm:1; // 50
+			uint32 mmin:2; // 51
+			uint32 noscissor:1; // 53
+			uint32 tpsm:4; // 54
+			uint32 aem:1; // 58
+			// TODO
+		};
+
+		struct
+		{
+			uint32 _pad1:24;
+			uint32 ababcd:8;
+			uint32 _pad2:2;
+			uint32 fb:2;
+			uint32 _pad3:1;
+			uint32 zb:2;
+		};
+
+		struct
+		{
+			uint32 lo;
+			uint32 hi;
+		};
+
+		uint64 key;
+
+		operator uint64() const { return key; }
+
+		bool IsSolidRect() const
+		{
+			return prim == GS_SPRITE_CLASS
+				&& iip == 0
+				&& tfx == TFX_NONE
+				&& abe == 0
+				&& ztst <= 1
+				&& atst <= 1
+				&& date == 0
+				&& fge == 0;
+		}
+	};
+
+	__aligned(struct, 32) TFXParameter
+	{
+		GSVector4i scissor;
+		GSVector4i dimx; // 4x4 signed char
+		uint32 fbp, zbp, bw;
+		uint32 fm, zm;
+		uint32 fog; // rgb
+		uint8 aref, afix;
+		uint8 ta0, ta1;
+		uint32 tbp[7], tbw[7];
+		int minu, maxu, minv, maxv; // umsk, ufix, vmsk, vfix
+		int lod; // lcm == 1
+		int mxl;
+		float l; // TEX1.L * -0x10000
+		float k; // TEX1.K * 0x10000
+		uint32 clut[256];
+	};
+
+	class TFXJob
+	{
+	public:
+		struct { int x, y, z, w; } rect;
+		TFXSelector sel;
+		uint32 ib_start;
+		uint32 prim_count;
+		GSVector4i* src_pages; // read by any texture level
+		GSVector4i* dst_pages; // f/z writes to it
+		uint32 fbp, zbp, bw;
+#ifdef DEBUG
+		TFXParameter* pb;
+#endif
+		TFXJob();
+		virtual ~TFXJob();
+
+		GSVector4i* GetSrcPages();
+		GSVector4i* GetDstPages();
+	};
+
+	class CL
+	{
+		std::string kernel_str;
+		std::map<uint32, cl::Kernel> prim_map;
+		std::map<uint32, cl::Kernel> tile_map;
+		std::map<uint64, cl::Kernel> tfx_map;
+
+		cl::Kernel Build(const char* entry, ostringstream& opt);
+		void AddDefs(ostringstream& opt);
+
+	public:
+		std::vector<OCLDeviceDesc> devs;
+		cl::Context context;
+		cl::CommandQueue queue[3];
+		cl::Buffer vm;
+		cl::Buffer tex;
+		struct { cl::Buffer buff[2]; size_t head, tail, size; unsigned char* ptr; void* mapped_ptr; } vb, ib, pb;
+		cl::Buffer env;
+		cl::CommandQueue* wq;
+		int wqidx;
+		uint32 WIs;
+		int version;
+
+	public:
+		CL();
+		virtual ~CL();
+
+		cl::Kernel& GetPrimKernel(const PrimSelector& sel);
+		cl::Kernel& GetTileKernel(const TileSelector& sel);
+		cl::Kernel& GetTFXKernel(const TFXSelector& sel);
+
+		void Map();
+		void Unmap();
+	};
+
+	CL m_cl;
+	std::list<shared_ptr<TFXJob>> m_jobs;
+	uint32 m_vb_start;
+	uint32 m_vb_count;
+	uint32 m_pb_start;
+	uint32 m_pb_count;
+	bool m_synced;
+
+	void Enqueue();
+	void EnqueueTFX(std::list<shared_ptr<TFXJob>>& jobs, uint32 bin_count, const cl_uchar4& bin_dim);
+	void UpdateTextureCache(TFXJob* job);
+	void InvalidateTextureCache(TFXJob* job);
+
+protected:
+	GSTexture* m_texture[2];
+	uint8* m_output;
+	
+	GSVector4i m_rw_pages[2][4]; // pages that may be read or modified by the rendering queue, f/z rw, tex r
+	GSVector4i m_tc_pages[4]; // invalidated texture cache pages (split this into 8:24?)
+	GSVector4i m_tmp_pages[4]; // TODO: this should be block level, too many overlaps inside pages with render targets
+
+	void Reset();
+	void VSync(int field);
+	void ResetDevice();
+	GSTexture* GetOutput(int i);
+
+	void Draw();
+	void Sync(int reason);
+	void InvalidateVideoMem(const GIFRegBITBLTBUF& BITBLTBUF, const GSVector4i& r);
+	void InvalidateLocalMem(const GIFRegBITBLTBUF& BITBLTBUF, const GSVector4i& r, bool clut = false);
+
+	bool SetupParameter(TFXJob* job, TFXParameter* pb, GSVertexCL* vertex, size_t vertex_count, const uint32* index, size_t index_count);
+
+public:
+	GSRendererCL();
+	virtual ~GSRendererCL();
+};
+
+#endif

plugins/GSdx/GSRendererSW.cpp

@@ -429,6 +429,15 @@ void GSRendererSW::Draw()
 
 	GSVector4i scissor = GSVector4i(context->scissor.in);
 	GSVector4i bbox = GSVector4i(m_vt.m_min.p.floor().xyxy(m_vt.m_max.p.ceil()));
+
+	// points and lines may have zero area bbox (single line: 0, 0 - 256, 0)
+
+	if(m_vt.m_primclass == GS_POINT_CLASS || m_vt.m_primclass == GS_LINE_CLASS)
+	{
+		if(bbox.x == bbox.z) bbox.z++;
+		if(bbox.y == bbox.w) bbox.w++;
+	}
+
 	GSVector4i r = bbox.rintersect(scissor);
 
 	scissor.z = std::min<int>(scissor.z, (int)context->FRAME.FBW * 64); // TODO: find a game that overflows and check which one is the right behaviour
@@ -973,7 +982,7 @@ bool GSRendererSW::GetScanlineGlobalData(SharedData* data)
 	gd.sel.zpsm = 3;
 	gd.sel.atst = ATST_ALWAYS;
 	gd.sel.tfx = TFX_NONE;
-	gd.sel.ababcd = 255;
+	gd.sel.ababcd = 0xff;
 	gd.sel.prim = primclass;
 
 	uint32 fm = context->FRAME.FBMSK;
@@ -1101,7 +1110,7 @@ bool GSRendererSW::GetScanlineGlobalData(SharedData* data)
 				gd.sel.mmin = (context->TEX1.MMIN & 1) + 1; // 1: round, 2: tri
 				gd.sel.lcm = context->TEX1.LCM;
 
-				int mxl = (std::min<int>((int)context->TEX1.MXL, 6) << 16);
+				int mxl = std::min<int>((int)context->TEX1.MXL, 6) << 16;
 				int k = context->TEX1.K << 12;
 
 				if((int)m_vt.m_lod.x >= (int)context->TEX1.MXL)

plugins/GSdx/GSSettingsDlg.cpp

@@ -31,6 +31,16 @@ GSSettingsDlg::GSSettingsDlg(bool isOpen2)
 	: GSDialog(isOpen2 ? IDD_CONFIG2 : IDD_CONFIG)
 	, m_IsOpen2(isOpen2)
 {
+	list<OCLDeviceDesc> ocldevs;
+
+	GSUtil::GetDeviceDescs(ocldevs);
+
+	int index = 0;
+
+	for(auto dev : ocldevs)
+	{
+		m_ocl_devs.push_back(GSSetting(index++, dev.name.c_str(), ""));
+	}
 }
 
 void GSSettingsDlg::OnInit()
@@ -40,11 +50,15 @@ void GSSettingsDlg::OnInit()
 	m_modes.clear();
 
 	CComPtr<IDirect3D9> d3d9;
+
 	d3d9.Attach(Direct3DCreate9(D3D_SDK_VERSION));
 
 	CComPtr<IDXGIFactory1> dxgi_factory;
-	if (GSUtil::CheckDXGI())
+	
+	if(GSUtil::CheckDXGI())
+	{
 		CreateDXGIFactory1(__uuidof(IDXGIFactory1), (void**)&dxgi_factory);
+	}
 
 	if(!m_IsOpen2)
 	{
@@ -81,51 +95,49 @@ void GSSettingsDlg::OnInit()
 	adapters.push_back(Adapter("Default Hardware Device", "default", GSUtil::CheckDirect3D11Level(NULL, D3D_DRIVER_TYPE_HARDWARE)));
 	adapters.push_back(Adapter("Reference Device", "ref", GSUtil::CheckDirect3D11Level(NULL, D3D_DRIVER_TYPE_REFERENCE)));
 
-	if (dxgi_factory)
+	if(dxgi_factory)
 	{
-		for (int i = 0;; i++)
+		for(int i = 0;; i++)
 		{
 			CComPtr<IDXGIAdapter1> adapter;
-			if (S_OK != dxgi_factory->EnumAdapters1(i, &adapter))
+
+			if(S_OK != dxgi_factory->EnumAdapters1(i, &adapter))
 				break;
+
 			DXGI_ADAPTER_DESC1 desc;
+			
 			HRESULT hr = adapter->GetDesc1(&desc);
-			if (S_OK == hr)
+			
+			if(S_OK == hr)
 			{
 				D3D_FEATURE_LEVEL level = GSUtil::CheckDirect3D11Level(adapter, D3D_DRIVER_TYPE_UNKNOWN);
-// GSDX isn't unicode!?
+				// GSDX isn't unicode!?
 #if 1
-				int size = WideCharToMultiByte(CP_ACP, 0,
+				int size = WideCharToMultiByte(CP_ACP, 0, desc.Description, sizeof(desc.Description), NULL, 0, NULL, NULL);
-					desc.Description, sizeof(desc.Description),
-					NULL, 0,
-					NULL, NULL);
 				char *buf = new char[size];
-				WideCharToMultiByte(CP_ACP, 0,
+				WideCharToMultiByte(CP_ACP, 0, desc.Description, sizeof(desc.Description), buf, size, NULL, NULL);
-					desc.Description, sizeof(desc.Description),
-					buf, size,
-					NULL, NULL);
 				adapters.push_back(Adapter(buf, GSAdapter(desc), level));
-				delete [] buf;
+				delete[] buf;
 #else
 				adapters.push_back(Adapter(desc.Description, GSAdapter(desc), level));
 #endif
 			}
 		}
 	}
-	else if (d3d9)
+	else if(d3d9)
 	{
 		int n = d3d9->GetAdapterCount();
-		for (int i = 0; i < n; i++)
+		for(int i = 0; i < n; i++)
 		{
 			D3DADAPTER_IDENTIFIER9 desc;
-			if (D3D_OK != d3d9->GetAdapterIdentifier(i, 0, &desc))
+
+			if(D3D_OK != d3d9->GetAdapterIdentifier(i, 0, &desc))
 				break;
-// GSDX isn't unicode!?
+
+			// GSDX isn't unicode!?
 #if 0
 			wchar_t buf[sizeof desc.Description * sizeof(WCHAR)];
-			MultiByteToWideChar(CP_ACP /* I have no idea if this is right */, 0,
+			MultiByteToWideChar(CP_ACP /* I have no idea if this is right */, 0, desc.Description, sizeof(desc.Description), buf, sizeof buf / sizeof *buf);
-				desc.Description, sizeof(desc.Description),
-				buf, sizeof buf / sizeof *buf);
 			adapters.push_back(Adapter(buf, GSAdapter(desc), (D3D_FEATURE_LEVEL)0));
 #else
 			adapters.push_back(Adapter(desc.Description, GSAdapter(desc), (D3D_FEATURE_LEVEL)0));
@@ -135,17 +147,37 @@ void GSSettingsDlg::OnInit()
 
 	std::string adapter_setting = theApp.GetConfig("Adapter", "default");
 	vector<GSSetting> adapter_settings;
-	unsigned adapter_sel = 0;
+	unsigned int adapter_sel = 0;
 
-	for (unsigned i = 0; i < adapters.size(); i++)
+	for(unsigned int i = 0; i < adapters.size(); i++)
+	{
+		if(adapters[i].id == adapter_setting)
 		{
-		if (adapters[i].id == adapter_setting)
 			adapter_sel = i;
+		}
+
 		adapter_settings.push_back(GSSetting(i, adapters[i].name.c_str(), ""));
 	}
 
+	std::string ocldev = theApp.GetConfig("ocldev", "");
+
+	unsigned int ocl_sel = 0;
+
+	for(unsigned int i = 0; i < m_ocl_devs.size(); i++)
+	{
+		if(ocldev == m_ocl_devs[i].name)
+		{
+			ocl_sel = i;
+
+			break;
+		}
+	}
+
 	ComboBoxInit(IDC_ADAPTER, adapter_settings, adapter_sel);
+	ComboBoxInit(IDC_OPENCL_DEVICE, m_ocl_devs, ocl_sel);
+
 	UpdateRenderers();
+	
 	ComboBoxInit(IDC_INTERLACE, theApp.m_gs_interlace, theApp.GetConfig("Interlace", 7)); // 7 = "auto", detects interlace based on SMODE2 register
 	ComboBoxInit(IDC_ASPECTRATIO, theApp.m_gs_aspectratio, theApp.GetConfig("AspectRatio", 1));
 	ComboBoxInit(IDC_UPSCALE_MULTIPLIER, theApp.m_gs_upscale_multiplier, theApp.GetConfig("upscale_multiplier", 1));
@@ -233,6 +265,11 @@ bool GSSettingsDlg::OnCommand(HWND hWnd, UINT id, UINT code)
 				theApp.SetConfig("Adapter", adapters[(int)data].id.c_str());
 			}
 
+			if(ComboBoxGetSelData(IDC_OPENCL_DEVICE, data))
+			{
+				theApp.SetConfig("ocldev", m_ocl_devs[(int)data].name.c_str());
+			}
+
 			if(!m_IsOpen2 && ComboBoxGetSelData(IDC_RESOLUTION, data))
 			{
 				const D3DDISPLAYMODE* mode = (D3DDISPLAYMODE*)data;
@@ -266,7 +303,7 @@ bool GSSettingsDlg::OnCommand(HWND hWnd, UINT id, UINT code)
 				theApp.SetConfig("upscale_multiplier", 1);
 			}
 
-			if (ComboBoxGetSelData(IDC_AFCOMBO, data))
+			if(ComboBoxGetSelData(IDC_AFCOMBO, data))
 			{
 				theApp.SetConfig("MaxAnisotropy", (int)data);
 			}
@@ -329,17 +366,20 @@ void GSSettingsDlg::UpdateRenderers()
 	{
 		GSSetting r = theApp.m_gs_renderers[i];
 
-		if(i >= 3 && i <= 5)
+		if(r.id >= 3 && r.id <= 5 || r.id == 15)
 		{
 			if(level < D3D_FEATURE_LEVEL_10_0) continue;
 
-			r.name = std::string("Direct3D") + (level >= D3D_FEATURE_LEVEL_11_0 ? "11" : "10");
+			r.name += (level >= D3D_FEATURE_LEVEL_11_0 ? "11" : "10");
 		}
 
 		renderers.push_back(r);
-		if (r.id == renderer_setting)
+
+		if(r.id == renderer_setting)
+		{
 			renderer_sel = renderer_setting;
 		}
+	}
 
 	ComboBoxInit(IDC_RENDERER, renderers, renderer_sel);
 }
@@ -357,16 +397,19 @@ void GSSettingsDlg::UpdateControls()
 
 	if(ComboBoxGetSelData(IDC_RENDERER, i))
 	{
-		bool dx9 = (i / 3) == 0;
+		bool dx9 = i >= 0 && i <= 2 || i == 14;
-		bool dx11 = (i / 3) == 1;
+		bool dx11 = i >= 3 && i <= 5 || i == 15;
-		bool ogl = (i / 3) == 4;
+		bool ogl = i >= 12 && i <= 13 || i == 17;
-		bool hw = (i % 3) == 0;
+		bool hw = i == 0 || i == 3 || i == 12;
-		//bool sw = (i % 3) == 1;
+		//bool sw = i == 1 || i == 4 || i == 10 || i == 13;
+		bool ocl = i >= 14 && i <= 17;
+
 		bool native = !!IsDlgButtonChecked(m_hWnd, IDC_NATIVERES);
 
 		ShowWindow(GetDlgItem(m_hWnd, IDC_LOGO9), dx9 ? SW_SHOW : SW_HIDE);
 		ShowWindow(GetDlgItem(m_hWnd, IDC_LOGO11), dx11 ? SW_SHOW : SW_HIDE);
 
+		EnableWindow(GetDlgItem(m_hWnd, IDC_OPENCL_DEVICE), ocl);
 		EnableWindow(GetDlgItem(m_hWnd, IDC_WINDOWED), dx9);
 		EnableWindow(GetDlgItem(m_hWnd, IDC_RESX), hw && !native && scaling == 1);
 		EnableWindow(GetDlgItem(m_hWnd, IDC_RESX_EDIT), hw && !native && scaling == 1);
@@ -607,13 +650,13 @@ bool GSHacksDlg::OnMessage(UINT message, WPARAM wParam, LPARAM lParam)
 				break;
 			case IDC_SPRITEHACK:
 				helpstr = "Sprite Hack\n\nHelps getting rid of black inner lines in some filtered sprites."
-						  " Half option is the preferred one. Use it for Mana Khemia or Ar tonelico for example."
+						  " Half option is the preferred one. Use it for Mana Khemia or Ar Tonelico for example."
 						  " Full can be used for Tales of Destiny.";
 				break;
 			case IDC_WILDHACK:
 				helpstr = "WildArms\n\nLowers the GS precision to avoid gaps between pixels when"
 						  " upscaling. Full option fixes the text on WildArms games, while Half option might improve portraits"
-						  " in Ar tonelico.";
+						  " in Ar Tonelico.";
 				break;
 			case IDC_MSAACB:
 			case IDC_STATIC_MSAA:

plugins/GSdx/GSSettingsDlg.h

@@ -69,16 +69,21 @@ public:
 class GSSettingsDlg : public GSDialog
 {
 	list<D3DDISPLAYMODE> m_modes;
+
 	struct Adapter
 	{
 		std::string name;
 		std::string id;
 		D3D_FEATURE_LEVEL level;
-		Adapter(const std::string &n, const std::string &i, const D3D_FEATURE_LEVEL &l)
+		Adapter(const std::string &n, const std::string &i, const D3D_FEATURE_LEVEL &l) : name(n), id(i), level(l) {}
-			: name(n), id(i), level(l) {}
 	};
+	
 	std::vector<const Adapter> adapters;
+
+	vector<GSSetting> m_ocl_devs;
+
 	bool m_IsOpen2;
+	uint32 m_lastValidMsaa; // used to revert to previous dialog value if the user changed to invalid one, or lesser one and canceled
 
 	void UpdateRenderers();
 	void UpdateControls();
@@ -87,8 +92,6 @@ protected:
 	void OnInit();
 	bool OnCommand(HWND hWnd, UINT id, UINT code);
 
-	uint32 m_lastValidMsaa; // used to revert to previous dialog value if the user changed to invalid one, or lesser one and canceled
-
 	// Shade Boost
 	GSShadeBostDlg ShadeBoostDlg;
 	GSHacksDlg HacksDlg;

plugins/GSdx/GSState.cpp

@@ -47,6 +47,10 @@ GSState::GSState()
 	s_savez = !!theApp.GetConfig("savez", 0);
 	s_saven = theApp.GetConfig("saven", 0);
 
+	//s_dump = 1;
+	//s_save = 1;
+	//s_savez = 1;
+
 	UserHacks_AggressiveCRC = !!theApp.GetConfig("UserHacks", 0) ? theApp.GetConfig("UserHacks_AggressiveCRC", 0) : 0;
 	UserHacks_DisableCrcHacks = !!theApp.GetConfig("UserHacks", 0) ? theApp.GetConfig( "UserHacks_DisableCrcHacks", 0 ) : 0;
 	UserHacks_WildHack = !!theApp.GetConfig("UserHacks", 0) ? theApp.GetConfig("UserHacks_WildHack", 0) : 0;
@@ -1551,7 +1555,8 @@ void GSState::Read(uint8* mem, int len)
 		return;
 	}
 
-	if (!m_init_read_fifo_supported) {
+	if(!m_init_read_fifo_supported)
+	{
 		if(m_tr.x == sx && m_tr.y == sy)
 		{
 			InvalidateLocalMem(m_env.BITBLTBUF, GSVector4i(sx, sy, sx + w, sy + h));
@@ -2316,20 +2321,20 @@ void GSState::GrowVertexBuffer()
 	GSVertex* vertex = (GSVertex*)_aligned_malloc(sizeof(GSVertex) * maxcount, 32);
 	uint32* index = (uint32*)_aligned_malloc(sizeof(uint32) * maxcount * 3, 32); // worst case is slightly less than vertex number * 3
 
-	if (!vertex || !index)
+	if(vertex == NULL || index == NULL)
 	{
 		printf("GSdx: failed to allocate %d bytes for verticles and %d for indices.\n", sizeof(GSVertex) * maxcount, sizeof(uint32) * maxcount * 3);
 		throw GSDXError();
 	}
 
-	if (m_vertex.buff != NULL)
+	if(m_vertex.buff != NULL)
 	{
 		memcpy(vertex, m_vertex.buff, sizeof(GSVertex) * m_vertex.tail);
 
 		_aligned_free(m_vertex.buff);
 	}
 
-	if (m_index.buff != NULL)
+	if(m_index.buff != NULL)
 	{
 		memcpy(index, m_index.buff, sizeof(uint32) * m_index.tail);
 		

plugins/GSdx/GSUtil.cpp

@@ -94,6 +94,7 @@ static class GSUtilMaps
 public:
 	uint8 PrimClassField[8];
 	uint8 VertexCountField[8];
+	uint8 ClassVertexCountField[4];
 	uint32 CompatibleBitsField[64][2];
 	uint32 SharedBitsField[64][2];
 
@@ -117,6 +118,11 @@ public:
 		VertexCountField[GS_SPRITE] = 2;
 		VertexCountField[GS_INVALID] = 1;
 
+		ClassVertexCountField[GS_POINT_CLASS] = 1;
+		ClassVertexCountField[GS_LINE_CLASS] = 2;
+		ClassVertexCountField[GS_TRIANGLE_CLASS] = 3;
+		ClassVertexCountField[GS_SPRITE_CLASS] = 2;
+
 		memset(CompatibleBitsField, 0, sizeof(CompatibleBitsField));
 
 		for(int i = 0; i < 64; i++)
@@ -163,6 +169,11 @@ int GSUtil::GetVertexCount(uint32 prim)
 	return s_maps.VertexCountField[prim];
 }
 
+int GSUtil::GetClassVertexCount(uint32 primclass)
+{
+	return s_maps.ClassVertexCountField[primclass];
+}
+
 const uint32* GSUtil::HasSharedBitsPtr(uint32 dpsm)
 {
 	return s_maps.SharedBitsField[dpsm];
@@ -215,6 +226,102 @@ bool GSUtil::CheckSSE()
 	return true;
 }
 
+#define OCL_PROGRAM_VERSION 1
+
+#ifdef ENABLE_OPENCL
+void GSUtil::GetDeviceDescs(list<OCLDeviceDesc>& dl)
+{
+	dl.clear();
+
+	try
+	{
+		std::vector<cl::Platform> platforms;
+
+		cl::Platform::get(&platforms);
+
+		for(auto& p : platforms)
+		{
+			std::string platform_vendor = p.getInfo<CL_PLATFORM_VENDOR>();
+
+			std::vector<cl::Device> ds;
+
+			p.getDevices(CL_DEVICE_TYPE_ALL, &ds);
+
+			for(auto& device : ds)
+			{
+				string type;
+
+				switch(device.getInfo<CL_DEVICE_TYPE>())
+				{
+				case CL_DEVICE_TYPE_GPU: type = "GPU"; break;
+				case CL_DEVICE_TYPE_CPU: type = "CPU"; break;
+				}
+
+				if(type.empty()) continue;
+
+				std::string version = device.getInfo<CL_DEVICE_OPENCL_C_VERSION>();
+
+				int major = 0;
+				int minor = 0;
+
+				if(!type.empty() && sscanf(version.c_str(), "OpenCL C %d.%d", &major, &minor) == 2 && major == 1 && minor >= 1 || major > 1)
+				{
+					OCLDeviceDesc desc;
+
+					desc.device = device;
+					desc.name = GetDeviceUniqueName(device);
+					desc.version = major * 100 + minor * 10;
+
+					// TODO: linux
+
+					char* buff = new char[MAX_PATH + 1];
+					GetTempPath(MAX_PATH, buff);
+					desc.tmppath = string(buff) + "/" + desc.name;
+
+					WIN32_FIND_DATA FindFileData;
+					HANDLE hFind = FindFirstFile(desc.tmppath.c_str(), &FindFileData);
+					if(hFind != INVALID_HANDLE_VALUE) FindClose(hFind);
+					else CreateDirectory(desc.tmppath.c_str(), NULL);
+
+					sprintf(buff, "/%d", OCL_PROGRAM_VERSION);
+					desc.tmppath += buff;
+					delete[] buff;
+
+					hFind = FindFirstFile(desc.tmppath.c_str(), &FindFileData);
+					if(hFind != INVALID_HANDLE_VALUE) FindClose(hFind);
+					else CreateDirectory(desc.tmppath.c_str(), NULL);
+
+					dl.push_back(desc);
+				}
+			}
+		}
+	}
+	catch(cl::Error err)
+	{
+		printf("%s (%d)\n", err.what(), err.err());
+	}
+}
+
+string GSUtil::GetDeviceUniqueName(cl::Device& device)
+{
+	std::string vendor = device.getInfo<CL_DEVICE_VENDOR>();
+	std::string name = device.getInfo<CL_DEVICE_NAME>();
+	std::string version = device.getInfo<CL_DEVICE_OPENCL_C_VERSION>();
+
+	string type;
+
+	switch(device.getInfo<CL_DEVICE_TYPE>())
+	{
+	case CL_DEVICE_TYPE_GPU: type = "GPU"; break;
+	case CL_DEVICE_TYPE_CPU: type = "CPU"; break;
+	}
+
+	version.erase(version.find_last_not_of(' ') + 1);
+
+	return vendor + " " + name + " " + version + " " + type;
+}
+#endif
+
 #ifdef _WINDOWS
 
 bool GSUtil::CheckDirectX()

plugins/GSdx/GSUtil.h

@@ -23,6 +23,16 @@
 
 #include "GS.h"
 
+struct OCLDeviceDesc
+{
+#ifdef ENABLE_OPENCL
+	cl::Device device;
+#endif
+	string name;
+	int version;
+	string tmppath;
+};
+
 class GSUtil
 {
 public:
@@ -30,6 +40,7 @@ public:
 
 	static GS_PRIM_CLASS GetPrimClass(uint32 prim);
 	static int GetVertexCount(uint32 prim);
+	static int GetClassVertexCount(uint32 primclass);
 
 	static const uint32* HasSharedBitsPtr(uint32 dpsm);
 	static bool HasSharedBits(uint32 spsm, const uint32* ptr);
@@ -39,6 +50,11 @@ public:
 
 	static bool CheckSSE();
 
+#ifdef ENABLE_OPENCL
+	static void GetDeviceDescs(list<OCLDeviceDesc>& dl);
+	static string GetDeviceUniqueName(cl::Device& device);
+#endif
+
 #ifdef _WINDOWS
 
 	static bool CheckDirectX();

plugins/GSdx/GSdx.cpp

@@ -41,8 +41,29 @@ BOOL APIENTRY DllMain(HMODULE hModule, DWORD ul_reason_for_call, LPVOID lpReserv
 	return TRUE;
 }
 
+bool GSdxApp::LoadResource(int id, vector<unsigned char>& buff, const char* type)
+{
+	buff.clear();
+	HRSRC hRsrc = FindResource((HMODULE)s_hModule, MAKEINTRESOURCE(id), type != NULL ? type : RT_RCDATA);
+	if(!hRsrc) return false;
+	HGLOBAL hGlobal = ::LoadResource((HMODULE)s_hModule, hRsrc);
+	if(!hGlobal) return false;
+	DWORD size = SizeofResource((HMODULE)s_hModule, hRsrc);
+	if(!size) return false;
+	buff.resize(size);
+	memcpy(buff.data(), LockResource(hGlobal), size);
+	return true;
+}
+
 #else
 
+bool GSdxApp::LoadResource(int id, vector<unsigned char>& buff, const char* type)
+{
+	buff.clear();
+	printf("LoadResource not implemented\n");
+	return false;
+}
+
 size_t GSdxApp::GetPrivateProfileString(const char* lpAppName, const char* lpKeyName, const char* lpDefault, char* lpReturnedString, size_t nSize, const char* lpFileName)
 {
 	BuildConfigurationMap(lpFileName);
@@ -108,20 +129,18 @@ GSdxApp::GSdxApp()
 
 	m_gs_renderers.push_back(GSSetting(0, "Direct3D9", "Hardware"));
 	m_gs_renderers.push_back(GSSetting(1, "Direct3D9", "Software"));
+	m_gs_renderers.push_back(GSSetting(14, "Direct3D9", "OpenCL"));
 	m_gs_renderers.push_back(GSSetting(2, "Direct3D9", "Null"));
-	m_gs_renderers.push_back(GSSetting(3, "Direct3D%d    ", "Hardware"));
+	m_gs_renderers.push_back(GSSetting(3, "Direct3D", "Hardware"));
-	m_gs_renderers.push_back(GSSetting(4, "Direct3D%d    ", "Software"));
+	m_gs_renderers.push_back(GSSetting(4, "Direct3D", "Software"));
-	m_gs_renderers.push_back(GSSetting(5, "Direct3D%d    ", "Null"));
+	m_gs_renderers.push_back(GSSetting(15, "Direct3D", "OpenCL"));
-#ifdef _LINUX
+	m_gs_renderers.push_back(GSSetting(5, "Direct3D", "Null"));
-	// note: SDL was removed. We keep those bits for compatibility of the renderer
-	// position in the linux dialog.
-	m_gs_renderers.push_back(GSSetting(7, "SDL 1.3", "Software"));
-	m_gs_renderers.push_back(GSSetting(8, "SDL 1.3", "Null"));
-#endif
 	m_gs_renderers.push_back(GSSetting(10, "Null", "Software"));
+	m_gs_renderers.push_back(GSSetting(16, "Null", "OpenCL"));
 	m_gs_renderers.push_back(GSSetting(11, "Null", "Null"));
 	m_gs_renderers.push_back(GSSetting(12, "OpenGL", "Hardware"));
 	m_gs_renderers.push_back(GSSetting(13, "OpenGL", "Software"));
+	m_gs_renderers.push_back(GSSetting(17, "OpenGL", "OpenCL"));
 
 	m_gs_interlace.push_back(GSSetting(0, "None", ""));
 	m_gs_interlace.push_back(GSSetting(1, "Weave tff", "saw-tooth"));

plugins/GSdx/GSdx.h

@@ -39,6 +39,7 @@ public:
 #ifdef _WINDOWS
  	HMODULE GetModuleHandle() {return (HMODULE)GetModuleHandlePtr();}
 #endif
+
 #ifdef _LINUX
 	void BuildConfigurationMap(const char* lpFileName);
 	void ReloadConfig();
@@ -48,6 +49,8 @@ public:
 	int GetPrivateProfileInt(const char* lpAppName, const char* lpKeyName, int nDefault, const char* lpFileName);
 #endif
 
+	bool LoadResource(int id, vector<unsigned char>& buff, const char* type = NULL);
+
 	string GetConfig(const char* entry, const char* value);
 	void SetConfig(const char* entry, const char* value);
 	int GetConfig(const char* entry, int value);

plugins/GSdx/GSdx.rc

@@ -51,9 +51,11 @@ BEGIN
     "#include ""res/tfx.fx""\r\n"
     "#include ""res/convert.fx""\r\n"
     "#include ""res/interlace.fx""\r\n"
-    "#include ""res/merge.fx""\r\0"
+    "#include ""res/merge.fx""\r\n"
-    "#include ""res/fxaa.fx""\r\0"
+    "#include ""res/fxaa.fx""\r\n"
-    "#include ""res/shadeboost.fx""\r\0"
+    "#include ""res/cs.fx""\r\n"
+    "#include ""res/shadeboost.fx""\r\n"
+    "#include ""res/tfx.cl""\r\0"
 END
 
 #endif    // APSTUDIO_INVOKED
@@ -64,13 +66,14 @@ END
 // RCDATA
 //
 
-IDR_CONVERT_FX          RCDATA                  "res\\convert.fx"
 IDR_TFX_FX              RCDATA                  "res\\tfx.fx"
-IDR_MERGE_FX            RCDATA                  "res\\merge.fx"
+IDR_CONVERT_FX          RCDATA                  "res\\convert.fx"
 IDR_INTERLACE_FX        RCDATA                  "res\\interlace.fx"
+IDR_MERGE_FX            RCDATA                  "res\\merge.fx"
 IDR_FXAA_FX             RCDATA                  "res\\fxaa.fx"
 IDR_CS_FX               RCDATA                  "res\\cs.fx"
 IDR_SHADEBOOST_FX       RCDATA                  "res\\shadeboost.fx"
+IDR_TFX_CL              RCDATA                  "res\\tfx.cl"
 
 /////////////////////////////////////////////////////////////////////////////
 //
@@ -227,49 +230,51 @@ BEGIN
     CONTROL         "Windowed",IDC_WINDOWED,"Button",BS_AUTOCHECKBOX | WS_TABSTOP,129,157,49,10
 END
 
-IDD_CONFIG2 DIALOGEX 0, 0, 187, 360
+IDD_CONFIG2 DIALOGEX 0, 0, 187, 370
 STYLE DS_SETFONT | DS_MODALFRAME | DS_FIXEDSYS | WS_POPUP | WS_CAPTION | WS_SYSMENU
 CAPTION "Settings..."
 FONT 8, "MS Shell Dlg", 400, 0, 0x1
 BEGIN
     CONTROL         IDB_LOGO10,IDC_LOGO11,"Static",SS_BITMAP | SS_CENTERIMAGE,6,6,173,42
-    DEFPUSHBUTTON   "OK",IDOK,40,336,50,14
+    DEFPUSHBUTTON   "OK",IDOK,40,346,50,14
     LTEXT           "Renderer:",IDC_STATIC,6,72,34,8
     COMBOBOX        IDC_RENDERER,70,70,111,118,CBS_DROPDOWNLIST | WS_VSCROLL | WS_TABSTOP
-    LTEXT           "Interlacing (F5):",IDC_STATIC,6,87,53,8
+    LTEXT           "Interlacing (F5):",IDC_STATIC,6,101,53,8
-    COMBOBOX        IDC_INTERLACE,70,85,111,98,CBS_DROPDOWNLIST | WS_VSCROLL | WS_TABSTOP
+    COMBOBOX        IDC_INTERLACE,70,99,111,98,CBS_DROPDOWNLIST | WS_VSCROLL | WS_TABSTOP
-    LTEXT           "Custom Resolution:",IDC_STATIC,26,149,65,8
+    LTEXT           "Custom Resolution:",IDC_STATIC,26,163,65,8
-    EDITTEXT        IDC_RESX_EDIT,92,147,35,13,ES_AUTOHSCROLL | ES_NUMBER
+    EDITTEXT        IDC_RESX_EDIT,92,161,35,13,ES_AUTOHSCROLL | ES_NUMBER
-    CONTROL         "",IDC_RESX,"msctls_updown32",UDS_SETBUDDYINT | UDS_ALIGNRIGHT | UDS_AUTOBUDDY | UDS_ARROWKEYS | UDS_NOTHOUSANDS,120,147,11,14
+    CONTROL         "",IDC_RESX,"msctls_updown32",UDS_SETBUDDYINT | UDS_ALIGNRIGHT | UDS_AUTOBUDDY | UDS_ARROWKEYS | UDS_NOTHOUSANDS,120,161,11,14
-    EDITTEXT        IDC_RESY_EDIT,130,147,35,13,ES_AUTOHSCROLL | ES_NUMBER
+    EDITTEXT        IDC_RESY_EDIT,130,161,35,13,ES_AUTOHSCROLL | ES_NUMBER
-    CONTROL         "",IDC_RESY,"msctls_updown32",UDS_SETBUDDYINT | UDS_ALIGNRIGHT | UDS_AUTOBUDDY | UDS_ARROWKEYS | UDS_NOTHOUSANDS,154,147,11,14
+    CONTROL         "",IDC_RESY,"msctls_updown32",UDS_SETBUDDYINT | UDS_ALIGNRIGHT | UDS_AUTOBUDDY | UDS_ARROWKEYS | UDS_NOTHOUSANDS,154,161,11,14
-    CONTROL         "Native",IDC_NATIVERES,"Button",BS_AUTOCHECKBOX | WS_TABSTOP,92,120,33,10
+    CONTROL         "Native",IDC_NATIVERES,"Button",BS_AUTOCHECKBOX | WS_TABSTOP,92,134,33,10
-    LTEXT           "Extra rendering threads:",IDC_STATIC,11,289,80,8
+    LTEXT           "Extra rendering threads:",IDC_STATIC,11,303,80,8
-    EDITTEXT        IDC_SWTHREADS_EDIT,94,287,35,13,ES_AUTOHSCROLL | ES_NUMBER
+    EDITTEXT        IDC_SWTHREADS_EDIT,94,301,35,13,ES_AUTOHSCROLL | ES_NUMBER
-    CONTROL         "",IDC_SWTHREADS,"msctls_updown32",UDS_SETBUDDYINT | UDS_ALIGNRIGHT | UDS_AUTOBUDDY | UDS_ARROWKEYS | UDS_NOTHOUSANDS,129,278,11,14
+    CONTROL         "",IDC_SWTHREADS,"msctls_updown32",UDS_SETBUDDYINT | UDS_ALIGNRIGHT | UDS_AUTOBUDDY | UDS_ARROWKEYS | UDS_NOTHOUSANDS,129,292,11,14
-    COMBOBOX        IDC_UPSCALE_MULTIPLIER,92,132,74,98,CBS_DROPDOWNLIST | WS_VSCROLL | WS_TABSTOP
+    COMBOBOX        IDC_UPSCALE_MULTIPLIER,92,146,74,98,CBS_DROPDOWNLIST | WS_VSCROLL | WS_TABSTOP
-    LTEXT           "Or use Scaling:",IDC_STATIC,38,134,49,8
+    LTEXT           "Or use Scaling:",IDC_STATIC,38,148,49,8
-    LTEXT           "Original PS2 Resolution:",IDC_STATIC,10,120,80,8
+    LTEXT           "Original PS2 Resolution:",IDC_STATIC,10,134,80,8
-    CONTROL         "Edge Anti-aliasing (AA1)",IDC_AA1,"Button",BS_AUTOCHECKBOX | WS_TABSTOP,10,305,93,10
+    CONTROL         "Edge Anti-aliasing (AA1)",IDC_AA1,"Button",BS_AUTOCHECKBOX | WS_TABSTOP,10,319,93,10
-    PUSHBUTTON      "Cancel",IDCANCEL,95,336,50,14
+    PUSHBUTTON      "Cancel",IDCANCEL,95,346,50,14
     CONTROL         IDB_LOGO9,IDC_LOGO9,"Static",SS_BITMAP | SS_CENTERIMAGE,6,6,175,44
-    GROUPBOX        "D3D Internal Resolution (can cause glitches)",IDC_STATIC,6,102,175,64,BS_CENTER
+    GROUPBOX        "D3D Internal Resolution (can cause glitches)",IDC_STATIC,6,116,175,64,BS_CENTER
-    GROUPBOX        "Software Mode Settings",IDC_STATIC,6,276,175,50,BS_CENTER
+    GROUPBOX        "Software Mode Settings",IDC_STATIC,6,290,175,50,BS_CENTER
-    GROUPBOX        "Hardware Mode Settings",IDC_STATIC,6,200,175,74,BS_CENTER
+    GROUPBOX        "Hardware Mode Settings",IDC_STATIC,6,214,175,74,BS_CENTER
-    CONTROL         "Logarithmic Z",IDC_LOGZ,"Button",BS_AUTOCHECKBOX | WS_TABSTOP,92,213,58,10
+    CONTROL         "Logarithmic Z",IDC_LOGZ,"Button",BS_AUTOCHECKBOX | WS_TABSTOP,92,227,58,10
-    CONTROL         "Alpha Correction (FBA)",IDC_FBA,"Button",BS_AUTOCHECKBOX | WS_TABSTOP,92,229,87,10
+    CONTROL         "Alpha Correction (FBA)",IDC_FBA,"Button",BS_AUTOCHECKBOX | WS_TABSTOP,92,243,87,10
-    CONTROL         "Allow 8-Bit Textures",IDC_PALTEX,"Button",BS_AUTOCHECKBOX | WS_TABSTOP,10,229,82,10
+    CONTROL         "Allow 8-Bit Textures",IDC_PALTEX,"Button",BS_AUTOCHECKBOX | WS_TABSTOP,10,243,82,10
-    CONTROL         "Texture Filtering",IDC_FILTER,"Button",BS_AUTO3STATE | WS_TABSTOP,10,213,67,10
+    CONTROL         "Texture Filtering",IDC_FILTER,"Button",BS_AUTO3STATE | WS_TABSTOP,10,227,67,10
-    CONTROL         "Enable Shade Boost",IDC_SHADEBOOST,"Button",BS_AUTOCHECKBOX | WS_TABSTOP,10,172,79,10
+    CONTROL         "Enable Shade Boost",IDC_SHADEBOOST,"Button",BS_AUTOCHECKBOX | WS_TABSTOP,10,186,79,10
-    PUSHBUTTON      "Settings...",IDC_SHADEBUTTON,92,169,75,14
+    PUSHBUTTON      "Settings...",IDC_SHADEBUTTON,92,183,75,14
-    CONTROL         "Enable HW Hacks",IDC_HACKS_ENABLED,"Button",BS_AUTOCHECKBOX | WS_TABSTOP,10,245,71,10
+    CONTROL         "Enable HW Hacks",IDC_HACKS_ENABLED,"Button",BS_AUTOCHECKBOX | WS_TABSTOP,10,259,71,10
-    PUSHBUTTON      "Configure...",IDC_HACKSBUTTON,92,242,75,14
+    PUSHBUTTON      "Configure...",IDC_HACKSBUTTON,92,256,75,14
     LTEXT           "Adapter:",IDC_STATIC,6,57,30,8
     COMBOBOX        IDC_ADAPTER,70,55,111,118,CBS_DROPDOWNLIST | WS_VSCROLL | WS_TABSTOP
-    CONTROL         "Enable FXAA",IDC_FXAA,"Button",BS_AUTOCHECKBOX | WS_TABSTOP,10,187,80,10
+    CONTROL         "Enable FXAA",IDC_FXAA,"Button",BS_AUTOCHECKBOX | WS_TABSTOP,10,201,80,10
-    CONTROL         "Enable FX Shader",IDC_SHADER_FX,"Button",BS_AUTOCHECKBOX | WS_TABSTOP,92,187,80,10
+    CONTROL         "Enable FX Shader",IDC_SHADER_FX,"Button",BS_AUTOCHECKBOX | WS_TABSTOP,92,201,80,10
-    CONTROL         "Anisotropic Filtering",IDC_ANISOTROPIC,"Button",BS_AUTOCHECKBOX | WS_TABSTOP,10,260,77,8
+    CONTROL         "Anisotropic Filtering",IDC_ANISOTROPIC,"Button",BS_AUTOCHECKBOX | WS_TABSTOP,10,274,77,8
-    COMBOBOX        IDC_AFCOMBO,93,258,35,30,CBS_DROPDOWNLIST | WS_VSCROLL | WS_TABSTOP
+    COMBOBOX        IDC_AFCOMBO,93,272,35,30,CBS_DROPDOWNLIST | WS_VSCROLL | WS_TABSTOP
+    LTEXT           "OpenCL Device:",IDC_STATIC,6,86,52,8
+    COMBOBOX        IDC_OPENCL_DEVICE,70,84,111,118,CBS_DROPDOWNLIST | WS_VSCROLL | WS_TABSTOP
 END
 
 
@@ -335,7 +340,7 @@ BEGIN
         VERTGUIDE, 11
         VERTGUIDE, 87
         TOPMARGIN, 6
-        BOTTOMMARGIN, 335
+        BOTTOMMARGIN, 360
     END
 END
 #endif    // APSTUDIO_INVOKED
@@ -394,6 +399,10 @@ END
 #include "res/convert.fx"
 #include "res/interlace.fx"
 #include "res/merge.fx"
+#include "res/fxaa.fx"
+#include "res/cs.fx"
+#include "res/shadeboost.fx"
+#include "res/tfx.cl"
 
 /////////////////////////////////////////////////////////////////////////////
 #endif    // not APSTUDIO_INVOKED

plugins/GSdx/GSdx_vs2013.vcxproj

@@ -374,6 +374,7 @@
     <Link>
       <ModuleDefinitionFile>.\GSdx.def</ModuleDefinitionFile>
       <TargetMachine>MachineX86</TargetMachine>
+      <AdditionalLibraryDirectories>$(SvnRootDir)\deps\$(Platform)\Debug;%(AdditionalLibraryDirectories)</AdditionalLibraryDirectories>
     </Link>
   </ItemDefinitionGroup>
   <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Debug SSE2|x64'">
@@ -392,6 +393,7 @@
     <Link>
       <ModuleDefinitionFile>.\GSdx.def</ModuleDefinitionFile>
       <TargetMachine>MachineX86</TargetMachine>
+      <AdditionalLibraryDirectories>$(SvnRootDir)\deps\$(Platform)\Release;%(AdditionalLibraryDirectories)</AdditionalLibraryDirectories>
     </Link>
   </ItemDefinitionGroup>
   <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Release SSE2|x64'">
@@ -410,6 +412,7 @@
     <Link>
       <ModuleDefinitionFile>.\GSdx.def</ModuleDefinitionFile>
       <TargetMachine>MachineX86</TargetMachine>
+      <AdditionalLibraryDirectories>$(SvnRootDir)\deps\$(Platform)\Release;%(AdditionalLibraryDirectories)</AdditionalLibraryDirectories>
     </Link>
   </ItemDefinitionGroup>
   <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Release SSSE3|x64'">
@@ -428,6 +431,7 @@
     <Link>
       <ModuleDefinitionFile>.\GSdx.def</ModuleDefinitionFile>
       <TargetMachine>MachineX86</TargetMachine>
+      <AdditionalLibraryDirectories>$(SvnRootDir)\deps\$(Platform)\Debug;%(AdditionalLibraryDirectories)</AdditionalLibraryDirectories>
     </Link>
   </ItemDefinitionGroup>
   <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Debug SSSE3|x64'">
@@ -446,6 +450,7 @@
     <Link>
       <ModuleDefinitionFile>.\GSdx.def</ModuleDefinitionFile>
       <TargetMachine>MachineX86</TargetMachine>
+      <AdditionalLibraryDirectories>$(SvnRootDir)\deps\$(Platform)\Debug;%(AdditionalLibraryDirectories)</AdditionalLibraryDirectories>
     </Link>
   </ItemDefinitionGroup>
   <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Debug AVX|Win32'">
@@ -455,6 +460,7 @@
     <Link>
       <ModuleDefinitionFile>.\GSdx.def</ModuleDefinitionFile>
       <TargetMachine>MachineX86</TargetMachine>
+      <AdditionalLibraryDirectories>$(SvnRootDir)\deps\$(Platform)\Debug;%(AdditionalLibraryDirectories)</AdditionalLibraryDirectories>
     </Link>
   </ItemDefinitionGroup>
   <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Debug AVX2|Win32'">
@@ -464,6 +470,7 @@
     <Link>
       <ModuleDefinitionFile>.\GSdx.def</ModuleDefinitionFile>
       <TargetMachine>MachineX86</TargetMachine>
+      <AdditionalLibraryDirectories>$(SvnRootDir)\deps\$(Platform)\Debug;%(AdditionalLibraryDirectories)</AdditionalLibraryDirectories>
     </Link>
   </ItemDefinitionGroup>
   <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Debug SSE4|x64'">
@@ -509,6 +516,7 @@
     <Link>
       <ModuleDefinitionFile>.\GSdx.def</ModuleDefinitionFile>
       <TargetMachine>MachineX86</TargetMachine>
+      <AdditionalLibraryDirectories>$(SvnRootDir)\deps\$(Platform)\Release;%(AdditionalLibraryDirectories)</AdditionalLibraryDirectories>
     </Link>
   </ItemDefinitionGroup>
   <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Release AVX|Win32'">
@@ -518,6 +526,7 @@
     <Link>
       <ModuleDefinitionFile>.\GSdx.def</ModuleDefinitionFile>
       <TargetMachine>MachineX86</TargetMachine>
+      <AdditionalLibraryDirectories>$(SvnRootDir)\deps\$(Platform)\Release;%(AdditionalLibraryDirectories)</AdditionalLibraryDirectories>
     </Link>
   </ItemDefinitionGroup>
   <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Release AVX2|Win32'">
@@ -527,6 +536,7 @@
     <Link>
       <ModuleDefinitionFile>.\GSdx.def</ModuleDefinitionFile>
       <TargetMachine>MachineX86</TargetMachine>
+      <AdditionalLibraryDirectories>$(SvnRootDir)\deps\$(Platform)\Release;%(AdditionalLibraryDirectories)</AdditionalLibraryDirectories>
     </Link>
   </ItemDefinitionGroup>
   <ItemDefinitionGroup Condition="'$(Configuration)|$(Platform)'=='Release SSE4|x64'">
@@ -687,6 +697,7 @@
       <AssemblerOutput Condition="'$(Configuration)|$(Platform)'=='Release AVX2|Win32'">AssemblyAndSourceCode</AssemblerOutput>
     </ClCompile>
     <ClCompile Include="GSRenderer.cpp" />
+    <ClCompile Include="GSRendererCL.cpp" />
     <ClCompile Include="GSRendererCS.cpp" />
     <ClCompile Include="GSRendererDX.cpp" />
     <ClCompile Include="GSRendererDX11.cpp" />
@@ -1970,6 +1981,7 @@
     <ClInclude Include="GSPerfMon.h" />
     <ClInclude Include="GSRasterizer.h" />
     <ClInclude Include="GSRenderer.h" />
+    <ClInclude Include="GSRendererCL.h" />
     <ClInclude Include="GSRendererCS.h" />
     <ClInclude Include="GSRendererDX.h" />
     <ClInclude Include="GSRendererDX11.h" />
@@ -2057,6 +2069,7 @@
     <None Include="res\interlace.fx" />
     <None Include="res\merge.fx" />
     <None Include="res\shadeboost.fx" />
+    <None Include="res\tfx.cl" />
     <None Include="res\tfx.fx" />
     <None Include="baseclasses\activex.rcv" />
     <None Include="baseclasses\activex.ver" />

plugins/GSdx/GSdx_vs2013.vcxproj.filters

@@ -348,6 +348,9 @@
     <ClCompile Include="GSSetupPrimCodeGenerator.x86.avx2.cpp">
       <Filter>Source Files</Filter>
     </ClCompile>
+    <ClCompile Include="GSRendererCL.cpp">
+      <Filter>Source Files</Filter>
+    </ClCompile>
   </ItemGroup>
   <ItemGroup>
     <ClInclude Include="GLLoader.h">
@@ -707,6 +710,9 @@
     <ClInclude Include="targetver.h">
       <Filter>Header Files</Filter>
     </ClInclude>
+    <ClInclude Include="GSRendererCL.h">
+      <Filter>Header Files</Filter>
+    </ClInclude>
   </ItemGroup>
   <ItemGroup>
     <None Include="res\logo10.bmp">
@@ -737,10 +743,13 @@
     <None Include="res\fxaa.fx">
       <Filter>Shaders</Filter>
     </None>
+    <None Include="res\shadeboost.fx">
+      <Filter>Shaders</Filter>
+    </None>
     <None Include="res\cs.fx">
       <Filter>Shaders</Filter>
     </None>
-    <None Include="res\shadeboost.fx">
+    <None Include="res\tfx.cl">
       <Filter>Shaders</Filter>
     </None>
   </ItemGroup>

plugins/GSdx/config.h

@@ -45,3 +45,7 @@
 
 // Output stencil to a color buffer
 //#define ENABLE_OGL_STENCIL_DEBUG
+
+#ifdef _WINDOWS
+#define ENABLE_OPENCL
+#endif

plugins/GSdx/res/cs.fx

@@ -1,3 +1,5 @@
+#ifdef SHADER_MODEL // make safe to include in resource file to enforce dependency
+
 #ifndef VS_TME
 #define VS_TME 1
 #define VS_FST 1
@@ -381,3 +383,5 @@ void ps_main1(GS_OUTPUT input)
 	WritePixel(addr.x, c, PS_FPSM);
 	WritePixel(addr.y, z, PS_ZPSM);
 }
+
+#endif

plugins/GSdx/resource.h

@@ -71,6 +71,7 @@
 #define IDC_SHADER_FX                   2088
 #define IDC_ANISOTROPIC                 2089
 #define IDC_AFCOMBO                     2090
+#define IDC_OPENCL_DEVICE               2091
 #define IDC_COLORSPACE                  3000
 #define IDR_CONVERT_FX                  10000
 #define IDR_TFX_FX                      10001
@@ -81,6 +82,7 @@
 #define IDR_CS_FX                       10006
 #define IDD_SHADEBOOST                  10007
 #define IDR_SHADEBOOST_FX               10009
+#define IDR_TFX_CL                      10010
 #define IDD_HACKS                       10011
 #define IDC_STATIC                      -1
 
@@ -90,7 +92,7 @@
 #ifndef APSTUDIO_READONLY_SYMBOLS
 #define _APS_NEXT_RESOURCE_VALUE        10012
 #define _APS_NEXT_COMMAND_VALUE         32771
-#define _APS_NEXT_CONTROL_VALUE         2091
+#define _APS_NEXT_CONTROL_VALUE         2092
 #define _APS_NEXT_SYMED_VALUE           5000
 #endif
 #endif

plugins/GSdx/stdafx.h

@@ -44,12 +44,22 @@
 #include <comutil.h>
 #include "../../common/include/comptr.h"
 
+
 #define D3DCOLORWRITEENABLE_RGBA (D3DCOLORWRITEENABLE_RED | D3DCOLORWRITEENABLE_GREEN | D3DCOLORWRITEENABLE_BLUE | D3DCOLORWRITEENABLE_ALPHA)
 #define D3D11_SHADER_MACRO D3D10_SHADER_MACRO
 #define ID3D11Blob ID3D10Blob
 
 #endif
 
+
+#ifdef ENABLE_OPENCL
+
+#define CL_USE_DEPRECATED_OPENCL_1_1_APIS
+#define __CL_ENABLE_EXCEPTIONS
+#include <CL/cl.hpp>
+
+#endif
+
 // put these into vc9/common7/ide/usertype.dat to have them highlighted
 
 typedef unsigned char uint8;

plugins/GSdx/vsprops/common.props

@@ -14,16 +14,16 @@
       <WarningLevel>Level4</WarningLevel>
       <DebugInformationFormat>ProgramDatabase</DebugInformationFormat>
       <DisableSpecificWarnings>4996;4995;4324;4100;4101;4201;4556;4127;4512;%(DisableSpecificWarnings)</DisableSpecificWarnings>
-      <AdditionalIncludeDirectories>$(DXSDK_DIR)include;$(VTUNE_AMPLIFIER_XE_2013_DIR)include;$(SolutionDir)3rdparty;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
+      <AdditionalIncludeDirectories>$(DXSDK_DIR)include;$(VTUNE_AMPLIFIER_XE_2015_DIR)include;$(SolutionDir)3rdparty;$(SolutionDir)3rdparty\opencl;%(AdditionalIncludeDirectories)</AdditionalIncludeDirectories>
       <MultiProcessorCompilation>true</MultiProcessorCompilation>
     </ClCompile>
     <Link>
-      <AdditionalDependencies>d3d11.lib;d3dx11.lib;d3d10_1.lib;d3dx10.lib;d3d9.lib;d3dx9.lib;dxgi.lib;dxguid.lib;winmm.lib;strmiids.lib;xinput.lib;opengl32.lib;comsuppw.lib;%(AdditionalDependencies)</AdditionalDependencies>
+      <AdditionalDependencies>d3d11.lib;d3dx11.lib;d3d10_1.lib;d3dx10.lib;d3d9.lib;d3dx9.lib;dxgi.lib;dxguid.lib;winmm.lib;strmiids.lib;xinput.lib;opengl32.lib;opencl.lib;comsuppw.lib;%(AdditionalDependencies)</AdditionalDependencies>
       <DelayLoadDLLs>d3d9.dll;d3dx9_43.dll;d3d11.dll;d3dx11_43.dll;dxgi.dll;opengl32.dll;%(DelayLoadDLLs)</DelayLoadDLLs>
       <GenerateDebugInformation>true</GenerateDebugInformation>
       <SubSystem>Windows</SubSystem>
       <RandomizedBaseAddress>false</RandomizedBaseAddress>
-      <AdditionalLibraryDirectories>$(VTUNE_AMPLIFIER_XE_2013_DIR)lib32;%(AdditionalLibraryDirectories)</AdditionalLibraryDirectories>
+      <AdditionalLibraryDirectories>$(VTUNE_AMPLIFIER_XE_2015_DIR)lib32;%(AdditionalLibraryDirectories)</AdditionalLibraryDirectories>
     </Link>
     <PostBuildEvent>
       <Command>.\postBuild.cmd "$(TargetPath)" "$(TargetName)" $(TargetExt) $(PcsxSubsection)</Command>