clover: Implement support for the ICD extension.

Tested-by: Tom Stellard <[email protected]>

17 files changed, 1046 insertions, 78 deletions

diff --git a/src/gallium/state_trackers/clover/Makefile.sources b/src/gallium/state_trackers/clover/Makefile.sources
index b8231e2cbd9..e55167733a4 100644
--- a/src/gallium/state_trackers/clover/Makefile.sources
+++ b/src/gallium/state_trackers/clover/Makefile.sources
@@ -41,6 +41,8 @@ CPP_SOURCES := \
 	core/module.hpp \
 	core/module.cpp \
 	api/util.hpp \
+	api/dispatch.hpp \
+	api/dispatch.cpp \
 	api/platform.cpp \
 	api/device.cpp \
 	api/context.cpp \
diff --git a/src/gallium/state_trackers/clover/api/context.cpp b/src/gallium/state_trackers/clover/api/context.cpp
index c6c71ceb92f..7b020a63bb7 100644
--- a/src/gallium/state_trackers/clover/api/context.cpp
+++ b/src/gallium/state_trackers/clover/api/context.cpp
@@ -26,7 +26,7 @@
 
 using namespace clover;
 
-PUBLIC cl_context
+CLOVER_API cl_context
 clCreateContext(const cl_context_properties *d_props, cl_uint num_devs,
                 const cl_device_id *d_devs,
                 void (CL_CALLBACK *pfn_notify)(const char *, const void *,
@@ -53,7 +53,7 @@ clCreateContext(const cl_context_properties *d_props, cl_uint num_devs,
    return NULL;
 }
 
-PUBLIC cl_context
+CLOVER_API cl_context
 clCreateContextFromType(const cl_context_properties *d_props,
                         cl_device_type type,
                         void (CL_CALLBACK *pfn_notify)(
@@ -79,7 +79,7 @@ clCreateContextFromType(const cl_context_properties *d_props,
    return NULL;
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clRetainContext(cl_context d_ctx) try {
    obj(d_ctx).retain();
    return CL_SUCCESS;
@@ -88,7 +88,7 @@ clRetainContext(cl_context d_ctx) try {
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clReleaseContext(cl_context d_ctx) try {
    if (obj(d_ctx).release())
       delete pobj(d_ctx);
@@ -99,7 +99,7 @@ clReleaseContext(cl_context d_ctx) try {
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clGetContextInfo(cl_context d_ctx, cl_context_info param,
                  size_t size, void *r_buf, size_t *r_size) try {
    property_buffer buf { r_buf, size, r_size };
diff --git a/src/gallium/state_trackers/clover/api/device.cpp b/src/gallium/state_trackers/clover/api/device.cpp
index dd7ef408f22..7bc8d0a7142 100644
--- a/src/gallium/state_trackers/clover/api/device.cpp
+++ b/src/gallium/state_trackers/clover/api/device.cpp
@@ -26,7 +26,7 @@
 
 using namespace clover;
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clGetDeviceIDs(cl_platform_id d_platform, cl_device_type device_type,
                cl_uint num_entries, cl_device_id *rd_devices,
                cl_uint *rnum_devices) try {
@@ -62,7 +62,7 @@ clGetDeviceIDs(cl_platform_id d_platform, cl_device_type device_type,
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clGetDeviceInfo(cl_device_id d_dev, cl_device_info param,
                 size_t size, void *r_buf, size_t *r_size) try {
    property_buffer buf { r_buf, size, r_size };
diff --git a/src/gallium/state_trackers/clover/api/dispatch.cpp b/src/gallium/state_trackers/clover/api/dispatch.cpp
new file mode 100644
index 00000000000..746372c0d64
--- /dev/null
+++ b/src/gallium/state_trackers/clover/api/dispatch.cpp
@@ -0,0 +1,147 @@
+//
+// Copyright 2013 Francisco Jerez
+//
+// Permission is hereby granted, free of charge, to any person obtaining a
+// copy of this software and associated documentation files (the "Software"),
+// to deal in the Software without restriction, including without limitation
+// the rights to use, copy, modify, merge, publish, distribute, sublicense,
+// and/or sell copies of the Software, and to permit persons to whom the
+// Software is 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 Software.
+//
+// THE SOFTWARE IS 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 SOFTWARE OR THE USE OR
+// OTHER DEALINGS IN THE SOFTWARE.
+//
+
+#include "api/dispatch.hpp"
+
+namespace clover {
+   const _cl_icd_dispatch _dispatch = {
+      clGetPlatformIDs,
+      GetPlatformInfo,
+      clGetDeviceIDs,
+      clGetDeviceInfo,
+      clCreateContext,
+      clCreateContextFromType,
+      clRetainContext,
+      clReleaseContext,
+      clGetContextInfo,
+      clCreateCommandQueue,
+      clRetainCommandQueue,
+      clReleaseCommandQueue,
+      clGetCommandQueueInfo,
+      NULL, // clSetCommandQueueProperty
+      clCreateBuffer,
+      clCreateImage2D,
+      clCreateImage3D,
+      clRetainMemObject,
+      clReleaseMemObject,
+      clGetSupportedImageFormats,
+      clGetMemObjectInfo,
+      clGetImageInfo,
+      clCreateSampler,
+      clRetainSampler,
+      clReleaseSampler,
+      clGetSamplerInfo,
+      clCreateProgramWithSource,
+      clCreateProgramWithBinary,
+      clRetainProgram,
+      clReleaseProgram,
+      clBuildProgram,
+      clUnloadCompiler,
+      clGetProgramInfo,
+      clGetProgramBuildInfo,
+      clCreateKernel,
+      clCreateKernelsInProgram,
+      clRetainKernel,
+      clReleaseKernel,
+      clSetKernelArg,
+      clGetKernelInfo,
+      clGetKernelWorkGroupInfo,
+      clWaitForEvents,
+      clGetEventInfo,
+      clRetainEvent,
+      clReleaseEvent,
+      clGetEventProfilingInfo,
+      clFlush,
+      clFinish,
+      clEnqueueReadBuffer,
+      clEnqueueWriteBuffer,
+      clEnqueueCopyBuffer,
+      clEnqueueReadImage,
+      clEnqueueWriteImage,
+      clEnqueueCopyImage,
+      clEnqueueCopyImageToBuffer,
+      clEnqueueCopyBufferToImage,
+      clEnqueueMapBuffer,
+      clEnqueueMapImage,
+      clEnqueueUnmapMemObject,
+      clEnqueueNDRangeKernel,
+      clEnqueueTask,
+      clEnqueueNativeKernel,
+      clEnqueueMarker,
+      clEnqueueWaitForEvents,
+      clEnqueueBarrier,
+      GetExtensionFunctionAddress,
+      NULL, // clCreateFromGLBuffer
+      NULL, // clCreateFromGLTexture2D
+      NULL, // clCreateFromGLTexture3D
+      NULL, // clCreateFromGLRenderbuffer
+      NULL, // clGetGLObjectInfo
+      NULL, // clGetGLTextureInfo
+      NULL, // clEnqueueAcquireGLObjects
+      NULL, // clEnqueueReleaseGLObjects
+      NULL, // clGetGLContextInfoKHR
+      NULL, // clGetDeviceIDsFromD3D10KHR
+      NULL, // clCreateFromD3D10BufferKHR
+      NULL, // clCreateFromD3D10Texture2DKHR
+      NULL, // clCreateFromD3D10Texture3DKHR
+      NULL, // clEnqueueAcquireD3D10ObjectsKHR
+      NULL, // clEnqueueReleaseD3D10ObjectsKHR
+      clSetEventCallback,
+      clCreateSubBuffer,
+      clSetMemObjectDestructorCallback,
+      clCreateUserEvent,
+      clSetUserEventStatus,
+      clEnqueueReadBufferRect,
+      clEnqueueWriteBufferRect,
+      clEnqueueCopyBufferRect,
+      NULL, // clCreateSubDevicesEXT
+      NULL, // clRetainDeviceEXT
+      NULL, // clReleaseDeviceEXT
+      NULL, // clCreateEventFromGLsyncKHR
+      NULL, // clCreateSubDevices
+      NULL, // clRetainDevice
+      NULL, // clReleaseDevice
+      NULL, // clCreateImage
+      NULL, // clCreateProgramWithBuiltInKernels
+      NULL, // clCompileProgram
+      NULL, // clLinkProgram
+      NULL, // clUnloadPlatformCompiler
+      NULL, // clGetKernelArgInfo
+      NULL, // clEnqueueFillBuffer
+      NULL, // clEnqueueFillImage
+      NULL, // clEnqueueMigrateMemObjects
+      NULL, // clEnqueueMarkerWithWaitList
+      NULL, // clEnqueueBarrierWithWaitList
+      NULL, // clGetExtensionFunctionAddressForPlatform
+      NULL, // clCreateFromGLTexture
+      NULL, // clGetDeviceIDsFromD3D11KHR
+      NULL, // clCreateFromD3D11BufferKHR
+      NULL, // clCreateFromD3D11Texture2DKHR
+      NULL, // clCreateFromD3D11Texture3DKHR
+      NULL, // clCreateFromDX9MediaSurfaceKHR
+      NULL, // clEnqueueAcquireD3D11ObjectsKHR
+      NULL, // clEnqueueReleaseD3D11ObjectsKHR
+      NULL, // clGetDeviceIDsFromDX9MediaAdapterKHR
+      NULL, // clEnqueueAcquireDX9MediaSurfacesKHR
+      NULL // clEnqueueReleaseDX9MediaSurfacesKHR
+   };
+}
diff --git a/src/gallium/state_trackers/clover/api/dispatch.hpp b/src/gallium/state_trackers/clover/api/dispatch.hpp
new file mode 100644
index 00000000000..9dfb2b9d022
--- /dev/null
+++ b/src/gallium/state_trackers/clover/api/dispatch.hpp
@@ -0,0 +1,759 @@
+//
+// Copyright 2013 Francisco Jerez
+//
+// Permission is hereby granted, free of charge, to any person obtaining a
+// copy of this software and associated documentation files (the "Software"),
+// to deal in the Software without restriction, including without limitation
+// the rights to use, copy, modify, merge, publish, distribute, sublicense,
+// and/or sell copies of the Software, and to permit persons to whom the
+// Software is 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 Software.
+//
+// THE SOFTWARE IS 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 SOFTWARE OR THE USE OR
+// OTHER DEALINGS IN THE SOFTWARE.
+//
+
+#ifndef API_DISPATCH_HPP
+#define API_DISPATCH_HPP
+
+#include "CL/cl.h"
+#include "CL/cl_ext.h"
+#include "CL/cl_gl.h"
+
+///
+/// OpenCL ICD vendor dispatch table.
+///
+/// The entry point ordering should always be in agreement with
+/// Khronos' ICD loader.
+///
+struct _cl_icd_dispatch {
+   CL_API_ENTRY cl_int (CL_API_CALL *clGetPlatformIDs)(
+      cl_uint num_entries,
+      cl_platform_id *platforms,
+      cl_uint *num_platforms);
+
+   CL_API_ENTRY 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_ENTRY 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_ENTRY 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_ENTRY 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_ENTRY 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_ENTRY cl_int (CL_API_CALL *clRetainContext)(
+      cl_context context);
+
+   CL_API_ENTRY cl_int (CL_API_CALL *clReleaseContext)(
+      cl_context context);
+
+   CL_API_ENTRY 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_ENTRY cl_command_queue (CL_API_CALL *clCreateCommandQueue)(
+      cl_context context,
+      cl_device_id device,
+      cl_command_queue_properties properties,
+      cl_int *errcode_ret);
+
+   CL_API_ENTRY cl_int (CL_API_CALL *clRetainCommandQueue)(
+      cl_command_queue command_queue);
+
+   CL_API_ENTRY cl_int (CL_API_CALL *clReleaseCommandQueue)(
+      cl_command_queue command_queue);
+
+   CL_API_ENTRY 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_ENTRY cl_int (CL_API_CALL *clSetCommandQueueProperty)(
+      cl_command_queue command_queue,
+      cl_command_queue_properties properties,
+      cl_bool enable,
+      cl_command_queue_properties *old_properties);
+
+   CL_API_ENTRY 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_ENTRY 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);
+
+   CL_API_ENTRY 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);
+
+   CL_API_ENTRY cl_int (CL_API_CALL *clRetainMemObject)(
+      cl_mem memobj);
+
+   CL_API_ENTRY cl_int (CL_API_CALL *clReleaseMemObject)(
+      cl_mem memobj);
+
+   CL_API_ENTRY 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_ENTRY 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_ENTRY 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_ENTRY 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);
+
+   CL_API_ENTRY cl_int (CL_API_CALL *clRetainSampler)(
+      cl_sampler sampler);
+
+   CL_API_ENTRY cl_int (CL_API_CALL *clReleaseSampler)(
+      cl_sampler sampler);
+
+   CL_API_ENTRY 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_ENTRY 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_ENTRY 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_ENTRY cl_int (CL_API_CALL *clRetainProgram)(
+      cl_program program);
+
+   CL_API_ENTRY cl_int (CL_API_CALL *clReleaseProgram)(
+      cl_program program);
+
+   CL_API_ENTRY 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, void *),
+      void *user_data);
+
+   CL_API_ENTRY cl_int (CL_API_CALL *clUnloadCompiler)(
+      void);
+
+   CL_API_ENTRY 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_ENTRY 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_ENTRY cl_kernel (CL_API_CALL *clCreateKernel)(
+      cl_program program,
+      const char *kernel_name,
+      cl_int *errcode_ret);
+
+   CL_API_ENTRY cl_int (CL_API_CALL *clCreateKernelsInProgram)(
+      cl_program program,
+      cl_uint num_kernels,
+      cl_kernel *kernels,
+      cl_uint *num_kernels_ret);
+
+   CL_API_ENTRY cl_int (CL_API_CALL *clRetainKernel)(
+      cl_kernel kernel);
+
+   CL_API_ENTRY cl_int (CL_API_CALL *clReleaseKernel)(
+      cl_kernel kernel);
+
+   CL_API_ENTRY cl_int (CL_API_CALL *clSetKernelArg)(
+      cl_kernel kernel,
+      cl_uint arg_index,
+      size_t arg_size,
+      const void *arg_value);
+
+   CL_API_ENTRY 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_ENTRY 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_ENTRY cl_int (CL_API_CALL *clWaitForEvents)(
+      cl_uint num_events,
+      const cl_event *event_list);
+
+   CL_API_ENTRY 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_ENTRY cl_int (CL_API_CALL *clRetainEvent)(
+      cl_event event);
+
+   CL_API_ENTRY cl_int (CL_API_CALL *clReleaseEvent)(
+      cl_event event);
+
+   CL_API_ENTRY 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_ENTRY cl_int (CL_API_CALL *clFlush)(
+      cl_command_queue command_queue);
+
+   CL_API_ENTRY cl_int (CL_API_CALL *clFinish)(
+      cl_command_queue command_queue);
+
+   CL_API_ENTRY cl_int (CL_API_CALL *clEnqueueReadBuffer)(
+      cl_command_queue command_queue,
+      cl_mem buffer,
+      cl_bool blocking_read,
+      size_t offset,
+      size_t cb,
+      void *ptr,
+      cl_uint num_events_in_wait_list,
+      const cl_event *event_wait_list,
+      cl_event *event);
+
+   CL_API_ENTRY cl_int (CL_API_CALL *clEnqueueWriteBuffer)(
+      cl_command_queue command_queue,
+      cl_mem buffer,
+      cl_bool blocking_write,
+      size_t offset,
+      size_t cb,
+      const void *ptr,
+      cl_uint num_events_in_wait_list,
+      const cl_event *event_wait_list,
+      cl_event *event);
+
+   CL_API_ENTRY 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 cb,
+      cl_uint num_events_in_wait_list,
+      const cl_event *event_wait_list,
+      cl_event *event);
+
+   CL_API_ENTRY 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_ENTRY 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_ENTRY 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_ENTRY 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_ENTRY 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_ENTRY 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 cb,
+      cl_uint num_events_in_wait_list,
+      const cl_event *event_wait_list,
+      cl_event *event,
+      cl_int *errcode_ret);
+
+   CL_API_ENTRY 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_ENTRY 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_ENTRY 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_ENTRY 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);
+
+   CL_API_ENTRY 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_ENTRY cl_int (CL_API_CALL *clEnqueueMarker)(
+      cl_command_queue command_queue,
+      cl_event *event);
+
+   CL_API_ENTRY cl_int (CL_API_CALL *clEnqueueWaitForEvents)(
+      cl_command_queue command_queue,
+      cl_uint num_events,
+      const cl_event *event_list);
+
+   CL_API_ENTRY cl_int (CL_API_CALL *clEnqueueBarrier)(
+      cl_command_queue command_queue);
+
+   CL_API_ENTRY void *(CL_API_CALL *clGetExtensionFunctionAddress)(
+      const char *function_name);
+
+   CL_API_ENTRY cl_mem (CL_API_CALL *clCreateFromGLBuffer)(
+      cl_context context,
+      cl_mem_flags flags,
+      cl_GLuint bufobj,
+      int *errcode_ret);
+
+   CL_API_ENTRY 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_API_ENTRY 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_API_ENTRY cl_mem (CL_API_CALL *clCreateFromGLRenderbuffer)(
+      cl_context context,
+      cl_mem_flags flags,
+      cl_GLuint renderbuffer,
+      cl_int *errcode_ret);
+
+   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_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_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_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_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);
+
+   void *clGetDeviceIDsFromD3D10KHR;
+   void *clCreateFromD3D10BufferKHR;
+   void *clCreateFromD3D10Texture2DKHR;
+   void *clCreateFromD3D10Texture3DKHR;
+   void *clEnqueueAcquireD3D10ObjectsKHR;
+   void *clEnqueueReleaseD3D10ObjectsKHR;
+
+   CL_API_ENTRY cl_int (CL_API_CALL *clSetEventCallback)(
+      cl_event event,
+      cl_int type,
+      void (CL_CALLBACK *pfn_notify)(cl_event, cl_int, void *),
+      void *user_data);
+
+   CL_API_ENTRY 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_ENTRY cl_int (CL_API_CALL *clSetMemObjectDestructorCallback)(
+      cl_mem memobj,
+      void (CL_CALLBACK *pfn_notify)(cl_mem, void *),
+      void *user_data);
+
+   CL_API_ENTRY cl_event (CL_API_CALL *clCreateUserEvent)(
+      cl_context context,
+      cl_int *errcode_ret);
+
+   CL_API_ENTRY cl_int (CL_API_CALL *clSetUserEventStatus)(
+      cl_event event,
+      cl_int status);
+
+   CL_API_ENTRY cl_int (CL_API_CALL *clEnqueueReadBufferRect)(
+      cl_command_queue command_queue,
+      cl_mem buffer,
+      cl_bool blocking_read,
+      const size_t *buffer_origin,
+      const size_t *host_origin,
+      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_ENTRY cl_int (CL_API_CALL *clEnqueueWriteBufferRect)(
+      cl_command_queue command_queue,
+      cl_mem buffer,
+      cl_bool blocking_read,
+      const size_t *buffer_origin,
+      const size_t *host_origin,
+      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_ENTRY 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_ENTRY cl_int (CL_API_CALL *clCreateSubDevicesEXT)(
+      cl_device_id in_device,
+      const cl_device_partition_property_ext *partition_properties,
+      cl_uint num_entries,
+      cl_device_id *out_devices,
+      cl_uint *num_devices);
+
+   CL_API_ENTRY cl_int (CL_API_CALL *clRetainDeviceEXT)(
+      cl_device_id device);
+
+   CL_API_ENTRY cl_int (CL_API_CALL *clReleaseDeviceEXT)(
+      cl_device_id device);
+
+   CL_API_ENTRY cl_event (CL_API_CALL *clCreateEventFromGLsyncKHR)(
+      cl_context context,
+      cl_GLsync sync,
+      cl_int *errcode_ret);
+
+   void *clCreateSubDevices;
+
+   CL_API_ENTRY cl_int (CL_API_CALL *clRetainDevice)(
+      cl_device_id device);
+
+   CL_API_ENTRY cl_int (CL_API_CALL *clReleaseDevice)(
+      cl_device_id device);
+
+   void *clCreateImage;
+
+   CL_API_ENTRY 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_ENTRY 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, void *),
+      void *user_data);
+
+   CL_API_ENTRY 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, void *),
+      void *user_data,
+      cl_int *errcode_ret);
+
+   CL_API_ENTRY cl_int (CL_API_CALL *clUnloadPlatformCompiler)(
+      cl_platform_id platform);
+
+   void *clGetKernelArgInfo;
+
+   CL_API_ENTRY 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 cb,
+      cl_uint num_events_in_wait_list,
+      const cl_event *event_wait_list,
+      cl_event *event);
+
+   CL_API_ENTRY cl_int (CL_API_CALL *clEnqueueFillImage)(
+      cl_command_queue command_queue,
+      cl_mem image,
+      const void *fill_color,
+      const size_t origin[3],
+      const size_t region[3],
+      cl_uint num_events_in_wait_list,
+      const cl_event *event_wait_list,
+      cl_event *event);
+
+   void *clEnqueueMigrateMemObjects;
+
+   CL_API_ENTRY 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_ENTRY 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_ENTRY void *(CL_API_CALL *clGetExtensionFunctionAddressForPlatform)(
+      cl_platform_id platform,
+      const char *function_name);
+
+   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);
+
+   void *clGetDeviceIDsFromD3D11KHR;
+   void *clCreateFromD3D11BufferKHR;
+   void *clCreateFromD3D11Texture2DKHR;
+   void *clCreateFromD3D11Texture3DKHR;
+   void *clCreateFromDX9MediaSurfaceKHR;
+   void *clEnqueueAcquireD3D11ObjectsKHR;
+   void *clEnqueueReleaseD3D11ObjectsKHR;
+   void *clGetDeviceIDsFromDX9MediaAdapterKHR;
+   void *clEnqueueAcquireDX9MediaSurfacesKHR;
+   void *clEnqueueReleaseDX9MediaSurfacesKHR;
+};
+
+namespace clover {
+   extern const _cl_icd_dispatch _dispatch;
+
+   cl_int
+   GetPlatformInfo(cl_platform_id d_platform, cl_platform_info param,
+                   size_t size, void *r_buf, size_t *r_size);
+
+   void *
+   GetExtensionFunctionAddress(const char *p_name);
+
+   cl_int
+   IcdGetPlatformIDsKHR(cl_uint num_entries, cl_platform_id *rd_platforms,
+                        cl_uint *rnum_platforms);
+}
+
+#endif
diff --git a/src/gallium/state_trackers/clover/api/event.cpp b/src/gallium/state_trackers/clover/api/event.cpp
index 4ed96b06313..1aff2c50238 100644
--- a/src/gallium/state_trackers/clover/api/event.cpp
+++ b/src/gallium/state_trackers/clover/api/event.cpp
@@ -25,7 +25,7 @@
 
 using namespace clover;
 
