[Beignet] [PATCH 4/8] Implement KHR ICD extension
Zhigang Gong
zhigang.gong at linux.intel.com
Wed Apr 17 00:51:55 PDT 2013
From: Simon Richter <Simon.Richter at hogyros.de>
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 at imag.fr> and
Vincent Danjean <Vincent.Danjean at 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 at linux.intel.com>
---
CMake/FindOCLIcd.cmake | 23 +++++++
CMakeLists.txt | 7 ++
src/CMakeLists.txt | 8 +++
src/cl_api.c | 9 ++-
src/cl_command_queue.c | 2 +
src/cl_command_queue.h | 1 +
src/cl_context.c | 2 +
src/cl_context.h | 2 +
src/cl_device_id.c | 4 ++
src/cl_device_id.h | 1 +
src/cl_event.h | 1 +
src/cl_extensions.c | 9 +++
src/cl_extensions.h | 3 +
src/cl_kernel.c | 2 +
src/cl_kernel.h | 1 +
src/cl_khr_icd.c | 175 +++++++++++++++++++++++++++++++++++++++++++++++++
src/cl_khr_icd.h | 30 +++++++++
src/cl_mem.c | 2 +
src/cl_mem.h | 1 +
src/cl_platform_id.c | 7 ++
src/cl_platform_id.h | 4 ++
src/cl_program.c | 2 +
src/cl_program.h | 1 +
src/cl_sampler.c | 2 +
src/cl_sampler.h | 1 +
25 files changed, 299 insertions(+), 1 deletion(-)
create mode 100644 CMake/FindOCLIcd.cmake
create mode 100644 src/cl_khr_icd.c
create mode 100644 src/cl_khr_icd.h
diff --git a/CMake/FindOCLIcd.cmake b/CMake/FindOCLIcd.cmake
new file mode 100644
index 0000000..076f00e
--- /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 7f37be9..9e565f3 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 140a864..2d15b90 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 c784d07..c39ef83 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 37e78b4..a22884f 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 1e2bcc1..6387ae1 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 d902537..4a1925c 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 d9f2fe4..5dff2ef 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 8d47aa5..9f8e6ad 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 b7ba6b3..610eaf6 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 879357c..23378e8 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 7d1031f..052b589 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 5a49cd6..51eb8e0 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 356a8a7..bbd4438 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 e444f3b..dd98fb3 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 0000000..5f0180a
--- /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 0000000..6c8b9f4
--- /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 3a8cfdd..e89aafa 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 6992454..8e7a2dd 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 cd95747..f216aaa 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,14 @@
.JOIN(FIELD,_sz) = sizeof(STRING) + 1,
static struct _cl_platform_id intel_platform_data = {
+#ifdef HAS_OCLIcd
+ .dispatch = &cl_khr_icd_dispatch,
+#endif
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 +108,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 +120,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 84fd0ef..edd3aae 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 ecffb00..0c48ef3 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 fd00621..161d858 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 fd88a77..d3e61da 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 800de4c..da9a488 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 */
--
1.7.11.7
More information about the Beignet
mailing list