[Mesa-dev] [PATCH 2/2] clover: Implement the ICD
Tom Stellard
tom at stellard.net
Fri Jul 12 09:10:16 PDT 2013
From: Tom Stellard <thomas.stellard at amd.com>
Use the --enable-opencl-icd option to enable clover to be used with
an icd loader.
---
configure.ac | 21 ++
src/gallium/state_trackers/clover/Makefile.am | 10 +-
src/gallium/state_trackers/clover/api/context.cpp | 18 +-
src/gallium/state_trackers/clover/api/device.cpp | 12 +-
src/gallium/state_trackers/clover/api/event.cpp | 56 +++--
src/gallium/state_trackers/clover/api/icd.cpp | 160 +++++++++++++
src/gallium/state_trackers/clover/api/icd.hpp | 258 +++++++++++++++++++++
src/gallium/state_trackers/clover/api/kernel.cpp | 56 +++--
src/gallium/state_trackers/clover/api/memory.cpp | 50 ++--
src/gallium/state_trackers/clover/api/platform.cpp | 61 ++++-
src/gallium/state_trackers/clover/api/program.cpp | 55 +++--
src/gallium/state_trackers/clover/api/queue.cpp | 23 +-
src/gallium/state_trackers/clover/api/sampler.cpp | 18 +-
src/gallium/state_trackers/clover/api/transfer.cpp | 107 ++++++---
src/gallium/state_trackers/clover/core/context.cpp | 10 +-
src/gallium/state_trackers/clover/core/context.hpp | 12 +-
src/gallium/state_trackers/clover/core/device.cpp | 71 ++++--
src/gallium/state_trackers/clover/core/device.hpp | 28 ++-
src/gallium/state_trackers/clover/core/event.cpp | 31 +--
src/gallium/state_trackers/clover/core/event.hpp | 30 +--
src/gallium/state_trackers/clover/core/format.cpp | 4 +-
src/gallium/state_trackers/clover/core/format.hpp | 5 +-
src/gallium/state_trackers/clover/core/kernel.cpp | 116 ++++-----
src/gallium/state_trackers/clover/core/kernel.hpp | 7 +-
src/gallium/state_trackers/clover/core/memory.cpp | 20 +-
src/gallium/state_trackers/clover/core/memory.hpp | 19 +-
.../state_trackers/clover/core/platform.cpp | 3 +-
.../state_trackers/clover/core/platform.hpp | 7 +-
src/gallium/state_trackers/clover/core/program.cpp | 20 +-
src/gallium/state_trackers/clover/core/program.hpp | 9 +-
src/gallium/state_trackers/clover/core/queue.cpp | 11 +-
src/gallium/state_trackers/clover/core/queue.hpp | 21 +-
.../state_trackers/clover/core/resource.hpp | 2 +-
src/gallium/state_trackers/clover/core/sampler.cpp | 14 +-
src/gallium/state_trackers/clover/core/sampler.hpp | 10 +-
src/gallium/targets/opencl/Makefile.am | 22 +-
36 files changed, 1015 insertions(+), 362 deletions(-)
create mode 100644 src/gallium/state_trackers/clover/api/icd.cpp
create mode 100644 src/gallium/state_trackers/clover/api/icd.hpp
diff --git a/configure.ac b/configure.ac
index 3a0cd77..3ee0643 100644
--- a/configure.ac
+++ b/configure.ac
@@ -612,6 +612,12 @@ AC_ARG_ENABLE([opencl],
@<:@default=no@:>@])],
[],
[enable_opencl=no])
+AC_ARG_ENABLE([opencl_icd],
+ [AS_HELP_STRING([--enable-opencl-icd],
+ [Build an OpenCL library that can be loaded by an ICD implementation
+ @<:@default=auto@:>@])],
+ [enable_opencl_icd="$enableval"],
+ [enable_opencl_icd=no])
AC_ARG_ENABLE([xlib_glx],
[AS_HELP_STRING([--enable-xlib-glx],
[make GLX library Xlib-based instead of DRI-based @<:@default=disabled@:>@])],
@@ -1375,8 +1381,23 @@ if test "x$enable_opencl" = xyes; then
GALLIUM_STATE_TRACKERS_DIRS="$GALLIUM_STATE_TRACKERS_DIRS clover"
GALLIUM_TARGET_DIRS="$GALLIUM_TARGET_DIRS opencl"
enable_gallium_loader=yes
+
+ if test "x$enable_opencl_icd" = xyes; then
+ PKG_CHECK_MODULES([OCL_ICD], [ocl-icd],
+ [], [AC_MSG_ERROR([ocl-icd not found.
+ If you want to use an ICD loader (recommended), please
+ install it using a distro provided package, or from source
+ (https://forge.imag.fr/projects/ocl-icd/). If you do not want to use an ICD
+ loader, then configure mesa with --disable-opencl-icd.)])])
+ OPENCL_LIBNAME="MesaOpenCL"
+ else
+ OPENCL_LIBNAME="OpenCL"
+ fi
+
fi
AM_CONDITIONAL(HAVE_CLOVER, test "x$enable_opencl" = xyes)
+AM_CONDITIONAL(HAVE_CLOVER_ICD, test "x$enable_opencl_icd" = xyes)
+AC_SUBST([OPENCL_LIBNAME])
dnl
dnl Gallium configuration
diff --git a/src/gallium/state_trackers/clover/Makefile.am b/src/gallium/state_trackers/clover/Makefile.am
index b4c197a..ca8275e 100644
--- a/src/gallium/state_trackers/clover/Makefile.am
+++ b/src/gallium/state_trackers/clover/Makefile.am
@@ -32,6 +32,13 @@ libclllvm_la_SOURCES = \
libclover_la_CXXFLAGS = \
-std=c++0x
+if HAVE_CLOVER_ICD
+
+libclover_la_CXXFLAGS += -DUSE_ICD
+
+endif
+
+
libclover_la_LIBADD = \
libcltgsi.la libclllvm.la
@@ -74,7 +81,8 @@ libclover_la_SOURCES = \
api/sampler.cpp \
api/event.cpp \
api/program.cpp \
- api/kernel.cpp
+ api/kernel.cpp \
+ api/icd.cpp
cldir = $(includedir)/CL
cl_HEADERS = \
diff --git a/src/gallium/state_trackers/clover/api/context.cpp b/src/gallium/state_trackers/clover/api/context.cpp
index 3717441..725ad1c 100644
--- a/src/gallium/state_trackers/clover/api/context.cpp
+++ b/src/gallium/state_trackers/clover/api/context.cpp
@@ -20,6 +20,7 @@
// OTHER DEALINGS IN THE SOFTWARE.
//
+#include "api/icd.hpp"
#include "api/util.hpp"
#include "core/context.hpp"
@@ -46,9 +47,8 @@ clCreateContext(const cl_context_properties *props, cl_uint num_devs,
}
ret_error(errcode_ret, CL_SUCCESS);
- return new context(
- property_vector(mprops),
- std::vector<cl_device_id>(devs, devs + num_devs));
+ return WRAP_ICD_CONTEXT(new context(
+ property_vector(mprops), device::unwrap_icd_list(devs, num_devs)));
} catch(error &e) {
ret_error(errcode_ret, e);
@@ -82,7 +82,8 @@ clCreateContextFromType(const cl_context_properties *props,
}
PUBLIC cl_int
-clRetainContext(cl_context ctx) {
+clRetainContext(cl_context _ctx) {
+ UNWRAP_ICD_PARAM_CONTEXT(ctx);
if (!ctx)
return CL_INVALID_CONTEXT;
@@ -91,7 +92,8 @@ clRetainContext(cl_context ctx) {
}
PUBLIC cl_int
-clReleaseContext(cl_context ctx) {
+clReleaseContext(cl_context _ctx) {
+ UNWRAP_ICD_PARAM_CONTEXT(ctx);
if (!ctx)
return CL_INVALID_CONTEXT;
@@ -102,8 +104,9 @@ clReleaseContext(cl_context ctx) {
}
PUBLIC cl_int
-clGetContextInfo(cl_context ctx, cl_context_info param,
+clGetContextInfo(cl_context _ctx, cl_context_info param,
size_t size, void *buf, size_t *size_ret) {
+ UNWRAP_ICD_PARAM_CONTEXT(ctx);
if (!ctx)
return CL_INVALID_CONTEXT;
@@ -115,7 +118,8 @@ clGetContextInfo(cl_context ctx, cl_context_info param,
return scalar_property<cl_uint>(buf, size, size_ret, ctx->devs.size());
case CL_CONTEXT_DEVICES:
- return vector_property<cl_device_id>(buf, size, size_ret, ctx->devs);
+ return vector_property<cl_device_id>(buf, size, size_ret,
+ clover::device::wrap_icd_list(ctx->devs));
case CL_CONTEXT_PROPERTIES:
return vector_property<cl_context_properties>(buf, size, size_ret,
diff --git a/src/gallium/state_trackers/clover/api/device.cpp b/src/gallium/state_trackers/clover/api/device.cpp
index e132656..669a186 100644
--- a/src/gallium/state_trackers/clover/api/device.cpp
+++ b/src/gallium/state_trackers/clover/api/device.cpp
@@ -20,6 +20,7 @@
// OTHER DEALINGS IN THE SOFTWARE.
//
+#include "api/icd.hpp"
#include "api/util.hpp"
#include "core/platform.hpp"
#include "core/device.hpp"
@@ -27,9 +28,10 @@
using namespace clover;
PUBLIC cl_int
-clGetDeviceIDs(cl_platform_id platform, cl_device_type device_type,
+clGetDeviceIDs(cl_platform_id _platform, cl_device_type device_type,
cl_uint num_entries, cl_device_id *devices,
cl_uint *num_devices) {
+ UNWRAP_ICD_PARAM_PLATFORM(platform);
std::vector<cl_device_id> devs;
if ((!num_entries && devices) ||
@@ -41,7 +43,7 @@ clGetDeviceIDs(cl_platform_id platform, cl_device_type device_type,
if (((device_type & CL_DEVICE_TYPE_DEFAULT) &&
&dev == &platform->front()) ||
(device_type & dev.type()))
- devs.push_back(&dev);
+ devs.push_back(WRAP_ICD_DEVICE(&dev));
}
if (devs.empty())
@@ -59,8 +61,9 @@ clGetDeviceIDs(cl_platform_id platform, cl_device_type device_type,
}
PUBLIC cl_int
-clGetDeviceInfo(cl_device_id dev, cl_device_info param,
+clGetDeviceInfo(cl_device_id _dev, cl_device_info param,
size_t size, void *buf, size_t *size_ret) {
+ UNWRAP_ICD_PARAM_DEVICE(dev);
if (!dev)
return CL_INVALID_DEVICE;
@@ -226,8 +229,7 @@ clGetDeviceInfo(cl_device_id dev, cl_device_info param,
return string_property(buf, size, size_ret, "");
case CL_DEVICE_PLATFORM:
- return scalar_property<cl_platform_id>(buf, size, size_ret,
- &dev->platform);
+ return scalar_property<platform*>(buf, size, size_ret, &dev->platform);
case CL_DEVICE_HOST_UNIFIED_MEMORY:
return scalar_property<cl_bool>(buf, size, size_ret, CL_TRUE);
diff --git a/src/gallium/state_trackers/clover/api/event.cpp b/src/gallium/state_trackers/clover/api/event.cpp
index 39a647b..cd3159f 100644
--- a/src/gallium/state_trackers/clover/api/event.cpp
+++ b/src/gallium/state_trackers/clover/api/event.cpp
@@ -20,18 +20,20 @@
// OTHER DEALINGS IN THE SOFTWARE.
//
+#include "api/icd.hpp"
#include "api/util.hpp"
#include "core/event.hpp"
using namespace clover;
PUBLIC cl_event
-clCreateUserEvent(cl_context ctx, cl_int *errcode_ret) try {
+clCreateUserEvent(cl_context _ctx, cl_int *errcode_ret) try {
+ UNWRAP_ICD_PARAM_CONTEXT(ctx)
if (!ctx)
throw error(CL_INVALID_CONTEXT);
ret_error(errcode_ret, CL_SUCCESS);
- return new soft_event(*ctx, {}, false);
+ return WRAP_ICD_EVENT(new soft_event(*ctx, {}, false));
} catch(error &e) {
ret_error(errcode_ret, e);
@@ -39,7 +41,8 @@ clCreateUserEvent(cl_context ctx, cl_int *errcode_ret) try {
}
PUBLIC cl_int
-clSetUserEventStatus(cl_event ev, cl_int status) {
+clSetUserEventStatus(cl_event _ev, cl_int status) {
+ UNWRAP_ICD_PARAM_EVENT(ev);
if (!dynamic_cast<soft_event *>(ev))
return CL_INVALID_EVENT;
@@ -62,11 +65,12 @@ clWaitForEvents(cl_uint num_evs, const cl_event *evs) try {
if (!num_evs || !evs)
throw error(CL_INVALID_VALUE);
- std::for_each(evs, evs + num_evs, [&](const cl_event ev) {
+ std::for_each(evs, evs + num_evs, [&](const cl_event _ev) {
+ UNWRAP_ICD_PARAM_EVENT(ev);
if (!ev)
throw error(CL_INVALID_EVENT);
- if (&ev->ctx != &evs[0]->ctx)
+ if (&ev->ctx != &UNWRAP_ICD_OBJECT(evs[0])->ctx)
throw error(CL_INVALID_CONTEXT);
if (ev->status() < 0)
@@ -76,7 +80,8 @@ clWaitForEvents(cl_uint num_evs, const cl_event *evs) try {
// Create a temporary soft event that depends on all the events in
// the wait list
ref_ptr<soft_event> sev = transfer(
- new soft_event(evs[0]->ctx, { evs, evs + num_evs }, true));
+ new soft_event(UNWRAP_ICD_OBJECT(evs[0])->ctx,
+ { evs, evs + num_evs }, true));
// ...and wait on it.
sev->wait();
@@ -88,17 +93,18 @@ clWaitForEvents(cl_uint num_evs, const cl_event *evs) try {
}
PUBLIC cl_int
-clGetEventInfo(cl_event ev, cl_event_info param,
+clGetEventInfo(cl_event _ev, cl_event_info param,
size_t size, void *buf, size_t *size_ret) {
+ UNWRAP_ICD_PARAM_EVENT(ev);
if (!ev)
return CL_INVALID_EVENT;
switch (param) {
case CL_EVENT_COMMAND_QUEUE:
- return scalar_property<cl_command_queue>(buf, size, size_ret, ev->queue());
+ return scalar_property<clover::command_queue*>(buf, size, size_ret, ev->queue());
case CL_EVENT_CONTEXT:
- return scalar_property<cl_context>(buf, size, size_ret, &ev->ctx);
+ return scalar_property<clover::context*>(buf, size, size_ret, &ev->ctx);
case CL_EVENT_COMMAND_TYPE:
return scalar_property<cl_command_type>(buf, size, size_ret, ev->command());
@@ -115,10 +121,11 @@ clGetEventInfo(cl_event ev, cl_event_info param,
}
PUBLIC cl_int
-clSetEventCallback(cl_event ev, cl_int type,
+clSetEventCallback(cl_event _ev, cl_int type,
void (CL_CALLBACK *pfn_event_notify)(cl_event, cl_int,
void *),
void *user_data) try {
+ UNWRAP_ICD_PARAM_EVENT(ev);
if (!ev)
throw error(CL_INVALID_EVENT);
@@ -128,10 +135,10 @@ clSetEventCallback(cl_event ev, cl_int type,
// Create a temporary soft event that depends on ev, with
// pfn_event_notify as completion action.
ref_ptr<soft_event> sev = transfer(
- new soft_event(ev->ctx, { ev }, true,
+ new soft_event(ev->ctx, { _ev }, true,
[=](event &) {
ev->wait();
- pfn_event_notify(ev, ev->status(), user_data);
+ pfn_event_notify(_ev, ev->status(), user_data);
}));
return CL_SUCCESS;
@@ -141,7 +148,8 @@ clSetEventCallback(cl_event ev, cl_int type,
}
PUBLIC cl_int
-clRetainEvent(cl_event ev) {
+clRetainEvent(cl_event _ev) {
+ UNWRAP_ICD_PARAM_EVENT(ev);
if (!ev)
return CL_INVALID_EVENT;
@@ -150,7 +158,8 @@ clRetainEvent(cl_event ev) {
}
PUBLIC cl_int
-clReleaseEvent(cl_event ev) {
+clReleaseEvent(cl_event _ev) {
+ UNWRAP_ICD_PARAM_EVENT(ev);
if (!ev)
return CL_INVALID_EVENT;
@@ -161,14 +170,15 @@ clReleaseEvent(cl_event ev) {
}
PUBLIC cl_int
-clEnqueueMarker(cl_command_queue q, cl_event *ev) try {
+clEnqueueMarker(cl_command_queue _q, cl_event *ev) try {
+ UNWRAP_ICD_PARAM_COMMAND_QUEUE(q);
if (!q)
throw error(CL_INVALID_COMMAND_QUEUE);
if (!ev)
throw error(CL_INVALID_VALUE);
- *ev = new hard_event(*q, CL_COMMAND_MARKER, {});
+ *ev = WRAP_ICD_EVENT(new hard_event(*q, CL_COMMAND_MARKER, {}));
return CL_SUCCESS;
@@ -177,7 +187,8 @@ clEnqueueMarker(cl_command_queue q, cl_event *ev) try {
}
PUBLIC cl_int
-clEnqueueBarrier(cl_command_queue q) {
+clEnqueueBarrier(cl_command_queue _q) {
+ UNWRAP_ICD_PARAM_COMMAND_QUEUE(q);
if (!q)
return CL_INVALID_COMMAND_QUEUE;
@@ -186,15 +197,17 @@ clEnqueueBarrier(cl_command_queue q) {
}
PUBLIC cl_int
-clEnqueueWaitForEvents(cl_command_queue q, cl_uint num_evs,
+clEnqueueWaitForEvents(cl_command_queue _q, cl_uint num_evs,
const cl_event *evs) try {
+ UNWRAP_ICD_PARAM_COMMAND_QUEUE(q);
if (!q)
throw error(CL_INVALID_COMMAND_QUEUE);
if (!num_evs || !evs)
throw error(CL_INVALID_VALUE);
- std::for_each(evs, evs + num_evs, [&](const cl_event ev) {
+ std::for_each(evs, evs + num_evs, [&](const cl_event _ev) {
+ UNWRAP_ICD_PARAM_EVENT(ev);
if (!ev)
throw error(CL_INVALID_EVENT);
@@ -215,13 +228,14 @@ clEnqueueWaitForEvents(cl_command_queue q, cl_uint num_evs,
}
PUBLIC cl_int
-clGetEventProfilingInfo(cl_event ev, cl_profiling_info param,
+clGetEventProfilingInfo(cl_event _ev, cl_profiling_info param,
size_t size, void *buf, size_t *size_ret) {
return CL_PROFILING_INFO_NOT_AVAILABLE;
}
PUBLIC cl_int
-clFinish(cl_command_queue q) try {
+clFinish(cl_command_queue _q) try {
+ UNWRAP_ICD_PARAM_COMMAND_QUEUE(q);
if (!q)
throw error(CL_INVALID_COMMAND_QUEUE);
diff --git a/src/gallium/state_trackers/clover/api/icd.cpp b/src/gallium/state_trackers/clover/api/icd.cpp
new file mode 100644
index 0000000..fcc01bb
--- /dev/null
+++ b/src/gallium/state_trackers/clover/api/icd.cpp
@@ -0,0 +1,160 @@
+//
+// Copyright 2013 Advanced Micro Devices Inc.
+//
+// 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.
+//
+//
+// Author: Tom Stellard <thomas.stellard at amd.com>
+//
+
+#include "api/icd.hpp"
+#include "core/context.hpp"
+#include "core/device.hpp"
+#include "core/memory.hpp"
+#include "core/platform.hpp"
+#include "core/program.hpp"
+#include "core/queue.hpp"
+
+ICD_CLASS_IMPL(_cl_command_queue, command_queue)
+ICD_CLASS_IMPL(_cl_context, context)
+ICD_CLASS_IMPL(_cl_device_id, device)
+ICD_CLASS_IMPL(_cl_event, event)
+ICD_CLASS_IMPL(_cl_kernel, kernel)
+ICD_CLASS_IMPL(_cl_mem, memory_obj)
+ICD_CLASS_IMPL(_cl_program, program)
+ICD_CLASS_IMPL(_cl_platform_id, platform)
+ICD_CLASS_IMPL(_cl_sampler, sampler)
+
+cl_icd_dispatch clover_icd_dispatch =
+#ifndef USE_ICD
+ 0
+#else
+{
+ &CLOVER_API(clGetPlatformIDs),
+ &CLOVER_API(clGetPlatformInfo),
+ &CLOVER_API(clGetDeviceIDs),
+ &CLOVER_API(clGetDeviceInfo),
+ &CLOVER_API(clCreateContext),
+ &CLOVER_API(clCreateContextFromType),
+ &CLOVER_API(clRetainContext),
+ &CLOVER_API(clReleaseContext),
+ &CLOVER_API(clGetContextInfo),
+ &CLOVER_API(clCreateCommandQueue),
+ &CLOVER_API(clRetainCommandQueue),
+ &CLOVER_API(clReleaseCommandQueue),
+ &CLOVER_API(clGetCommandQueueInfo),
+ NULL, // &CLOVER_API(clSetCommandQueueProperty),
+ &CLOVER_API(clCreateBuffer),
+ NULL, // &CLOVER_API(clCreateImage2D),
+ NULL, // &CLOVER_API(clCreateImage3D),
+ &CLOVER_API(clRetainMemObject),
+ &CLOVER_API(clReleaseMemObject),
+ &CLOVER_API(clGetSupportedImageFormats),
+ &CLOVER_API(clGetMemObjectInfo),
+ &CLOVER_API(clGetImageInfo),
+ &CLOVER_API(clCreateSampler),
+ &CLOVER_API(clRetainSampler),
+ &CLOVER_API(clReleaseSampler),
+ &CLOVER_API(clGetSamplerInfo),
+ &CLOVER_API(clCreateProgramWithSource),
+ &CLOVER_API(clCreateProgramWithBinary),
+ &CLOVER_API(clRetainProgram),
+ &CLOVER_API(clReleaseProgram),
+ &CLOVER_API(clBuildProgram),
+ NULL, // &CLOVER_API(clUnloadCompiler),
+ &CLOVER_API(clGetProgramInfo),
+ &CLOVER_API(clGetProgramBuildInfo),
+ &CLOVER_API(clCreateKernel),
+ &CLOVER_API(clCreateKernelsInProgram),
+ &CLOVER_API(clRetainKernel),
+ &CLOVER_API(clReleaseKernel),
+ &CLOVER_API(clSetKernelArg),
+ &CLOVER_API(clGetKernelInfo),
+ &CLOVER_API(clGetKernelWorkGroupInfo),
+ &CLOVER_API(clWaitForEvents),
+ &CLOVER_API(clGetEventInfo),
+ &CLOVER_API(clRetainEvent),
+ &CLOVER_API(clReleaseEvent),
+ &CLOVER_API(clGetEventProfilingInfo),
+ &CLOVER_API(clFlush),
+ &CLOVER_API(clFinish),
+ &CLOVER_API(clEnqueueReadBuffer),
+ &CLOVER_API(clEnqueueWriteBuffer),
+ &CLOVER_API(clEnqueueCopyBuffer),
+ &CLOVER_API(clEnqueueReadImage),
+ &CLOVER_API(clEnqueueWriteImage),
+ &CLOVER_API(clEnqueueCopyImage),
+ &CLOVER_API(clEnqueueCopyImageToBuffer),
+ &CLOVER_API(clEnqueueCopyBufferToImage),
+ &CLOVER_API(clEnqueueMapBuffer),
+ &CLOVER_API(clEnqueueMapImage),
+ &CLOVER_API(clEnqueueUnmapMemObject),
+ &CLOVER_API(clEnqueueNDRangeKernel),
+ &CLOVER_API(clEnqueueTask),
+ &CLOVER_API(clEnqueueNativeKernel),
+ NULL, // &CLOVER_API(clEnqueueMarker),
+ NULL, // &CLOVER_API(clEnqueueWaitForEvents),
+ NULL, // &CLOVER_API(clEnqueueBarrier),
+ clGetExtensionFunctionAddress,
+ NULL, // &CLOVER_API(clCreateFromGLBuffer),
+ NULL, // &CLOVER_API(clCreateFromGLTexture2D),
+ NULL, // &CLOVER_API(clCreateFromGLTexture3D),
+ NULL, // &CLOVER_API(clCreateFromGLRenderbuffer),
+ NULL, // &CLOVER_API(clGetGLObjectInfo),
+ NULL, // &CLOVER_API(clGetGLTextureInfo),
+ NULL, // &CLOVER_API(clEnqueueAcquireGLObjects),
+ NULL, // &CLOVER_API(clEnqueueReleaseGLObjects),
+ NULL, // &CLOVER_API(clGetGLContextInfoKHR),
+ NULL,
+ NULL,
+ NULL,
+ NULL,
+ NULL,
+ NULL,
+ &CLOVER_API(clSetEventCallback),
+ &CLOVER_API(clCreateSubBuffer),
+ &CLOVER_API(clSetMemObjectDestructorCallback),
+ &CLOVER_API(clCreateUserEvent),
+ &CLOVER_API(clSetUserEventStatus),
+ &CLOVER_API(clEnqueueReadBufferRect),
+ &CLOVER_API(clEnqueueWriteBufferRect),
+ &CLOVER_API(clEnqueueCopyBufferRect),
+ NULL, // &CLOVER_API(clCreateSubDevicesEXT),
+ NULL, // &CLOVER_API(clRetainDeviceEXT),
+ NULL, // &CLOVER_API(clReleaseDeviceEXT),
+ NULL,
+ NULL, // &CLOVER_API(clCreateSubDevices),
+ NULL, // &CLOVER_API(clRetainDevice),
+ NULL, // &CLOVER_API(clReleaseDevice),
+ NULL, // &CLOVER_API(clCreateImage),
+ NULL, // &CLOVER_API(clCreateProgramWithBuiltInKernels),
+ NULL, // &CLOVER_API(clCompileProgram),
+ NULL, // &CLOVER_API(clLinkProgram),
+ NULL, // &CLOVER_API(clUnloadPlatformCompiler),
+ NULL, // &CLOVER_API(clGetKernelArgInfo),
+ NULL, // &CLOVER_API(clEnqueueFillBuffer),
+ NULL, // &CLOVER_API(clEnqueueFillImage),
+ NULL, // &CLOVER_API(clEnqueueMigrateMemObjects),
+ NULL, // &CLOVER_API(clEnqueueMarkerWithWaitList),
+ NULL, // &CLOVER_API(clEnqueueBarrierWithWaitList),
+ NULL, // &CLOVER_API(clGetExtensionFunctionAddressForPlatform),
+ NULL, // &CLOVER_API(clCreateFromGLTexture),
+}
+#endif // USE_ICD
+;
diff --git a/src/gallium/state_trackers/clover/api/icd.hpp b/src/gallium/state_trackers/clover/api/icd.hpp
new file mode 100644
index 0000000..c81c3e0
--- /dev/null
+++ b/src/gallium/state_trackers/clover/api/icd.hpp
@@ -0,0 +1,258 @@
+//
+// Copyright 2013 Advanced Micro Devices Inc.
+//
+// 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.
+//
+//
+// Author: Tom Stellard <thomas.stellard at amd.com>
+//
+
+#ifndef __CL_ICD_HPP__
+#define __CL_ICD_HPP__
+
+namespace clover {
+ class command_queue;
+ class context;
+ class event;
+ class device;
+ class kernel;
+ class memory_obj;
+ class platform;
+ class program;
+ class sampler;
+}
+
+#ifdef USE_ICD
+
+typedef struct _cl_icd_dispatch cl_icd_dispatch;
+
+#else // USE_ICD
+
+typedef unsigned cl_icd_dispatch;
+
+#endif // USE_ICD
+
+extern cl_icd_dispatch clover_icd_dispatch;
+
+#define ICD_CLASS_DECL(name, internal_class) \
+struct name { \
+ cl_icd_dispatch *dispatch; \
+ clover::internal_class *__obj; \
+ name(clover::internal_class *obj); \
+};
+
+#define ICD_CLASS_IMPL(name, internal_class) \
+name::name(clover::internal_class *obj) : dispatch(&clover_icd_dispatch), \
+ __obj(obj) { } \
+
+namespace clover {
+
+template <class T>
+class icd_container {
+
+public:
+ icd_container(T icd_obj) : __icd_obj(icd_obj) { }
+ virtual ~icd_container() {
+ delete __icd_obj;
+ }
+ T icd_obj() { return __icd_obj; }
+private:
+ T __icd_obj;
+};
+
+}
+
+ICD_CLASS_DECL(_cl_command_queue, command_queue)
+ICD_CLASS_DECL(_cl_context, context)
+ICD_CLASS_DECL(_cl_device_id, device)
+ICD_CLASS_DECL(_cl_event, event)
+ICD_CLASS_DECL(_cl_kernel, kernel)
+ICD_CLASS_DECL(_cl_mem, memory_obj)
+ICD_CLASS_DECL(_cl_program, program)
+ICD_CLASS_DECL(_cl_platform_id, platform)
+ICD_CLASS_DECL(_cl_sampler, sampler)
+
+#define UNWRAP_ICD_OBJECT(obj) (obj)->__obj
+#define UNWRAP_ICD_PARAM(type, name) \
+ type *name = _##name->__obj;
+#define WRAP_ICD_OBJECT(obj) (obj)->icd_obj()
+#define WRAP_ICD_RET_OBJECT(clover_type, cl_type, out, in) \
+ clover::clover_type *_##out; \
+ ret_object(&_##out, in); \
+ if (out) \
+ *out = WRAP_ICD_OBJECT((_##out));
+
+
+#define UNWRAP_ICD_PARAM_COMMAND_QUEUE(name) \
+ UNWRAP_ICD_PARAM(clover::command_queue, name)
+#define UNWRAP_ICD_PARAM_CONTEXT(name) \
+ UNWRAP_ICD_PARAM(clover::context, name)
+#define UNWRAP_ICD_PARAM_DEVICE(name) \
+ UNWRAP_ICD_PARAM(clover::device, name)
+#define UNWRAP_ICD_PARAM_EVENT(name) \
+ UNWRAP_ICD_PARAM(clover::event, name)
+#define UNWRAP_ICD_PARAM_KERNEL(name) \
+ UNWRAP_ICD_PARAM(clover::kernel, name)
+#define UNWRAP_ICD_PARAM_MEM(name) \
+ UNWRAP_ICD_PARAM(clover::memory_obj, name)
+#define UNWRAP_ICD_PARAM_PLATFORM(name) \
+ UNWRAP_ICD_PARAM(clover::platform, name)
+#define UNWRAP_ICD_PARAM_PROGRAM(name) \
+ UNWRAP_ICD_PARAM(clover::program, name)
+#define UNWRAP_ICD_PARAM_SAMPLER(name) \
+ UNWRAP_ICD_PARAM(clover::sampler, name)
+
+#define WRAP_ICD_CONTEXT(ctx) WRAP_ICD_OBJECT(ctx)
+#define WRAP_ICD_DEVICE(dev) WRAP_ICD_OBJECT(dev)
+#define WRAP_ICD_EVENT(event) WRAP_ICD_OBJECT(event)
+#define WRAP_ICD_KERNEL(kernel) WRAP_ICD_OBJECT(kernel)
+#define WRAP_ICD_MEM(mem) WRAP_ICD_OBJECT(mem)
+#define WRAP_ICD_PROGRAM(prog) WRAP_ICD_OBJECT(prog)
+#define WRAP_ICD_SAMPLER(sampler) WRAP_ICD_OBJECT(sampler)
+
+#define WRAP_ICD_RET_OBJECT_EVENT(out, in) \
+ WRAP_ICD_RET_OBJECT(event, _cl_event, out, in)
+
+
+//===----------------------------------------------------------------------===//
+// Function prototypes
+//===----------------------------------------------------------------------===//
+
+#ifndef USE_ICD
+
+#define CLOVER_API(function) function
+
+#else
+
+#define CLOVER_API(function) CLOVER##function
+
+#define clGetPlatformIDs CLOVER_API(clGetPlatformIDs)
+#define clGetPlatformInfo CLOVER_API(clGetPlatformInfo)
+#define clGetDeviceIDs CLOVER_API(clGetDeviceIDs)
+#define clGetDeviceInfo CLOVER_API(clGetDeviceInfo)
+#define clCreateContext CLOVER_API(clCreateContext)
+#define clCreateContextFromType CLOVER_API(clCreateContextFromType)
+#define clRetainContext CLOVER_API(clRetainContext)
+#define clReleaseContext CLOVER_API(clReleaseContext)
+#define clGetContextInfo CLOVER_API(clGetContextInfo)
+#define clCreateCommandQueue CLOVER_API(clCreateCommandQueue)
+#define clRetainCommandQueue CLOVER_API(clRetainCommandQueue)
+#define clReleaseCommandQueue CLOVER_API(clReleaseCommandQueue)
+#define clGetCommandQueueInfo CLOVER_API(clGetCommandQueueInfo)
+#define clSetCommandQueueProperty CLOVER_API(clSetCommandQueueProperty)
+#define clCreateBuffer CLOVER_API(clCreateBuffer)
+#define clCreateImage2D CLOVER_API(clCreateImage2D)
+#define clCreateImage3D CLOVER_API(clCreateImage3D)
+#define clRetainMemObject CLOVER_API(clRetainMemObject)
+#define clReleaseMemObject CLOVER_API(clReleaseMemObject)
+#define clGetSupportedImageFormats CLOVER_API(clGetSupportedImageFormats)
+#define clGetMemObjectInfo CLOVER_API(clGetMemObjectInfo)
+#define clGetImageInfo CLOVER_API(clGetImageInfo)
+#define clCreateSampler CLOVER_API(clCreateSampler)
+#define clRetainSampler CLOVER_API(clRetainSampler)
+#define clReleaseSampler CLOVER_API(clReleaseSampler)
+#define clGetSamplerInfo CLOVER_API(clGetSamplerInfo)
+#define clCreateProgramWithSource CLOVER_API(clCreateProgramWithSource)
+#define clCreateProgramWithBinary CLOVER_API(clCreateProgramWithBinary)
+#define clRetainProgram CLOVER_API(clRetainProgram)
+#define clReleaseProgram CLOVER_API(clReleaseProgram)
+#define clBuildProgram CLOVER_API(clBuildProgram)
+#define clUnloadCompiler CLOVER_API(clUnloadCompiler)
+#define clGetProgramInfo CLOVER_API(clGetProgramInfo)
+#define clGetProgramBuildInfo CLOVER_API(clGetProgramBuildInfo)
+#define clCreateKernel CLOVER_API(clCreateKernel)
+#define clCreateKernelsInProgram CLOVER_API(clCreateKernelsInProgram)
+#define clRetainKernel CLOVER_API(clRetainKernel)
+#define clReleaseKernel CLOVER_API(clReleaseKernel)
+#define clSetKernelArg CLOVER_API(clSetKernelArg)
+#define clGetKernelInfo CLOVER_API(clGetKernelInfo)
+#define clGetKernelWorkGroupInfo CLOVER_API(clGetKernelWorkGroupInfo)
+#define clWaitForEvents CLOVER_API(clWaitForEvents)
+#define clGetEventInfo CLOVER_API(clGetEventInfo)
+#define clRetainEvent CLOVER_API(clRetainEvent)
+#define clReleaseEvent CLOVER_API(clReleaseEvent)
+#define clGetEventProfilingInfo CLOVER_API(clGetEventProfilingInfo)
+#define clFlush CLOVER_API(clFlush)
+#define clFinish CLOVER_API(clFinish)
+#define clEnqueueReadBuffer CLOVER_API(clEnqueueReadBuffer)
+#define clEnqueueWriteBuffer CLOVER_API(clEnqueueWriteBuffer)
+#define clEnqueueCopyBuffer CLOVER_API(clEnqueueCopyBuffer)
+#define clEnqueueReadImage CLOVER_API(clEnqueueReadImage)
+#define clEnqueueWriteImage CLOVER_API(clEnqueueWriteImage)
+#define clEnqueueCopyImage CLOVER_API(clEnqueueCopyImage)
+#define clEnqueueCopyImageToBuffer CLOVER_API(clEnqueueCopyImageToBuffer)
+#define clEnqueueCopyBufferToImage CLOVER_API(clEnqueueCopyBufferToImage)
+#define clEnqueueMapBuffer CLOVER_API(clEnqueueMapBuffer)
+#define clEnqueueMapImage CLOVER_API(clEnqueueMapImage)
+#define clEnqueueUnmapMemObject CLOVER_API(clEnqueueUnmapMemObject)
+#define clEnqueueNDRangeKernel CLOVER_API(clEnqueueNDRangeKernel)
+#define clEnqueueTask CLOVER_API(clEnqueueTask)
+#define clEnqueueNativeKernel CLOVER_API(clEnqueueNativeKernel)
+#define clEnqueueMarker CLOVER_API(clEnqueueMarker)
+#define clEnqueueWaitForEvents CLOVER_API(clEnqueueWaitForEvents)
+#define clEnqueueBarrier CLOVER_API(clEnqueueBarrier)
+#define clCreateFromGLBuffer CLOVER_API(clCreateFromGLBuffer)
+#define clCreateFromGLTexture2D CLOVER_API(clCreateFromGLTexture2D)
+#define clCreateFromGLTexture3D CLOVER_API(clCreateFromGLTexture3D)
+#define clCreateFromGLRenderbuffer CLOVER_API(clCreateFromGLRenderbuffer)
+#define clGetGLObjectInfo CLOVER_API(clGetGLObjectInfo)
+#define clGetGLTextureInfo CLOVER_API(clGetGLTextureInfo)
+#define clEnqueueAcquireGLObjects CLOVER_API(clEnqueueAcquireGLObjects)
+#define clEnqueueReleaseGLObjects CLOVER_API(clEnqueueReleaseGLObjects)
+#define clGetGLContextInfoKHR CLOVER_API(clGetGLContextInfoKHR)
+
+
+
+
+
+
+#define clSetEventCallback CLOVER_API(clSetEventCallback)
+#define clCreateSubBuffer CLOVER_API(clCreateSubBuffer)
+#define clSetMemObjectDestructorCallback CLOVER_API(clSetMemObjectDestructorCallback)
+#define clCreateUserEvent CLOVER_API(clCreateUserEvent)
+#define clSetUserEventStatus CLOVER_API(clSetUserEventStatus)
+#define clEnqueueReadBufferRect CLOVER_API(clEnqueueReadBufferRect)
+#define clEnqueueWriteBufferRect CLOVER_API(clEnqueueWriteBufferRect)
+#define clEnqueueCopyBufferRect CLOVER_API(clEnqueueCopyBufferRect)
+#define clCreateSubDevicesEXT CLOVER_API(clCreateSubDevicesEXT)
+#define clRetainDeviceEXT CLOVER_API(clRetainDeviceEXT)
+#define clReleaseDeviceEXT CLOVER_API(clReleaseDeviceEXT)
+
+#define clCreateSubDevices CLOVER_API(clCreateSubDevices)
+#define clRetainDevice CLOVER_API(clRetainDevice)
+#define clReleaseDevice CLOVER_API(clReleaseDevice)
+#define clCreateImage CLOVER_API(clCreateImage)
+#define clCreateProgramWithBuiltInKernels CLOVER_API(clCreateProgramWithBuiltInKernels)
+#define clCompileProgram CLOVER_API(clCompileProgram)
+#define clLinkProgram CLOVER_API(clLinkProgram)
+#define clUnloadPlatformCompiler CLOVER_API(clUnloadPlatformCompiler)
+#define clGetKernelArgInfo CLOVER_API(clGetKernelArgInfo)
+#define clEnqueueFillBuffer CLOVER_API(clEnqueueFillBuffer)
+#define clEnqueueFillImage CLOVER_API(clEnqueueFillImage)
+#define clEnqueueMigrateMemObjects CLOVER_API(clEnqueueMigrateMemObjects)
+#define clEnqueueMarkerWithWaitList CLOVER_API(clEnqueueMarkerWithWaitList)
+#define clEnqueueBarrierWithWaitList CLOVER_API(clEnqueueBarrierWithWaitList)
+#define clCreateFromGLTexture CLOVER_API(clCreateFromGLTexture)
+
+#include <CL/cl.h>
+#include <ocl_icd.h>
+
+#endif
+
+#endif // __CL_ICD_HPP__
diff --git a/src/gallium/state_trackers/clover/api/kernel.cpp b/src/gallium/state_trackers/clover/api/kernel.cpp
index 13113a2..31dabfa 100644
--- a/src/gallium/state_trackers/clover/api/kernel.cpp
+++ b/src/gallium/state_trackers/clover/api/kernel.cpp
@@ -20,6 +20,7 @@
// OTHER DEALINGS IN THE SOFTWARE.
//
+#include "api/icd.hpp"
#include "api/util.hpp"
#include "core/kernel.hpp"
#include "core/event.hpp"
@@ -27,8 +28,9 @@
using namespace clover;
PUBLIC cl_kernel
-clCreateKernel(cl_program prog, const char *name,
+clCreateKernel(cl_program _prog, const char *name,
cl_int *errcode_ret) try {
+ UNWRAP_ICD_PARAM_PROGRAM(prog);
if (!prog)
throw error(CL_INVALID_PROGRAM);
@@ -41,7 +43,8 @@ clCreateKernel(cl_program prog, const char *name,
auto sym = prog->binaries().begin()->second.sym(name);
ret_error(errcode_ret, CL_SUCCESS);
- return new kernel(*prog, name, { sym.args.begin(), sym.args.end() });
+ return WRAP_ICD_KERNEL(new kernel(*prog, name, { sym.args.begin(),
+ sym.args.end() }));
} catch (module::noent_error &e) {
ret_error(errcode_ret, CL_INVALID_KERNEL_NAME);
@@ -53,8 +56,9 @@ clCreateKernel(cl_program prog, const char *name,
}
PUBLIC cl_int
-clCreateKernelsInProgram(cl_program prog, cl_uint count,
+clCreateKernelsInProgram(cl_program _prog, cl_uint count,
cl_kernel *kerns, cl_uint *count_ret) {
+ UNWRAP_ICD_PARAM_PROGRAM(prog);
if (!prog)
throw error(CL_INVALID_PROGRAM);
@@ -69,8 +73,9 @@ clCreateKernelsInProgram(cl_program prog, cl_uint count,
if (kerns)
std::transform(syms.begin(), syms.end(), kerns,
[=](const module::symbol &sym) {
- return new kernel(*prog, compat::string(sym.name),
- { sym.args.begin(), sym.args.end() });
+ return WRAP_ICD_KERNEL(new kernel(*prog,
+ compat::string(sym.name),
+ { sym.args.begin(), sym.args.end() }));
});
if (count_ret)
@@ -80,7 +85,8 @@ clCreateKernelsInProgram(cl_program prog, cl_uint count,
}
PUBLIC cl_int
-clRetainKernel(cl_kernel kern) {
+clRetainKernel(cl_kernel _kern) {
+ UNWRAP_ICD_PARAM_KERNEL(kern);
if (!kern)
return CL_INVALID_KERNEL;
@@ -89,7 +95,8 @@ clRetainKernel(cl_kernel kern) {
}
PUBLIC cl_int
-clReleaseKernel(cl_kernel kern) {
+clReleaseKernel(cl_kernel _kern) {
+ UNWRAP_ICD_PARAM_KERNEL(kern);
if (!kern)
return CL_INVALID_KERNEL;
@@ -100,8 +107,9 @@ clReleaseKernel(cl_kernel kern) {
}
PUBLIC cl_int
-clSetKernelArg(cl_kernel kern, cl_uint idx, size_t size,
+clSetKernelArg(cl_kernel _kern, cl_uint idx, size_t size,
const void *value) try {
+ UNWRAP_ICD_PARAM_KERNEL(kern);
if (!kern)
throw error(CL_INVALID_KERNEL);
@@ -117,8 +125,9 @@ clSetKernelArg(cl_kernel kern, cl_uint idx, size_t size,
}
PUBLIC cl_int
-clGetKernelInfo(cl_kernel kern, cl_kernel_info param,
+clGetKernelInfo(cl_kernel _kern, cl_kernel_info param,
size_t size, void *buf, size_t *size_ret) {
+ UNWRAP_ICD_PARAM_KERNEL(kern);
if (!kern)
return CL_INVALID_KERNEL;
@@ -135,11 +144,11 @@ clGetKernelInfo(cl_kernel kern, cl_kernel_info param,
kern->ref_count());
case CL_KERNEL_CONTEXT:
- return scalar_property<cl_context>(buf, size, size_ret,
+ return scalar_property<clover::context *>(buf, size, size_ret,
&kern->prog.ctx);
case CL_KERNEL_PROGRAM:
- return scalar_property<cl_program>(buf, size, size_ret,
+ return scalar_property<clover::program *>(buf, size, size_ret,
&kern->prog);
default:
@@ -148,9 +157,11 @@ clGetKernelInfo(cl_kernel kern, cl_kernel_info param,
}
PUBLIC cl_int
-clGetKernelWorkGroupInfo(cl_kernel kern, cl_device_id dev,
+clGetKernelWorkGroupInfo(cl_kernel _kern, cl_device_id _dev,
cl_kernel_work_group_info param,
size_t size, void *buf, size_t *size_ret) {
+ UNWRAP_ICD_PARAM_KERNEL(kern);
+ UNWRAP_ICD_PARAM_DEVICE(dev);
if (!kern)
return CL_INVALID_KERNEL;
@@ -188,7 +199,7 @@ namespace {
/// Common argument checking shared by kernel invocation commands.
///
void
- kernel_validate(cl_command_queue q, cl_kernel kern,
+ kernel_validate(clover::command_queue *q, clover::kernel *kern,
cl_uint dims, const size_t *grid_offset,
const size_t *grid_size, const size_t *block_size,
cl_uint num_deps, const cl_event *deps,
@@ -200,7 +211,8 @@ namespace {
throw error(CL_INVALID_KERNEL);
if (&kern->prog.ctx != &q->ctx ||
- any_of([&](const cl_event ev) {
+ any_of([&](const cl_event _ev) {
+ UNWRAP_ICD_PARAM_EVENT(ev);
return &ev->ctx != &q->ctx;
}, deps, deps + num_deps))
throw error(CL_INVALID_CONTEXT);
@@ -246,7 +258,7 @@ namespace {
/// Common event action shared by kernel invocation commands.
///
std::function<void (event &)>
- kernel_op(cl_command_queue q, cl_kernel kern,
+ kernel_op(clover::command_queue *q, clover::kernel *kern,
const std::vector<size_t> &grid_offset,
const std::vector<size_t> &grid_size,
const std::vector<size_t> &block_size) {
@@ -269,11 +281,13 @@ namespace {
}
PUBLIC cl_int
-clEnqueueNDRangeKernel(cl_command_queue q, cl_kernel kern,
+clEnqueueNDRangeKernel(cl_command_queue _q, cl_kernel _kern,
cl_uint dims, const size_t *pgrid_offset,
const size_t *pgrid_size, const size_t *pblock_size,
cl_uint num_deps, const cl_event *deps,
cl_event *ev) try {
+ UNWRAP_ICD_PARAM_COMMAND_QUEUE(q);
+ UNWRAP_ICD_PARAM_KERNEL(kern);
auto grid_offset = opt_vector(pgrid_offset, dims, 0);
auto grid_size = opt_vector(pgrid_size, dims, 1);
auto block_size = opt_vector(pblock_size, dims, 1);
@@ -285,7 +299,7 @@ clEnqueueNDRangeKernel(cl_command_queue q, cl_kernel kern,
*q, CL_COMMAND_NDRANGE_KERNEL, { deps, deps + num_deps },
kernel_op(q, kern, grid_offset, grid_size, block_size));
- ret_object(ev, hev);
+ WRAP_ICD_RET_OBJECT_EVENT(ev, hev);
return CL_SUCCESS;
} catch(error &e) {
@@ -293,9 +307,11 @@ clEnqueueNDRangeKernel(cl_command_queue q, cl_kernel kern,
}
PUBLIC cl_int
-clEnqueueTask(cl_command_queue q, cl_kernel kern,
+clEnqueueTask(cl_command_queue _q, cl_kernel _kern,
cl_uint num_deps, const cl_event *deps,
cl_event *ev) try {
+ UNWRAP_ICD_PARAM_COMMAND_QUEUE(q);
+ UNWRAP_ICD_PARAM_KERNEL(kern);
const std::vector<size_t> grid_offset = { 0 };
const std::vector<size_t> grid_size = { 1 };
const std::vector<size_t> block_size = { 1 };
@@ -307,7 +323,7 @@ clEnqueueTask(cl_command_queue q, cl_kernel kern,
*q, CL_COMMAND_TASK, { deps, deps + num_deps },
kernel_op(q, kern, grid_offset, grid_size, block_size));
- ret_object(ev, hev);
+ WRAP_ICD_RET_OBJECT_EVENT(ev, hev);
return CL_SUCCESS;
} catch(error &e) {
@@ -315,7 +331,7 @@ clEnqueueTask(cl_command_queue q, cl_kernel kern,
}
PUBLIC cl_int
-clEnqueueNativeKernel(cl_command_queue q, void (*func)(void *),
+clEnqueueNativeKernel(cl_command_queue _q, void (*func)(void *),
void *args, size_t args_size,
cl_uint obj_count, const cl_mem *obj_list,
const void **obj_args, cl_uint num_deps,
diff --git a/src/gallium/state_trackers/clover/api/memory.cpp b/src/gallium/state_trackers/clover/api/memory.cpp
index 63b3619..7bbca07 100644
--- a/src/gallium/state_trackers/clover/api/memory.cpp
+++ b/src/gallium/state_trackers/clover/api/memory.cpp
@@ -20,6 +20,7 @@
// OTHER DEALINGS IN THE SOFTWARE.
//
+#include "api/icd.hpp"
#include "api/util.hpp"
#include "core/memory.hpp"
#include "core/format.hpp"
@@ -27,8 +28,9 @@
using namespace clover;
PUBLIC cl_mem
-clCreateBuffer(cl_context ctx, cl_mem_flags flags, size_t size,
+clCreateBuffer(cl_context _ctx, cl_mem_flags flags, size_t size,
void *host_ptr, cl_int *errcode_ret) try {
+ UNWRAP_ICD_PARAM_CONTEXT(ctx);
if (!ctx)
throw error(CL_INVALID_CONTEXT);
@@ -45,7 +47,7 @@ clCreateBuffer(cl_context ctx, cl_mem_flags flags, size_t size,
throw error(CL_INVALID_VALUE);
ret_error(errcode_ret, CL_SUCCESS);
- return new root_buffer(*ctx, flags, size, host_ptr);
+ return WRAP_ICD_MEM(new root_buffer(*ctx, flags, size, host_ptr));
} catch (error &e) {
ret_error(errcode_ret, e);
@@ -53,8 +55,9 @@ clCreateBuffer(cl_context ctx, cl_mem_flags flags, size_t size,
}
PUBLIC cl_mem
-clCreateSubBuffer(cl_mem obj, cl_mem_flags flags, cl_buffer_create_type op,
+clCreateSubBuffer(cl_mem _obj, cl_mem_flags flags, cl_buffer_create_type op,
const void *op_info, cl_int *errcode_ret) try {
+ UNWRAP_ICD_PARAM_MEM(obj);
root_buffer *parent = dynamic_cast<root_buffer *>(obj);
if (!parent)
@@ -79,7 +82,8 @@ clCreateSubBuffer(cl_mem obj, cl_mem_flags flags, cl_buffer_create_type op,
throw error(CL_INVALID_BUFFER_SIZE);
ret_error(errcode_ret, CL_SUCCESS);
- return new sub_buffer(*parent, flags, reg->origin, reg->size);
+ return WRAP_ICD_MEM(new sub_buffer(*parent, flags, reg->origin,
+ reg->size));
} else {
throw error(CL_INVALID_VALUE);
@@ -91,10 +95,11 @@ clCreateSubBuffer(cl_mem obj, cl_mem_flags flags, cl_buffer_create_type op,
}
PUBLIC cl_mem
-clCreateImage2D(cl_context ctx, cl_mem_flags flags,
+clCreateImage2D(cl_context _ctx, cl_mem_flags flags,
const cl_image_format *format,
size_t width, size_t height, size_t row_pitch,
void *host_ptr, cl_int *errcode_ret) try {
+ UNWRAP_ICD_PARAM_CONTEXT(ctx);
if (!ctx)
throw error(CL_INVALID_CONTEXT);
@@ -117,8 +122,8 @@ clCreateImage2D(cl_context ctx, cl_mem_flags flags,
throw error(CL_IMAGE_FORMAT_NOT_SUPPORTED);
ret_error(errcode_ret, CL_SUCCESS);
- return new image2d(*ctx, flags, format, width, height,
- row_pitch, host_ptr);
+ return WRAP_ICD_MEM(new image2d(*ctx, flags, format, width, height,
+ row_pitch, host_ptr));
} catch (error &e) {
ret_error(errcode_ret, e);
@@ -126,11 +131,12 @@ clCreateImage2D(cl_context ctx, cl_mem_flags flags,
}
PUBLIC cl_mem
-clCreateImage3D(cl_context ctx, cl_mem_flags flags,
+clCreateImage3D(cl_context _ctx, cl_mem_flags flags,
const cl_image_format *format,
size_t width, size_t height, size_t depth,
size_t row_pitch, size_t slice_pitch,
void *host_ptr, cl_int *errcode_ret) try {
+ UNWRAP_ICD_PARAM_CONTEXT(ctx);
if (!ctx)
throw error(CL_INVALID_CONTEXT);
@@ -153,8 +159,8 @@ clCreateImage3D(cl_context ctx, cl_mem_flags flags,
throw error(CL_IMAGE_FORMAT_NOT_SUPPORTED);
ret_error(errcode_ret, CL_SUCCESS);
- return new image3d(*ctx, flags, format, width, height, depth,
- row_pitch, slice_pitch, host_ptr);
+ return WRAP_ICD_MEM(new image3d(*ctx, flags, format, width, height, depth,
+ row_pitch, slice_pitch, host_ptr));
} catch (error &e) {
ret_error(errcode_ret, e);
@@ -162,9 +168,10 @@ clCreateImage3D(cl_context ctx, cl_mem_flags flags,
}
PUBLIC cl_int
-clGetSupportedImageFormats(cl_context ctx, cl_mem_flags flags,
+clGetSupportedImageFormats(cl_context _ctx, cl_mem_flags flags,
cl_mem_object_type type, cl_uint count,
cl_image_format *buf, cl_uint *count_ret) try {
+ UNWRAP_ICD_PARAM_CONTEXT(ctx);
if (!ctx)
throw error(CL_INVALID_CONTEXT);
@@ -191,8 +198,9 @@ clGetSupportedImageFormats(cl_context ctx, cl_mem_flags flags,
}
PUBLIC cl_int
-clGetMemObjectInfo(cl_mem obj, cl_mem_info param,
+clGetMemObjectInfo(cl_mem _obj, cl_mem_info param,
size_t size, void *buf, size_t *size_ret) {
+ UNWRAP_ICD_PARAM_MEM(obj);
if (!obj)
return CL_INVALID_MEM_OBJECT;
@@ -217,11 +225,11 @@ clGetMemObjectInfo(cl_mem obj, cl_mem_info param,
return scalar_property<cl_uint>(buf, size, size_ret, obj->ref_count());
case CL_MEM_CONTEXT:
- return scalar_property<cl_context>(buf, size, size_ret, &obj->ctx);
+ return scalar_property<clover::context*>(buf, size, size_ret, &obj->ctx);
case CL_MEM_ASSOCIATED_MEMOBJECT: {
sub_buffer *sub = dynamic_cast<sub_buffer *>(obj);
- return scalar_property<cl_mem>(buf, size, size_ret,
+ return scalar_property<clover::memory_obj*>(buf, size, size_ret,
(sub ? &sub->parent : NULL));
}
case CL_MEM_OFFSET: {
@@ -235,8 +243,9 @@ clGetMemObjectInfo(cl_mem obj, cl_mem_info param,
}
PUBLIC cl_int
-clGetImageInfo(cl_mem obj, cl_image_info param,
+clGetImageInfo(cl_mem _obj, cl_image_info param,
size_t size, void *buf, size_t *size_ret) {
+ UNWRAP_ICD_PARAM_MEM(obj);
image *img = dynamic_cast<image *>(obj);
if (!img)
return CL_INVALID_MEM_OBJECT;
@@ -270,7 +279,8 @@ clGetImageInfo(cl_mem obj, cl_image_info param,
}
PUBLIC cl_int
-clRetainMemObject(cl_mem obj) {
+clRetainMemObject(cl_mem _obj) {
+ UNWRAP_ICD_PARAM_MEM(obj);
if (!obj)
return CL_INVALID_MEM_OBJECT;
@@ -279,7 +289,8 @@ clRetainMemObject(cl_mem obj) {
}
PUBLIC cl_int
-clReleaseMemObject(cl_mem obj) {
+clReleaseMemObject(cl_mem _obj) {
+ UNWRAP_ICD_PARAM_MEM(obj);
if (!obj)
return CL_INVALID_MEM_OBJECT;
@@ -290,16 +301,17 @@ clReleaseMemObject(cl_mem obj) {
}
PUBLIC cl_int
-clSetMemObjectDestructorCallback(cl_mem obj,
+clSetMemObjectDestructorCallback(cl_mem _obj,
void (CL_CALLBACK *pfn_notify)(cl_mem, void *),
void *user_data) {
+ UNWRAP_ICD_PARAM_MEM(obj);
if (!obj)
return CL_INVALID_MEM_OBJECT;
if (!pfn_notify)
return CL_INVALID_VALUE;
- obj->destroy_notify([=]{ pfn_notify(obj, user_data); });
+ obj->destroy_notify([=]{ pfn_notify(_obj, user_data); });
return CL_SUCCESS;
}
diff --git a/src/gallium/state_trackers/clover/api/platform.cpp b/src/gallium/state_trackers/clover/api/platform.cpp
index bfac67b..e514350 100644
--- a/src/gallium/state_trackers/clover/api/platform.cpp
+++ b/src/gallium/state_trackers/clover/api/platform.cpp
@@ -20,12 +20,19 @@
// OTHER DEALINGS IN THE SOFTWARE.
//
+#include "api/icd.hpp"
#include "api/util.hpp"
+#include "core/context.hpp"
+#include "core/device.hpp"
+#include "core/event.hpp"
+#include "core/kernel.hpp"
+#include "core/memory.hpp"
#include "core/platform.hpp"
+#include "core/program.hpp"
using namespace clover;
-static platform __platform;
+static clover::platform __platform;
PUBLIC cl_int
clGetPlatformIDs(cl_uint num_entries, cl_platform_id *platforms,
@@ -37,15 +44,16 @@ clGetPlatformIDs(cl_uint num_entries, cl_platform_id *platforms,
if (num_platforms)
*num_platforms = 1;
if (platforms)
- *platforms = &__platform;
+ *platforms = WRAP_ICD_OBJECT(&__platform);
return CL_SUCCESS;
}
PUBLIC cl_int
-clGetPlatformInfo(cl_platform_id platform, cl_platform_info param_name,
- size_t size, void *buf, size_t *size_ret) {
- if (platform != &__platform)
+CLOVER_API(clGetPlatformInfo)(cl_platform_id platform,
+ cl_platform_info param_name,
+ size_t size, void *buf, size_t *size_ret) {
+ if (platform != WRAP_ICD_OBJECT(&__platform))
return CL_INVALID_PLATFORM;
switch (param_name) {
@@ -59,13 +67,54 @@ clGetPlatformInfo(cl_platform_id platform, cl_platform_info param_name,
case CL_PLATFORM_NAME:
return string_property(buf, size, size_ret, "Default");
+#ifdef USE_ICD
+ case CL_PLATFORM_ICD_SUFFIX_KHR:
+#endif
case CL_PLATFORM_VENDOR:
return string_property(buf, size, size_ret, "Mesa");
case CL_PLATFORM_EXTENSIONS:
- return string_property(buf, size, size_ret, "");
+ return string_property(buf, size, size_ret, "cl_khr_icd");
default:
return CL_INVALID_VALUE;
}
}
+
+#ifdef USE_ICD
+
+#undef clGetPlatformInfo
+
+PUBLIC cl_int
+clGetPlatformInfo(cl_platform_id platform, cl_platform_info param_name,
+ size_t size, void *buf, size_t *size_ret) {
+ return CLOVER_API(clGetPlatformInfo)(platform, param_name, size, buf,
+ size_ret);
+}
+
+PUBLIC cl_int
+clIcdGetPlatformIDsKHR(cl_uint num_entries, cl_platform_id *platforms,
+ cl_uint *num_platforms) {
+ assert(std::is_standard_layout<_cl_platform_id>());
+ assert(std::is_standard_layout<_cl_device_id>());
+ assert(std::is_standard_layout<_cl_context>());
+ assert(std::is_standard_layout<_cl_command_queue>());
+ assert(std::is_standard_layout<_cl_mem>());
+ assert(std::is_standard_layout<_cl_program>());
+ assert(std::is_standard_layout<_cl_kernel>());
+ assert(std::is_standard_layout<_cl_event>());
+ assert(std::is_standard_layout<_cl_sampler>());
+ return clGetPlatformIDs(num_entries, platforms, num_platforms);
+}
+
+PUBLIC void*
+clGetExtensionFunctionAddress(const char *funcname) {
+ if (strcmp(funcname, "clIcdGetPlatformIDsKHR") == 0) {
+ return (void*)&clIcdGetPlatformIDsKHR;
+ } else if (strcmp(funcname, "clGetPlatformInfo") == 0) {
+ return (void*)&clGetPlatformInfo;
+ }
+ return NULL;
+}
+
+#endif
diff --git a/src/gallium/state_trackers/clover/api/program.cpp b/src/gallium/state_trackers/clover/api/program.cpp
index 6e5ca08..3b8b6bd 100644
--- a/src/gallium/state_trackers/clover/api/program.cpp
+++ b/src/gallium/state_trackers/clover/api/program.cpp
@@ -20,15 +20,17 @@
// OTHER DEALINGS IN THE SOFTWARE.
//
+#include "api/icd.hpp"
#include "api/util.hpp"
#include "core/program.hpp"
using namespace clover;
PUBLIC cl_program
-clCreateProgramWithSource(cl_context ctx, cl_uint count,
+clCreateProgramWithSource(cl_context _ctx, cl_uint count,
const char **strings, const size_t *lengths,
cl_int *errcode_ret) try {
+ UNWRAP_ICD_PARAM_CONTEXT(ctx);
std::string source;
if (!ctx)
@@ -46,7 +48,7 @@ clCreateProgramWithSource(cl_context ctx, cl_uint count,
// ...and create a program object for them.
ret_error(errcode_ret, CL_SUCCESS);
- return new program(*ctx, source);
+ return WRAP_ICD_PROGRAM(new program(*ctx, source));
} catch (error &e) {
ret_error(errcode_ret, e);
@@ -54,17 +56,19 @@ clCreateProgramWithSource(cl_context ctx, cl_uint count,
}
PUBLIC cl_program
-clCreateProgramWithBinary(cl_context ctx, cl_uint count,
+clCreateProgramWithBinary(cl_context _ctx, cl_uint count,
const cl_device_id *devs, const size_t *lengths,
const unsigned char **binaries, cl_int *status_ret,
cl_int *errcode_ret) try {
+ UNWRAP_ICD_PARAM_CONTEXT(ctx);
if (!ctx)
throw error(CL_INVALID_CONTEXT);
if (!count || !devs || !lengths || !binaries)
throw error(CL_INVALID_VALUE);
- if (any_of([&](const cl_device_id dev) {
+ if (any_of([&](const cl_device_id _dev) {
+ UNWRAP_ICD_PARAM_DEVICE(dev);
return !ctx->has_device(dev);
}, devs, devs + count))
throw error(CL_INVALID_DEVICE);
@@ -102,9 +106,11 @@ clCreateProgramWithBinary(cl_context ctx, cl_uint count,
// initialize a program object with them.
ret_error(errcode_ret, CL_SUCCESS);
- return new program(*ctx, { devs, devs + count },
- map(values<cl_int, module>,
- modules.begin(), modules.end()));
+ std::vector<clover::device*> clover_devs =
+ clover::device::unwrap_icd_list(devs, count);
+ return WRAP_ICD_PROGRAM(new program(*ctx, { clover_devs.begin(), clover_devs.end() },
+ map(values<cl_int, module>,
+ modules.begin(), modules.end())));
} catch (error &e) {
ret_error(errcode_ret, e);
@@ -112,7 +118,8 @@ clCreateProgramWithBinary(cl_context ctx, cl_uint count,
}
PUBLIC cl_int
-clRetainProgram(cl_program prog) {
+clRetainProgram(cl_program _prog) {
+ UNWRAP_ICD_PARAM_PROGRAM(prog);
if (!prog)
return CL_INVALID_PROGRAM;
@@ -121,7 +128,8 @@ clRetainProgram(cl_program prog) {
}
PUBLIC cl_int
-clReleaseProgram(cl_program prog) {
+clReleaseProgram(cl_program _prog) {
+ UNWRAP_ICD_PARAM_PROGRAM(prog);
if (!prog)
return CL_INVALID_PROGRAM;
@@ -132,26 +140,30 @@ clReleaseProgram(cl_program prog) {
}
PUBLIC cl_int
-clBuildProgram(cl_program prog, cl_uint count, const cl_device_id *devs,
+clBuildProgram(cl_program _prog, cl_uint count, const cl_device_id *_devs,
const char *opts, void (*pfn_notify)(cl_program, void *),
void *user_data) try {
+ UNWRAP_ICD_PARAM_PROGRAM(prog);
if (!prog)
throw error(CL_INVALID_PROGRAM);
- if (bool(count) != bool(devs) ||
+ if (bool(count) != bool(_devs) ||
(!pfn_notify && user_data))
throw error(CL_INVALID_VALUE);
if (!opts)
opts = "";
- if (devs) {
- if (any_of([&](const cl_device_id dev) {
+ if (_devs) {
+ if (any_of([&](const cl_device_id _dev) {
+ UNWRAP_ICD_PARAM_DEVICE(dev);
return !prog->ctx.has_device(dev);
- }, devs, devs + count))
+ }, _devs, _devs + count))
throw error(CL_INVALID_DEVICE);
- prog->build({ devs, devs + count }, opts);
+ std::vector<clover::device *>devs =
+ clover::device::unwrap_icd_list(_devs, count);
+ prog->build({ devs.begin(), devs.end() }, opts);
} else {
prog->build(prog->ctx.devs, opts);
}
@@ -168,8 +180,9 @@ clUnloadCompiler() {
}
PUBLIC cl_int
-clGetProgramInfo(cl_program prog, cl_program_info param,
+clGetProgramInfo(cl_program _prog, cl_program_info param,
size_t size, void *buf, size_t *size_ret) {
+ UNWRAP_ICD_PARAM_PROGRAM(prog);
if (!prog)
return CL_INVALID_PROGRAM;
@@ -179,15 +192,15 @@ clGetProgramInfo(cl_program prog, cl_program_info param,
prog->ref_count());
case CL_PROGRAM_CONTEXT:
- return scalar_property<cl_context>(buf, size, size_ret,
- &prog->ctx);
+ return scalar_property<clover::context *>(buf, size, size_ret,
+ &prog->ctx);
case CL_PROGRAM_NUM_DEVICES:
return scalar_property<cl_uint>(buf, size, size_ret,
prog->binaries().size());
case CL_PROGRAM_DEVICES:
- return vector_property<cl_device_id>(
+ return vector_property<clover::device *>(
buf, size, size_ret,
map(keys<device *, module>,
prog->binaries().begin(), prog->binaries().end()));
@@ -223,9 +236,11 @@ clGetProgramInfo(cl_program prog, cl_program_info param,
}
PUBLIC cl_int
-clGetProgramBuildInfo(cl_program prog, cl_device_id dev,
+clGetProgramBuildInfo(cl_program _prog, cl_device_id _dev,
cl_program_build_info param,
size_t size, void *buf, size_t *size_ret) {
+ UNWRAP_ICD_PARAM_PROGRAM(prog);
+ UNWRAP_ICD_PARAM_DEVICE(dev);
if (!prog)
return CL_INVALID_PROGRAM;
diff --git a/src/gallium/state_trackers/clover/api/queue.cpp b/src/gallium/state_trackers/clover/api/queue.cpp
index 9c7d7e6..9297761 100644
--- a/src/gallium/state_trackers/clover/api/queue.cpp
+++ b/src/gallium/state_trackers/clover/api/queue.cpp
@@ -20,15 +20,18 @@
// OTHER DEALINGS IN THE SOFTWARE.
//
+#include "api/icd.hpp"
#include "api/util.hpp"
#include "core/queue.hpp"
using namespace clover;
PUBLIC cl_command_queue
-clCreateCommandQueue(cl_context ctx, cl_device_id dev,
+clCreateCommandQueue(cl_context _ctx, cl_device_id _dev,
cl_command_queue_properties props,
cl_int *errcode_ret) try {
+ UNWRAP_ICD_PARAM_CONTEXT(ctx);
+ UNWRAP_ICD_PARAM_DEVICE(dev);
if (!ctx)
throw error(CL_INVALID_CONTEXT);
@@ -40,7 +43,7 @@ clCreateCommandQueue(cl_context ctx, cl_device_id dev,
throw error(CL_INVALID_VALUE);
ret_error(errcode_ret, CL_SUCCESS);
- return new command_queue(*ctx, *dev, props);
+ return new _cl_command_queue(new command_queue(*ctx, *dev, props));
} catch (error &e) {
ret_error(errcode_ret, e);
@@ -48,7 +51,8 @@ clCreateCommandQueue(cl_context ctx, cl_device_id dev,
}
PUBLIC cl_int
-clRetainCommandQueue(cl_command_queue q) {
+clRetainCommandQueue(cl_command_queue _q) {
+ UNWRAP_ICD_PARAM_COMMAND_QUEUE(q);
if (!q)
return CL_INVALID_COMMAND_QUEUE;
@@ -57,7 +61,8 @@ clRetainCommandQueue(cl_command_queue q) {
}
PUBLIC cl_int
-clReleaseCommandQueue(cl_command_queue q) {
+clReleaseCommandQueue(cl_command_queue _q) {
+ UNWRAP_ICD_PARAM_COMMAND_QUEUE(q);
if (!q)
return CL_INVALID_COMMAND_QUEUE;
@@ -68,17 +73,18 @@ clReleaseCommandQueue(cl_command_queue q) {
}
PUBLIC cl_int
-clGetCommandQueueInfo(cl_command_queue q, cl_command_queue_info param,
+clGetCommandQueueInfo(cl_command_queue _q, cl_command_queue_info param,
size_t size, void *buf, size_t *size_ret) {
+ UNWRAP_ICD_PARAM_COMMAND_QUEUE(q);
if (!q)
return CL_INVALID_COMMAND_QUEUE;
switch (param) {
case CL_QUEUE_CONTEXT:
- return scalar_property<cl_context>(buf, size, size_ret, &q->ctx);
+ return scalar_property<clover::context*>(buf, size, size_ret, &q->ctx);
case CL_QUEUE_DEVICE:
- return scalar_property<cl_device_id>(buf, size, size_ret, &q->dev);
+ return scalar_property<clover::device*>(buf, size, size_ret, &q->dev);
case CL_QUEUE_REFERENCE_COUNT:
return scalar_property<cl_uint>(buf, size, size_ret, q->ref_count());
@@ -93,7 +99,8 @@ clGetCommandQueueInfo(cl_command_queue q, cl_command_queue_info param,
}
PUBLIC cl_int
-clFlush(cl_command_queue q) {
+clFlush(cl_command_queue _q) {
+ UNWRAP_ICD_PARAM_COMMAND_QUEUE(q);
if (!q)
return CL_INVALID_COMMAND_QUEUE;
diff --git a/src/gallium/state_trackers/clover/api/sampler.cpp b/src/gallium/state_trackers/clover/api/sampler.cpp
index 4f5ea2f..d0206dc 100644
--- a/src/gallium/state_trackers/clover/api/sampler.cpp
+++ b/src/gallium/state_trackers/clover/api/sampler.cpp
@@ -20,20 +20,23 @@
// OTHER DEALINGS IN THE SOFTWARE.
//
+#include "api/icd.hpp"
#include "api/util.hpp"
#include "core/sampler.hpp"
using namespace clover;
PUBLIC cl_sampler
-clCreateSampler(cl_context ctx, cl_bool norm_mode,
+clCreateSampler(cl_context _ctx, cl_bool norm_mode,
cl_addressing_mode addr_mode, cl_filter_mode filter_mode,
cl_int *errcode_ret) try {
+ UNWRAP_ICD_PARAM_CONTEXT(ctx)
if (!ctx)
throw error(CL_INVALID_CONTEXT);
ret_error(errcode_ret, CL_SUCCESS);
- return new sampler(*ctx, norm_mode, addr_mode, filter_mode);
+ return WRAP_ICD_SAMPLER(new sampler(*ctx, norm_mode, addr_mode,
+ filter_mode));
} catch (error &e) {
ret_error(errcode_ret, e);
@@ -41,7 +44,8 @@ clCreateSampler(cl_context ctx, cl_bool norm_mode,
}
PUBLIC cl_int
-clRetainSampler(cl_sampler s) {
+clRetainSampler(cl_sampler _s) {
+ UNWRAP_ICD_PARAM_SAMPLER(s);
if (!s)
throw error(CL_INVALID_SAMPLER);
@@ -50,7 +54,8 @@ clRetainSampler(cl_sampler s) {
}
PUBLIC cl_int
-clReleaseSampler(cl_sampler s) {
+clReleaseSampler(cl_sampler _s) {
+ UNWRAP_ICD_PARAM_SAMPLER(s);
if (!s)
throw error(CL_INVALID_SAMPLER);
@@ -61,8 +66,9 @@ clReleaseSampler(cl_sampler s) {
}
PUBLIC cl_int
-clGetSamplerInfo(cl_sampler s, cl_sampler_info param,
+clGetSamplerInfo(cl_sampler _s, cl_sampler_info param,
size_t size, void *buf, size_t *size_ret) {
+ UNWRAP_ICD_PARAM_SAMPLER(s);
if (!s)
throw error(CL_INVALID_SAMPLER);
@@ -71,7 +77,7 @@ clGetSamplerInfo(cl_sampler s, cl_sampler_info param,
return scalar_property<cl_uint>(buf, size, size_ret, s->ref_count());
case CL_SAMPLER_CONTEXT:
- return scalar_property<cl_context>(buf, size, size_ret, &s->ctx);
+ return scalar_property<clover::context *>(buf, size, size_ret, &s->ctx);
case CL_SAMPLER_NORMALIZED_COORDS:
return scalar_property<cl_bool>(buf, size, size_ret, s->norm_mode());
diff --git a/src/gallium/state_trackers/clover/api/transfer.cpp b/src/gallium/state_trackers/clover/api/transfer.cpp
index 1615d6f..26005db 100644
--- a/src/gallium/state_trackers/clover/api/transfer.cpp
+++ b/src/gallium/state_trackers/clover/api/transfer.cpp
@@ -22,6 +22,7 @@
#include <cstring>
+#include "api/icd.hpp"
#include "api/util.hpp"
#include "core/event.hpp"
#include "core/resource.hpp"
@@ -35,7 +36,8 @@ namespace {
/// Common argument checking shared by memory transfer commands.
///
void
- validate_base(cl_command_queue q, cl_uint num_deps, const cl_event *deps) {
+ validate_base(clover::command_queue *q, cl_uint num_deps,
+ const cl_event *deps) {
if (!q)
throw error(CL_INVALID_COMMAND_QUEUE);
@@ -44,7 +46,7 @@ namespace {
throw error(CL_INVALID_EVENT_WAIT_LIST);
if (any_of([&](const cl_event ev) {
- return &ev->ctx != &q->ctx;
+ return &ev->__obj->ctx != &q->ctx;
}, deps, deps + num_deps))
throw error(CL_INVALID_CONTEXT);
}
@@ -54,7 +56,7 @@ namespace {
/// transfer commands.
///
void
- validate_obj(cl_command_queue q, cl_mem obj) {
+ validate_obj(clover::command_queue *q, clover::memory_obj *obj) {
if (!obj)
throw error(CL_INVALID_MEM_OBJECT);
@@ -71,7 +73,7 @@ namespace {
template<> struct __map<void *> {
static void *
- get(cl_command_queue q, void *obj, cl_map_flags flags,
+ get(clover::command_queue *q, void *obj, cl_map_flags flags,
size_t offset, size_t size) {
return (char *)obj + offset;
}
@@ -79,7 +81,7 @@ namespace {
template<> struct __map<const void *> {
static const void *
- get(cl_command_queue q, const void *obj, cl_map_flags flags,
+ get(clover::command_queue *q, const void *obj, cl_map_flags flags,
size_t offset, size_t size) {
return (const char *)obj + offset;
}
@@ -87,7 +89,7 @@ namespace {
template<> struct __map<memory_obj *> {
static mapping
- get(cl_command_queue q, memory_obj *obj, cl_map_flags flags,
+ get(clover::command_queue *q, memory_obj *obj, cl_map_flags flags,
size_t offset, size_t size) {
return { *q, obj->resource(q), flags, true, { offset }, { size, 1, 1 }};
}
@@ -99,7 +101,7 @@ namespace {
///
template<typename T, typename S>
std::function<void (event &)>
- soft_copy_op(cl_command_queue q,
+ soft_copy_op(clover::command_queue *q,
T dst_obj, const point &dst_orig, const point &dst_pitch,
S src_obj, const point &src_orig, const point &src_pitch,
const point ®ion) {
@@ -125,7 +127,7 @@ namespace {
///
template<typename T, typename S>
std::function<void (event &)>
- hard_copy_op(cl_command_queue q, T dst_obj, const point &dst_orig,
+ hard_copy_op(clover::command_queue *q, T dst_obj, const point &dst_orig,
S src_obj, const point &src_orig, const point ®ion) {
return [=](event &) {
dst_obj->resource(q).copy(*q, dst_orig, region,
@@ -135,10 +137,12 @@ namespace {
}
PUBLIC cl_int
-clEnqueueReadBuffer(cl_command_queue q, cl_mem obj, cl_bool blocking,
+clEnqueueReadBuffer(cl_command_queue _q, cl_mem _obj, cl_bool blocking,
size_t offset, size_t size, void *ptr,
cl_uint num_deps, const cl_event *deps,
cl_event *ev) try {
+ UNWRAP_ICD_PARAM_COMMAND_QUEUE(q);
+ UNWRAP_ICD_PARAM_MEM(obj);
validate_base(q, num_deps, deps);
validate_obj(q, obj);
@@ -152,7 +156,7 @@ clEnqueueReadBuffer(cl_command_queue q, cl_mem obj, cl_bool blocking,
obj, { offset }, { 1 },
{ size, 1, 1 }));
- ret_object(ev, hev);
+ WRAP_ICD_RET_OBJECT_EVENT(ev, hev);
return CL_SUCCESS;
} catch (error &e) {
@@ -160,10 +164,12 @@ clEnqueueReadBuffer(cl_command_queue q, cl_mem obj, cl_bool blocking,
}
PUBLIC cl_int
-clEnqueueWriteBuffer(cl_command_queue q, cl_mem obj, cl_bool blocking,
+clEnqueueWriteBuffer(cl_command_queue _q, cl_mem _obj, cl_bool blocking,
size_t offset, size_t size, const void *ptr,
cl_uint num_deps, const cl_event *deps,
cl_event *ev) try {
+ UNWRAP_ICD_PARAM_COMMAND_QUEUE(q);
+ UNWRAP_ICD_PARAM_MEM(obj);
validate_base(q, num_deps, deps);
validate_obj(q, obj);
@@ -177,7 +183,7 @@ clEnqueueWriteBuffer(cl_command_queue q, cl_mem obj, cl_bool blocking,
ptr, { 0 }, { 1 },
{ size, 1, 1 }));
- ret_object(ev, hev);
+ WRAP_ICD_RET_OBJECT_EVENT(ev, hev);
return CL_SUCCESS;
} catch (error &e) {
@@ -185,7 +191,7 @@ clEnqueueWriteBuffer(cl_command_queue q, cl_mem obj, cl_bool blocking,
}
PUBLIC cl_int
-clEnqueueReadBufferRect(cl_command_queue q, cl_mem obj, cl_bool blocking,
+clEnqueueReadBufferRect(cl_command_queue _q, cl_mem _obj, cl_bool blocking,
const size_t *obj_origin, const size_t *host_origin,
const size_t *region,
size_t obj_row_pitch, size_t obj_slice_pitch,
@@ -193,6 +199,8 @@ clEnqueueReadBufferRect(cl_command_queue q, cl_mem obj, cl_bool blocking,
void *ptr,
cl_uint num_deps, const cl_event *deps,
cl_event *ev) try {
+ UNWRAP_ICD_PARAM_COMMAND_QUEUE(q);
+ UNWRAP_ICD_PARAM_MEM(obj);
validate_base(q, num_deps, deps);
validate_obj(q, obj);
@@ -208,7 +216,7 @@ clEnqueueReadBufferRect(cl_command_queue q, cl_mem obj, cl_bool blocking,
{ 1, obj_row_pitch, obj_slice_pitch },
region));
- ret_object(ev, hev);
+ WRAP_ICD_RET_OBJECT_EVENT(ev, hev);
return CL_SUCCESS;
} catch (error &e) {
@@ -216,7 +224,7 @@ clEnqueueReadBufferRect(cl_command_queue q, cl_mem obj, cl_bool blocking,
}
PUBLIC cl_int
-clEnqueueWriteBufferRect(cl_command_queue q, cl_mem obj, cl_bool blocking,
+clEnqueueWriteBufferRect(cl_command_queue _q, cl_mem _obj, cl_bool blocking,
const size_t *obj_origin, const size_t *host_origin,
const size_t *region,
size_t obj_row_pitch, size_t obj_slice_pitch,
@@ -224,6 +232,8 @@ clEnqueueWriteBufferRect(cl_command_queue q, cl_mem obj, cl_bool blocking,
const void *ptr,
cl_uint num_deps, const cl_event *deps,
cl_event *ev) try {
+ UNWRAP_ICD_PARAM_COMMAND_QUEUE(q);
+ UNWRAP_ICD_PARAM_MEM(obj);
validate_base(q, num_deps, deps);
validate_obj(q, obj);
@@ -239,7 +249,7 @@ clEnqueueWriteBufferRect(cl_command_queue q, cl_mem obj, cl_bool blocking,
{ 1, host_row_pitch, host_slice_pitch },
region));
- ret_object(ev, hev);
+ WRAP_ICD_RET_OBJECT_EVENT(ev, hev);
return CL_SUCCESS;
} catch (error &e) {
@@ -247,10 +257,13 @@ clEnqueueWriteBufferRect(cl_command_queue q, cl_mem obj, cl_bool blocking,
}
PUBLIC cl_int
-clEnqueueCopyBuffer(cl_command_queue q, cl_mem src_obj, cl_mem dst_obj,
+clEnqueueCopyBuffer(cl_command_queue _q, cl_mem _src_obj, cl_mem _dst_obj,
size_t src_offset, size_t dst_offset, size_t size,
cl_uint num_deps, const cl_event *deps,
cl_event *ev) try {
+ UNWRAP_ICD_PARAM_COMMAND_QUEUE(q);
+ UNWRAP_ICD_PARAM_MEM(src_obj);
+ UNWRAP_ICD_PARAM_MEM(dst_obj);
validate_base(q, num_deps, deps);
validate_obj(q, src_obj);
validate_obj(q, dst_obj);
@@ -261,7 +274,7 @@ clEnqueueCopyBuffer(cl_command_queue q, cl_mem src_obj, cl_mem dst_obj,
src_obj, { src_offset },
{ size, 1, 1 }));
- ret_object(ev, hev);
+ WRAP_ICD_RET_OBJECT_EVENT(ev, hev);
return CL_SUCCESS;
} catch (error &e) {
@@ -269,13 +282,16 @@ clEnqueueCopyBuffer(cl_command_queue q, cl_mem src_obj, cl_mem dst_obj,
}
PUBLIC cl_int
-clEnqueueCopyBufferRect(cl_command_queue q, cl_mem src_obj, cl_mem dst_obj,
+clEnqueueCopyBufferRect(cl_command_queue _q, cl_mem _src_obj, cl_mem _dst_obj,
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_deps, const cl_event *deps,
cl_event *ev) try {
+ UNWRAP_ICD_PARAM_COMMAND_QUEUE(q);
+ UNWRAP_ICD_PARAM_MEM(src_obj);
+ UNWRAP_ICD_PARAM_MEM(dst_obj);
validate_base(q, num_deps, deps);
validate_obj(q, src_obj);
validate_obj(q, dst_obj);
@@ -289,7 +305,7 @@ clEnqueueCopyBufferRect(cl_command_queue q, cl_mem src_obj, cl_mem dst_obj,
{ 1, src_row_pitch, src_slice_pitch },
region));
- ret_object(ev, hev);
+ WRAP_ICD_RET_OBJECT_EVENT(ev, hev);
return CL_SUCCESS;
} catch (error &e) {
@@ -297,11 +313,13 @@ clEnqueueCopyBufferRect(cl_command_queue q, cl_mem src_obj, cl_mem dst_obj,
}
PUBLIC cl_int
-clEnqueueReadImage(cl_command_queue q, cl_mem obj, cl_bool blocking,
+clEnqueueReadImage(cl_command_queue _q, cl_mem _obj, cl_bool blocking,
const size_t *origin, const size_t *region,
size_t row_pitch, size_t slice_pitch, void *ptr,
cl_uint num_deps, const cl_event *deps,
cl_event *ev) try {
+ UNWRAP_ICD_PARAM_COMMAND_QUEUE(q);
+ UNWRAP_ICD_PARAM_MEM(obj);
image *img = dynamic_cast<image *>(obj);
validate_base(q, num_deps, deps);
@@ -319,7 +337,7 @@ clEnqueueReadImage(cl_command_queue q, cl_mem obj, cl_bool blocking,
{ 1, img->row_pitch(), img->slice_pitch() },
region));
- ret_object(ev, hev);
+ WRAP_ICD_RET_OBJECT_EVENT(ev, hev);
return CL_SUCCESS;
} catch (error &e) {
@@ -327,11 +345,13 @@ clEnqueueReadImage(cl_command_queue q, cl_mem obj, cl_bool blocking,
}
PUBLIC cl_int
-clEnqueueWriteImage(cl_command_queue q, cl_mem obj, cl_bool blocking,
+clEnqueueWriteImage(cl_command_queue _q, cl_mem _obj, cl_bool blocking,
const size_t *origin, const size_t *region,
size_t row_pitch, size_t slice_pitch, const void *ptr,
cl_uint num_deps, const cl_event *deps,
cl_event *ev) try {
+ UNWRAP_ICD_PARAM_COMMAND_QUEUE(q);
+ UNWRAP_ICD_PARAM_MEM(obj);
image *img = dynamic_cast<image *>(obj);
validate_base(q, num_deps, deps);
@@ -349,7 +369,7 @@ clEnqueueWriteImage(cl_command_queue q, cl_mem obj, cl_bool blocking,
{ 1, row_pitch, slice_pitch },
region));
- ret_object(ev, hev);
+ WRAP_ICD_RET_OBJECT_EVENT(ev, hev);
return CL_SUCCESS;
} catch (error &e) {
@@ -357,11 +377,14 @@ clEnqueueWriteImage(cl_command_queue q, cl_mem obj, cl_bool blocking,
}
PUBLIC cl_int
-clEnqueueCopyImage(cl_command_queue q, cl_mem src_obj, cl_mem dst_obj,
+clEnqueueCopyImage(cl_command_queue _q, cl_mem _src_obj, cl_mem _dst_obj,
const size_t *src_origin, const size_t *dst_origin,
const size_t *region,
cl_uint num_deps, const cl_event *deps,
cl_event *ev) try {
+ UNWRAP_ICD_PARAM_COMMAND_QUEUE(q);
+ UNWRAP_ICD_PARAM_MEM(src_obj);
+ UNWRAP_ICD_PARAM_MEM(dst_obj);
image *src_img = dynamic_cast<image *>(src_obj);
image *dst_img = dynamic_cast<image *>(dst_obj);
@@ -373,7 +396,7 @@ clEnqueueCopyImage(cl_command_queue q, cl_mem src_obj, cl_mem dst_obj,
*q, CL_COMMAND_COPY_IMAGE, { deps, deps + num_deps },
hard_copy_op(q, dst_obj, dst_origin, src_obj, src_origin, region));
- ret_object(ev, hev);
+ WRAP_ICD_RET_OBJECT_EVENT(ev, hev);
return CL_SUCCESS;
} catch (error &e) {
@@ -381,11 +404,14 @@ clEnqueueCopyImage(cl_command_queue q, cl_mem src_obj, cl_mem dst_obj,
}
PUBLIC cl_int
-clEnqueueCopyImageToBuffer(cl_command_queue q, cl_mem src_obj, cl_mem dst_obj,
+clEnqueueCopyImageToBuffer(cl_command_queue _q, cl_mem _src_obj, cl_mem _dst_obj,
const size_t *src_origin, const size_t *region,
size_t dst_offset,
cl_uint num_deps, const cl_event *deps,
cl_event *ev) try {
+ UNWRAP_ICD_PARAM_COMMAND_QUEUE(q);
+ UNWRAP_ICD_PARAM_MEM(src_obj);
+ UNWRAP_ICD_PARAM_MEM(dst_obj);
image *src_img = dynamic_cast<image *>(src_obj);
validate_base(q, num_deps, deps);
@@ -401,7 +427,7 @@ clEnqueueCopyImageToBuffer(cl_command_queue q, cl_mem src_obj, cl_mem dst_obj,
{ 1, src_img->row_pitch(), src_img->slice_pitch() },
region));
- ret_object(ev, hev);
+ WRAP_ICD_RET_OBJECT_EVENT(ev, hev);
return CL_SUCCESS;
} catch (error &e) {
@@ -409,11 +435,14 @@ clEnqueueCopyImageToBuffer(cl_command_queue q, cl_mem src_obj, cl_mem dst_obj,
}
PUBLIC cl_int
-clEnqueueCopyBufferToImage(cl_command_queue q, cl_mem src_obj, cl_mem dst_obj,
+clEnqueueCopyBufferToImage(cl_command_queue _q, cl_mem _src_obj, cl_mem _dst_obj,
size_t src_offset,
const size_t *dst_origin, const size_t *region,
cl_uint num_deps, const cl_event *deps,
cl_event *ev) try {
+ UNWRAP_ICD_PARAM_COMMAND_QUEUE(q);
+ UNWRAP_ICD_PARAM_MEM(src_obj);
+ UNWRAP_ICD_PARAM_MEM(dst_obj);
image *dst_img = dynamic_cast<image *>(dst_obj);
validate_base(q, num_deps, deps);
@@ -429,7 +458,7 @@ clEnqueueCopyBufferToImage(cl_command_queue q, cl_mem src_obj, cl_mem dst_obj,
{ 0, 0, 0 },
region));
- ret_object(ev, hev);
+ WRAP_ICD_RET_OBJECT_EVENT(ev, hev);
return CL_SUCCESS;
} catch (error &e) {
@@ -437,10 +466,12 @@ clEnqueueCopyBufferToImage(cl_command_queue q, cl_mem src_obj, cl_mem dst_obj,
}
PUBLIC void *
-clEnqueueMapBuffer(cl_command_queue q, cl_mem obj, cl_bool blocking,
+clEnqueueMapBuffer(cl_command_queue _q, cl_mem _obj, cl_bool blocking,
cl_map_flags flags, size_t offset, size_t size,
cl_uint num_deps, const cl_event *deps,
cl_event *ev, cl_int *errcode_ret) try {
+ UNWRAP_ICD_PARAM_COMMAND_QUEUE(q);
+ UNWRAP_ICD_PARAM_MEM(obj);
validate_base(q, num_deps, deps);
validate_obj(q, obj);
@@ -450,7 +481,7 @@ clEnqueueMapBuffer(cl_command_queue q, cl_mem obj, cl_bool blocking,
void *map = obj->resource(q).add_map(
*q, flags, blocking, { offset }, { size });
- ret_object(ev, new hard_event(*q, CL_COMMAND_MAP_BUFFER,
+ WRAP_ICD_RET_OBJECT_EVENT(ev, new hard_event(*q, CL_COMMAND_MAP_BUFFER,
{ deps, deps + num_deps }));
ret_error(errcode_ret, CL_SUCCESS);
return map;
@@ -461,12 +492,14 @@ clEnqueueMapBuffer(cl_command_queue q, cl_mem obj, cl_bool blocking,
}
PUBLIC void *
-clEnqueueMapImage(cl_command_queue q, cl_mem obj, cl_bool blocking,
+clEnqueueMapImage(cl_command_queue _q, cl_mem _obj, cl_bool blocking,
cl_map_flags flags,
const size_t *origin, const size_t *region,
size_t *row_pitch, size_t *slice_pitch,
cl_uint num_deps, const cl_event *deps,
cl_event *ev, cl_int *errcode_ret) try {
+ UNWRAP_ICD_PARAM_COMMAND_QUEUE(q);
+ UNWRAP_ICD_PARAM_MEM(obj);
image *img = dynamic_cast<image *>(obj);
validate_base(q, num_deps, deps);
@@ -475,7 +508,7 @@ clEnqueueMapImage(cl_command_queue q, cl_mem obj, cl_bool blocking,
void *map = obj->resource(q).add_map(
*q, flags, blocking, origin, region);
- ret_object(ev, new hard_event(*q, CL_COMMAND_MAP_IMAGE,
+ WRAP_ICD_RET_OBJECT_EVENT(ev, new hard_event(*q, CL_COMMAND_MAP_IMAGE,
{ deps, deps + num_deps }));
ret_error(errcode_ret, CL_SUCCESS);
return map;
@@ -486,9 +519,11 @@ clEnqueueMapImage(cl_command_queue q, cl_mem obj, cl_bool blocking,
}
PUBLIC cl_int
-clEnqueueUnmapMemObject(cl_command_queue q, cl_mem obj, void *ptr,
+clEnqueueUnmapMemObject(cl_command_queue _q, cl_mem _obj, void *ptr,
cl_uint num_deps, const cl_event *deps,
cl_event *ev) try {
+ UNWRAP_ICD_PARAM_COMMAND_QUEUE(q);
+ UNWRAP_ICD_PARAM_MEM(obj);
validate_base(q, num_deps, deps);
validate_obj(q, obj);
@@ -498,7 +533,7 @@ clEnqueueUnmapMemObject(cl_command_queue q, cl_mem obj, void *ptr,
obj->resource(q).del_map(ptr);
});
- ret_object(ev, hev);
+ WRAP_ICD_RET_OBJECT_EVENT(ev, hev);
return CL_SUCCESS;
} catch (error &e) {
diff --git a/src/gallium/state_trackers/clover/core/context.cpp b/src/gallium/state_trackers/clover/core/context.cpp
index d6817bf..a748b17 100644
--- a/src/gallium/state_trackers/clover/core/context.cpp
+++ b/src/gallium/state_trackers/clover/core/context.cpp
@@ -22,16 +22,16 @@
#include <algorithm>
+#include "api/icd.hpp"
#include "core/context.hpp"
using namespace clover;
-_cl_context::_cl_context(const std::vector<cl_context_properties> &props,
- const std::vector<device *> &devs) :
- devs(devs), __props(props) {
-}
+context::context(const std::vector<cl_context_properties> &props,
+ const std::vector<clover::device *> &devs) :
+ icd_container(new _cl_context(this)), devs(devs), __props(props) { }
bool
-_cl_context::has_device(clover::device *dev) const {
+context::has_device(clover::device *dev) const {
return std::count(devs.begin(), devs.end(), dev);
}
diff --git a/src/gallium/state_trackers/clover/core/context.hpp b/src/gallium/state_trackers/clover/core/context.hpp
index 9d19b15..4b0b4c6 100644
--- a/src/gallium/state_trackers/clover/core/context.hpp
+++ b/src/gallium/state_trackers/clover/core/context.hpp
@@ -27,14 +27,13 @@
#include "core/device.hpp"
namespace clover {
- typedef struct _cl_context context;
-}
+ struct device;
-struct _cl_context : public clover::ref_counter {
+struct context : icd_container<cl_context>, public clover::ref_counter {
public:
- _cl_context(const std::vector<cl_context_properties> &props,
+ context(const std::vector<cl_context_properties> &props,
const std::vector<clover::device *> &devs);
- _cl_context(const _cl_context &ctx) = delete;
+ context(const _cl_context &ctx) = delete;
bool has_device(clover::device *dev) const;
@@ -42,10 +41,11 @@ public:
return __props;
}
- const std::vector<clover::device *> devs;
+ std::vector<clover::device *> devs;
private:
std::vector<cl_context_properties> __props;
};
+} // End namespace clover
#endif
diff --git a/src/gallium/state_trackers/clover/core/device.cpp b/src/gallium/state_trackers/clover/core/device.cpp
index 95a422c..c327b30 100644
--- a/src/gallium/state_trackers/clover/core/device.cpp
+++ b/src/gallium/state_trackers/clover/core/device.cpp
@@ -20,6 +20,7 @@
// OTHER DEALINGS IN THE SOFTWARE.
//
+#include "api/icd.hpp"
#include "core/device.hpp"
#include "pipe/p_screen.h"
#include "pipe/p_state.h"
@@ -38,29 +39,30 @@ namespace {
}
}
-_cl_device_id::_cl_device_id(clover::platform &platform,
- pipe_loader_device *ldev) :
+device::device(clover::platform &platform, pipe_loader_device *ldev) :
+ icd_container(new _cl_device_id(this)),
platform(platform), ldev(ldev) {
pipe = pipe_loader_create_screen(ldev, PIPE_SEARCH_DIR);
if (!pipe || !pipe->get_param(pipe, PIPE_CAP_COMPUTE))
throw error(CL_INVALID_DEVICE);
}
-_cl_device_id::_cl_device_id(_cl_device_id &&dev) :
+device::device(device &&dev) :
+ icd_container(new _cl_device_id(this)),
platform(dev.platform), pipe(dev.pipe), ldev(dev.ldev) {
dev.pipe = NULL;
dev.ldev = NULL;
}
-_cl_device_id::~_cl_device_id() {
+device::~device() {
if (pipe)
pipe->destroy(pipe);
if (ldev)
pipe_loader_release(&ldev, 1);
}
-_cl_device_id &
-_cl_device_id::operator=(_cl_device_id dev) {
+device &
+device::operator=(device dev) {
assert(&platform == &dev.platform);
std::swap(pipe, dev.pipe);
@@ -70,7 +72,7 @@ _cl_device_id::operator=(_cl_device_id dev) {
}
cl_device_type
-_cl_device_id::type() const {
+device::type() const {
switch (ldev->type) {
case PIPE_LOADER_DEVICE_SOFTWARE:
return CL_DEVICE_TYPE_CPU;
@@ -83,7 +85,7 @@ _cl_device_id::type() const {
}
cl_uint
-_cl_device_id::vendor_id() const {
+device::vendor_id() const {
switch (ldev->type) {
case PIPE_LOADER_DEVICE_SOFTWARE:
return 0;
@@ -96,99 +98,118 @@ _cl_device_id::vendor_id() const {
}
size_t
-_cl_device_id::max_images_read() const {
+device::max_images_read() const {
return PIPE_MAX_SHADER_RESOURCES;
}
size_t
-_cl_device_id::max_images_write() const {
+device::max_images_write() const {
return PIPE_MAX_SHADER_RESOURCES;
}
cl_uint
-_cl_device_id::max_image_levels_2d() const {
+device::max_image_levels_2d() const {
return pipe->get_param(pipe, PIPE_CAP_MAX_TEXTURE_2D_LEVELS);
}
cl_uint
-_cl_device_id::max_image_levels_3d() const {
+device::max_image_levels_3d() const {
return pipe->get_param(pipe, PIPE_CAP_MAX_TEXTURE_3D_LEVELS);
}
cl_uint
-_cl_device_id::max_samplers() const {
+device::max_samplers() const {
return pipe->get_shader_param(pipe, PIPE_SHADER_COMPUTE,
PIPE_SHADER_CAP_MAX_TEXTURE_SAMPLERS);
}
cl_ulong
-_cl_device_id::max_mem_global() const {
+device::max_mem_global() const {
return get_compute_param<uint64_t>(pipe,
PIPE_COMPUTE_CAP_MAX_GLOBAL_SIZE)[0];
}
cl_ulong
-_cl_device_id::max_mem_local() const {
+device::max_mem_local() const {
return get_compute_param<uint64_t>(pipe,
PIPE_COMPUTE_CAP_MAX_LOCAL_SIZE)[0];
}
cl_ulong
-_cl_device_id::max_mem_input() const {
+device::max_mem_input() const {
return get_compute_param<uint64_t>(pipe,
PIPE_COMPUTE_CAP_MAX_INPUT_SIZE)[0];
}
cl_ulong
-_cl_device_id::max_const_buffer_size() const {
+device::max_const_buffer_size() const {
return pipe->get_shader_param(pipe, PIPE_SHADER_COMPUTE,
PIPE_SHADER_CAP_MAX_CONSTS) * 16;
}
cl_uint
-_cl_device_id::max_const_buffers() const {
+device::max_const_buffers() const {
return pipe->get_shader_param(pipe, PIPE_SHADER_COMPUTE,
PIPE_SHADER_CAP_MAX_CONST_BUFFERS);
}
size_t
-_cl_device_id::max_threads_per_block() const {
+device::max_threads_per_block() const {
return get_compute_param<uint64_t>(
pipe, PIPE_COMPUTE_CAP_MAX_THREADS_PER_BLOCK)[0];
}
cl_ulong
-_cl_device_id::max_mem_alloc_size() const {
+device::max_mem_alloc_size() const {
return get_compute_param<uint64_t>(pipe,
PIPE_COMPUTE_CAP_MAX_MEM_ALLOC_SIZE)[0];
}
std::vector<size_t>
-_cl_device_id::max_block_size() const {
+device::max_block_size() const {
auto v = get_compute_param<uint64_t>(pipe, PIPE_COMPUTE_CAP_MAX_BLOCK_SIZE);
return { v.begin(), v.end() };
}
std::string
-_cl_device_id::device_name() const {
+device::device_name() const {
return pipe->get_name(pipe);
}
std::string
-_cl_device_id::vendor_name() const {
+device::vendor_name() const {
return pipe->get_vendor(pipe);
}
enum pipe_shader_ir
-_cl_device_id::ir_format() const {
+device::ir_format() const {
return (enum pipe_shader_ir) pipe->get_shader_param(pipe,
PIPE_SHADER_COMPUTE,
PIPE_SHADER_CAP_PREFERRED_IR);
}
std::string
-_cl_device_id::ir_target() const {
+device::ir_target() const {
std::vector<char> target = get_compute_param<char>(pipe,
PIPE_COMPUTE_CAP_IR_TARGET);
return { target.data() };
}
+
+std::vector<clover::device *>
+device::unwrap_icd_list(const cl_device_id *devs, unsigned count) {
+ std::vector<clover::device *> out;
+ for (unsigned i = 0; i < count; ++i) {
+ out.push_back(UNWRAP_ICD_OBJECT(devs[i]));
+ }
+ return out;
+}
+
+std::vector<cl_device_id>
+device::wrap_icd_list(const std::vector<clover::device *> &list) {
+ std::vector<cl_device_id> icd_list;
+ for (std::vector<clover::device*>::const_iterator i = list.begin(),
+ e = list.end(); i != e; ++i) {
+ icd_list.push_back((*i)->icd_obj());
+ }
+ return icd_list;
+}
diff --git a/src/gallium/state_trackers/clover/core/device.hpp b/src/gallium/state_trackers/clover/core/device.hpp
index 465af9f..69748d8 100644
--- a/src/gallium/state_trackers/clover/core/device.hpp
+++ b/src/gallium/state_trackers/clover/core/device.hpp
@@ -26,25 +26,26 @@
#include <set>
#include <vector>
+#include "api/icd.hpp"
#include "core/base.hpp"
#include "core/format.hpp"
#include "pipe-loader/pipe_loader.h"
namespace clover {
- typedef struct _cl_device_id device;
- typedef struct _cl_platform_id platform;
+ struct context;
+ struct platform;
+ class command_queue;
class root_resource;
class hard_event;
-}
-struct _cl_device_id {
+struct device : icd_container <cl_device_id> {
public:
- _cl_device_id(clover::platform &platform, pipe_loader_device *ldev);
- _cl_device_id(_cl_device_id &&dev);
- _cl_device_id(const _cl_device_id &dev) = delete;
- ~_cl_device_id();
+ device(clover::platform &platform, pipe_loader_device *ldev);
+ device(device &&dev);
+ device(const device &dev) = delete;
+ ~device();
- _cl_device_id &operator=(_cl_device_id dev);
+ device &operator=(device dev);
cl_device_type type() const;
cl_uint vendor_id() const;
@@ -66,12 +67,16 @@ public:
std::string vendor_name() const;
enum pipe_shader_ir ir_format() const;
std::string ir_target() const;
+ static std::vector<clover::device *>
+ unwrap_icd_list(const cl_device_id *devs, unsigned count);
+ static std::vector<cl_device_id> wrap_icd_list(
+ const std::vector<clover::device *> &list);
- friend struct _cl_command_queue;
+ friend struct clover::command_queue;
friend class clover::root_resource;
friend class clover::hard_event;
friend std::set<cl_image_format>
- clover::supported_formats(cl_context, cl_mem_object_type);
+ clover::supported_formats(clover::context*, cl_mem_object_type);
clover::platform &platform;
@@ -80,4 +85,5 @@ private:
pipe_loader_device *ldev;
};
+} // End namespace clover
#endif
diff --git a/src/gallium/state_trackers/clover/core/event.cpp b/src/gallium/state_trackers/clover/core/event.cpp
index 93d3b58..c6f4f15 100644
--- a/src/gallium/state_trackers/clover/core/event.cpp
+++ b/src/gallium/state_trackers/clover/core/event.cpp
@@ -20,25 +20,26 @@
// OTHER DEALINGS IN THE SOFTWARE.
//
+#include "api/icd.hpp"
#include "core/event.hpp"
#include "pipe/p_screen.h"
using namespace clover;
-_cl_event::_cl_event(clover::context &ctx,
- std::vector<clover::event *> deps,
+event::event(clover::context &ctx,
+ std::vector<cl_event> deps,
action action_ok, action action_fail) :
- ctx(ctx), __status(0), wait_count(1),
+ icd_container(new _cl_event(this)), ctx(ctx), __status(0), wait_count(1),
action_ok(action_ok), action_fail(action_fail) {
for (auto ev : deps)
- ev->chain(this);
+ UNWRAP_ICD_OBJECT(ev)->chain(this);
}
-_cl_event::~_cl_event() {
+event::~event() {
}
void
-_cl_event::trigger() {
+event::trigger() {
if (!--wait_count) {
action_ok(*this);
@@ -50,7 +51,7 @@ _cl_event::trigger() {
}
void
-_cl_event::abort(cl_int status) {
+event::abort(cl_int status) {
__status = status;
action_fail(*this);
@@ -61,12 +62,12 @@ _cl_event::abort(cl_int status) {
}
bool
-_cl_event::signalled() const {
+event::signalled() const {
return !wait_count;
}
void
-_cl_event::chain(clover::event *ev) {
+event::chain(clover::event *ev) {
if (wait_count) {
ev->wait_count++;
__chain.push_back(ev);
@@ -75,8 +76,8 @@ _cl_event::chain(clover::event *ev) {
}
hard_event::hard_event(clover::command_queue &q, cl_command_type command,
- std::vector<clover::event *> deps, action action) :
- _cl_event(q.ctx, deps, action, [](event &ev){}),
+ std::vector<cl_event> deps, action action) :
+ event(q.ctx, deps, action, [](clover::event &ev){}),
__queue(q), __command(command), __fence(NULL) {
q.sequence(this);
trigger();
@@ -104,7 +105,7 @@ hard_event::status() const {
return CL_COMPLETE;
}
-cl_command_queue
+clover::command_queue*
hard_event::queue() const {
return &__queue;
}
@@ -133,9 +134,9 @@ hard_event::fence(pipe_fence_handle *fence) {
}
soft_event::soft_event(clover::context &ctx,
- std::vector<clover::event *> deps,
+ std::vector<cl_event> deps,
bool __trigger, action action) :
- _cl_event(ctx, deps, action, action) {
+ event(ctx, deps, action, action) {
if (__trigger)
trigger();
}
@@ -155,7 +156,7 @@ soft_event::status() const {
return CL_COMPLETE;
}
-cl_command_queue
+clover::command_queue*
soft_event::queue() const {
return NULL;
}
diff --git a/src/gallium/state_trackers/clover/core/event.hpp b/src/gallium/state_trackers/clover/core/event.hpp
index eb81953..d931af3 100644
--- a/src/gallium/state_trackers/clover/core/event.hpp
+++ b/src/gallium/state_trackers/clover/core/event.hpp
@@ -29,8 +29,8 @@
#include "core/queue.hpp"
namespace clover {
- typedef struct _cl_event event;
-}
+
+class command_queue;
///
/// Class that represents a task that might be executed asynchronously
@@ -49,20 +49,20 @@ namespace clover {
/// the status() method, and it can be waited for completion using the
/// wait() method.
///
-struct _cl_event : public clover::ref_counter {
+struct event : icd_container<cl_event>, public clover::ref_counter {
public:
typedef std::function<void (clover::event &)> action;
- _cl_event(clover::context &ctx, std::vector<clover::event *> deps,
+ event(clover::context &ctx, std::vector<cl_event> deps,
action action_ok, action action_fail);
- virtual ~_cl_event();
+ virtual ~event();
void trigger();
void abort(cl_int status);
bool signalled() const;
virtual cl_int status() const = 0;
- virtual cl_command_queue queue() const = 0;
+ virtual clover::command_queue *queue() const = 0;
virtual cl_command_type command() const = 0;
virtual void wait() const = 0;
@@ -81,7 +81,6 @@ private:
std::vector<clover::ref_ptr<clover::event>> __chain;
};
-namespace clover {
///
/// Class that represents a task executed by a command queue.
///
@@ -97,16 +96,16 @@ namespace clover {
class hard_event : public event {
public:
hard_event(clover::command_queue &q, cl_command_type command,
- std::vector<clover::event *> deps,
- action action = [](event &){});
+ std::vector<cl_event> deps,
+ action action = [](clover::event &){});
~hard_event();
virtual cl_int status() const;
- virtual cl_command_queue queue() const;
+ virtual clover::command_queue *queue() const;
virtual cl_command_type command() const;
virtual void wait() const;
- friend class ::_cl_command_queue;
+ friend class clover::command_queue;
private:
virtual void fence(pipe_fence_handle *fence);
@@ -125,14 +124,15 @@ namespace clover {
///
class soft_event : public event {
public:
- soft_event(clover::context &ctx, std::vector<clover::event *> deps,
- bool trigger, action action = [](event &){});
+ soft_event(clover::context &ctx, std::vector<cl_event> deps,
+ bool trigger, action action = [](clover::event &){});
virtual cl_int status() const;
- virtual cl_command_queue queue() const;
+ virtual clover::command_queue *queue() const;
virtual cl_command_type command() const;
virtual void wait() const;
};
-}
+
+} // End namespace clover
#endif
diff --git a/src/gallium/state_trackers/clover/core/format.cpp b/src/gallium/state_trackers/clover/core/format.cpp
index d0c0153..61c9727 100644
--- a/src/gallium/state_trackers/clover/core/format.cpp
+++ b/src/gallium/state_trackers/clover/core/format.cpp
@@ -145,7 +145,7 @@ namespace clover {
}
std::set<cl_image_format>
- supported_formats(cl_context ctx, cl_mem_object_type type) {
+ supported_formats(clover::context *ctx, cl_mem_object_type type) {
std::set<cl_image_format> s;
pipe_texture_target target = translate_target(type);
unsigned bindings = (PIPE_BIND_SAMPLER_VIEW |
@@ -155,7 +155,7 @@ namespace clover {
for (auto f : formats) {
if (std::all_of(ctx->devs.begin(), ctx->devs.end(),
- [=](const device *dev) {
+ [=](const clover::device *dev) {
return dev->pipe->is_format_supported(
dev->pipe, f.second, target, 1, bindings);
}))
diff --git a/src/gallium/state_trackers/clover/core/format.hpp b/src/gallium/state_trackers/clover/core/format.hpp
index 30b79fd..f31ab14 100644
--- a/src/gallium/state_trackers/clover/core/format.hpp
+++ b/src/gallium/state_trackers/clover/core/format.hpp
@@ -30,6 +30,9 @@
#include "pipe/p_format.h"
namespace clover {
+
+ class context;
+
pipe_texture_target translate_target(cl_mem_object_type type);
pipe_format translate_format(const cl_image_format &format);
@@ -37,7 +40,7 @@ namespace clover {
/// Return all the image formats supported by a given context for
/// the given memory object type.
///
- std::set<cl_image_format> supported_formats(cl_context ctx,
+ std::set<cl_image_format> supported_formats(clover::context *ctx,
cl_mem_object_type type);
}
diff --git a/src/gallium/state_trackers/clover/core/kernel.cpp b/src/gallium/state_trackers/clover/core/kernel.cpp
index 68e4137..e938280 100644
--- a/src/gallium/state_trackers/clover/core/kernel.cpp
+++ b/src/gallium/state_trackers/clover/core/kernel.cpp
@@ -23,13 +23,14 @@
#include "core/kernel.hpp"
#include "core/resource.hpp"
#include "pipe/p_context.h"
+#include "api/icd.hpp"
using namespace clover;
-_cl_kernel::_cl_kernel(clover::program &prog,
+kernel::kernel(clover::program &prog,
const std::string &name,
const std::vector<clover::module::argument> &args) :
- prog(prog), __name(name), exec(*this) {
+ icd_container(new _cl_kernel(this)), prog(prog), __name(name), exec(*this) {
for (auto arg : args) {
if (arg.type == module::argument::scalar)
this->args.emplace_back(new scalar_argument(arg.size));
@@ -61,7 +62,7 @@ pad_vector(clover::command_queue &q, const V &v, T x) {
}
void
-_cl_kernel::launch(clover::command_queue &q,
+kernel::launch(clover::command_queue &q,
const std::vector<size_t> &grid_offset,
const std::vector<size_t> &grid_size,
const std::vector<size_t> &block_size) {
@@ -93,7 +94,7 @@ _cl_kernel::launch(clover::command_queue &q,
}
size_t
-_cl_kernel::mem_local() const {
+kernel::mem_local() const {
size_t sz = 0;
for (auto &arg : args) {
@@ -105,42 +106,42 @@ _cl_kernel::mem_local() const {
}
size_t
-_cl_kernel::mem_private() const {
+kernel::mem_private() const {
return 0;
}
size_t
-_cl_kernel::max_block_size() const {
+kernel::max_block_size() const {
return SIZE_MAX;
}
const std::string &
-_cl_kernel::name() const {
+kernel::name() const {
return __name;
}
std::vector<size_t>
-_cl_kernel::block_size() const {
+kernel::block_size() const {
return { 0, 0, 0 };
}
const clover::module &
-_cl_kernel::module(const clover::command_queue &q) const {
+kernel::module(const clover::command_queue &q) const {
return prog.binaries().find(&q.dev)->second;
}
-_cl_kernel::exec_context::exec_context(clover::kernel &kern) :
+kernel::exec_context::exec_context(clover::kernel &kern) :
kern(kern), q(NULL), mem_local(0), st(NULL) {
}
-_cl_kernel::exec_context::~exec_context() {
+kernel::exec_context::~exec_context() {
if (st)
q->pipe->delete_compute_state(q->pipe, st);
}
void *
-_cl_kernel::exec_context::bind(clover::command_queue *__q) {
+kernel::exec_context::bind(clover::command_queue *__q) {
std::swap(q, __q);
for (auto &arg : kern.args)
@@ -163,7 +164,7 @@ _cl_kernel::exec_context::bind(clover::command_queue *__q) {
}
void
-_cl_kernel::exec_context::unbind() {
+kernel::exec_context::unbind() {
for (auto &arg : kern.args)
arg->unbind(*this);
@@ -176,26 +177,26 @@ _cl_kernel::exec_context::unbind() {
mem_local = 0;
}
-_cl_kernel::argument::argument(size_t size) :
+kernel::argument::argument(size_t size) :
__size(size), __set(false) {
}
bool
-_cl_kernel::argument::set() const {
+kernel::argument::set() const {
return __set;
}
size_t
-_cl_kernel::argument::storage() const {
+kernel::argument::storage() const {
return 0;
}
-_cl_kernel::scalar_argument::scalar_argument(size_t size) :
+kernel::scalar_argument::scalar_argument(size_t size) :
argument(size) {
}
void
-_cl_kernel::scalar_argument::set(size_t size, const void *value) {
+kernel::scalar_argument::set(size_t size, const void *value) {
if (size != __size)
throw error(CL_INVALID_ARG_SIZE);
@@ -204,24 +205,25 @@ _cl_kernel::scalar_argument::set(size_t size, const void *value) {
}
void
-_cl_kernel::scalar_argument::bind(exec_context &ctx) {
+kernel::scalar_argument::bind(exec_context &ctx) {
ctx.input.insert(ctx.input.end(), v.begin(), v.end());
}
void
-_cl_kernel::scalar_argument::unbind(exec_context &ctx) {
+kernel::scalar_argument::unbind(exec_context &ctx) {
}
-_cl_kernel::global_argument::global_argument(size_t size) :
+kernel::global_argument::global_argument(size_t size) :
argument(size) {
}
void
-_cl_kernel::global_argument::set(size_t size, const void *value) {
- if (size != sizeof(cl_mem))
+kernel::global_argument::set(size_t size, const void *value) {
+ if (size != sizeof(clover::memory_obj *))
throw error(CL_INVALID_ARG_SIZE);
- obj = dynamic_cast<clover::buffer *>(*(cl_mem *)value);
+ cl_mem mem = *(cl_mem*)value;
+ obj = dynamic_cast<clover::buffer *>(UNWRAP_ICD_OBJECT(mem));
if (!obj)
throw error(CL_INVALID_MEM_OBJECT);
@@ -229,7 +231,7 @@ _cl_kernel::global_argument::set(size_t size, const void *value) {
}
void
-_cl_kernel::global_argument::bind(exec_context &ctx) {
+kernel::global_argument::bind(exec_context &ctx) {
size_t offset = ctx.input.size();
size_t idx = ctx.g_buffers.size();
@@ -243,20 +245,20 @@ _cl_kernel::global_argument::bind(exec_context &ctx) {
}
void
-_cl_kernel::global_argument::unbind(exec_context &ctx) {
+kernel::global_argument::unbind(exec_context &ctx) {
}
-_cl_kernel::local_argument::local_argument() :
+kernel::local_argument::local_argument() :
argument(sizeof(uint32_t)) {
}
size_t
-_cl_kernel::local_argument::storage() const {
+kernel::local_argument::storage() const {
return __storage;
}
void
-_cl_kernel::local_argument::set(size_t size, const void *value) {
+kernel::local_argument::set(size_t size, const void *value) {
if (value)
throw error(CL_INVALID_ARG_VALUE);
@@ -265,7 +267,7 @@ _cl_kernel::local_argument::set(size_t size, const void *value) {
}
void
-_cl_kernel::local_argument::bind(exec_context &ctx) {
+kernel::local_argument::bind(exec_context &ctx) {
size_t offset = ctx.input.size();
size_t ptr = ctx.mem_local;
@@ -276,19 +278,20 @@ _cl_kernel::local_argument::bind(exec_context &ctx) {
}
void
-_cl_kernel::local_argument::unbind(exec_context &ctx) {
+kernel::local_argument::unbind(exec_context &ctx) {
}
-_cl_kernel::constant_argument::constant_argument() :
+kernel::constant_argument::constant_argument() :
argument(sizeof(uint32_t)) {
}
void
-_cl_kernel::constant_argument::set(size_t size, const void *value) {
- if (size != sizeof(cl_mem))
+kernel::constant_argument::set(size_t size, const void *value) {
+ if (size != sizeof(clover::memory_obj))
throw error(CL_INVALID_ARG_SIZE);
- obj = dynamic_cast<clover::buffer *>(*(cl_mem *)value);
+ cl_mem mem = *(cl_mem*)value;
+ obj = dynamic_cast<clover::buffer *>(UNWRAP_ICD_OBJECT(mem));
if (!obj)
throw error(CL_INVALID_MEM_OBJECT);
@@ -296,7 +299,7 @@ _cl_kernel::constant_argument::set(size_t size, const void *value) {
}
void
-_cl_kernel::constant_argument::bind(exec_context &ctx) {
+kernel::constant_argument::bind(exec_context &ctx) {
size_t offset = ctx.input.size();
size_t idx = ctx.resources.size();
@@ -308,20 +311,21 @@ _cl_kernel::constant_argument::bind(exec_context &ctx) {
}
void
-_cl_kernel::constant_argument::unbind(exec_context &ctx) {
+kernel::constant_argument::unbind(exec_context &ctx) {
obj->resource(ctx.q).unbind_surface(*ctx.q, st);
}
-_cl_kernel::image_rd_argument::image_rd_argument() :
+kernel::image_rd_argument::image_rd_argument() :
argument(sizeof(uint32_t)) {
}
void
-_cl_kernel::image_rd_argument::set(size_t size, const void *value) {
- if (size != sizeof(cl_mem))
+kernel::image_rd_argument::set(size_t size, const void *value) {
+ if (size != sizeof(clover::memory_obj))
throw error(CL_INVALID_ARG_SIZE);
- obj = dynamic_cast<clover::image *>(*(cl_mem *)value);
+ cl_mem mem = *(cl_mem*)value;
+ obj = dynamic_cast<clover::image *>(UNWRAP_ICD_OBJECT(mem));
if (!obj)
throw error(CL_INVALID_MEM_OBJECT);
@@ -329,7 +333,7 @@ _cl_kernel::image_rd_argument::set(size_t size, const void *value) {
}
void
-_cl_kernel::image_rd_argument::bind(exec_context &ctx) {
+kernel::image_rd_argument::bind(exec_context &ctx) {
size_t offset = ctx.input.size();
size_t idx = ctx.sviews.size();
@@ -341,20 +345,21 @@ _cl_kernel::image_rd_argument::bind(exec_context &ctx) {
}
void
-_cl_kernel::image_rd_argument::unbind(exec_context &ctx) {
+kernel::image_rd_argument::unbind(exec_context &ctx) {
obj->resource(ctx.q).unbind_sampler_view(*ctx.q, st);
}
-_cl_kernel::image_wr_argument::image_wr_argument() :
+kernel::image_wr_argument::image_wr_argument() :
argument(sizeof(uint32_t)) {
}
void
-_cl_kernel::image_wr_argument::set(size_t size, const void *value) {
- if (size != sizeof(cl_mem))
+kernel::image_wr_argument::set(size_t size, const void *value) {
+ if (size != sizeof(clover::memory_obj))
throw error(CL_INVALID_ARG_SIZE);
- obj = dynamic_cast<clover::image *>(*(cl_mem *)value);
+ cl_mem mem = *(cl_mem*)value;
+ obj = dynamic_cast<clover::image *>(UNWRAP_ICD_OBJECT(mem));
if (!obj)
throw error(CL_INVALID_MEM_OBJECT);
@@ -362,7 +367,7 @@ _cl_kernel::image_wr_argument::set(size_t size, const void *value) {
}
void
-_cl_kernel::image_wr_argument::bind(exec_context &ctx) {
+kernel::image_wr_argument::bind(exec_context &ctx) {
size_t offset = ctx.input.size();
size_t idx = ctx.resources.size();
@@ -374,25 +379,26 @@ _cl_kernel::image_wr_argument::bind(exec_context &ctx) {
}
void
-_cl_kernel::image_wr_argument::unbind(exec_context &ctx) {
+kernel::image_wr_argument::unbind(exec_context &ctx) {
obj->resource(ctx.q).unbind_surface(*ctx.q, st);
}
-_cl_kernel::sampler_argument::sampler_argument() :
+kernel::sampler_argument::sampler_argument() :
argument(0) {
}
void
-_cl_kernel::sampler_argument::set(size_t size, const void *value) {
- if (size != sizeof(cl_sampler))
+kernel::sampler_argument::set(size_t size, const void *value) {
+ if (size != sizeof(clover::sampler))
throw error(CL_INVALID_ARG_SIZE);
- obj = *(cl_sampler *)value;
+ cl_sampler sampler = *(cl_sampler*)value;
+ obj = UNWRAP_ICD_OBJECT(sampler);
__set = true;
}
void
-_cl_kernel::sampler_argument::bind(exec_context &ctx) {
+kernel::sampler_argument::bind(exec_context &ctx) {
size_t idx = ctx.samplers.size();
ctx.samplers.resize(idx + 1);
@@ -400,6 +406,6 @@ _cl_kernel::sampler_argument::bind(exec_context &ctx) {
}
void
-_cl_kernel::sampler_argument::unbind(exec_context &ctx) {
+kernel::sampler_argument::unbind(exec_context &ctx) {
obj->unbind(*ctx.q, st);
}
diff --git a/src/gallium/state_trackers/clover/core/kernel.hpp b/src/gallium/state_trackers/clover/core/kernel.hpp
index fd64f3e..c03705f 100644
--- a/src/gallium/state_trackers/clover/core/kernel.hpp
+++ b/src/gallium/state_trackers/clover/core/kernel.hpp
@@ -32,11 +32,9 @@
#include "pipe/p_state.h"
namespace clover {
- typedef struct _cl_kernel kernel;
class argument;
-}
-struct _cl_kernel : public clover::ref_counter {
+struct kernel : icd_container<cl_kernel>, public clover::ref_counter {
private:
///
/// Class containing all the state required to execute a compute
@@ -94,7 +92,7 @@ public:
bool __set;
};
- _cl_kernel(clover::program &prog,
+ kernel(clover::program &prog,
const std::string &name,
const std::vector<clover::module::argument> &args);
@@ -211,4 +209,5 @@ private:
exec_context exec;
};
+} // End namespace clover
#endif
diff --git a/src/gallium/state_trackers/clover/core/memory.cpp b/src/gallium/state_trackers/clover/core/memory.cpp
index 63050b3..c8afdd9 100644
--- a/src/gallium/state_trackers/clover/core/memory.cpp
+++ b/src/gallium/state_trackers/clover/core/memory.cpp
@@ -21,12 +21,14 @@
//
#include "core/memory.hpp"
+#include "core/queue.hpp"
#include "core/resource.hpp"
using namespace clover;
-_cl_mem::_cl_mem(clover::context &ctx, cl_mem_flags flags,
+memory_obj::memory_obj(clover::context &ctx, cl_mem_flags flags,
size_t size, void *host_ptr) :
+ icd_container(new _cl_mem(this)),
ctx(ctx), __flags(flags),
__size(size), __host_ptr(host_ptr),
__destroy_notify([]{}) {
@@ -34,27 +36,27 @@ _cl_mem::_cl_mem(clover::context &ctx, cl_mem_flags flags,
data.append((char *)host_ptr, size);
}
-_cl_mem::~_cl_mem() {
+memory_obj::~memory_obj() {
__destroy_notify();
}
void
-_cl_mem::destroy_notify(std::function<void ()> f) {
+memory_obj::destroy_notify(std::function<void ()> f) {
__destroy_notify = f;
}
cl_mem_flags
-_cl_mem::flags() const {
+memory_obj::flags() const {
return __flags;
}
size_t
-_cl_mem::size() const {
+memory_obj::size() const {
return __size;
}
void *
-_cl_mem::host_ptr() const {
+memory_obj::host_ptr() const {
return __host_ptr;
}
@@ -74,7 +76,7 @@ root_buffer::root_buffer(clover::context &ctx, cl_mem_flags flags,
}
clover::resource &
-root_buffer::resource(cl_command_queue q) {
+root_buffer::resource(clover::command_queue *q) {
// Create a new resource if there's none for this device yet.
if (!resources.count(&q->dev)) {
auto r = (!resources.empty() ?
@@ -97,7 +99,7 @@ sub_buffer::sub_buffer(clover::root_buffer &parent, cl_mem_flags flags,
}
clover::resource &
-sub_buffer::resource(cl_command_queue q) {
+sub_buffer::resource(clover::command_queue *q) {
// Create a new resource if there's none for this device yet.
if (!resources.count(&q->dev)) {
auto r = new sub_resource(parent.resource(q), { offset() });
@@ -125,7 +127,7 @@ image::image(clover::context &ctx, cl_mem_flags flags,
}
clover::resource &
-image::resource(cl_command_queue q) {
+image::resource(clover::command_queue *q) {
// Create a new resource if there's none for this device yet.
if (!resources.count(&q->dev)) {
auto r = (!resources.empty() ?
diff --git a/src/gallium/state_trackers/clover/core/memory.hpp b/src/gallium/state_trackers/clover/core/memory.hpp
index c71a03c..f36f562 100644
--- a/src/gallium/state_trackers/clover/core/memory.hpp
+++ b/src/gallium/state_trackers/clover/core/memory.hpp
@@ -31,23 +31,21 @@
#include "core/queue.hpp"
namespace clover {
- typedef struct _cl_mem memory_obj;
class resource;
class sub_resource;
-}
-struct _cl_mem : public clover::ref_counter {
+struct memory_obj : icd_container<cl_mem>, public clover::ref_counter {
protected:
- _cl_mem(clover::context &ctx, cl_mem_flags flags,
+ memory_obj(clover::context &ctx, cl_mem_flags flags,
size_t size, void *host_ptr);
- _cl_mem(const _cl_mem &obj) = delete;
+ memory_obj(const memory_obj &obj) = delete;
public:
- virtual ~_cl_mem();
+ virtual ~memory_obj();
virtual cl_mem_object_type type() const = 0;
- virtual clover::resource &resource(cl_command_queue q) = 0;
+ virtual clover::resource &resource(clover::command_queue *q) = 0;
void destroy_notify(std::function<void ()> f);
cl_mem_flags flags() const;
@@ -66,7 +64,6 @@ protected:
std::string data;
};
-namespace clover {
struct buffer : public memory_obj {
protected:
buffer(clover::context &ctx, cl_mem_flags flags,
@@ -81,7 +78,7 @@ namespace clover {
root_buffer(clover::context &ctx, cl_mem_flags flags,
size_t size, void *host_ptr);
- virtual clover::resource &resource(cl_command_queue q);
+ virtual clover::resource &resource(clover::command_queue *q);
private:
std::map<clover::device *,
@@ -93,7 +90,7 @@ namespace clover {
sub_buffer(clover::root_buffer &parent, cl_mem_flags flags,
size_t offset, size_t size);
- virtual clover::resource &resource(cl_command_queue q);
+ virtual clover::resource &resource(clover::command_queue *q);
size_t offset() const;
clover::root_buffer &parent;
@@ -113,7 +110,7 @@ namespace clover {
void *host_ptr);
public:
- virtual clover::resource &resource(cl_command_queue q);
+ virtual clover::resource &resource(clover::command_queue *q);
cl_image_format format() const;
size_t width() const;
size_t height() const;
diff --git a/src/gallium/state_trackers/clover/core/platform.cpp b/src/gallium/state_trackers/clover/core/platform.cpp
index 6d002e7..a77639d 100644
--- a/src/gallium/state_trackers/clover/core/platform.cpp
+++ b/src/gallium/state_trackers/clover/core/platform.cpp
@@ -20,11 +20,12 @@
// OTHER DEALINGS IN THE SOFTWARE.
//
+#include "api/icd.hpp"
#include "core/platform.hpp"
using namespace clover;
-_cl_platform_id::_cl_platform_id() {
+platform::platform() : icd_container(new _cl_platform_id(this)){
int n = pipe_loader_probe(NULL, 0);
std::vector<pipe_loader_device *> ldevs(n);
diff --git a/src/gallium/state_trackers/clover/core/platform.hpp b/src/gallium/state_trackers/clover/core/platform.hpp
index 831a742..51c5f86 100644
--- a/src/gallium/state_trackers/clover/core/platform.hpp
+++ b/src/gallium/state_trackers/clover/core/platform.hpp
@@ -29,14 +29,12 @@
#include "core/device.hpp"
namespace clover {
- typedef struct _cl_platform_id platform;
-}
-struct _cl_platform_id {
+struct platform : icd_container <cl_platform_id> {
public:
typedef std::vector<clover::device>::iterator iterator;
- _cl_platform_id();
+ platform();
///
/// Container of all compute devices that are available in the platform.
@@ -63,4 +61,5 @@ protected:
std::vector<clover::device> devs;
};
+} // End namespace clover
#endif
diff --git a/src/gallium/state_trackers/clover/core/program.cpp b/src/gallium/state_trackers/clover/core/program.cpp
index e85f27a..3712766 100644
--- a/src/gallium/state_trackers/clover/core/program.cpp
+++ b/src/gallium/state_trackers/clover/core/program.cpp
@@ -25,15 +25,15 @@
using namespace clover;
-_cl_program::_cl_program(clover::context &ctx,
+program::program(clover::context &ctx,
const std::string &source) :
- ctx(ctx), __source(source) {
+ icd_container(new _cl_program(this)), ctx(ctx), __source(source) {
}
-_cl_program::_cl_program(clover::context &ctx,
+program::program(clover::context &ctx,
const std::vector<clover::device *> &devs,
const std::vector<clover::module> &binaries) :
- ctx(ctx) {
+ icd_container(new _cl_program(this)), ctx(ctx) {
for_each([&](clover::device *dev, const clover::module &bin) {
__binaries.insert({ dev, bin });
},
@@ -41,7 +41,7 @@ _cl_program::_cl_program(clover::context &ctx,
}
void
-_cl_program::build(const std::vector<clover::device *> &devs,
+program::build(const std::vector<clover::device *> &devs,
const char *opts) {
for (auto dev : devs) {
@@ -67,26 +67,26 @@ _cl_program::build(const std::vector<clover::device *> &devs,
}
const std::string &
-_cl_program::source() const {
+program::source() const {
return __source;
}
const std::map<clover::device *, clover::module> &
-_cl_program::binaries() const {
+program::binaries() const {
return __binaries;
}
cl_build_status
-_cl_program::build_status(clover::device *dev) const {
+program::build_status(clover::device *dev) const {
return __binaries.count(dev) ? CL_BUILD_SUCCESS : CL_BUILD_NONE;
}
std::string
-_cl_program::build_opts(clover::device *dev) const {
+program::build_opts(clover::device *dev) const {
return __opts.count(dev) ? __opts.find(dev)->second : "";
}
std::string
-_cl_program::build_log(clover::device *dev) const {
+program::build_log(clover::device *dev) const {
return __logs.count(dev) ? __logs.find(dev)->second : "";
}
diff --git a/src/gallium/state_trackers/clover/core/program.hpp b/src/gallium/state_trackers/clover/core/program.hpp
index a1452df..b71aa50 100644
--- a/src/gallium/state_trackers/clover/core/program.hpp
+++ b/src/gallium/state_trackers/clover/core/program.hpp
@@ -30,14 +30,12 @@
#include "core/module.hpp"
namespace clover {
- typedef struct _cl_program program;
-}
-struct _cl_program : public clover::ref_counter {
+struct program : icd_container<cl_program>, public clover::ref_counter {
public:
- _cl_program(clover::context &ctx,
+ program(clover::context &ctx,
const std::string &source);
- _cl_program(clover::context &ctx,
+ program(clover::context &ctx,
const std::vector<clover::device *> &devs,
const std::vector<clover::module> &binaries);
@@ -59,4 +57,5 @@ private:
std::string __source;
};
+} // End namespace clover
#endif
diff --git a/src/gallium/state_trackers/clover/core/queue.cpp b/src/gallium/state_trackers/clover/core/queue.cpp
index 0b1c494..4130fb2 100644
--- a/src/gallium/state_trackers/clover/core/queue.cpp
+++ b/src/gallium/state_trackers/clover/core/queue.cpp
@@ -29,20 +29,21 @@
using namespace clover;
-_cl_command_queue::_cl_command_queue(context &ctx, device &dev,
+command_queue::command_queue(context &ctx, device &dev,
cl_command_queue_properties props) :
- ctx(ctx), dev(dev), __props(props) {
+ icd_container(new _cl_command_queue(this)), ctx(ctx), dev(dev),
+ __props(props) {
pipe = dev.pipe->context_create(dev.pipe, NULL);
if (!pipe)
throw error(CL_INVALID_DEVICE);
}
-_cl_command_queue::~_cl_command_queue() {
+command_queue::~command_queue() {
pipe->destroy(pipe);
}
void
-_cl_command_queue::flush() {
+command_queue::flush() {
pipe_screen *screen = dev.pipe;
pipe_fence_handle *fence = NULL;
@@ -61,7 +62,7 @@ _cl_command_queue::flush() {
}
void
-_cl_command_queue::sequence(clover::hard_event *ev) {
+command_queue::sequence(clover::hard_event *ev) {
if (!queued_events.empty())
queued_events.back()->chain(ev);
diff --git a/src/gallium/state_trackers/clover/core/queue.hpp b/src/gallium/state_trackers/clover/core/queue.hpp
index eee8527..8b62847 100644
--- a/src/gallium/state_trackers/clover/core/queue.hpp
+++ b/src/gallium/state_trackers/clover/core/queue.hpp
@@ -28,18 +28,22 @@
#include "pipe/p_context.h"
namespace clover {
- typedef struct _cl_command_queue command_queue;
+ struct context;
+ struct device;
+ struct kernel;
+ class root_resource;
class resource;
class mapping;
+ class sampler;
class hard_event;
-}
-struct _cl_command_queue : public clover::ref_counter {
+struct command_queue : icd_container<cl_command_queue>,
+ public clover::ref_counter {
public:
- _cl_command_queue(clover::context &ctx, clover::device &dev,
+ command_queue(clover::context &ctx, clover::device &dev,
cl_command_queue_properties props);
- _cl_command_queue(const _cl_command_queue &q) = delete;
- ~_cl_command_queue();
+ command_queue(const command_queue &q) = delete;
+ ~command_queue();
void flush();
@@ -54,8 +58,8 @@ public:
friend class clover::root_resource;
friend class clover::mapping;
friend class clover::hard_event;
- friend struct _cl_sampler;
- friend struct _cl_kernel;
+ friend struct clover::sampler;
+ friend struct clover::kernel;
private:
/// Serialize a hardware event with respect to the previous ones,
@@ -69,4 +73,5 @@ private:
std::vector<event_ptr> queued_events;
};
+} // End namespace clover
#endif
diff --git a/src/gallium/state_trackers/clover/core/resource.hpp b/src/gallium/state_trackers/clover/core/resource.hpp
index f5a4deb..a25bfa3 100644
--- a/src/gallium/state_trackers/clover/core/resource.hpp
+++ b/src/gallium/state_trackers/clover/core/resource.hpp
@@ -57,7 +57,7 @@ namespace clover {
friend class sub_resource;
friend class mapping;
- friend struct ::_cl_kernel;
+ friend struct clover::kernel;
protected:
resource(clover::device &dev, clover::memory_obj &obj);
diff --git a/src/gallium/state_trackers/clover/core/sampler.cpp b/src/gallium/state_trackers/clover/core/sampler.cpp
index 55a66a1..83e1e1c 100644
--- a/src/gallium/state_trackers/clover/core/sampler.cpp
+++ b/src/gallium/state_trackers/clover/core/sampler.cpp
@@ -25,30 +25,30 @@
using namespace clover;
-_cl_sampler::_cl_sampler(clover::context &ctx, bool norm_mode,
+sampler::sampler(clover::context &ctx, bool norm_mode,
cl_addressing_mode addr_mode,
cl_filter_mode filter_mode) :
- ctx(ctx), __norm_mode(norm_mode),
+ icd_container(new _cl_sampler(this)), ctx(ctx), __norm_mode(norm_mode),
__addr_mode(addr_mode), __filter_mode(filter_mode) {
}
bool
-_cl_sampler::norm_mode() {
+sampler::norm_mode() {
return __norm_mode;
}
cl_addressing_mode
-_cl_sampler::addr_mode() {
+sampler::addr_mode() {
return __addr_mode;
}
cl_filter_mode
-_cl_sampler::filter_mode() {
+sampler::filter_mode() {
return __filter_mode;
}
void *
-_cl_sampler::bind(clover::command_queue &q) {
+sampler::bind(clover::command_queue &q) {
struct pipe_sampler_state info {};
info.normalized_coords = norm_mode();
@@ -68,6 +68,6 @@ _cl_sampler::bind(clover::command_queue &q) {
}
void
-_cl_sampler::unbind(clover::command_queue &q, void *st) {
+sampler::unbind(clover::command_queue &q, void *st) {
q.pipe->delete_sampler_state(q.pipe, st);
}
diff --git a/src/gallium/state_trackers/clover/core/sampler.hpp b/src/gallium/state_trackers/clover/core/sampler.hpp
index 850240e..a8a4f2e 100644
--- a/src/gallium/state_trackers/clover/core/sampler.hpp
+++ b/src/gallium/state_trackers/clover/core/sampler.hpp
@@ -27,12 +27,10 @@
#include "core/queue.hpp"
namespace clover {
- typedef struct _cl_sampler sampler;
-}
-struct _cl_sampler : public clover::ref_counter {
+struct sampler : icd_container<cl_sampler>, public clover::ref_counter {
public:
- _cl_sampler(clover::context &ctx, bool norm_mode,
+ sampler(clover::context &ctx, bool norm_mode,
cl_addressing_mode addr_mode, cl_filter_mode filter_mode);
bool norm_mode();
@@ -41,7 +39,7 @@ public:
clover::context &ctx;
- friend class _cl_kernel;
+ friend class clover::kernel;
private:
void *bind(clover::command_queue &q);
@@ -52,4 +50,6 @@ private:
cl_filter_mode __filter_mode;
};
+} // End namespace clover
+
#endif
diff --git a/src/gallium/targets/opencl/Makefile.am b/src/gallium/targets/opencl/Makefile.am
index 46bb29f..dcba8bb 100644
--- a/src/gallium/targets/opencl/Makefile.am
+++ b/src/gallium/targets/opencl/Makefile.am
@@ -1,12 +1,11 @@
AUTOMAKE_OPTIONS = subdir-objects
-lib_LTLIBRARIES = libOpenCL.la
+lib_LTLIBRARIES = lib at OPENCL_LIBNAME@.la
-libOpenCL_la_LDFLAGS = \
+lib at OPENCL_LIBNAME@_la_LDFLAGS = \
$(LLVM_LDFLAGS) \
-version-number 1:0
-
-libOpenCL_la_LIBADD = \
+lib at OPENCL_LIBNAME@_la_LIBADD = \
$(top_builddir)/src/gallium/auxiliary/pipe-loader/libpipe_loader.la \
$(top_builddir)/src/gallium/winsys/sw/null/libws_null.la \
$(top_builddir)/src/gallium/state_trackers/clover/libclover.la \
@@ -29,13 +28,20 @@ libOpenCL_la_LIBADD = \
$(LLVM_LIBS)
-libOpenCL_la_SOURCES =
+lib at OPENCL_LIBNAME@_la_SOURCES =
# Force usage of a C++ linker
-nodist_EXTRA_libOpenCL_la_SOURCES = dummy.cpp
+nodist_EXTRA_lib at 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: libOpenCL.la
+all-local: lib at OPENCL_LIBNAME@.la
$(MKDIR_P) $(top_builddir)/$(LIB_DIR)
- ln -f .libs/libOpenCL.so* $(top_builddir)/$(LIB_DIR)/
+ ln -f .libs/lib at OPENCL_LIBNAME@.so* $(top_builddir)/$(LIB_DIR)/
--
1.7.11.4
More information about the mesa-dev
mailing list