--- /dev/null
+#
+# 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)
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)
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(
#include "cl_utils.h"
#include "CL/cl.h"
+#include "CL/cl_ext.h"
#include "CL/cl_intel.h"
#include <stdio.h>
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;
}
#include "cl_utils.h"
#include "cl_alloc.h"
#include "cl_driver.h"
+#include "cl_khr_icd.h"
#include <assert.h>
#include <stdio.h>
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;
/* 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 */
#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"
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;
#include "cl_internals.h"
#include "cl_driver.h"
#include "CL/cl.h"
+#include "cl_khr_icd.h"
#include <stdint.h>
#include <pthread.h>
/* 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 */
#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>
#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},
};
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},
/* 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},
/* 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;
#define __CL_EVENT_H__
struct _cl_event {
+ DEFINE_ICD(dispatch)
};
#endif /* __CL_EVENT_H__ */
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
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);
#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 {
#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"
{
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;
/* 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 */
--- /dev/null
+/*
+ * 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
+};
+
--- /dev/null
+/*
+ * 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
#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"
/* 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;
/* 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 */
#include "cl_internals.h"
#include "cl_utils.h"
#include "CL/cl.h"
+#include "CL/cl_ext.h"
#include <stdlib.h>
#include <string.h>
.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
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;
}
}
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;
}
}
#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;
};
#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"
/* 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;
/* 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 */
#include "cl_sampler.h"
#include "cl_utils.h"
#include "cl_alloc.h"
+#include "cl_khr_icd.h"
#include <assert.h>
/* 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;
/* 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 */