clover: Import OpenCL state tracker.

11 files changed, 2417 insertions, 0 deletions

diff --git a/src/gallium/state_trackers/clover/api/context.cpp b/src/gallium/state_trackers/clover/api/context.cpp
new file mode 100644
index 00000000000..c8d668933e5
--- /dev/null
+++ b/src/gallium/state_trackers/clover/api/context.cpp
@@ -0,0 +1,120 @@
+//
+// Copyright 2012 Francisco Jerez
+//
+// 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 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.
+//
+
+#include "api/util.hpp"
+#include "core/context.hpp"
+
+using namespace clover;
+
+PUBLIC cl_context
+clCreateContext(const cl_context_properties *props, cl_uint num_devs,
+                const cl_device_id *devs,
+                void (CL_CALLBACK *pfn_notify)(const char *, const void *,
+                                               size_t, void *),
+                void *user_data, cl_int *errcode_ret) try {
+   auto mprops = property_map(props);
+
+   if (!devs || !num_devs ||
+       (!pfn_notify && user_data))
+      throw error(CL_INVALID_VALUE);
+
+   if (any_of(is_zero<cl_device_id>(), devs, devs + num_devs))
+      throw error(CL_INVALID_DEVICE);
+
+   for (auto p : mprops) {
+      if (!(p.first == CL_CONTEXT_PLATFORM &&
+            (cl_platform_id)p.second == NULL))
+         throw error(CL_INVALID_PROPERTY);
+   }
+
+   ret_error(errcode_ret, CL_SUCCESS);
+   return new context(
+      property_vector(mprops),
+      std::vector<cl_device_id>(devs, devs + num_devs));
+
+} catch(error &e) {
+   ret_error(errcode_ret, e);
+   return NULL;
+}
+
+PUBLIC cl_context
+clCreateContextFromType(const cl_context_properties *props,
+                        cl_device_type type,
+                        void (CL_CALLBACK *pfn_notify)(
+                           const char *, const void *, size_t, void *),
+                        void *user_data, cl_int *errcode_ret) {
+   cl_device_id dev;
+   cl_int ret;
+
+   ret = clGetDeviceIDs(0, type, 1, &dev, 0);
+   if (ret) {
+      ret_error(errcode_ret, ret);
+      return NULL;
+   }
+
+   return clCreateContext(props, 1, &dev, pfn_notify, user_data, errcode_ret);
+}
+
+PUBLIC cl_int
+clRetainContext(cl_context ctx) {
+   if (!ctx)
+      return CL_INVALID_CONTEXT;
+
+   ctx->retain();
+   return CL_SUCCESS;
+}
+
+PUBLIC cl_int
+clReleaseContext(cl_context ctx) {
+   if (!ctx)
+      return CL_INVALID_CONTEXT;
+
+   if (ctx->release())
+      delete ctx;
+
+   return CL_SUCCESS;
+}
+
+PUBLIC cl_int
+clGetContextInfo(cl_context ctx, cl_context_info param,
+                 size_t size, void *buf, size_t *size_ret) {
+   if (!ctx)
+      return CL_INVALID_CONTEXT;
+
+   switch (param) {
+   case CL_CONTEXT_REFERENCE_COUNT:
+      return scalar_property<cl_uint>(buf, size, size_ret, ctx->ref_count());
+
+   case CL_CONTEXT_NUM_DEVICES:
+      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);
+
+   case CL_CONTEXT_PROPERTIES:
+      return vector_property<cl_context_properties>(buf, size, size_ret,
+                                                    ctx->props());
+
+   default:
+      return CL_INVALID_VALUE;
+   }
+}
diff --git a/src/gallium/state_trackers/clover/api/device.cpp b/src/gallium/state_trackers/clover/api/device.cpp
new file mode 100644
index 00000000000..03767519aaf
--- /dev/null
+++ b/src/gallium/state_trackers/clover/api/device.cpp
@@ -0,0 +1,262 @@
+//
+// Copyright 2012 Francisco Jerez
+//
+// 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 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.
+//
+
+#include "api/util.hpp"
+#include "core/device.hpp"
+
+using namespace clover;
+
+static device_registry registry;
+
+PUBLIC cl_int
+clGetDeviceIDs(cl_platform_id platform, cl_device_type device_type,
+               cl_uint num_entries, cl_device_id *devices,
+               cl_uint *num_devices) {
+   std::vector<cl_device_id> devs;
+
+   if (platform != NULL)
+      return CL_INVALID_PLATFORM;
+
+   if ((!num_entries && devices) ||
+       (!num_devices && !devices))
+      return CL_INVALID_VALUE;
+
+   // Collect matching devices
+   for (device &dev : registry) {
+      if (((device_type & CL_DEVICE_TYPE_DEFAULT) &&
+           &dev == &registry.front()) ||
+          (device_type & dev.type()))
+         devs.push_back(&dev);
+   }
+
+   if (devs.empty())
+      return CL_DEVICE_NOT_FOUND;
+
+   // ...and return the requested data.
+   if (num_devices)
+      *num_devices = devs.size();
+   if (devices)
+      std::copy_n(devs.begin(),
+                  std::min((cl_uint)devs.size(), num_entries),
+                  devices);
+
+   return CL_SUCCESS;
+}
+
+PUBLIC cl_int
+clGetDeviceInfo(cl_device_id dev, cl_device_info param,
+                size_t size, void *buf, size_t *size_ret) {
+   if (!dev)
+      return CL_INVALID_DEVICE;
+
+   switch (param) {
+   case CL_DEVICE_TYPE:
+      return scalar_property<cl_device_type>(buf, size, size_ret, dev->type());
+
+   case CL_DEVICE_VENDOR_ID:
+      return scalar_property<cl_uint>(buf, size, size_ret, dev->vendor_id());
+
+   case CL_DEVICE_MAX_COMPUTE_UNITS:
+      return scalar_property<cl_uint>(buf, size, size_ret, 1);
+
+   case CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS:
+      return scalar_property<cl_uint>(buf, size, size_ret,
+                                      dev->max_block_size().size());
+
+   case CL_DEVICE_MAX_WORK_ITEM_SIZES:
+      return vector_property<size_t>(buf, size, size_ret,
+                                     dev->max_block_size());
+
+   case CL_DEVICE_MAX_WORK_GROUP_SIZE:
+      return scalar_property<size_t>(buf, size, size_ret, SIZE_MAX);
+
+   case CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR:
+      return scalar_property<cl_uint>(buf, size, size_ret, 16);
+
+   case CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT:
+      return scalar_property<cl_uint>(buf, size, size_ret, 8);
+
+   case CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT:
+      return scalar_property<cl_uint>(buf, size, size_ret, 4);
+
+   case CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG:
+      return scalar_property<cl_uint>(buf, size, size_ret, 2);
+
+   case CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT:
+      return scalar_property<cl_uint>(buf, size, size_ret, 4);
+
+   case CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE:
+      return scalar_property<cl_uint>(buf, size, size_ret, 2);
+
+   case CL_DEVICE_PREFERRED_VECTOR_WIDTH_HALF:
+      return scalar_property<cl_uint>(buf, size, size_ret, 0);
+
+   case CL_DEVICE_MAX_CLOCK_FREQUENCY:
+      return scalar_property<cl_uint>(buf, size, size_ret, 0);
+
+   case CL_DEVICE_ADDRESS_BITS:
+      return scalar_property<cl_uint>(buf, size, size_ret, 32);
+
+   case CL_DEVICE_MAX_READ_IMAGE_ARGS:
+      return scalar_property<cl_uint>(buf, size, size_ret,
+                                      dev->max_images_read());
+
+   case CL_DEVICE_MAX_WRITE_IMAGE_ARGS:
+      return scalar_property<cl_uint>(buf, size, size_ret,
+                                      dev->max_images_write());
+
+   case CL_DEVICE_MAX_MEM_ALLOC_SIZE:
+      return scalar_property<cl_ulong>(buf, size, size_ret, 0);
+
+   case CL_DEVICE_IMAGE2D_MAX_WIDTH:
+   case CL_DEVICE_IMAGE2D_MAX_HEIGHT:
+      return scalar_property<size_t>(buf, size, size_ret,
+                                     1 << dev->max_image_levels_2d());
+
+   case CL_DEVICE_IMAGE3D_MAX_WIDTH:
+   case CL_DEVICE_IMAGE3D_MAX_HEIGHT:
+   case CL_DEVICE_IMAGE3D_MAX_DEPTH:
+      return scalar_property<size_t>(buf, size, size_ret,
+                                     1 << dev->max_image_levels_3d());
+
+   case CL_DEVICE_IMAGE_SUPPORT:
+      return scalar_property<cl_bool>(buf, size, size_ret, CL_TRUE);
+
+   case CL_DEVICE_MAX_PARAMETER_SIZE:
+      return scalar_property<size_t>(buf, size, size_ret,
+                                     dev->max_mem_input());
+
+   case CL_DEVICE_MAX_SAMPLERS:
+      return scalar_property<cl_uint>(buf, size, size_ret,
+                                      dev->max_samplers());
+
+   case CL_DEVICE_MEM_BASE_ADDR_ALIGN:
+   case CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE:
+      return scalar_property<cl_uint>(buf, size, size_ret, 128);
+
+   case CL_DEVICE_SINGLE_FP_CONFIG:
+      return scalar_property<cl_device_fp_config>(buf, size, size_ret,
+         CL_FP_DENORM | CL_FP_INF_NAN | CL_FP_ROUND_TO_NEAREST);
+
+   case CL_DEVICE_GLOBAL_MEM_CACHE_TYPE:
+      return scalar_property<cl_device_mem_cache_type>(buf, size, size_ret,
+                                                       CL_NONE);
+
+   case CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE:
+      return scalar_property<cl_uint>(buf, size, size_ret, 0);
+
+   case CL_DEVICE_GLOBAL_MEM_CACHE_SIZE:
+      return scalar_property<cl_ulong>(buf, size, size_ret, 0);
+
+   case CL_DEVICE_GLOBAL_MEM_SIZE:
+      return scalar_property<cl_ulong>(buf, size, size_ret,
+                                       dev->max_mem_global());
+
+   case CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE:
+      return scalar_property<cl_ulong>(buf, size, size_ret,
+                                       dev->max_const_buffer_size());
+
+   case CL_DEVICE_MAX_CONSTANT_ARGS:
+      return scalar_property<cl_uint>(buf, size, size_ret,
+                                      dev->max_const_buffers());
+
+   case CL_DEVICE_LOCAL_MEM_TYPE:
+      return scalar_property<cl_device_local_mem_type>(buf, size, size_ret,
+                                                       CL_LOCAL);
+
+   case CL_DEVICE_LOCAL_MEM_SIZE:
+      return scalar_property<cl_ulong>(buf, size, size_ret,
+                                       dev->max_mem_local());
+
+   case CL_DEVICE_ERROR_CORRECTION_SUPPORT:
+      return scalar_property<cl_bool>(buf, size, size_ret, CL_FALSE);
+
+   case CL_DEVICE_PROFILING_TIMER_RESOLUTION:
+      return scalar_property<size_t>(buf, size, size_ret, 0);
+
+   case CL_DEVICE_ENDIAN_LITTLE:
+      return scalar_property<cl_bool>(buf, size, size_ret, CL_TRUE);
+
+   case CL_DEVICE_AVAILABLE:
+   case CL_DEVICE_COMPILER_AVAILABLE:
+      return scalar_property<cl_bool>(buf, size, size_ret, CL_TRUE);
+
+   case CL_DEVICE_EXECUTION_CAPABILITIES:
+      return scalar_property<cl_device_exec_capabilities>(buf, size, size_ret,
+                                                          CL_EXEC_KERNEL);
+
+   case CL_DEVICE_QUEUE_PROPERTIES:
+      return scalar_property<cl_command_queue_properties>(buf, size, size_ret,
+         CL_QUEUE_PROFILING_ENABLE);
+
+   case CL_DEVICE_NAME:
+      return string_property(buf, size, size_ret, dev->device_name());
+
+   case CL_DEVICE_VENDOR:
+      return string_property(buf, size, size_ret, dev->vendor_name());
+
+   case CL_DRIVER_VERSION:
+      return string_property(buf, size, size_ret, MESA_VERSION);
+
+   case CL_DEVICE_PROFILE:
+      return string_property(buf, size, size_ret, "FULL_PROFILE");
+
+   case CL_DEVICE_VERSION:
+      return string_property(buf, size, size_ret, "OpenCL 1.1 MESA " MESA_VERSION);
+
+   case CL_DEVICE_EXTENSIONS:
+      return string_property(buf, size, size_ret, "");
+
+   case CL_DEVICE_PLATFORM:
+      return scalar_property<cl_platform_id>(buf, size, size_ret, NULL);
+
+   case CL_DEVICE_HOST_UNIFIED_MEMORY:
+      return scalar_property<cl_bool>(buf, size, size_ret, CL_TRUE);
+
+   case CL_DEVICE_NATIVE_VECTOR_WIDTH_CHAR:
+      return scalar_property<cl_uint>(buf, size, size_ret, 16);
+
+   case CL_DEVICE_NATIVE_VECTOR_WIDTH_SHORT:
+      return scalar_property<cl_uint>(buf, size, size_ret, 8);
+
+   case CL_DEVICE_NATIVE_VECTOR_WIDTH_INT:
+      return scalar_property<cl_uint>(buf, size, size_ret, 4);
+
+   case CL_DEVICE_NATIVE_VECTOR_WIDTH_LONG:
+      return scalar_property<cl_uint>(buf, size, size_ret, 2);
+
+   case CL_DEVICE_NATIVE_VECTOR_WIDTH_FLOAT:
+      return scalar_property<cl_uint>(buf, size, size_ret, 4);
+
+   case CL_DEVICE_NATIVE_VECTOR_WIDTH_DOUBLE:
+      return scalar_property<cl_uint>(buf, size, size_ret, 2);
+
+   case CL_DEVICE_NATIVE_VECTOR_WIDTH_HALF:
+      return scalar_property<cl_uint>(buf, size, size_ret, 0);
+
+   case CL_DEVICE_OPENCL_C_VERSION:
+      return string_property(buf, size, size_ret, "OpenCL C 1.1");
+
+   default:
+      return CL_INVALID_VALUE;
+   }
+}
diff --git a/src/gallium/state_trackers/clover/api/event.cpp b/src/gallium/state_trackers/clover/api/event.cpp
new file mode 100644
index 00000000000..d6c37f6aef2
--- /dev/null
+++ b/src/gallium/state_trackers/clover/api/event.cpp
@@ -0,0 +1,239 @@
+//
+// Copyright 2012 Francisco Jerez
+//
+// 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 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.
+//
+
+#include "api/util.hpp"
+#include "core/event.hpp"
+
+using namespace clover;
+
+PUBLIC cl_event
+clCreateUserEvent(cl_context ctx, cl_int *errcode_ret) try {
+   if (!ctx)
+      throw error(CL_INVALID_CONTEXT);
+
+   ret_error(errcode_ret, CL_SUCCESS);
+   return new soft_event(*ctx, {}, false);
+
+} catch(error &e) {
+   ret_error(errcode_ret, e);
+   return NULL;
+}
+
+PUBLIC cl_int
+clSetUserEventStatus(cl_event ev, cl_int status) {
+   if (!dynamic_cast<soft_event *>(ev))
+      return CL_INVALID_EVENT;
+
+   if (status > 0)
+      return CL_INVALID_VALUE;
+
+   if (ev->status() <= 0)
+      return CL_INVALID_OPERATION;
+
+   if (status)
+      ev->abort(status);
+   else
+      ev->trigger();
+
+   return CL_SUCCESS;
+}
+
+PUBLIC cl_int
+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) {
+         if (!ev)
+            throw error(CL_INVALID_EVENT);
+
+         if (&ev->ctx != &evs[0]->ctx)
+            throw error(CL_INVALID_CONTEXT);
+
+         if (ev->status() < 0)
+            throw error(CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST);
+      });
+
+   // 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));
+
+   // ...and wait on it.
+   sev->wait();
+
+   return CL_SUCCESS;
+
+} catch(error &e) {
+   return e.get();
+}
+
+PUBLIC cl_int
+clGetEventInfo(cl_event ev, cl_event_info param,
+               size_t size, void *buf, size_t *size_ret) {
+   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());
+
+   case CL_EVENT_CONTEXT:
+      return scalar_property<cl_context>(buf, size, size_ret, &ev->ctx);
+
+   case CL_EVENT_COMMAND_TYPE:
+      return scalar_property<cl_command_type>(buf, size, size_ret, ev->command());
+
+   case CL_EVENT_COMMAND_EXECUTION_STATUS:
+      return scalar_property<cl_int>(buf, size, size_ret, ev->status());
+
+   case CL_EVENT_REFERENCE_COUNT:
+      return scalar_property<cl_uint>(buf, size, size_ret, ev->ref_count());
+
+   default:
+      return CL_INVALID_VALUE;
+   }
+}
+
+PUBLIC cl_int
+clSetEventCallback(cl_event ev, cl_int type,
+                   void (CL_CALLBACK *pfn_event_notify)(cl_event, cl_int,
+                                                        void *),
+                   void *user_data) try {
+   if (!ev)
+      throw error(CL_INVALID_EVENT);
+
+   if (!pfn_event_notify || type != CL_COMPLETE)
+      throw error(CL_INVALID_VALUE);
+
+   // 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,
+                     [=](event &) {
+                        ev->wait();
+                        pfn_event_notify(ev, ev->status(), user_data);
+                     }));
+
+   return CL_SUCCESS;
+
+} catch(error &e) {
+   return e.get();
+}
+
+PUBLIC cl_int
+clRetainEvent(cl_event ev) {
+   if (!ev)
+      return CL_INVALID_EVENT;
+
+   ev->retain();
+   return CL_SUCCESS;
+}
+
+PUBLIC cl_int
+clReleaseEvent(cl_event ev) {
+   if (!ev)
+      return CL_INVALID_EVENT;
+
+   if (ev->release())
+      delete ev;
+
+   return CL_SUCCESS;
+}
+
+PUBLIC cl_int
+clEnqueueMarker(cl_command_queue q, cl_event *ev) try {
+   if (!q)
+      throw error(CL_INVALID_COMMAND_QUEUE);
+
+   if (!ev)
+      throw error(CL_INVALID_VALUE);
+
+   *ev = new hard_event(*q, CL_COMMAND_MARKER, {});
+
+   return CL_SUCCESS;
+
+} catch(error &e) {
+   return e.get();
+}
+
+PUBLIC cl_int
+clEnqueueBarrier(cl_command_queue q) {
+   if (!q)
+      return CL_INVALID_COMMAND_QUEUE;
+
+   // No need to do anything, q preserves data ordering strictly.
+   return CL_SUCCESS;
+}
+
+PUBLIC cl_int
+clEnqueueWaitForEvents(cl_command_queue q, cl_uint num_evs,
+                       const cl_event *evs) try {
+   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) {
+         if (!ev)
+            throw error(CL_INVALID_EVENT);
+
+         if (&ev->ctx != &q->ctx)
+            throw error(CL_INVALID_CONTEXT);
+      });
+
+   // Create a hard event that depends on the events in the wait list:
+   // subsequent commands in the same queue will be implicitly
+   // serialized with respect to it -- hard events always are.
+   ref_ptr<hard_event> hev = transfer(
+      new hard_event(*q, 0, { evs, evs + num_evs }));
+
+   return CL_SUCCESS;
+
+} catch(error &e) {
+   return e.get();
+}
+
+PUBLIC cl_int
+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 {
+   if (!q)
+      throw error(CL_INVALID_COMMAND_QUEUE);
+
+   // Create a temporary hard event -- it implicitly depends on all
+   // the previously queued hard events.
+   ref_ptr<hard_event> hev = transfer(new hard_event(*q, 0, { }));
+
+   // And wait on it.
+   hev->wait();
+
+   return CL_SUCCESS;
+
+} catch(error &e) {
+   return e.get();
+}
diff --git a/src/gallium/state_trackers/clover/api/kernel.cpp b/src/gallium/state_trackers/clover/api/kernel.cpp
new file mode 100644
index 00000000000..44eeb277127
--- /dev/null
+++ b/src/gallium/state_trackers/clover/api/kernel.cpp
@@ -0,0 +1,318 @@
+//
+// Copyright 2012 Francisco Jerez
+//
+// 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 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.
+//
+
+#include "api/util.hpp"
+#include "core/kernel.hpp"
+#include "core/event.hpp"
+
+using namespace clover;
+
+PUBLIC cl_kernel
+clCreateKernel(cl_program prog, const char *name,
+               cl_int *errcode_ret) try {
+   if (!prog)
+      throw error(CL_INVALID_PROGRAM);
+
+   if (!name)
+      throw error(CL_INVALID_VALUE);
+
+   if (prog->binaries().empty())
+      throw error(CL_INVALID_PROGRAM_EXECUTABLE);
+
+   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() });
+
+} catch (module::noent_error &e) {
+   ret_error(errcode_ret, CL_INVALID_KERNEL_NAME);
+   return NULL;
+
+} catch(error &e) {
+   ret_error(errcode_ret, e);
+   return NULL;
+}
+
+PUBLIC cl_int
+clCreateKernelsInProgram(cl_program prog, cl_uint count,
+                         cl_kernel *kerns, cl_uint *count_ret) {
+   if (!prog)
+      throw error(CL_INVALID_PROGRAM);
+
+   if (prog->binaries().empty())
+      throw error(CL_INVALID_PROGRAM_EXECUTABLE);
+
+   auto &syms = prog->binaries().begin()->second.syms;
+
+   if (kerns && count < syms.size())
+      throw error(CL_INVALID_VALUE);
+
+   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() });
+                     });
+
+   if (count_ret)
+      *count_ret = syms.size();
+
+   return CL_SUCCESS;
+}
+
+PUBLIC cl_int
+clRetainKernel(cl_kernel kern) {
+   if (!kern)
+      return CL_INVALID_KERNEL;
+
+   kern->retain();
+   return CL_SUCCESS;
+}
+
+PUBLIC cl_int
+clReleaseKernel(cl_kernel kern) {
+   if (!kern)
+      return CL_INVALID_KERNEL;
+
+   if (kern->release())
+      delete kern;
+
+   return CL_SUCCESS;
+}
+
+PUBLIC cl_int
+clSetKernelArg(cl_kernel kern, cl_uint idx, size_t size,
+               const void *value) try {
+   if (!kern)
+      throw error(CL_INVALID_KERNEL);
+
+   if (idx >= kern->args.size())
+      throw error(CL_INVALID_ARG_INDEX);
+
+   kern->args[idx]->set(size, value);
+
+   return CL_SUCCESS;
+
+} catch(error &e) {
+   return e.get();
+}
+
+PUBLIC cl_int
+clGetKernelInfo(cl_kernel kern, cl_kernel_info param,
+                size_t size, void *buf, size_t *size_ret) {
+   if (!kern)
+      return CL_INVALID_KERNEL;
+
+   switch (param) {
+   case CL_KERNEL_FUNCTION_NAME:
+      return string_property(buf, size, size_ret, kern->name());
+
+   case CL_KERNEL_NUM_ARGS:
+      return scalar_property<cl_uint>(buf, size, size_ret,
+                                      kern->args.size());
+
+   case CL_KERNEL_REFERENCE_COUNT:
+      return scalar_property<cl_uint>(buf, size, size_ret,
+                                      kern->ref_count());
+
+   case CL_KERNEL_CONTEXT:
+      return scalar_property<cl_context>(buf, size, size_ret,
+                                         &kern->prog.ctx);
+
+   case CL_KERNEL_PROGRAM:
+      return scalar_property<cl_program>(buf, size, size_ret,
+                                         &kern->prog);
+
+   default:
+      return CL_INVALID_VALUE;
+   }
+}
+
+PUBLIC cl_int
+clGetKernelWorkGroupInfo(cl_kernel kern, cl_device_id dev,
+                         cl_kernel_work_group_info param,
+                         size_t size, void *buf, size_t *size_ret) {
+   if (!kern)
+      return CL_INVALID_KERNEL;
+
+   if ((!dev && kern->prog.binaries().size() != 1) ||
+       (dev && !kern->prog.binaries().count(dev)))
+      return CL_INVALID_DEVICE;
+
+   switch (param) {
+   case CL_KERNEL_WORK_GROUP_SIZE:
+      return scalar_property<size_t>(buf, size, size_ret,
+                                     kern->max_block_size());
+
+   case CL_KERNEL_COMPILE_WORK_GROUP_SIZE:
+      return vector_property<size_t>(buf, size, size_ret,
+                                     kern->block_size());
+
+   case CL_KERNEL_LOCAL_MEM_SIZE:
+      return scalar_property<cl_ulong>(buf, size, size_ret,
+                                       kern->mem_local());
+
+   case CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE:
+      return scalar_property<size_t>(buf, size, size_ret, 1);
+
+   case CL_KERNEL_PRIVATE_MEM_SIZE:
+      return scalar_property<cl_ulong>(buf, size, size_ret,
+                                       kern->mem_private());
+
+   default:
+      return CL_INVALID_VALUE;
+   }
+}
+
+namespace {
+   ///
+   /// Common argument checking shared by kernel invocation commands.
+   ///
+   void
+   kernel_validate(cl_command_queue q, cl_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,
+                   cl_event *ev) {
+      if (!q)
+         throw error(CL_INVALID_COMMAND_QUEUE);
+
+      if (!kern)
+         throw error(CL_INVALID_KERNEL);
+
+      if (&kern->prog.ctx != &q->ctx ||
+          any_of([&](const cl_event ev) {
+                return &ev->ctx != &q->ctx;
+             }, deps, deps + num_deps))
+         throw error(CL_INVALID_CONTEXT);
+
+      if (bool(num_deps) != bool(deps) ||
+          any_of(is_zero<cl_event>(), deps, deps + num_deps))
+         throw error(CL_INVALID_EVENT_WAIT_LIST);
+
+      if (any_of([](std::unique_ptr<kernel::argument> &arg) {
+               return !arg->set();
+            }, kern->args.begin(), kern->args.end()))
+         throw error(CL_INVALID_KERNEL_ARGS);
+
+      if (!kern->prog.binaries().count(&q->dev))
+         throw error(CL_INVALID_PROGRAM_EXECUTABLE);
+
+      if (dims < 1 || dims > q->dev.max_block_size().size())
+         throw error(CL_INVALID_WORK_DIMENSION);
+
+      if (!grid_size || any_of(is_zero<size_t>(), grid_size, grid_size + dims))
+         throw error(CL_INVALID_GLOBAL_WORK_SIZE);
+
+      if (block_size && any_of([](size_t b, size_t max) {
+               return b == 0 || b > max;
+            }, block_size, block_size + dims,
+            q->dev.max_block_size().begin()))
+         throw error(CL_INVALID_WORK_ITEM_SIZE);
+
+      if (block_size && any_of([](size_t b, size_t g) {
+               return g % b;
+            }, block_size, block_size + dims, grid_size))
+         throw error(CL_INVALID_WORK_GROUP_SIZE);
+   }
+
+   ///
+   /// Common event action shared by kernel invocation commands.
+   ///
+   std::function<void (event &)>
+   kernel_op(cl_command_queue q, cl_kernel kern,
+             const std::vector<size_t> &grid_offset,
+             const std::vector<size_t> &grid_size,
+             const std::vector<size_t> &block_size) {
+      const std::vector<size_t> reduced_grid_size = map(
+         std::divides<size_t>(), grid_size.begin(), grid_size.end(),
+         block_size.begin());
+
+      return [=](event &) {
+         kern->launch(*q, grid_offset, reduced_grid_size, block_size);
+      };
+   }
+
+   template<typename T, typename S>
+   std::vector<T>
+   opt_vector(const T *p, S n) {
+      if (p)
+         return { p, p + n };
+      else
+         return { n };
+   }
+}
+
+PUBLIC cl_int
+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 {
+   const std::vector<size_t> grid_offset = opt_vector(pgrid_offset, dims);
+   const std::vector<size_t> grid_size = opt_vector(pgrid_size, dims);
+   const std::vector<size_t> block_size = opt_vector(pblock_size, dims);
+
+   kernel_validate(q, kern, dims, pgrid_offset, pgrid_size, pblock_size,
+                   num_deps, deps, ev);
+
+   hard_event *hev = new hard_event(
+      *q, CL_COMMAND_NDRANGE_KERNEL, { deps, deps + num_deps },
+      kernel_op(q, kern, grid_offset, grid_size, block_size));
+
+   ret_object(ev, hev);
+   return CL_SUCCESS;
+
+} catch(error &e) {
+   return e.get();
+}
+
+PUBLIC cl_int
+clEnqueueTask(cl_command_queue q, cl_kernel kern,
+              cl_uint num_deps, const cl_event *deps,
+              cl_event *ev) try {
+   const std::vector<size_t> grid_offset = { 0 };
+   const std::vector<size_t> grid_size = { 1 };
+   const std::vector<size_t> block_size = { 1 };
+
+   kernel_validate(q, kern, 1, grid_offset.data(), grid_size.data(),
+                   block_size.data(), num_deps, deps, ev);
+
+   hard_event *hev = new hard_event(
+      *q, CL_COMMAND_TASK, { deps, deps + num_deps },
+      kernel_op(q, kern, grid_offset, grid_size, block_size));
+
+   ret_object(ev, hev);
+   return CL_SUCCESS;
+
+} catch(error &e) {
+   return e.get();
+}
+
+PUBLIC cl_int
+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,
+                      const cl_event *deps, cl_event *ev) {
+   return CL_INVALID_OPERATION;
+}
diff --git a/src/gallium/state_trackers/clover/api/memory.cpp b/src/gallium/state_trackers/clover/api/memory.cpp
new file mode 100644
index 00000000000..1b1ae73796f
--- /dev/null
+++ b/src/gallium/state_trackers/clover/api/memory.cpp
@@ -0,0 +1,305 @@
+//
+// Copyright 2012 Francisco Jerez
+//
+// 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 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.
+//
+
+#include "api/util.hpp"
+#include "core/memory.hpp"
+#include "core/format.hpp"
+
+using namespace clover;
+
+PUBLIC cl_mem
+clCreateBuffer(cl_context ctx, cl_mem_flags flags, size_t size,
+               void *host_ptr, cl_int *errcode_ret) try {
+   if (!ctx)
+      throw error(CL_INVALID_CONTEXT);
+
+   if (bool(host_ptr) != bool(flags & (CL_MEM_USE_HOST_PTR |
+                                       CL_MEM_COPY_HOST_PTR)))
+      throw error(CL_INVALID_HOST_PTR);
+
+   if (!size)
+      throw error(CL_INVALID_BUFFER_SIZE);
+
+   if (flags & ~(CL_MEM_WRITE_ONLY | CL_MEM_READ_ONLY |
+                 CL_MEM_USE_HOST_PTR | CL_MEM_ALLOC_HOST_PTR |
+                 CL_MEM_COPY_HOST_PTR))
+      throw error(CL_INVALID_VALUE);
+
+   ret_error(errcode_ret, CL_SUCCESS);
+   return new root_buffer(*ctx, flags, size, host_ptr);
+
+} catch (error &e) {
+   ret_error(errcode_ret, e);
+   return NULL;
+}
+
+PUBLIC cl_mem
+clCreateSubBuffer(cl_mem obj, cl_mem_flags flags, cl_buffer_create_type op,
+                  const void *op_info, cl_int *errcode_ret) try {
+   root_buffer *parent = dynamic_cast<root_buffer *>(obj);
+
+   if (!parent)
+      throw error(CL_INVALID_MEM_OBJECT);
+
+   if ((flags & (CL_MEM_USE_HOST_PTR |
+                 CL_MEM_ALLOC_HOST_PTR |
+                 CL_MEM_COPY_HOST_PTR)) ||
+       (~flags & parent->flags() & (CL_MEM_READ_ONLY |
+                                    CL_MEM_WRITE_ONLY)))
+      throw error(CL_INVALID_VALUE);
+
+   if (op == CL_BUFFER_CREATE_TYPE_REGION) {
+      const cl_buffer_region *reg = (const cl_buffer_region *)op_info;
+
+      if (!reg ||
+          reg->origin > parent->size() ||
+          reg->origin + reg->size > parent->size())
+         throw error(CL_INVALID_VALUE);
+
+      if (!reg->size)
+         throw error(CL_INVALID_BUFFER_SIZE);
+
+      ret_error(errcode_ret, CL_SUCCESS);
+      return new sub_buffer(*parent, flags, reg->origin, reg->size);
+
+   } else {
+      throw error(CL_INVALID_VALUE);
+   }
+
+} catch (error &e) {
+   ret_error(errcode_ret, e);
+   return NULL;
+}
+
+PUBLIC cl_mem
+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 {
+   if (!ctx)
+      throw error(CL_INVALID_CONTEXT);
+
+   if (flags & ~(CL_MEM_WRITE_ONLY | CL_MEM_READ_ONLY |
+                 CL_MEM_USE_HOST_PTR | CL_MEM_ALLOC_HOST_PTR |
+                 CL_MEM_COPY_HOST_PTR))
+      throw error(CL_INVALID_VALUE);
+
+   if (!format)
+      throw error(CL_INVALID_IMAGE_FORMAT_DESCRIPTOR);
+
+   if (width < 1 || height < 1)
+      throw error(CL_INVALID_IMAGE_SIZE);
+
+   if (bool(host_ptr) != bool(flags & (CL_MEM_USE_HOST_PTR |
+                                       CL_MEM_COPY_HOST_PTR)))
+      throw error(CL_INVALID_HOST_PTR);
+
+   if (!supported_formats(ctx, CL_MEM_OBJECT_IMAGE2D).count(*format))
+      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);
+
+} catch (error &e) {
+   ret_error(errcode_ret, e);
+   return NULL;
+}
+
+PUBLIC cl_mem
+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 {
+   if (!ctx)
+      throw error(CL_INVALID_CONTEXT);
+
+   if (flags & ~(CL_MEM_WRITE_ONLY | CL_MEM_READ_ONLY |
+                 CL_MEM_USE_HOST_PTR | CL_MEM_ALLOC_HOST_PTR |
+                 CL_MEM_COPY_HOST_PTR))
+      throw error(CL_INVALID_VALUE);
+
+   if (!format)
+      throw error(CL_INVALID_IMAGE_FORMAT_DESCRIPTOR);
+
+   if (width < 1 || height < 1 || depth < 2)
+      throw error(CL_INVALID_IMAGE_SIZE);
+
+   if (bool(host_ptr) != bool(flags & (CL_MEM_USE_HOST_PTR |
+                                       CL_MEM_COPY_HOST_PTR)))
+      throw error(CL_INVALID_HOST_PTR);
+
+   if (!supported_formats(ctx, CL_MEM_OBJECT_IMAGE3D).count(*format))
+      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);
+
+} catch (error &e) {
+   ret_error(errcode_ret, e);
+   return NULL;
+}
+
+PUBLIC cl_int
+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 {
+   if (!ctx)
+      throw error(CL_INVALID_CONTEXT);
+
+   if (flags & ~(CL_MEM_WRITE_ONLY | CL_MEM_READ_ONLY |
+                 CL_MEM_USE_HOST_PTR | CL_MEM_ALLOC_HOST_PTR |
+                 CL_MEM_COPY_HOST_PTR))
+      throw error(CL_INVALID_VALUE);
+
+   if (!count && buf)
+      throw error(CL_INVALID_VALUE);
+
+   auto formats = supported_formats(ctx, type);
+
+   if (buf)
+      std::copy_n(formats.begin(), std::min((cl_uint)formats.size(), count),
+                  buf);
+   if (count_ret)
+      *count_ret = formats.size();
+
+   return CL_SUCCESS;
+
+} catch (error &e) {
+   return e.get();
+}
+
+PUBLIC cl_int
+clGetMemObjectInfo(cl_mem obj, cl_mem_info param,
+                   size_t size, void *buf, size_t *size_ret) {
+   if (!obj)
+      return CL_INVALID_MEM_OBJECT;
+
+   switch (param) {
+   case CL_MEM_TYPE:
+      return scalar_property<cl_mem_object_type>(buf, size, size_ret,
+                                                 obj->type());
+
+   case CL_MEM_FLAGS:
+      return scalar_property<cl_mem_flags>(buf, size, size_ret, obj->flags());
+
+   case CL_MEM_SIZE:
+      return scalar_property<size_t>(buf, size, size_ret, obj->size());
+
+   case CL_MEM_HOST_PTR:
+      return scalar_property<void *>(buf, size, size_ret, obj->host_ptr());
+
+   case CL_MEM_MAP_COUNT:
+      return scalar_property<cl_uint>(buf, size, size_ret, 0);
+
+   case CL_MEM_REFERENCE_COUNT:
+      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);
+
+   case CL_MEM_ASSOCIATED_MEMOBJECT: {
+      sub_buffer *sub = dynamic_cast<sub_buffer *>(obj);
+      return scalar_property<cl_mem>(buf, size, size_ret,
+                                     (sub ? &sub->parent : NULL));
+   }
+   case CL_MEM_OFFSET: {
+      sub_buffer *sub = dynamic_cast<sub_buffer *>(obj);
+      return scalar_property<size_t>(buf, size, size_ret,
+                                     (sub ? sub->offset() : 0));
+   }
+   default:
+      return CL_INVALID_VALUE;
+   }
+}
+
+PUBLIC cl_int
+clGetImageInfo(cl_mem obj, cl_image_info param,
+               size_t size, void *buf, size_t *size_ret) {
+   image *img = dynamic_cast<image *>(obj);
+   if (!img)
+      return CL_INVALID_MEM_OBJECT;
+
+   switch (param) {
+   case CL_IMAGE_FORMAT:
+      return scalar_property<cl_image_format>(buf, size, size_ret,
+                                              img->format());
+
+   case CL_IMAGE_ELEMENT_SIZE:
+      return scalar_property<size_t>(buf, size, size_ret, 0);
+
+   case CL_IMAGE_ROW_PITCH:
+      return scalar_property<size_t>(buf, size, size_ret, img->row_pitch());
+
+   case CL_IMAGE_SLICE_PITCH:
+      return scalar_property<size_t>(buf, size, size_ret, img->slice_pitch());
+
+   case CL_IMAGE_WIDTH:
+      return scalar_property<size_t>(buf, size, size_ret, img->width());
+
+   case CL_IMAGE_HEIGHT:
+      return scalar_property<size_t>(buf, size, size_ret, img->height());
+
+   case CL_IMAGE_DEPTH:
+      return scalar_property<size_t>(buf, size, size_ret, img->depth());
+
+   default:
+      return CL_INVALID_VALUE;
+   }
+}
+
+PUBLIC cl_int
+clRetainMemObject(cl_mem obj) {
+   if (!obj)
+      return CL_INVALID_MEM_OBJECT;
+
+   obj->retain();
+   return CL_SUCCESS;
+}
+
+PUBLIC cl_int
+clReleaseMemObject(cl_mem obj) {
+   if (!obj)
+      return CL_INVALID_MEM_OBJECT;
+
+   if (obj->release())
+      delete obj;
+
+   return CL_SUCCESS;
+}
+
+PUBLIC cl_int
+clSetMemObjectDestructorCallback(cl_mem obj,
+                                 void (CL_CALLBACK *pfn_notify)(cl_mem, void *),
+                                 void *user_data) {
+   if (!obj)
+      return CL_INVALID_MEM_OBJECT;
+
+   if (!pfn_notify)
+      return CL_INVALID_VALUE;
+
+   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
new file mode 100644
index 00000000000..e5e80b85256
--- /dev/null
+++ b/src/gallium/state_trackers/clover/api/platform.cpp
@@ -0,0 +1,68 @@
+//
+// Copyright 2012 Francisco Jerez
+//
+// 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 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.
+//
+
+#include "api/util.hpp"
+
+using namespace clover;
+
+PUBLIC cl_int
+clGetPlatformIDs(cl_uint num_entries, cl_platform_id *platforms,
+                 cl_uint *num_platforms) {
+   if ((!num_entries && platforms) ||
+       (!num_platforms && !platforms))
+      return CL_INVALID_VALUE;
+
+   if (num_platforms)
+      *num_platforms = 1;
+   if (platforms)
+      *platforms = NULL;
+
+   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 != NULL)
+      return CL_INVALID_PLATFORM;
+
+   switch (param_name) {
+   case CL_PLATFORM_PROFILE:
+      return string_property(buf, size, size_ret, "FULL_PROFILE");
+
+   case CL_PLATFORM_VERSION:
+      return string_property(buf, size, size_ret,
+                             "OpenCL 1.1 MESA " MESA_VERSION);
+
+   case CL_PLATFORM_NAME:
+      return string_property(buf, size, size_ret, "Default");
+
+   case CL_PLATFORM_VENDOR:
+      return string_property(buf, size, size_ret, "Mesa");
+
+   case CL_PLATFORM_EXTENSIONS:
+      return string_property(buf, size, size_ret, "");
+
+   default:
+      return CL_INVALID_VALUE;
+   }
+}
diff --git a/src/gallium/state_trackers/clover/api/program.cpp b/src/gallium/state_trackers/clover/api/program.cpp
new file mode 100644
index 00000000000..e874c51ad7d
--- /dev/null
+++ b/src/gallium/state_trackers/clover/api/program.cpp
@@ -0,0 +1,241 @@
+//
+// Copyright 2012 Francisco Jerez
+//
+// 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 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.
+//
+
+#include "api/util.hpp"
+#include "core/program.hpp"
+
+using namespace clover;
+
+PUBLIC cl_program
+clCreateProgramWithSource(cl_context ctx, cl_uint count,
+                          const char **strings, const size_t *lengths,
+                          cl_int *errcode_ret) try {
+   std::string source;
+
+   if (!ctx)
+      throw error(CL_INVALID_CONTEXT);
+
+   if (!count || !strings ||
+       any_of(is_zero<const char *>(), strings, strings + count))
+      throw error(CL_INVALID_VALUE);
+
+   // Concatenate all the provided fragments together
+   for (unsigned i = 0; i < count; ++i)
+         source += (lengths && lengths[i] ?
+                    std::string(strings[i], strings[i] + lengths[i]) :
+                    std::string(strings[i]));
+
+   // ...and create a program object for them.
+   ret_error(errcode_ret, CL_SUCCESS);
+   return new program(*ctx, source);
+
+} catch (error &e) {
+   ret_error(errcode_ret, e);
+   return NULL;
+}
+
+PUBLIC cl_program
+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 {
+   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) {
+            return !ctx->has_device(dev);
+         }, devs, devs + count))
+      throw error(CL_INVALID_DEVICE);
+
+   // Deserialize the provided binaries,
+   auto modules = map(
+      [](const unsigned char *p, size_t l) -> std::pair<cl_int, module> {
+         if (!p || !l)
+            return { CL_INVALID_VALUE, {} };
+
+         try {
+            compat::istream::buffer_t bin(p, l);
+            compat::istream s(bin);
+
+            return { CL_SUCCESS, module::deserialize(s) };
+
+         } catch (compat::istream::error &e) {
+            return { CL_INVALID_BINARY, {} };
+         }
+      },
+      binaries, binaries + count, lengths);
+
+   // update the status array,
+   if (status_ret)
+      std::transform(modules.begin(), modules.end(), status_ret,
+                     keys<cl_int, module>);
+
+   if (any_of(key_equals<cl_int, module>(CL_INVALID_VALUE),
+              modules.begin(), modules.end()))
+      throw error(CL_INVALID_VALUE);
+
+   if (any_of(key_equals<cl_int, module>(CL_INVALID_BINARY),
+              modules.begin(), modules.end()))
+      throw error(CL_INVALID_BINARY);
+
+   // 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()));
+
+} catch (error &e) {
+   ret_error(errcode_ret, e);
+   return NULL;
+}
+
+PUBLIC cl_int
+clRetainProgram(cl_program prog) {
+   if (!prog)
+      return CL_INVALID_PROGRAM;
+
+   prog->retain();
+   return CL_SUCCESS;
+}
+
+PUBLIC cl_int
+clReleaseProgram(cl_program prog) {
+   if (!prog)
+      return CL_INVALID_PROGRAM;
+
+   if (prog->release())
+      delete prog;
+
+   return CL_SUCCESS;
+}
+
+PUBLIC cl_int
+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 {
+   if (!prog)
+      throw error(CL_INVALID_PROGRAM);
+
+   if (bool(count) != bool(devs) ||
+       (!pfn_notify && user_data))
+      throw error(CL_INVALID_VALUE);
+
+   if (any_of([&](const cl_device_id dev) {
+            return !prog->ctx.has_device(dev);
+         }, devs, devs + count))
+      throw error(CL_INVALID_DEVICE);
+
+   prog->build({ devs, devs + count });
+   return CL_SUCCESS;
+
+} catch (error &e) {
+   return e.get();
+}
+
+PUBLIC cl_int
+clUnloadCompiler() {
+   return CL_SUCCESS;
+}
+
+PUBLIC cl_int
+clGetProgramInfo(cl_program prog, cl_program_info param,
+                 size_t size, void *buf, size_t *size_ret) {
+   if (!prog)
+      return CL_INVALID_PROGRAM;
+
+   switch (param) {
+   case CL_PROGRAM_REFERENCE_COUNT:
+      return scalar_property<cl_uint>(buf, size, size_ret,
+                                      prog->ref_count());
+
+   case CL_PROGRAM_CONTEXT:
+      return scalar_property<cl_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>(
+         buf, size, size_ret,
+         map(keys<device *, module>,
+             prog->binaries().begin(), prog->binaries().end()));
+
+   case CL_PROGRAM_SOURCE:
+      return string_property(buf, size, size_ret, prog->source());
+
+   case CL_PROGRAM_BINARY_SIZES:
+      return vector_property<size_t>(
+         buf, size, size_ret,
+         map([](const std::pair<device *, module> &ent) {
+               compat::ostream::buffer_t bin;
+               compat::ostream s(bin);
+               ent.second.serialize(s);
+               return bin.size();
+            },
+            prog->binaries().begin(), prog->binaries().end()));
+
+   case CL_PROGRAM_BINARIES:
+      return matrix_property<unsigned char>(
+         buf, size, size_ret,
+         map([](const std::pair<device *, module> &ent) {
+               compat::ostream::buffer_t bin;
+               compat::ostream s(bin);
+               ent.second.serialize(s);
+               return bin;
+            },
+            prog->binaries().begin(), prog->binaries().end()));
+
+   default:
+      return CL_INVALID_VALUE;
+   }
+}
+
+PUBLIC cl_int
+clGetProgramBuildInfo(cl_program prog, cl_device_id dev,
+                      cl_program_build_info param,
+                      size_t size, void *buf, size_t *size_ret) {
+   if (!prog)
+      return CL_INVALID_PROGRAM;
+
+   if (!prog->ctx.has_device(dev))
+      return CL_INVALID_DEVICE;
+
+   switch (param) {
+   case CL_PROGRAM_BUILD_STATUS:
+      return scalar_property<cl_build_status>(buf, size, size_ret,
+                                              prog->build_status(dev));
+
+   case CL_PROGRAM_BUILD_OPTIONS:
+      return string_property(buf, size, size_ret, prog->build_opts(dev));
+
+   case CL_PROGRAM_BUILD_LOG:
+      return string_property(buf, size, size_ret, prog->build_log(dev));
+
+   default:
+      return CL_INVALID_VALUE;
+   }
+}
diff --git a/src/gallium/state_trackers/clover/api/queue.cpp b/src/gallium/state_trackers/clover/api/queue.cpp
new file mode 100644
index 00000000000..a7905bc4396
--- /dev/null
+++ b/src/gallium/state_trackers/clover/api/queue.cpp
@@ -0,0 +1,102 @@
+//
+// Copyright 2012 Francisco Jerez
+//
+// 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 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.
+//
+
+#include "api/util.hpp"
+#include "core/queue.hpp"
+
+using namespace clover;
+
+PUBLIC cl_command_queue
+clCreateCommandQueue(cl_context ctx, cl_device_id dev,
+                     cl_command_queue_properties props,
+                     cl_int *errcode_ret) try {
+   if (!ctx)
+      throw error(CL_INVALID_CONTEXT);
+
+   if (!ctx->has_device(dev))
+      throw error(CL_INVALID_DEVICE);
+
+   if (props & ~(CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE |
+                 CL_QUEUE_PROFILING_ENABLE))
+      throw error(CL_INVALID_VALUE);
+
+   ret_error(errcode_ret, CL_SUCCESS);
+   return new command_queue(*ctx, *dev, props);
+
+} catch (error &e) {
+   ret_error(errcode_ret, e);
+   return NULL;
+}
+
+PUBLIC cl_int
+clRetainCommandQueue(cl_command_queue q) {
+   if (!q)
+      return CL_INVALID_COMMAND_QUEUE;
+
+   q->retain();
+   return CL_SUCCESS;
+}
+
+PUBLIC cl_int
+clReleaseCommandQueue(cl_command_queue q) {
+   if (!q)
+      return CL_INVALID_COMMAND_QUEUE;
+
+   if (q->release())
+      delete q;
+
+   return CL_SUCCESS;
+}
+
+PUBLIC cl_int
+clGetCommandQueueInfo(cl_command_queue q, cl_command_queue_info param,
+                      size_t size, void *buf, size_t *size_ret) {
+   if (!q)
+      return CL_INVALID_COMMAND_QUEUE;
+
+   switch (param) {
+   case CL_QUEUE_CONTEXT:
+      return scalar_property<cl_context>(buf, size, size_ret, &q->ctx);
+
+   case CL_QUEUE_DEVICE:
+      return scalar_property<cl_device_id>(buf, size, size_ret, &q->dev);
+
+   case CL_QUEUE_REFERENCE_COUNT:
+      return scalar_property<cl_uint>(buf, size, size_ret, q->ref_count());
+
+   case CL_QUEUE_PROPERTIES:
+      return scalar_property<cl_command_queue_properties>(buf, size, size_ret,
+                                                          q->props());
+
+   default:
+      return CL_INVALID_VALUE;
+   }
+}
+
+PUBLIC cl_int
+clFlush(cl_command_queue q) {
+   if (!q)
+      return CL_INVALID_COMMAND_QUEUE;
+
+   q->flush();
+   return CL_SUCCESS;
+}
diff --git a/src/gallium/state_trackers/clover/api/sampler.cpp b/src/gallium/state_trackers/clover/api/sampler.cpp
new file mode 100644
index 00000000000..32ce22ef90f
--- /dev/null
+++ b/src/gallium/state_trackers/clover/api/sampler.cpp
@@ -0,0 +1,90 @@
+//
+// Copyright 2012 Francisco Jerez
+//
+// 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 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.
+//
+
+#include "api/util.hpp"
+#include "core/sampler.hpp"
+
+using namespace clover;
+
+PUBLIC cl_sampler
+clCreateSampler(cl_context ctx, cl_bool norm_mode,
+                cl_addressing_mode addr_mode, cl_filter_mode filter_mode,
+                cl_int *errcode_ret) try {
+   if (!ctx)
+      throw error(CL_INVALID_CONTEXT);
+
+   ret_error(errcode_ret, CL_SUCCESS);
+   return new sampler(*ctx, norm_mode, addr_mode, filter_mode);
+
+} catch (error &e) {
+   ret_error(errcode_ret, e);
+   return NULL;
+}
+
+PUBLIC cl_int
+clRetainSampler(cl_sampler s) {
+   if (!s)
+      throw error(CL_INVALID_SAMPLER);
+
+   s->retain();
+   return CL_SUCCESS;
+}
+
+PUBLIC cl_int
+clReleaseSampler(cl_sampler s) {
+   if (!s)
+      throw error(CL_INVALID_SAMPLER);
+
+   if (s->release())
+      delete s;
+
+   return CL_SUCCESS;
+}
+
+PUBLIC cl_int
+clGetSamplerInfo(cl_sampler s, cl_sampler_info param,
+                 size_t size, void *buf, size_t *size_ret) {
+   if (!s)
+      throw error(CL_INVALID_SAMPLER);
+
+   switch (param) {
+   case CL_SAMPLER_REFERENCE_COUNT:
+      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);
+
+   case CL_SAMPLER_NORMALIZED_COORDS:
+      return scalar_property<cl_bool>(buf, size, size_ret, s->norm_mode());
+
+   case CL_SAMPLER_ADDRESSING_MODE:
+      return scalar_property<cl_addressing_mode>(buf, size, size_ret,
+                                                 s->addr_mode());
+
+   case CL_SAMPLER_FILTER_MODE:
+      return scalar_property<cl_filter_mode>(buf, size, size_ret,
+                                             s->filter_mode());
+
+   default:
+      return CL_INVALID_VALUE;
+   }
+}
diff --git a/src/gallium/state_trackers/clover/api/transfer.cpp b/src/gallium/state_trackers/clover/api/transfer.cpp
new file mode 100644
index 00000000000..c67b75e8034
--- /dev/null
+++ b/src/gallium/state_trackers/clover/api/transfer.cpp
@@ -0,0 +1,506 @@
+//
+// Copyright 2012 Francisco Jerez
+//
+// 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 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.
+//
+
+#include <cstring>
+
+#include "api/util.hpp"
+#include "core/event.hpp"
+#include "core/resource.hpp"
+
+using namespace clover;
+
+namespace {
+   typedef resource::point point;
+
+   ///
+   /// Common argument checking shared by memory transfer commands.
+   ///
+   void
+   validate_base(cl_command_queue q, cl_uint num_deps, const cl_event *deps) {
+      if (!q)
+         throw error(CL_INVALID_COMMAND_QUEUE);
+
+      if (bool(num_deps) != bool(deps) ||
+          any_of(is_zero<cl_event>(), deps, deps + num_deps))
+         throw error(CL_INVALID_EVENT_WAIT_LIST);
+
+      if (any_of([&](const cl_event ev) {
+               return &ev->ctx != &q->ctx;
+            }, deps, deps + num_deps))
+         throw error(CL_INVALID_CONTEXT);
+   }
+
+   ///
+   /// Memory object-specific argument checking shared by most memory
+   /// transfer commands.
+   ///
+   void
+   validate_obj(cl_command_queue q, cl_mem obj) {
+      if (!obj)
+         throw error(CL_INVALID_MEM_OBJECT);
+
+      if (&obj->ctx != &q->ctx)
+         throw error(CL_INVALID_CONTEXT);
+   }
+
+   ///
+   /// Class that encapsulates the task of mapping an object of type
+   /// \a T.  The return value of get() should be implicitly
+   /// convertible to \a void *.
+   ///
+   template<typename T> struct __map;
+
+   template<> struct __map<void *> {
+      static void *
+      get(cl_command_queue q, void *obj, cl_map_flags flags,
+          size_t offset, size_t size) {
+         return (char *)obj + offset;
+      }
+   };
+
+   template<> struct __map<const void *> {
+      static const void *
+      get(cl_command_queue q, const void *obj, cl_map_flags flags,
+          size_t offset, size_t size) {
+         return (const char *)obj + offset;
+      }
+   };
+
+   template<> struct __map<memory_obj *> {
+      static mapping
+      get(cl_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 }};
+      }
+   };
+
+   ///
+   /// Software copy from \a src_obj to \a dst_obj.  They can be
+   /// either pointers or memory objects.
+   ///
+   template<typename T, typename S>
+   std::function<void (event &)>
+   soft_copy_op(cl_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 &region) {
+      return [=](event &) {
+         auto dst = __map<T>::get(q, dst_obj, CL_MAP_WRITE,
+                                  dst_pitch(dst_orig), dst_pitch(region));
+         auto src = __map<S>::get(q, src_obj, CL_MAP_READ,
+                                  src_pitch(src_orig), src_pitch(region));
+         point p;
+
+         for (p[2] = 0; p[2] < region[2]; ++p[2]) {
+            for (p[1] = 0; p[1] < region[1]; ++p[1]) {
+               std::memcpy(static_cast<char *>(dst) + dst_pitch(p),
+                           static_cast<const char *>(src) + src_pitch(p),
+                           src_pitch[0] * region[0]);
+            }
+         }
+      };
+   }
+
+   ///
+   /// Hardware copy from \a src_obj to \a dst_obj.
+   ///
+   template<typename T, typename S>
+   std::function<void (event &)>
+   hard_copy_op(cl_command_queue q, T dst_obj, const point &dst_orig,
+                S src_obj, const point &src_orig, const point &region) {
+      return [=](event &) {
+         dst_obj->resource(q).copy(*q, dst_orig, region,
+                                   src_obj->resource(q), src_orig);
+      };
+   }
+}
+
+PUBLIC cl_int
+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 {
+   validate_base(q, num_deps, deps);
+   validate_obj(q, obj);
+
+   if (!ptr || offset > obj->size() || offset + size > obj->size())
+      throw error(CL_INVALID_VALUE);
+
+   hard_event *hev = new hard_event(
+      *q, CL_COMMAND_READ_BUFFER, { deps, deps + num_deps },
+      soft_copy_op(q,
+                   ptr, { 0 }, { 1 },
+                   obj, { offset }, { 1 },
+                   { size, 1, 1 }));
+
+   ret_object(ev, hev);
+   return CL_SUCCESS;
+
+} catch (error &e) {
+   return e.get();
+}
+
+PUBLIC cl_int
+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 {
+   validate_base(q, num_deps, deps);
+   validate_obj(q, obj);
+
+   if (!ptr || offset > obj->size() || offset + size > obj->size())
+      throw error(CL_INVALID_VALUE);
+
+   hard_event *hev = new hard_event(
+      *q, CL_COMMAND_WRITE_BUFFER, { deps, deps + num_deps },
+      soft_copy_op(q,
+                   obj, { offset }, { 1 },
+                   ptr, { 0 }, { 1 },
+                   { size, 1, 1 }));
+
+   ret_object(ev, hev);
+   return CL_SUCCESS;
+
+} catch (error &e) {
+   return e.get();
+}
+
+PUBLIC cl_int
+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,
+                        size_t host_row_pitch, size_t host_slice_pitch,
+                        void *ptr,
+                        cl_uint num_deps, const cl_event *deps,
+                        cl_event *ev) try {
+   validate_base(q, num_deps, deps);
+   validate_obj(q, obj);
+
+   if (!ptr)
+      throw error(CL_INVALID_VALUE);
+
+   hard_event *hev = new hard_event(
+      *q, CL_COMMAND_READ_BUFFER_RECT, { deps, deps + num_deps },
+      soft_copy_op(q,
+                   ptr, host_origin,
+                   { 1, host_row_pitch, host_slice_pitch },
+                   obj, obj_origin,
+                   { 1, obj_row_pitch, obj_slice_pitch },
+                   region));
+
+   ret_object(ev, hev);
+   return CL_SUCCESS;
+
+} catch (error &e) {
+   return e.get();
+}
+
+PUBLIC cl_int
+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,
+                         size_t host_row_pitch, size_t host_slice_pitch,
+                         const void *ptr,
+                         cl_uint num_deps, const cl_event *deps,
+                         cl_event *ev) try {
+   validate_base(q, num_deps, deps);
+   validate_obj(q, obj);
+
+   if (!ptr)
+      throw error(CL_INVALID_VALUE);
+
+   hard_event *hev = new hard_event(
+      *q, CL_COMMAND_WRITE_BUFFER_RECT, { deps, deps + num_deps },
+      soft_copy_op(q,
+                   obj, obj_origin,
+                   { 1, obj_row_pitch, obj_slice_pitch },
+                   ptr, host_origin,
+                   { 1, host_row_pitch, host_slice_pitch },
+                   region));
+
+   ret_object(ev, hev);
+   return CL_SUCCESS;
+
+} catch (error &e) {
+   return e.get();
+}
+
+PUBLIC cl_int
+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 {
+   validate_base(q, num_deps, deps);
+   validate_obj(q, src_obj);
+   validate_obj(q, dst_obj);
+
+   hard_event *hev = new hard_event(
+      *q, CL_COMMAND_COPY_BUFFER, { deps, deps + num_deps },
+      hard_copy_op(q, dst_obj, { dst_offset },
+                   src_obj, { src_offset },
+                   { size, 1, 1 }));
+
+   ret_object(ev, hev);
+   return CL_SUCCESS;
+
+} catch (error &e) {
+   return e.get();
+}
+
+PUBLIC cl_int
+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 {
+   validate_base(q, num_deps, deps);
+   validate_obj(q, src_obj);
+   validate_obj(q, dst_obj);
+
+   hard_event *hev = new hard_event(
+      *q, CL_COMMAND_COPY_BUFFER_RECT, { deps, deps + num_deps },
+      soft_copy_op(q,
+                   dst_obj, dst_origin,
+                   { 1, dst_row_pitch, dst_slice_pitch },
+                   src_obj, src_origin,
+                   { 1, src_row_pitch, src_slice_pitch },
+                   region));
+
+   ret_object(ev, hev);
+   return CL_SUCCESS;
+
+} catch (error &e) {
+   return e.get();
+}
+
+PUBLIC cl_int
+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 {
+   image *img = dynamic_cast<image *>(obj);
+
+   validate_base(q, num_deps, deps);
+   validate_obj(q, img);
+
+   if (!ptr)
+      throw error(CL_INVALID_VALUE);
+
+   hard_event *hev = new hard_event(
+      *q, CL_COMMAND_READ_IMAGE, { deps, deps + num_deps },
+      soft_copy_op(q,
+                   ptr, {},
+                   { 1, row_pitch, slice_pitch },
+                   obj, origin,
+                   { 1, img->row_pitch(), img->slice_pitch() },
+                   region));
+
+   ret_object(ev, hev);
+   return CL_SUCCESS;
+
+} catch (error &e) {
+   return e.get();
+}
+
+PUBLIC cl_int
+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 {
+   image *img = dynamic_cast<image *>(obj);
+
+   validate_base(q, num_deps, deps);
+   validate_obj(q, img);
+
+   if (!ptr)
+      throw error(CL_INVALID_VALUE);
+
+   hard_event *hev = new hard_event(
+      *q, CL_COMMAND_WRITE_IMAGE, { deps, deps + num_deps },
+      soft_copy_op(q,
+                   obj, origin,
+                   { 1, img->row_pitch(), img->slice_pitch() },
+                   ptr, {},
+                   { 1, row_pitch, slice_pitch },
+                   region));
+
+   ret_object(ev, hev);
+   return CL_SUCCESS;
+
+} catch (error &e) {
+   return e.get();
+}
+
+PUBLIC cl_int
+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 {
+   image *src_img = dynamic_cast<image *>(src_obj);
+   image *dst_img = dynamic_cast<image *>(dst_obj);
+
+   validate_base(q, num_deps, deps);
+   validate_obj(q, src_img);
+   validate_obj(q, dst_img);
+
+   hard_event *hev = new hard_event(
+      *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);
+   return CL_SUCCESS;
+
+} catch (error &e) {
+   return e.get();
+}
+
+PUBLIC cl_int
+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 {
+   image *src_img = dynamic_cast<image *>(src_obj);
+
+   validate_base(q, num_deps, deps);
+   validate_obj(q, src_img);
+   validate_obj(q, dst_obj);
+
+   hard_event *hev = new hard_event(
+      *q, CL_COMMAND_COPY_IMAGE_TO_BUFFER, { deps, deps + num_deps },
+      soft_copy_op(q,
+                   dst_obj, { dst_offset },
+                   { 0, 0, 0 },
+                   src_obj, src_origin,
+                   { 1, src_img->row_pitch(), src_img->slice_pitch() },
+                   region));
+
+   ret_object(ev, hev);
+   return CL_SUCCESS;
+
+} catch (error &e) {
+   return e.get();
+}
+
+PUBLIC cl_int
+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 {
+   image *dst_img = dynamic_cast<image *>(src_obj);
+
+   validate_base(q, num_deps, deps);
+   validate_obj(q, src_obj);
+   validate_obj(q, dst_img);
+
+   hard_event *hev = new hard_event(
+      *q, CL_COMMAND_COPY_BUFFER_TO_IMAGE, { deps, deps + num_deps },
+      soft_copy_op(q,
+                   dst_obj, dst_origin,
+                   { 1, dst_img->row_pitch(), dst_img->slice_pitch() },
+                   src_obj, { src_offset },
+                   { 0, 0, 0 },
+                   region));
+
+   ret_object(ev, hev);
+   return CL_SUCCESS;
+
+} catch (error &e) {
+   return e.get();
+}
+
+PUBLIC void *
+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 {
+   validate_base(q, num_deps, deps);
+   validate_obj(q, obj);
+
+   if (offset > obj->size() || offset + size > obj->size())
+      throw error(CL_INVALID_VALUE);
+
+   void *map = obj->resource(q).add_map(
+      *q, flags, blocking, { offset }, { size });
+
+   ret_object(ev, new hard_event(*q, CL_COMMAND_MAP_BUFFER,
+                                 { deps, deps + num_deps }));
+   ret_error(errcode_ret, CL_SUCCESS);
+   return map;
+
+} catch (error &e) {
+   ret_error(errcode_ret, e);
+   return NULL;
+}
+
+PUBLIC void *
+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 {
+   image *img = dynamic_cast<image *>(obj);
+
+   validate_base(q, num_deps, deps);
+   validate_obj(q, img);
+
+   void *map = obj->resource(q).add_map(
+      *q, flags, blocking, origin, region);
+
+   ret_object(ev, new hard_event(*q, CL_COMMAND_MAP_IMAGE,
+                                 { deps, deps + num_deps }));
+   ret_error(errcode_ret, CL_SUCCESS);
+   return map;
+
+} catch (error &e) {
+   ret_error(errcode_ret, e);
+   return NULL;
+}
+
+PUBLIC cl_int
+clEnqueueUnmapMemObject(cl_command_queue q, cl_mem obj, void *ptr,
+                        cl_uint num_deps, const cl_event *deps,
+                        cl_event *ev) try {
+   validate_base(q, num_deps, deps);
+   validate_obj(q, obj);
+
+   hard_event *hev = new hard_event(
+      *q, CL_COMMAND_UNMAP_MEM_OBJECT, { deps, deps + num_deps },
+      [=](event &) {
+         obj->resource(q).del_map(ptr);
+      });
+
+   ret_object(ev, hev);
+   return CL_SUCCESS;
+
+} catch (error &e) {
+   return e.get();
+}
diff --git a/src/gallium/state_trackers/clover/api/util.hpp b/src/gallium/state_trackers/clover/api/util.hpp
new file mode 100644
index 00000000000..2f9ec1f6a10
--- /dev/null
+++ b/src/gallium/state_trackers/clover/api/util.hpp
@@ -0,0 +1,166 @@
+//
+// Copyright 2012 Francisco Jerez
+//
+// 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 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.
+//
+
+#ifndef __CL_UTIL_HPP__
+#define __CL_UTIL_HPP__
+
+#include <cstdint>
+#include <cstring>
+#include <algorithm>
+#include <map>
+
+#include "core/base.hpp"
+#include "pipe/p_compiler.h"
+
+namespace clover {
+   ///
+   /// Return a matrix (a container of containers) in \a buf with
+   /// argument and bounds checking.  Intended to be used by
+   /// implementations of \a clGetXXXInfo().
+   ///
+   template<typename T, typename V>
+   cl_int
+   matrix_property(void *buf, size_t size, size_t *size_ret, const V& v) {
+      if (buf && size < sizeof(T *) * v.size())
+         return CL_INVALID_VALUE;
+
+      if (size_ret)
+         *size_ret = sizeof(T *) * v.size();
+
+      if (buf)
+         for_each([](typename V::value_type src, T *dst) {
+               if (dst)
+                  std::copy(src.begin(), src.end(), dst);
+            },
+            v.begin(), v.end(), (T **)buf);
+
+      return CL_SUCCESS;
+   }
+
+   ///
+   /// Return a vector in \a buf with argument and bounds checking.
+   /// Intended to be used by implementations of \a clGetXXXInfo().
+   ///
+   template<typename T, typename V>
+   cl_int
+   vector_property(void *buf, size_t size, size_t *size_ret, const V& v) {
+      if (buf && size < sizeof(T) * v.size())
+         return CL_INVALID_VALUE;
+
+      if (size_ret)
+         *size_ret = sizeof(T) * v.size();
+      if (buf)
+         std::copy(v.begin(), v.end(), (T *)buf);
+
+      return CL_SUCCESS;
+   }
+
+   ///
+   /// Return a scalar in \a buf with argument and bounds checking.
+   /// Intended to be used by implementations of \a clGetXXXInfo().
+   ///
+   template<typename T>
+   cl_int
+   scalar_property(void *buf, size_t size, size_t *size_ret, T v) {
+      return vector_property<T>(buf, size, size_ret, std::vector<T>(1, v));
+   }
+
+   ///
+   /// Return a string in \a buf with argument and bounds checking.
+   /// Intended to be used by implementations of \a clGetXXXInfo().
+   ///
+   inline cl_int
+   string_property(void *buf, size_t size, size_t *size_ret,
+                   const std::string &v) {
+      if (buf && size < v.size() + 1)
+         return CL_INVALID_VALUE;
+
+      if (size_ret)
+         *size_ret = v.size() + 1;
+      if (buf)
+         std::strcpy((char *)buf, v.c_str());
+
+      return CL_SUCCESS;
+   }
+
+   ///
+   /// Convert a NULL-terminated property list into an std::map.
+   ///
+   template<typename T>
+   std::map<T, T>
+   property_map(const T *props) {
+      std::map<T, T> m;
+
+      while (props && *props) {
+         T key = *props++;
+         T value = *props++;
+
+         if (m.count(key))
+            throw clover::error(CL_INVALID_PROPERTY);
+
+         m.insert({ key, value });
+      }
+
+      return m;
+   }
+
+   ///
+   /// Convert an std::map into a NULL-terminated property list.
+   ///
+   template<typename T>
+   std::vector<T>
+   property_vector(const std::map<T, T> &m) {
+      std::vector<T> v;
+
+      for (auto &p : m) {
+         v.push_back(p.first);
+         v.push_back(p.second);
+      }
+
+      v.push_back(0);
+      return v;
+   }
+
+   ///
+   /// Return an error code in \a p if non-zero.
+   ///
+   inline void
+   ret_error(cl_int *p, const clover::error &e) {
+      if (p)
+         *p = e.get();
+   }
+
+   ///
+   /// Return a reference-counted object in \a p if non-zero.
+   /// Otherwise release object ownership.
+   ///
+   template<typename T, typename S>
+   void
+   ret_object(T p, S v) {
+      if (p)
+         *p = v;
+      else
+         v->release();
+   }
+}
+
+#endif