-PUBLIC cl_event
+CLOVER_API cl_event
 clCreateUserEvent(cl_context d_ctx, cl_int *r_errcode) try {
    auto &ctx = obj(d_ctx);
 
@@ -37,7 +37,7 @@ clCreateUserEvent(cl_context d_ctx, cl_int *r_errcode) try {
    return NULL;
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clSetUserEventStatus(cl_event d_ev, cl_int status) try {
    auto &sev = obj<soft_event>(d_ev);
 
@@ -58,7 +58,7 @@ clSetUserEventStatus(cl_event d_ev, cl_int status) try {
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clWaitForEvents(cl_uint num_evs, const cl_event *d_evs) try {
    auto evs = objs(d_evs, num_evs);
 
@@ -84,7 +84,7 @@ clWaitForEvents(cl_uint num_evs, const cl_event *d_evs) try {
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clGetEventInfo(cl_event d_ev, cl_event_info param,
                size_t size, void *r_buf, size_t *r_size) try {
    property_buffer buf { r_buf, size, r_size };
@@ -121,7 +121,7 @@ clGetEventInfo(cl_event d_ev, cl_event_info param,
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clSetEventCallback(cl_event d_ev, cl_int type,
                    void (CL_CALLBACK *pfn_notify)(cl_event, cl_int, void *),
                    void *user_data) try {
@@ -145,7 +145,7 @@ clSetEventCallback(cl_event d_ev, cl_int type,
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clRetainEvent(cl_event d_ev) try {
    obj(d_ev).retain();
    return CL_SUCCESS;
@@ -154,7 +154,7 @@ clRetainEvent(cl_event d_ev) try {
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clReleaseEvent(cl_event d_ev) try {
    if (obj(d_ev).release())
       delete pobj(d_ev);
@@ -165,7 +165,7 @@ clReleaseEvent(cl_event d_ev) try {
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clEnqueueMarker(cl_command_queue d_q, cl_event *rd_ev) try {
    auto &q = obj(d_q);
 
@@ -180,7 +180,7 @@ clEnqueueMarker(cl_command_queue d_q, cl_event *rd_ev) try {
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clEnqueueBarrier(cl_command_queue d_q) try {
    obj(d_q);
 
@@ -192,7 +192,7 @@ clEnqueueBarrier(cl_command_queue d_q) try {
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clEnqueueWaitForEvents(cl_command_queue d_q, cl_uint num_evs,
                        const cl_event *d_evs) try {
    auto &q = obj(d_q);
@@ -214,7 +214,7 @@ clEnqueueWaitForEvents(cl_command_queue d_q, cl_uint num_evs,
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clGetEventProfilingInfo(cl_event d_ev, cl_profiling_info param,
                         size_t size, void *r_buf, size_t *r_size) try {
    property_buffer buf { r_buf, size, r_size };
@@ -256,7 +256,7 @@ clGetEventProfilingInfo(cl_event d_ev, cl_profiling_info param,
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clFinish(cl_command_queue d_q) try {
    auto &q = obj(d_q);
 
diff --git a/src/gallium/state_trackers/clover/api/kernel.cpp b/src/gallium/state_trackers/clover/api/kernel.cpp
index e55a55b546e..d6129e63427 100644
--- a/src/gallium/state_trackers/clover/api/kernel.cpp
+++ b/src/gallium/state_trackers/clover/api/kernel.cpp
@@ -26,7 +26,7 @@
 
 using namespace clover;
 
-PUBLIC cl_kernel
+CLOVER_API cl_kernel
 clCreateKernel(cl_program d_prog, const char *name, cl_int *r_errcode) try {
    auto &prog = obj(d_prog);
 
@@ -47,7 +47,7 @@ clCreateKernel(cl_program d_prog, const char *name, cl_int *r_errcode) try {
    return NULL;
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clCreateKernelsInProgram(cl_program d_prog, cl_uint count,
                          cl_kernel *rd_kerns, cl_uint *r_count) try {
    auto &prog = obj(d_prog);
@@ -72,7 +72,7 @@ clCreateKernelsInProgram(cl_program d_prog, cl_uint count,
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clRetainKernel(cl_kernel d_kern) try {
    obj(d_kern).retain();
    return CL_SUCCESS;
@@ -81,7 +81,7 @@ clRetainKernel(cl_kernel d_kern) try {
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clReleaseKernel(cl_kernel d_kern) try {
    if (obj(d_kern).release())
       delete pobj(d_kern);
@@ -92,7 +92,7 @@ clReleaseKernel(cl_kernel d_kern) try {
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clSetKernelArg(cl_kernel d_kern, cl_uint idx, size_t size,
                const void *value) try {
    obj(d_kern).args().at(idx).set(size, value);
@@ -105,7 +105,7 @@ clSetKernelArg(cl_kernel d_kern, cl_uint idx, size_t size,
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clGetKernelInfo(cl_kernel d_kern, cl_kernel_info param,
                 size_t size, void *r_buf, size_t *r_size) try {
    property_buffer buf { r_buf, size, r_size };
@@ -142,7 +142,7 @@ clGetKernelInfo(cl_kernel d_kern, cl_kernel_info param,
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clGetKernelWorkGroupInfo(cl_kernel d_kern, cl_device_id d_dev,
                          cl_kernel_work_group_info param,
                          size_t size, void *r_buf, size_t *r_size) try {
@@ -243,7 +243,7 @@ namespace {
    }
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clEnqueueNDRangeKernel(cl_command_queue d_q, cl_kernel d_kern,
                        cl_uint dims, const size_t *d_grid_offset,
                        const size_t *d_grid_size, const size_t *d_block_size,
@@ -272,7 +272,7 @@ clEnqueueNDRangeKernel(cl_command_queue d_q, cl_kernel d_kern,
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clEnqueueTask(cl_command_queue d_q, cl_kernel d_kern,
               cl_uint num_deps, const cl_event *d_deps,
               cl_event *rd_ev) try {
@@ -295,7 +295,7 @@ clEnqueueTask(cl_command_queue d_q, cl_kernel d_kern,
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clEnqueueNativeKernel(cl_command_queue d_q, void (*func)(void *),
                       void *args, size_t args_size,
                       cl_uint num_mems, const cl_mem *d_mems,
diff --git a/src/gallium/state_trackers/clover/api/memory.cpp b/src/gallium/state_trackers/clover/api/memory.cpp
index b01a1bc5bd1..785a509b6d8 100644
--- a/src/gallium/state_trackers/clover/api/memory.cpp
+++ b/src/gallium/state_trackers/clover/api/memory.cpp
@@ -26,7 +26,7 @@
 
 using namespace clover;
 
-PUBLIC cl_mem
+CLOVER_API cl_mem
 clCreateBuffer(cl_context d_ctx, cl_mem_flags flags, size_t size,
                void *host_ptr, cl_int *r_errcode) try {
    auto &ctx = obj(d_ctx);
@@ -51,7 +51,7 @@ clCreateBuffer(cl_context d_ctx, cl_mem_flags flags, size_t size,
    return NULL;
 }
 
-PUBLIC cl_mem
+CLOVER_API cl_mem
 clCreateSubBuffer(cl_mem d_mem, cl_mem_flags flags,
                   cl_buffer_create_type op,
                   const void *op_info, cl_int *r_errcode) try {
@@ -87,7 +87,7 @@ clCreateSubBuffer(cl_mem d_mem, cl_mem_flags flags,
    return NULL;
 }
 
-PUBLIC cl_mem
+CLOVER_API cl_mem
 clCreateImage2D(cl_context d_ctx, cl_mem_flags flags,
                 const cl_image_format *format,
                 size_t width, size_t height, size_t row_pitch,
@@ -121,7 +121,7 @@ clCreateImage2D(cl_context d_ctx, cl_mem_flags flags,
    return NULL;
 }
 
-PUBLIC cl_mem
+CLOVER_API cl_mem
 clCreateImage3D(cl_context d_ctx, cl_mem_flags flags,
                 const cl_image_format *format,
                 size_t width, size_t height, size_t depth,
@@ -156,7 +156,7 @@ clCreateImage3D(cl_context d_ctx, cl_mem_flags flags,
    return NULL;
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clGetSupportedImageFormats(cl_context d_ctx, cl_mem_flags flags,
                            cl_mem_object_type type, cl_uint count,
                            cl_image_format *r_buf, cl_uint *r_count) try {
@@ -186,7 +186,7 @@ clGetSupportedImageFormats(cl_context d_ctx, cl_mem_flags flags,
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clGetMemObjectInfo(cl_mem d_mem, cl_mem_info param,
                    size_t size, void *r_buf, size_t *r_size) try {
    property_buffer buf { r_buf, size, r_size };
@@ -241,7 +241,7 @@ clGetMemObjectInfo(cl_mem d_mem, cl_mem_info param,
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clGetImageInfo(cl_mem d_mem, cl_image_info param,
                size_t size, void *r_buf, size_t *r_size) try {
    property_buffer buf { r_buf, size, r_size };
@@ -286,7 +286,7 @@ clGetImageInfo(cl_mem d_mem, cl_image_info param,
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clRetainMemObject(cl_mem d_mem) try {
    obj(d_mem).retain();
    return CL_SUCCESS;
@@ -295,7 +295,7 @@ clRetainMemObject(cl_mem d_mem) try {
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clReleaseMemObject(cl_mem d_mem) try {
    if (obj(d_mem).release())
       delete pobj(d_mem);
@@ -306,7 +306,7 @@ clReleaseMemObject(cl_mem d_mem) try {
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clSetMemObjectDestructorCallback(cl_mem d_mem,
                                  void (CL_CALLBACK *pfn_notify)(cl_mem, void *),
                                  void *user_data) try {
diff --git a/src/gallium/state_trackers/clover/api/platform.cpp b/src/gallium/state_trackers/clover/api/platform.cpp
index 985c2211027..81b08548158 100644
--- a/src/gallium/state_trackers/clover/api/platform.cpp
+++ b/src/gallium/state_trackers/clover/api/platform.cpp
@@ -29,7 +29,7 @@ namespace {
    platform _clover_platform;
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clGetPlatformIDs(cl_uint num_entries, cl_platform_id *rd_platforms,
                  cl_uint *rnum_platforms) {
    if ((!num_entries && rd_platforms) ||
@@ -44,9 +44,9 @@ clGetPlatformIDs(cl_uint num_entries, cl_platform_id *rd_platforms,
    return CL_SUCCESS;
 }
 
-PUBLIC cl_int
-clGetPlatformInfo(cl_platform_id d_platform, cl_platform_info param,
-                  size_t size, void *r_buf, size_t *r_size) try {
+cl_int
+clover::GetPlatformInfo(cl_platform_id d_platform, cl_platform_info param,
+                        size_t size, void *r_buf, size_t *r_size) try {
    property_buffer buf { r_buf, size, r_size };
 
    obj(d_platform);
@@ -69,7 +69,11 @@ clGetPlatformInfo(cl_platform_id d_platform, cl_platform_info param,
       break;
 
    case CL_PLATFORM_EXTENSIONS:
-      buf.as_string() = "";
+      buf.as_string() = "cl_khr_icd";
+      break;
+
+   case CL_PLATFORM_ICD_SUFFIX_KHR:
+      buf.as_string() = "MESA";
       break;
 
    default:
@@ -81,3 +85,36 @@ clGetPlatformInfo(cl_platform_id d_platform, cl_platform_info param,
 } catch (error &e) {
    return e.get();
 }
+
+void *
+clover::GetExtensionFunctionAddress(const char *p_name) {
+   std::string name { p_name };
+
+   if (name == "clIcdGetPlatformIDsKHR")
+      return reinterpret_cast<void *>(IcdGetPlatformIDsKHR);
+   else
+      return NULL;
+}
+
+cl_int
+clover::IcdGetPlatformIDsKHR(cl_uint num_entries, cl_platform_id *rd_platforms,
+                             cl_uint *rnum_platforms) {
+   return clGetPlatformIDs(num_entries, rd_platforms, rnum_platforms);
+}
+
+CLOVER_ICD_API cl_int
+clGetPlatformInfo(cl_platform_id d_platform, cl_platform_info param,
+                  size_t size, void *r_buf, size_t *r_size) {
+   return GetPlatformInfo(d_platform, param, size, r_buf, r_size);
+}
+
+CLOVER_ICD_API void *
+clGetExtensionFunctionAddress(const char *p_name) {
+   return GetExtensionFunctionAddress(p_name);
+}
+
+CLOVER_ICD_API cl_int
+clIcdGetPlatformIDsKHR(cl_uint num_entries, cl_platform_id *rd_platforms,
+                       cl_uint *rnum_platforms) {
+   return IcdGetPlatformIDsKHR(num_entries, rd_platforms, rnum_platforms);
+}
diff --git a/src/gallium/state_trackers/clover/api/program.cpp b/src/gallium/state_trackers/clover/api/program.cpp
index 01133813f98..70835762b46 100644
--- a/src/gallium/state_trackers/clover/api/program.cpp
+++ b/src/gallium/state_trackers/clover/api/program.cpp
@@ -25,7 +25,7 @@
 
 using namespace clover;
 
-PUBLIC cl_program
+CLOVER_API cl_program
 clCreateProgramWithSource(cl_context d_ctx, cl_uint count,
                           const char **strings, const size_t *lengths,
                           cl_int *r_errcode) try {
@@ -51,7 +51,7 @@ clCreateProgramWithSource(cl_context d_ctx, cl_uint count,
    return NULL;
 }
 
-PUBLIC cl_program
+CLOVER_API cl_program
 clCreateProgramWithBinary(cl_context d_ctx, cl_uint n,
                           const cl_device_id *d_devs,
                           const size_t *lengths,
@@ -106,7 +106,7 @@ clCreateProgramWithBinary(cl_context d_ctx, cl_uint n,
    return NULL;
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clRetainProgram(cl_program d_prog) try {
    obj(d_prog).retain();
    return CL_SUCCESS;
@@ -115,7 +115,7 @@ clRetainProgram(cl_program d_prog) try {
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clReleaseProgram(cl_program d_prog) try {
    if (obj(d_prog).release())
       delete pobj(d_prog);
@@ -126,7 +126,7 @@ clReleaseProgram(cl_program d_prog) try {
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clBuildProgram(cl_program d_prog, cl_uint num_devs,
                const cl_device_id *d_devs, const char *p_opts,
                void (*pfn_notify)(cl_program, void *),
@@ -152,12 +152,12 @@ clBuildProgram(cl_program d_prog, cl_uint num_devs,
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clUnloadCompiler() {
    return CL_SUCCESS;
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clGetProgramInfo(cl_program d_prog, cl_program_info param,
                  size_t size, void *r_buf, size_t *r_size) try {
    property_buffer buf { r_buf, size, r_size };
@@ -214,7 +214,7 @@ clGetProgramInfo(cl_program d_prog, cl_program_info param,
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clGetProgramBuildInfo(cl_program d_prog, cl_device_id d_dev,
                       cl_program_build_info param,
                       size_t size, void *r_buf, size_t *r_size) try {
diff --git a/src/gallium/state_trackers/clover/api/queue.cpp b/src/gallium/state_trackers/clover/api/queue.cpp
index 0fdb7356b35..1c251b01589 100644
--- a/src/gallium/state_trackers/clover/api/queue.cpp
+++ b/src/gallium/state_trackers/clover/api/queue.cpp
@@ -25,7 +25,7 @@
 
 using namespace clover;
 
-PUBLIC cl_command_queue
+CLOVER_API cl_command_queue
 clCreateCommandQueue(cl_context d_ctx, cl_device_id d_dev,
                      cl_command_queue_properties props,
                      cl_int *r_errcode) try {
@@ -47,7 +47,7 @@ clCreateCommandQueue(cl_context d_ctx, cl_device_id d_dev,
    return NULL;
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clRetainCommandQueue(cl_command_queue d_q) try {
    obj(d_q).retain();
    return CL_SUCCESS;
@@ -56,7 +56,7 @@ clRetainCommandQueue(cl_command_queue d_q) try {
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clReleaseCommandQueue(cl_command_queue d_q) try {
    if (obj(d_q).release())
       delete pobj(d_q);
@@ -67,7 +67,7 @@ clReleaseCommandQueue(cl_command_queue d_q) try {
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clGetCommandQueueInfo(cl_command_queue d_q, cl_command_queue_info param,
                       size_t size, void *r_buf, size_t *r_size) try {
    property_buffer buf { r_buf, size, r_size };
@@ -100,7 +100,7 @@ clGetCommandQueueInfo(cl_command_queue d_q, cl_command_queue_info param,
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clFlush(cl_command_queue d_q) try {
    obj(d_q).flush();
    return CL_SUCCESS;
diff --git a/src/gallium/state_trackers/clover/api/sampler.cpp b/src/gallium/state_trackers/clover/api/sampler.cpp
index 288df080eba..5a04d2b745d 100644
--- a/src/gallium/state_trackers/clover/api/sampler.cpp
+++ b/src/gallium/state_trackers/clover/api/sampler.cpp
@@ -25,7 +25,7 @@
 
 using namespace clover;
 
-PUBLIC cl_sampler
+CLOVER_API cl_sampler
 clCreateSampler(cl_context d_ctx, cl_bool norm_mode,
                 cl_addressing_mode addr_mode, cl_filter_mode filter_mode,
                 cl_int *r_errcode) try {
@@ -39,7 +39,7 @@ clCreateSampler(cl_context d_ctx, cl_bool norm_mode,
    return NULL;
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clRetainSampler(cl_sampler d_s) try {
    obj(d_s).retain();
    return CL_SUCCESS;
@@ -48,7 +48,7 @@ clRetainSampler(cl_sampler d_s) try {
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clReleaseSampler(cl_sampler d_s) try {
    if (obj(d_s).release())
       delete pobj(d_s);
@@ -59,7 +59,7 @@ clReleaseSampler(cl_sampler d_s) try {
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clGetSamplerInfo(cl_sampler d_s, cl_sampler_info param,
                  size_t size, void *r_buf, size_t *r_size) try {
    property_buffer buf { r_buf, size, r_size };
diff --git a/src/gallium/state_trackers/clover/api/transfer.cpp b/src/gallium/state_trackers/clover/api/transfer.cpp
index e6fa26ca4e4..b16912072f4 100644
--- a/src/gallium/state_trackers/clover/api/transfer.cpp
+++ b/src/gallium/state_trackers/clover/api/transfer.cpp
@@ -131,7 +131,7 @@ namespace {
    }
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clEnqueueReadBuffer(cl_command_queue d_q, cl_mem d_mem, cl_bool blocking,
                     size_t offset, size_t size, void *ptr,
                     cl_uint num_deps, const cl_event *d_deps,
@@ -159,7 +159,7 @@ clEnqueueReadBuffer(cl_command_queue d_q, cl_mem d_mem, cl_bool blocking,
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clEnqueueWriteBuffer(cl_command_queue d_q, cl_mem d_mem, cl_bool blocking,
                      size_t offset, size_t size, const void *ptr,
                      cl_uint num_deps, const cl_event *d_deps,
@@ -187,7 +187,7 @@ clEnqueueWriteBuffer(cl_command_queue d_q, cl_mem d_mem, cl_bool blocking,
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clEnqueueReadBufferRect(cl_command_queue d_q, cl_mem d_mem, cl_bool blocking,
                         const size_t *obj_origin,
                         const size_t *host_origin,
@@ -222,7 +222,7 @@ clEnqueueReadBufferRect(cl_command_queue d_q, cl_mem d_mem, cl_bool blocking,
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clEnqueueWriteBufferRect(cl_command_queue d_q, cl_mem d_mem, cl_bool blocking,
                          const size_t *obj_origin,
                          const size_t *host_origin,
@@ -257,7 +257,7 @@ clEnqueueWriteBufferRect(cl_command_queue d_q, cl_mem d_mem, cl_bool blocking,
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clEnqueueCopyBuffer(cl_command_queue d_q, cl_mem d_src_mem, cl_mem d_dst_mem,
                     size_t src_offset, size_t dst_offset, size_t size,
                     cl_uint num_deps, const cl_event *d_deps,
@@ -282,7 +282,7 @@ clEnqueueCopyBuffer(cl_command_queue d_q, cl_mem d_src_mem, cl_mem d_dst_mem,
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clEnqueueCopyBufferRect(cl_command_queue d_q, cl_mem d_src_mem,
                         cl_mem d_dst_mem,
                         const size_t *src_origin, const size_t *dst_origin,
@@ -314,7 +314,7 @@ clEnqueueCopyBufferRect(cl_command_queue d_q, cl_mem d_src_mem,
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clEnqueueReadImage(cl_command_queue d_q, cl_mem d_mem, cl_bool blocking,
                    const size_t *origin, const size_t *region,
                    size_t row_pitch, size_t slice_pitch, void *ptr,
@@ -345,7 +345,7 @@ clEnqueueReadImage(cl_command_queue d_q, cl_mem d_mem, cl_bool blocking,
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clEnqueueWriteImage(cl_command_queue d_q, cl_mem d_mem, cl_bool blocking,
                     const size_t *origin, const size_t *region,
                     size_t row_pitch, size_t slice_pitch, const void *ptr,
@@ -376,7 +376,7 @@ clEnqueueWriteImage(cl_command_queue d_q, cl_mem d_mem, cl_bool blocking,
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clEnqueueCopyImage(cl_command_queue d_q, cl_mem d_src_mem, cl_mem d_dst_mem,
                    const size_t *src_origin, const size_t *dst_origin,
                    const size_t *region,
@@ -403,7 +403,7 @@ clEnqueueCopyImage(cl_command_queue d_q, cl_mem d_src_mem, cl_mem d_dst_mem,
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clEnqueueCopyImageToBuffer(cl_command_queue d_q,
                            cl_mem d_src_mem, cl_mem d_dst_mem,
                            const size_t *src_origin, const size_t *region,
@@ -433,7 +433,7 @@ clEnqueueCopyImageToBuffer(cl_command_queue d_q,
    return e.get();
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clEnqueueCopyBufferToImage(cl_command_queue d_q,
                            cl_mem d_src_mem, cl_mem d_dst_mem,
                            size_t src_offset,
@@ -463,7 +463,7 @@ clEnqueueCopyBufferToImage(cl_command_queue d_q,
    return e.get();
 }
 
-PUBLIC void *
+CLOVER_API void *
 clEnqueueMapBuffer(cl_command_queue d_q, cl_mem d_mem, cl_bool blocking,
                    cl_map_flags flags, size_t offset, size_t size,
                    cl_uint num_deps, const cl_event *d_deps,
@@ -489,7 +489,7 @@ clEnqueueMapBuffer(cl_command_queue d_q, cl_mem d_mem, cl_bool blocking,
    return NULL;
 }
 
-PUBLIC void *
+CLOVER_API void *
 clEnqueueMapImage(cl_command_queue d_q, cl_mem d_mem, cl_bool blocking,
                   cl_map_flags flags,
                   const size_t *origin, const size_t *region,
@@ -514,7 +514,7 @@ clEnqueueMapImage(cl_command_queue d_q, cl_mem d_mem, cl_bool blocking,
    return NULL;
 }
 
-PUBLIC cl_int
+CLOVER_API cl_int
 clEnqueueUnmapMemObject(cl_command_queue d_q, cl_mem d_mem, void *ptr,
                         cl_uint num_deps, const cl_event *d_deps,
                         cl_event *rd_ev) try {
diff --git a/src/gallium/state_trackers/clover/api/util.hpp b/src/gallium/state_trackers/clover/api/util.hpp
index 88f6548f23f..60c8709cc05 100644
--- a/src/gallium/state_trackers/clover/api/util.hpp
+++ b/src/gallium/state_trackers/clover/api/util.hpp
@@ -29,6 +29,14 @@
 #include "core/property.hpp"
 #include "util/algorithm.hpp"
 
+#ifdef HAVE_CLOVER_ICD
+#define CLOVER_API
+#define CLOVER_ICD_API PUBLIC
+#else
+#define CLOVER_API PUBLIC
+#define CLOVER_ICD_API PUBLIC
+#endif
+
 namespace clover {
    ///
    /// Return an error code in \a p if non-zero.
diff --git a/src/gallium/state_trackers/clover/core/object.hpp b/src/gallium/state_trackers/clover/core/object.hpp
index ea71879512c..b8b3da81eb7 100644
--- a/src/gallium/state_trackers/clover/core/object.hpp
+++ b/src/gallium/state_trackers/clover/core/object.hpp
@@ -31,6 +31,7 @@
 
 #include "core/error.hpp"
 #include "core/property.hpp"
+#include "api/dispatch.hpp"
 
 ///
 /// Main namespace of the CL state tracker.
@@ -43,6 +44,13 @@ namespace clover {
    struct descriptor {
       typedef T object_type;
       typedef S descriptor_type;
+
+      descriptor() : dispatch(&_dispatch) {
+         static_assert(std::is_standard_layout<descriptor_type>::value,
+                       "ICD requires CL API objects to be standard layout.");
+      }
+
+      const _cl_icd_dispatch *dispatch;
    };
 
    struct default_tag;
@@ -57,7 +65,8 @@ namespace clover {
          static void
          validate(D *d) {
             auto o = static_cast<typename D::object_type *>(d);
-            if (!o || !dynamic_cast<object_type *>(o))
+            if (!o || o->dispatch != &_dispatch ||
+                !dynamic_cast<object_type *>(o))
                throw invalid_object_error<T>();
          }
 
@@ -74,7 +83,7 @@ namespace clover {
 
          static void
          validate(D *d) {
-            if (!d)
+            if (!d || d->dispatch != &_dispatch)
                throw invalid_object_error<object_type>();
          }
 
@@ -91,7 +100,7 @@ namespace clover {
 
          static void
          validate(D *d) {
-            if (!d)
+            if (!d || d->dispatch != &_dispatch)
                throw invalid_wait_list_error();
          }
 
diff --git a/src/gallium/targets/opencl/Makefile.am b/src/gallium/targets/opencl/Makefile.am
index cdc7500f315..653302c3a57 100644
--- a/src/gallium/targets/opencl/Makefile.am
+++ b/src/gallium/targets/opencl/Makefile.am
@@ -33,6 +33,11 @@ lib@OPENCL_LIBNAME@_la_SOURCES =
 # Force usage of a C++ linker
 nodist_EXTRA_lib@OPENCL_LIBNAME@_la_SOURCES = dummy.cpp
 
+if HAVE_CLOVER_ICD
+icddir = /etc/OpenCL/vendors/
+icd_DATA = mesa.icd
+endif
+
 # Provide compatibility with scripts for the old Mesa build system for
 # a while by putting a link to the driver into /lib of the build tree.
 all-local: lib@[email protected]
diff --git a/src/gallium/targets/opencl/mesa.icd b/src/gallium/targets/opencl/mesa.icd
new file mode 100644
index 00000000000..6a6a8706d7c
--- /dev/null
+++ b/src/gallium/targets/opencl/mesa.icd
@@ -0,0 +1 @@
+libMesaOpenCL.so