Implement KHR ICD extension

This adds a pointer to the dispatch table at the beginning of every object
of type

 - cl_command_queue
 - cl_context
 - cl_device_id
 - cl_event
 - cl_kernel
 - cl_mem
 - cl_platform_id
 - cl_program
 - cl_sampler

as required by the ICD specification. The layout of the dispatch table
comes from the OpenCL ICD loader by Brice Videau <brice.videau@imag.fr> and
Vincent Danjean <Vincent.Danjean@ens-lyon.org>.

To avoid dispatch table entries being overwritten with the ICD loader's
implementations of the CL functions (as would be the proper behaviour for
the ELF loader), the -Bsymbolic option is given to the linker.

Signed-off-by: Zhigang Gong <zhigang.gong@linux.intel.com>

25 files changed, 297 insertions, 1 deletions

diff --git a/CMake/FindOCLIcd.cmake b/CMake/FindOCLIcd.cmake
new file mode 100644
index 00000000..076f00e5
--- /dev/null
+++ b/CMake/FindOCLIcd.cmake
@@ -0,0 +1,23 @@
+#
+# Try to find ocl_icd library and include path.
+# Once done this will define
+#
+# OCLIcd_FOUND
+# OCLIcd_INCLUDE_PATH
+#
+
+FIND_PATH(OCLIcd_INCLUDE_PATH ocl_icd.h
+  ~/include/
+  /usr/include/
+  /usr/local/include/
+  /sw/include/
+  /opt/local/include/
+  DOC "The directory where ocl_icd.h resides")
+
+IF(OCLIcd_INCLUDE_PATH)
+  SET(OCLIcd_FOUND 1 CACHE STRING "Set to 1 if OCLIcd is found, 0 otherwise")
+ELSE(OCLIcd_INCLUDE_PATH)
+  SET(OCLIcd_FOUND 0 CACHE STRING "Set to 1 if OCLIcd is found, 0 otherwise")
+ENDIF(OCLIcd_INCLUDE_PATH)
+
+MARK_AS_ADVANCED(OCLIcd_FOUND)
diff --git a/CMakeLists.txt b/CMakeLists.txt
index cff4a67b..4b402ee4 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -130,6 +130,13 @@ ELSE(EGL_FOUND)
   MESSAGE(STATUS "Looking for EGL - not found")
 ENDIF(EGL_FOUND)
 
+Find_Package(OCLIcd)
+IF(OCLIcd_FOUND)
+  MESSAGE(STATUS "Looking for OCL ICD header file - found")
+ELSE(OCLIcd_FOUND)
+  MESSAGE(STATUS "Looking for OCL ICD header file - not found")
+ENDIF(OCLIcd_FOUND)
+
 ADD_SUBDIRECTORY(include)
 ADD_SUBDIRECTORY(backend)
 ADD_SUBDIRECTORY(src)
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
index 140a8641..2d15b90b 100644
--- a/src/CMakeLists.txt
+++ b/src/CMakeLists.txt
@@ -34,6 +34,14 @@ SET(CMAKE_CXX_FLAGS "-DHAS_EGL ${CMAKE_CXX_FLAGS}")
 SET(CMAKE_C_FLAGS "-DHAS_EGL ${CMAKE_C_FLAGS}")
 endif (EGL_FOUND)
 
