From 0118a372252180fb6b8c597dd0874fb31b477435 Mon Sep 17 00:00:00 2001 From: Simon Richter Date: Wed, 3 Apr 2013 20:32:45 +0200 Subject: [PATCH] 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 and Vincent Danjean . 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 --- 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 | 5 ++ 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, 297 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 cff4a67..4b402ee 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 @@ -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 #include @@ -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 #include @@ -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 @@ -30,6 +31,7 @@ #include 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 . + */ + +#include + +#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 . + */ + +#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..2f66064 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 #include @@ -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 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 @@ -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 */ -- 2.7.4