+if (OCLIcd_FOUND)
+set (OPENCL_SRC ${OPENCL_SRC} cl_khr_icd.c)
+SET(CMAKE_CXX_FLAGS "-DHAS_OCLIcd ${CMAKE_CXX_FLAGS}")
+SET(CMAKE_C_FLAGS "-DHAS_OCLIcd ${CMAKE_C_FLAGS}")
+endif (OCLIcd_FOUND)
+
+SET(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} -Wl,-Bsymbolic")
+
 link_directories (${LLVM_LIBRARY_DIR})
 add_library(cl SHARED ${OPENCL_SRC})
 target_link_libraries(
diff --git a/src/cl_api.c b/src/cl_api.c
index c784d07f..c39ef832 100644
--- a/src/cl_api.c
+++ b/src/cl_api.c
@@ -30,6 +30,7 @@
 #include "cl_utils.h"
 
 #include "CL/cl.h"
+#include "CL/cl_ext.h"
 #include "CL/cl_intel.h"
 
 #include <stdio.h>
@@ -1169,7 +1170,13 @@ clEnqueueBarrier(cl_command_queue  command_queue)
 void*
 clGetExtensionFunctionAddress(const char *func_name)
 {
-  /* No extensions supported at present */
+  if (func_name == NULL)
+    return NULL;
+#ifdef HAS_OCLIcd
+  /* cl_khr_icd */
+  if (strcmp("clIcdGetPlatformIDsKHR", func_name) == 0)
+    return (void *)clIcdGetPlatformIDsKHR;
+#endif
   return NULL;
 }
 
diff --git a/src/cl_command_queue.c b/src/cl_command_queue.c
index 37e78b46..a22884f8 100644
--- a/src/cl_command_queue.c
+++ b/src/cl_command_queue.c
@@ -26,6 +26,7 @@
 #include "cl_utils.h"
 #include "cl_alloc.h"
 #include "cl_driver.h"
+#include "cl_khr_icd.h"
 
 #include <assert.h>
 #include <stdio.h>
@@ -38,6 +39,7 @@ cl_command_queue_new(cl_context ctx)
 
   assert(ctx);
   TRY_ALLOC_NO_ERR (queue, CALLOC(struct _cl_command_queue));
+  SET_ICD(queue->dispatch)
   queue->magic = CL_MAGIC_QUEUE_HEADER;
   queue->ref_n = 1;
   queue->ctx = ctx;
diff --git a/src/cl_command_queue.h b/src/cl_command_queue.h
index 1e2bcc15..6387ae1b 100644
--- a/src/cl_command_queue.h
+++ b/src/cl_command_queue.h
@@ -29,6 +29,7 @@ struct intel_gpgpu;
 
 /* Basically, this is a (kind-of) batch buffer */
 struct _cl_command_queue {
+  DEFINE_ICD(dispatch)
   uint64_t magic;              /* To identify it as a command queue */
   volatile int ref_n;          /* We reference count this object */
   cl_context ctx;              /* Its parent context */
diff --git a/src/cl_context.c b/src/cl_context.c
index d9025372..4a1925c3 100644
--- a/src/cl_context.c
+++ b/src/cl_context.c
@@ -25,6 +25,7 @@
 #include "cl_alloc.h"
 #include "cl_utils.h"
 #include "cl_driver.h"
+#include "cl_khr_icd.h"
 
 #include "CL/cl.h"
 #include "CL/cl_gl.h"
@@ -154,6 +155,7 @@ cl_context_new(struct _cl_context_prop *props)
 
   TRY_ALLOC_NO_ERR (ctx, CALLOC(struct _cl_context));
   TRY_ALLOC_NO_ERR (ctx->drv, cl_driver_new(props));
+  SET_ICD(ctx->dispatch)
   ctx->props = *props;
   ctx->magic = CL_MAGIC_CONTEXT_HEADER;
   ctx->ref_n = 1;
diff --git a/src/cl_context.h b/src/cl_context.h
index d9f2fe4e..5dff2ef7 100644
--- a/src/cl_context.h
+++ b/src/cl_context.h
@@ -23,6 +23,7 @@
 #include "cl_internals.h"
 #include "cl_driver.h"
 #include "CL/cl.h"
+#include "cl_khr_icd.h"
 
 #include <stdint.h>
 #include <pthread.h>
@@ -52,6 +53,7 @@ struct _cl_context_prop {
 
 /* Encapsulate the whole device */
 struct _cl_context {
+  DEFINE_ICD(dispatch)
   uint64_t magic;                   /* To identify it as a context */
   volatile int ref_n;               /* We reference count this object */
   cl_driver drv;                    /* Handles HW or simulator */
diff --git a/src/cl_device_id.c b/src/cl_device_id.c
index 8d47aa50..9f8e6ade 100644
--- a/src/cl_device_id.c
+++ b/src/cl_device_id.c
@@ -23,6 +23,7 @@
 #include "cl_utils.h"
 #include "cl_driver.h"
 #include "cl_device_data.h"
+#include "cl_khr_icd.h"
 #include "CL/cl.h"
 
 #include <assert.h>
@@ -30,6 +31,7 @@
 #include <string.h>
 
 static struct _cl_device_id intel_ivb_gt2_device = {
+  INIT_ICD(dispatch)
   .max_compute_unit = 128,
   .max_thread_per_unit = 8,
   .max_work_item_sizes = {512, 512, 512},
@@ -41,6 +43,7 @@ static struct _cl_device_id intel_ivb_gt2_device = {
 };
 
 static struct _cl_device_id intel_ivb_gt1_device = {
+  INIT_ICD(dispatch)
   .max_compute_unit = 64,
   .max_thread_per_unit = 8,
   .max_work_item_sizes = {512, 512, 512},
@@ -53,6 +56,7 @@ static struct _cl_device_id intel_ivb_gt1_device = {
 
 /* XXX we clone IVB for HSW now */
 static struct _cl_device_id intel_hsw_device = {
+  INIT_ICD(dispatch)
   .max_compute_unit = 64,
   .max_thread_per_unit = 8,
   .max_work_item_sizes = {512, 512, 512},
diff --git a/src/cl_device_id.h b/src/cl_device_id.h
index b7ba6b3a..610eaf6a 100644
--- a/src/cl_device_id.h
+++ b/src/cl_device_id.h
@@ -22,6 +22,7 @@
 
 /* Store complete information about the device */
 struct _cl_device_id {
+  DEFINE_ICD(dispatch)
   cl_device_type device_type;
   cl_uint  vendor_id;
   cl_uint  max_compute_unit;
diff --git a/src/cl_event.h b/src/cl_event.h
index 879357c2..23378e88 100644
--- a/src/cl_event.h
+++ b/src/cl_event.h
@@ -21,6 +21,7 @@
 #define __CL_EVENT_H__
 
 struct _cl_event {
+  DEFINE_ICD(dispatch)
 };
 
 #endif /* __CL_EVENT_H__ */
diff --git a/src/cl_extensions.c b/src/cl_extensions.c
index 7d1031fb..052b5897 100644
--- a/src/cl_extensions.c
+++ b/src/cl_extensions.c
@@ -31,6 +31,14 @@ void check_basic_extension(cl_extensions_t *extensions)
       extensions->extensions[id].base.ext_enabled = 1;
 }
 
+void check_opt1_extension(cl_extensions_t *extensions)
+{
+  int id;
+  for(id = OPT1_EXT_START_ID; id <= OPT1_EXT_END_ID; id++)
+    if (id == EXT_ID(khr_icd))
+      extensions->extensions[id].base.ext_enabled = 1;
+}
+
 void
 check_gl_extension(cl_extensions_t *extensions) {
 #ifdef HAS_EGL
@@ -101,6 +109,7 @@ cl_intel_platform_extension_init(cl_platform_id intel_platform)
     return;
   }
   check_basic_extension(&intel_extensions);
+  check_opt1_extension(&intel_extensions);
   check_gl_extension(&intel_extensions);
   check_intel_extension(&intel_extensions);
   process_extension_str(&intel_extensions);
diff --git a/src/cl_extensions.h b/src/cl_extensions.h
index 5a49cd65..51eb8e05 100644
--- a/src/cl_extensions.h
+++ b/src/cl_extensions.h
@@ -52,10 +52,13 @@ cl_khr_extension_id_max
 
 #define BASE_EXT_START_ID EXT_ID(khr_global_int32_base_atomics)
 #define BASE_EXT_END_ID EXT_ID(khr_fp64)
+#define OPT1_EXT_START_ID EXT_ID(khr_int64_base_atomics)
+#define OPT1_EXT_END_ID EXT_ID(khr_icd)
 #define GL_EXT_START_ID EXT_ID(khr_gl_sharing)
 #define GL_EXT_END_ID EXT_ID(khr_gl_msaa_sharing)
 
 #define IS_BASE_EXTENSION(id)  (id >= BASE_EXT_START_ID && id <= BASE_EXT_END_ID)
+#define IS_OPT1_EXTENSION(id)  (id >= OPT1_EXT_START_ID && id <= OPT1_EXT_END_ID)
 #define IS_GL_EXTENSION(id)    (id >= GL_EXT_START_ID && id <= GL_EXT_END_ID)
 
 struct cl_extension_base {
diff --git a/src/cl_kernel.c b/src/cl_kernel.c
index 356a8a7d..bbd4438a 100644
--- a/src/cl_kernel.c
+++ b/src/cl_kernel.c
@@ -24,6 +24,7 @@
 #include "cl_mem.h"
 #include "cl_alloc.h"
 #include "cl_utils.h"
+#include "cl_khr_icd.h"
 #include "CL/cl.h"
 #include "cl_sampler.h"
 
@@ -64,6 +65,7 @@ cl_kernel_new(cl_program p)
 {
   cl_kernel k = NULL;
   TRY_ALLOC_NO_ERR (k, CALLOC(struct _cl_kernel));
+  SET_ICD(k->dispatch)
   k->ref_n = 1;
   k->magic = CL_MAGIC_KERNEL_HEADER;
   k->program = p;
diff --git a/src/cl_kernel.h b/src/cl_kernel.h
index e444f3b6..dd98fb34 100644
--- a/src/cl_kernel.h
+++ b/src/cl_kernel.h
@@ -43,6 +43,7 @@ typedef struct cl_argument {
 
 /* One OCL function */
 struct _cl_kernel {
+  DEFINE_ICD(dispatch)
   uint64_t magic;             /* To identify it as a kernel */
   volatile int ref_n;         /* We reference count this object */
   cl_buffer bo;               /* The code itself */
diff --git a/src/cl_khr_icd.c b/src/cl_khr_icd.c
new file mode 100644
index 00000000..5f0180a3
--- /dev/null
+++ b/src/cl_khr_icd.c
@@ -0,0 +1,175 @@
+/* 
+ * Copyright © 2013 Simon Richter
+ *
+ * This library is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation; either
+ * version 2 of the License, or (at your option) any later version.
+ *
+ * This library 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
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with this library. If not, see <http://www.gnu.org/licenses/>.
+ */
+
+#include <ocl_icd.h>
+
+#include "cl_platform_id.h"
+
+/* The interop functions are not implemented in Beignet */
+#define CL_GL_INTEROP(x) NULL
+/* OpenCL 1.2 is not implemented in Beignet */
+#define CL_1_2_NOTYET(x) NULL
+
+/** Return platform list through ICD interface
+ * This code is used only if a client is linked directly against the library
+ * instead of using the ICD loader. In this case, no other implementations
+ * should exist in the process address space, so the call is equivalent to
+ * clGetPlatformIDs().
+ *
+ * @param[in]   num_entries     Number of entries allocated in return buffer
+ * @param[out]  platforms       Platform identifiers supported by this implementation
+ * @param[out]  num_platforms   Number of platform identifiers returned
+ * @return      OpenCL error code
+ * @retval      CL_SUCCESS                      Successful execution
+ * @retval      CL_PLATFORM_NOT_FOUND_KHR       No platforms provided
+ * @retval      CL_INVALID_VALUE                Invalid parameters
+ */
+cl_int
+clIcdGetPlatformIDsKHR(cl_uint          num_entries,
+                 cl_platform_id * platforms,
+                 cl_uint *        num_platforms)
+{
+  return cl_get_platform_ids(num_entries, platforms, num_platforms);
+}
+
+struct _cl_icd_dispatch const cl_khr_icd_dispatch = {
+  clGetPlatformIDs,
+  clGetPlatformInfo,
+  clGetDeviceIDs,
+  clGetDeviceInfo,
+  clCreateContext,
+  clCreateContextFromType,
+  clRetainContext,
+  clReleaseContext,
+  clGetContextInfo,
+  clCreateCommandQueue,
+  clRetainCommandQueue,
+  clReleaseCommandQueue,
+  clGetCommandQueueInfo,
+  (void *) 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,
+  clGetExtensionFunctionAddress,
+  CL_GL_INTEROP(clCreateFromGLBuffer),
+  CL_GL_INTEROP(clCreateFromGLTexture2D),
+  CL_GL_INTEROP(clCreateFromGLTexture3D),
+  CL_GL_INTEROP(clCreateFromGLRenderbuffer),
+  CL_GL_INTEROP(clGetGLObjectInfo),
+  CL_GL_INTEROP(clGetGLTextureInfo),
+  CL_GL_INTEROP(clEnqueueAcquireGLObjects),
+  CL_GL_INTEROP(clEnqueueReleaseGLObjects),
+  CL_GL_INTEROP(clGetGLContextInfoKHR),
+  (void *) NULL,
+  (void *) NULL,
+  (void *) NULL,
+  (void *) NULL,
+  (void *) NULL,
+  (void *) NULL,
+  clSetEventCallback,
+  clCreateSubBuffer,
+  clSetMemObjectDestructorCallback,
+  clCreateUserEvent,
+  clSetUserEventStatus,
+  clEnqueueReadBufferRect,
+  clEnqueueWriteBufferRect,
+  clEnqueueCopyBufferRect,
+  CL_1_2_NOTYET(clCreateSubDevicesEXT),
+  CL_1_2_NOTYET(clRetainDeviceEXT),
+  CL_1_2_NOTYET(clReleaseDeviceEXT),
+#ifdef CL_VERSION_1_2
+  (void *) NULL,
+  CL_1_2_NOTYET(clCreateSubDevices),
+  CL_1_2_NOTYET(clRetainDevice),
+  CL_1_2_NOTYET(clReleaseDevice),
+  CL_1_2_NOTYET(clCreateImage),
+  CL_1_2_NOTYET(clCreateProgramWithBuiltInKernels),
+  CL_1_2_NOTYET(clCompileProgram),
+  CL_1_2_NOTYET(clLinkProgram),
+  CL_1_2_NOTYET(clUnloadPlatformCompiler),
+  CL_1_2_NOTYET(clGetKernelArgInfo),
+  CL_1_2_NOTYET(clEnqueueFillBuffer),
+  CL_1_2_NOTYET(clEnqueueFillImage),
+  CL_1_2_NOTYET(clEnqueueMigrateMemObjects),
+  CL_1_2_NOTYET(clEnqueueMarkerWithWaitList),
+  CL_1_2_NOTYET(clEnqueueBarrierWithWaitList),
+  CL_1_2_NOTYET(clGetExtensionFunctionAddressForPlatform),
+  CL_GL_INTEROP(clCreateFromGLTexture),
+  (void *) NULL,
+  (void *) NULL,
+  (void *) NULL,
+  (void *) NULL,
+  (void *) NULL,
+  (void *) NULL,
+  (void *) NULL,
+  (void *) NULL,
+  (void *) NULL,
+  (void *) NULL,
+  (void *) NULL,
+  (void *) NULL,
+  (void *) NULL
+#endif
+};
+
diff --git a/src/cl_khr_icd.h b/src/cl_khr_icd.h
new file mode 100644
index 00000000..6c8b9f4c
--- /dev/null
+++ b/src/cl_khr_icd.h
@@ -0,0 +1,30 @@
+/* 
+ * Copyright © 2013 Simon Richter
+ *
+ * This library is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU Lesser General Public
+ * License as published by the Free Software Foundation; either
+ * version 2 of the License, or (at your option) any later version.
+ *
+ * This library 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
+ * Lesser General Public License for more details.
+ *
+ * You should have received a copy of the GNU Lesser General Public
+ * License along with this library. If not, see <http://www.gnu.org/licenses/>.
+ */
+
+#ifdef HAS_OCLIcd
+
+#define SET_ICD(dispatch) \
+  dispatch = &cl_khr_icd_dispatch;
+#define INIT_ICD(member)  .member = &cl_khr_icd_dispatch,
+#define DEFINE_ICD(member) struct _cl_icd_dispatch const *member;
+
+extern struct _cl_icd_dispatch const cl_khr_icd_dispatch;
+#else
+#define SET_ICD(dispatch)
+#define INIT_ICD(member)
+#define DEFINE_ICD(member)
+#endif
diff --git a/src/cl_mem.c b/src/cl_mem.c
index 3a8cfdda..e89aafad 100644
--- a/src/cl_mem.c
+++ b/src/cl_mem.c
@@ -24,6 +24,7 @@
 #include "cl_alloc.h"
 #include "cl_device_id.h"
 #include "cl_driver.h"
+#include "cl_khr_icd.h"
 
 #include "CL/cl.h"
 #include "CL/cl_intel.h"
@@ -63,6 +64,7 @@ cl_mem_allocate(cl_context ctx,
 
   /* Allocate and inialize the structure itself */
   TRY_ALLOC (mem, CALLOC(struct _cl_mem));
+  SET_ICD(mem->dispatch)
   mem->ref_n = 1;
   mem->magic = CL_MAGIC_MEM_HEADER;
   mem->flags = flags;
diff --git a/src/cl_mem.h b/src/cl_mem.h
index 6992454c..8e7a2dd1 100644
--- a/src/cl_mem.h
+++ b/src/cl_mem.h
@@ -32,6 +32,7 @@ typedef enum cl_image_tiling {
 
 /* Used for buffers and images */
 struct _cl_mem {
+  DEFINE_ICD(dispatch)
   uint64_t magic;           /* To identify it as a memory object */
   volatile int ref_n;       /* This object is reference counted */
   cl_buffer bo;             /* Data in GPU memory */
diff --git a/src/cl_platform_id.c b/src/cl_platform_id.c
index cd957472..2f660641 100644
--- a/src/cl_platform_id.c
+++ b/src/cl_platform_id.c
@@ -21,6 +21,7 @@
 #include "cl_internals.h"
 #include "cl_utils.h"
 #include "CL/cl.h"
+#include "CL/cl_ext.h"
 
 #include <stdlib.h>
 #include <string.h>
@@ -30,10 +31,12 @@
     .JOIN(FIELD,_sz) = sizeof(STRING) + 1,
 
 static struct _cl_platform_id intel_platform_data = {
+  INIT_ICD(dispatch)
   DECL_INFO_STRING(profile, "FULL_PROFILE")
   DECL_INFO_STRING(version, OCL_VERSION_STRING)
   DECL_INFO_STRING(name, "Experiment Intel Gen OCL Driver")
   DECL_INFO_STRING(vendor, "Intel")
+  DECL_INFO_STRING(icd_suffix_khr, "Intel")
 };
 
 #undef DECL_INFO_STRING
@@ -103,6 +106,7 @@ cl_get_platform_info(cl_platform_id    platform,
       GET_FIELD_SZ (PLATFORM_NAME,       name);
       GET_FIELD_SZ (PLATFORM_VENDOR,     vendor);
       GET_FIELD_SZ (PLATFORM_EXTENSIONS, extensions);
+      GET_FIELD_SZ (PLATFORM_ICD_SUFFIX_KHR, icd_suffix_khr);
       default: return CL_INVALID_VALUE;
     }
   }
@@ -114,6 +118,7 @@ cl_get_platform_info(cl_platform_id    platform,
     DECL_FIELD (PLATFORM_NAME,       name);
     DECL_FIELD (PLATFORM_VENDOR,     vendor);
     DECL_FIELD (PLATFORM_EXTENSIONS, extensions);
+    DECL_FIELD (PLATFORM_ICD_SUFFIX_KHR, icd_suffix_khr);
     default: return CL_INVALID_VALUE;
   }
 }
diff --git a/src/cl_platform_id.h b/src/cl_platform_id.h
index 84fd0ef8..edd3aaeb 100644
--- a/src/cl_platform_id.h
+++ b/src/cl_platform_id.h
@@ -22,19 +22,23 @@
 
 #include "cl_internals.h"
 #include "cl_extensions.h"
+#include "cl_khr_icd.h"
 #include "CL/cl.h"
 
 struct _cl_platform_id {
+  DEFINE_ICD(dispatch)
   const char *profile;
   const char *version;
   const char *name;
   const char *vendor;
   char *extensions;
+  const char *icd_suffix_khr;
   size_t profile_sz;
   size_t version_sz;
   size_t name_sz;
   size_t vendor_sz;
   size_t extensions_sz;
+  size_t icd_suffix_khr_sz;
   struct cl_extensions *internal_extensions;
 };
 
diff --git a/src/cl_program.c b/src/cl_program.c
index ecffb001..0c48ef30 100644
--- a/src/cl_program.c
+++ b/src/cl_program.c
@@ -23,6 +23,7 @@
 #include "cl_context.h"
 #include "cl_alloc.h"
 #include "cl_utils.h"
+#include "cl_khr_icd.h"
 #include "CL/cl.h"
 #include "CL/cl_intel.h"
 
@@ -91,6 +92,7 @@ cl_program_new(cl_context ctx)
 
   /* Allocate the structure */
   TRY_ALLOC_NO_ERR (p, CALLOC(struct _cl_program));
+  SET_ICD(p->dispatch)
   p->ref_n = 1;
   p->magic = CL_MAGIC_PROGRAM_HEADER;
   p->ctx = ctx;
diff --git a/src/cl_program.h b/src/cl_program.h
index fd006210..161d8587 100644
--- a/src/cl_program.h
+++ b/src/cl_program.h
@@ -38,6 +38,7 @@ enum {
 
 /* This maps an OCL file containing some kernels */
 struct _cl_program {
+  DEFINE_ICD(dispatch)
   uint64_t magic;         /* To identify it as a program */
   volatile int ref_n;     /* We reference count this object */
   gbe_program opaque;     /* (Opaque) program as ouput by the compiler */
diff --git a/src/cl_sampler.c b/src/cl_sampler.c
index fd88a772..d3e61da6 100644
--- a/src/cl_sampler.c
+++ b/src/cl_sampler.c
@@ -21,6 +21,7 @@
 #include "cl_sampler.h"
 #include "cl_utils.h"
 #include "cl_alloc.h"
+#include "cl_khr_icd.h"
 
 #include <assert.h>
 
@@ -36,6 +37,7 @@ cl_sampler_new(cl_context ctx,
 
   /* Allocate and inialize the structure itself */
   TRY_ALLOC (sampler, CALLOC(struct _cl_sampler));
+  SET_ICD(sampler->dispatch)
   sampler->ref_n = 1;
   sampler->magic = CL_MAGIC_SAMPLER_HEADER;
   sampler->normalized_coords = normalized_coords;
diff --git a/src/cl_sampler.h b/src/cl_sampler.h
index 800de4c6..da9a488e 100644
--- a/src/cl_sampler.h
+++ b/src/cl_sampler.h
@@ -25,6 +25,7 @@
 
 /* How to access images */
 struct _cl_sampler {
+  DEFINE_ICD(dispatch)
   uint64_t magic;            /* To identify it as a sampler object */
   volatile int ref_n;        /* This object is reference counted */
   cl_sampler prev, next;     /* We chain the samplers in the allocator */