diff --git a/doc/make_constants.py b/doc/make_constants.py
index d5b444a5e6668ebf5ae863b866c75924e2f212a6..f9df2f36c2970b4861a2e28a9df99f214ca8b148 100644
--- a/doc/make_constants.py
+++ b/doc/make_constants.py
@@ -3,12 +3,26 @@ import pyopencl as cl
def doc_class(cls):
print ".. class :: %s" % cls.__name__
print
+ if cls.__name__.startswith("gl_"):
+ print " Only available when PyOpenCL is compiled with GL support. See :func:`have_gl`."
+ print
for i in sorted(dir(cls)):
- if not i.startswith("_"):
+ if not i.startswith("_") and not i == "to_string":
print " .. attribute :: %s" % i
+ print " .. method :: to_string(value)"
print
+ print " Returns a :class:`str` representing *value*."
+ print
+ print " .. versionadded:: 0.91"
+ print
+
+if not cl.have_gl():
+ print "***************************************************************"
+ print "WARNING: GL not supported. Incompleted docs will be generated."
+ print "***************************************************************"
print ".. This is an automatically generated file. DO NOT EDIT"
+print
for cls in cl.CONSTANT_CLASSES:
doc_class(cls)
diff --git a/doc/source/misc.rst b/doc/source/misc.rst
index 61abb7fb5bcd1cce980b0041eea4a28a376ea799..0b06f9426049a9202a2d117a095018b6de680bbe 100644
--- a/doc/source/misc.rst
+++ b/doc/source/misc.rst
@@ -7,12 +7,43 @@ Installation information is maintained collaboratively on the
Acknowledgments
===============
-* James Snyder provided a patch to make PyOpenCL work on OS X 10.6.
+* James Snyder provided patches to make PyOpenCL work on OS X 10.6.
* Roger Pau Monné supplied the example :file:`examples/benchmark-all.py`.
User-visible Changes
====================
+Version 0.91
+------------
+
+* Add :ref:`gl-interop`.
+* Add a test suite.
+* Fix numerous `get_info` bugs. (reports by David Garcia and the test suite)
+* Add :meth:`pyopencl.ImageFormat.__repr__`.
+* Add :meth:`pyopencl.addressing_mode.to_string` and colleagues.
+* The `pitch` arguments to
+ :func:`pyopencl.create_image_2d`,
+ :func:`pyopencl.create_image_3d`,
+ :func:`pyopencl.enqueue_read_image`, and
+ :func:`pyopencl.enqueue_write_image`
+ are now defaulted to zero. The argument order of `enqueue_{read,write}_image`
+ has changed for this reason.
+* :meth:`pyopencl.MemoryObject.get_image_info` now actually exists.
+* Add :attr:`pyopencl.MemoryObject.image.info`.
+* Fix API tracing.
+* Add constructor arguments to :class:`pyopencl.ImageFormat`.
+ (suggested by David Garcia)
+
+Version 0.90.4
+--------------
+
+* Add build fixes for Windows and OS X.
+
+Version 0.90.3
+--------------
+
+* Fix a GNU-ism in the C++ code of the wrapper.
+
Version 0.90.2
--------------
diff --git a/doc/source/reference.rst b/doc/source/reference.rst
index 8ea5f616ebd5087feced671ee2c6f3c2a550631a..caa215e3f34bc8eea8d60021042176bf9d432d9a 100644
--- a/doc/source/reference.rst
+++ b/doc/source/reference.rst
@@ -51,7 +51,9 @@ Platforms, Devices and Contexts
using *"=="* and *"!="*.
.. |buf-iface| replace:: must implement the Python buffer interface.
(e.g. by being an :class:`numpy.ndarray`)
-.. |enqueue-waitfor| replace:: Returns a new :class:`Event`.
+.. |enqueue-waitfor| replace:: Returns a new :class:`Event`. *wait_for*
+ may either be *None* or a list of :class:`Event` instances for
+ whose completion this command waits before starting exeuction.
.. function:: get_platforms()
@@ -188,18 +190,39 @@ Memory
See :class:`mem_info` for values of *param*.
+ .. method:: get_image_info(param)
+
+ See :class:`image_info` for values of *param*.
+
+ .. versionadded:: 0.91
+
.. attribute:: info
Lower case versions of the :class:`mem_info` constants
may be used as attributes on instances of this class
to directly query info attributes.
- .. method:: get_image_info(param)
+ .. attribute:: image.info
- See :class:`image_info` for values of *param*.
+ Lower case versions of the :class:`image_info` constants
+ may be used as attributes on the attribute `image` of this
+ class to directly query image info.
+
+ For example, you may use *img.image.depth* instead of
+ *img.get_image_info(pyopencl.image_info.DEPTH)*.
.. method:: release()
+ .. method:: get_gl_object_info()
+
+ Return a tuple *(obj_type, obj_name)*, where *obj_type* is one of the
+ :class:`gl_object_type` constants, and *obj_name* is the GL object
+ name.
+ Only available when PyOpenCL is compiled with GL support. See :func:`have_gl`.
+
+ .. method:: get_gl_texture_info(param)
+
+ See :class:`gl_texture_info` for values of *param*. Only available when PyOpenCL is compiled with GL support. See :func:`have_gl`.
|comparable|
Buffers
@@ -227,7 +250,10 @@ Buffers
Image Formats
^^^^^^^^^^^^^
-.. class:: ImageFormat
+.. class:: ImageFormat([channel_order, channel_type])
+
+ .. versionchanged:: 0.91
+ Constructor arguments added.
.. attribute:: channel_order
@@ -237,6 +263,12 @@ Image Formats
See :class:`channel_type` for possible values.
+ .. method:: __repr__
+
+ Returns a :class:`str` representation of the image format.
+
+ .. versionadded:: 0.91
+
.. function:: get_supported_image_formats(context, flags, image_type)
See :class:`mem_flags` for possible values of *flags*
@@ -245,24 +277,36 @@ Image Formats
Images
^^^^^^
-.. function:: create_image_2d(context, flags, format, width, height, pitch, host_buffer=None)
+.. function:: create_image_2d(context, flags, format, width, height, pitch=0, host_buffer=None)
See :class:`mem_flags` for possible values of *flags*.
Returns a new image-type :class:`MemoryObject`.
-.. function:: create_image_3d(context, flags, format, width, height, depth, row_pitch, slice_pitch, host_buffer=None)
+ .. versionchanged:: 0.91
+ *pitch* argument defaults to zero, moved.
+
+.. function:: create_image_3d(context, flags, format, width, height, depth, row_pitch=0, slice_pitch=0, host_buffer=None)
See :class:`mem_flags` for possible values of *flags*.
Returns a new image-type :class:`MemoryObject`.
-.. function:: enqueue_read_image(queue, mem, origin, region, row_pitch, slice_pitch, host_buffer, wait_for=None, is_blocking=False)
+ .. versionchanged:: 0.91
+ *pitch* arguments defaults to zero, moved.
+
+.. function:: enqueue_read_image(queue, mem, origin, region, host_buffer, row_pitch=0, slice_pitch=0, wait_for=None, is_blocking=False)
|enqueue-waitfor|
-.. function:: enqueue_write_image(queue, mem, origin, region, row_pitch, slice_pitch, host_buffer, wait_for=None, is_blocking=False)
+ .. versionchanged:: 0.91
+ *pitch* arguments defaults to zero, moved.
+
+.. function:: enqueue_write_image(queue, mem, origin, region, host_buffer, row_pitch=0, slice_pitch=0, wait_for=None, is_blocking=False)
|enqueue-waitfor|
+ .. versionchanged:: 0.91
+ *pitch* arguments defaults to zero, moved.
+
.. function:: enqueue_copy_image(queue, src, dest, src_origin, dest_origin, region, wait_for=None)
|enqueue-waitfor|
@@ -331,7 +375,7 @@ Programs and Kernels
may be used as attributes on instances of this class
to directly query info attributes.
- .. method:: get_build_info(param, device)
+ .. method:: get_build_info(device, param)
See :class:`program_build_info` for values of *param*.
@@ -353,6 +397,10 @@ Programs and Kernels
|comparable|
+ .. method:: all_kernels()
+
+ Returns a list of all :class:`Kernel` objects in the :class:`Program`.
+
.. function:: unload_compiler()
.. function:: create_program_with_source(context, src)
.. function:: create_program_with_binary(context, devices, binaries)
@@ -390,3 +438,52 @@ Programs and Kernels
.. function:: enqueue_task(queue, kernel, wait_for=None)
|enqueue-waitfor|
+
+.. _gl-interop:
+
+GL Interoperability
+-------------------
+
+Functionality in this section is only available when PyOpenCL is compiled
+with GL support. See :func:`have_gl`.
+
+.. versionadded:: 0.91
+
+.. function:: have_gl()
+
+ Return *True* if PyOpenCL was compiled with OpenGL interoperability, otherwise *False*.
+
+.. function:: create_from_gl_buffer(context, mem_flags, gl_buffer_obj)
+
+ See :class:`mem_flags` for values of *flags*.
+ Returns a new :class:`MemoryObject`.
+
+.. function:: create_from_gl_texture_2d(context, mem_flags, texture_target, miplevel, texture)
+
+ See :class:`mem_flags` for values of *flags*.
+ Returns a new :class:`MemoryObject`.
+
+.. function:: create_from_gl_texture_3d(context, mem_flags, texture_target, miplevel, texture)
+
+ See :class:`mem_flags` for values of *flags*.
+ Returns a new :class:`MemoryObject`.
+
+.. function:: create_from_gl_renderbuffer(context, mem_flags, gl_renderbuffer)
+
+ See :class:`mem_flags` for values of *flags*.
+ Returns a new :class:`MemoryObject`.
+
+.. function:: enqueue_acquire_gl_objects(queue, mem_objects, wait_for=None)
+
+ *mem_objects* is a list of :class:`MemoryObject` instances.
+ |enqueue-waitfor|
+
+.. function:: enqueue_release_gl_objects(queue, mem_objects, wait_for=None)
+
+ *mem_objects* is a list of :class:`MemoryObject` instances.
+ |enqueue-waitfor|
+
+.. seealso::
+
+ * :meth:`MemoryObject.get_gl_object_info`
+ * :meth:`MemoryObject.get_gl_texture_info`
diff --git a/examples/benchmark-all.py b/examples/benchmark-all.py
index 493dd48eb3f3b0f8895612537d8220beed35d001..93685006dc8817307c12025d5813f9c5485f1e6d 100644
--- a/examples/benchmark-all.py
+++ b/examples/benchmark-all.py
@@ -23,12 +23,6 @@ print "Execution time of test without OpenCL: ", time2 - time1, "s"
for platform in cl.get_platforms():
for device in platform.get_devices():
- dev_type = "unknown"
-
- for dev_type_candidate in dir(cl.device_type):
- if getattr(cl.device_type, dev_type_candidate) == device.type:
- dev_type = dev_type_candidate
-
print "==============================================================="
print "Platform name:", platform.name
print "Platform profile:", platform.profile
@@ -36,7 +30,7 @@ for platform in cl.get_platforms():
print "Platform version:", platform.version
print "---------------------------------------------------------------"
print "Device name:", device.name
- print "Device type: ", dev_type
+ print "Device type:", cl.device_type.to_string(device.type)
print "Device memory: ", device.global_mem_size//1024//1024, 'MB'
print "Device max clock speed:", device.max_clock_frequency, 'MHz'
print "Device compute units:", device.max_compute_units
diff --git a/pyopencl/__init__.py b/pyopencl/__init__.py
index a8a992051ab6aa51c8e763eef326b7762dcd7df6..2f1e03c6c91cd0cac2df56001b455108180447c5 100644
--- a/pyopencl/__init__.py
+++ b/pyopencl/__init__.py
@@ -20,6 +20,42 @@ def _add_functionality():
_cl.Kernel: _cl.kernel_info,
}
+ def to_string(cls, value):
+ for name in dir(cls):
+ if (not name.startswith("_") and getattr(cls, name) == value):
+ return name
+
+ raise ValueError("a name for value %d was not found in %s"
+ % (value, cls.__name__))
+
+ addressing_mode.to_string = classmethod(to_string)
+ channel_order.to_string = classmethod(to_string)
+ channel_type.to_string = classmethod(to_string)
+ command_execution_status.to_string = classmethod(to_string)
+ command_queue_info.to_string = classmethod(to_string)
+ command_queue_properties.to_string = classmethod(to_string)
+ context_info.to_string = classmethod(to_string)
+ context_properties.to_string = classmethod(to_string)
+ device_exec_capabilities.to_string = classmethod(to_string)
+ device_fp_config.to_string = classmethod(to_string)
+ device_info.to_string = classmethod(to_string)
+ device_local_mem_type.to_string = classmethod(to_string)
+ device_mem_cache_type.to_string = classmethod(to_string)
+ device_type.to_string = classmethod(to_string)
+ event_info.to_string = classmethod(to_string)
+ filter_mode.to_string = classmethod(to_string)
+ image_info.to_string = classmethod(to_string)
+ kernel_info.to_string = classmethod(to_string)
+ kernel_work_group_info.to_string = classmethod(to_string)
+ map_flags.to_string = classmethod(to_string)
+ mem_info.to_string = classmethod(to_string)
+ mem_object_type.to_string = classmethod(to_string)
+ platform_info.to_string = classmethod(to_string)
+ profiling_info.to_string = classmethod(to_string)
+ program_build_info.to_string = classmethod(to_string)
+ program_info.to_string = classmethod(to_string)
+ sampler_info.to_string = classmethod(to_string)
+
class ProfilingInfoGetter:
def __init__(self, event):
self.event = event
@@ -37,6 +73,23 @@ def _add_functionality():
_cl.Event.profile = property(ProfilingInfoGetter)
+ class ImageInfoGetter:
+ def __init__(self, mem):
+ self.mem = mem
+
+ def __getattr__(self, name):
+ info_cls = _cl.image_info
+
+ try:
+ inf_attr = getattr(info_cls, name.upper())
+ except AttributeError:
+ raise AttributeError("%s has no attribute '%s'"
+ % (type(self), name))
+ else:
+ return self.mem.get_image_info(inf_attr)
+
+ _cl.MemoryObject.image = property(ImageInfoGetter)
+
def make_getattr(info_cls):
def result(self, name):
try:
@@ -57,7 +110,11 @@ def _add_functionality():
pi_attr = getattr(program_info, attr.upper())
except AttributeError:
try:
- return Kernel(self, attr)
+ knl = Kernel(self, attr)
+ # Nvidia does not raise errors even for invalid names,
+ # but this will give an error if the kernel is invalid.
+ knl.num_args
+ return knl
except LogicError:
raise AttributeError("'%s' was not found as a program info attribute or as a kernel name"
% attr)
@@ -84,6 +141,13 @@ def _add_functionality():
Kernel.__call__ = kernel_call
+ def image_format_repr(self):
+ return "ImageFormat(%s, %s)" % (
+ channel_order.to_string(self.channel_order),
+ channel_type.to_string(self.channel_data_type))
+
+ ImageFormat.__repr__ = image_format_repr
+
def event_wait(self):
wait_for_events([self])
return self
diff --git a/pyopencl/version.py b/pyopencl/version.py
index 8029e54f53bf8d8b8406a185db72dd0fa3eef71b..5abc8d8af0686b44410279fda6687f4fb017ed26 100644
--- a/pyopencl/version.py
+++ b/pyopencl/version.py
@@ -1,5 +1,5 @@
-VERSION = (0, 90, 4)
-VERSION_STATUS = ""
+VERSION = (0, 91)
+VERSION_STATUS = "alpha"
VERSION_TEXT = ".".join(str(x) for x in VERSION) + VERSION_STATUS
diff --git a/setup.py b/setup.py
index 9098a3fed908b6cc47142a7db6ca2a5d8d15a84a..5543bded8b7a076d8eede4683ce6fe0f01629f8a 100644
--- a/setup.py
+++ b/setup.py
@@ -13,6 +13,7 @@ def get_config_schema():
BoostLibraries("thread"),
Switch("CL_TRACE", False, "Enable OpenCL API tracing"),
+ Switch("CL_ENABLE_GL", False, "Enable OpenCL<->OpenGL interoperability"),
Switch("SHIPPED_CL_HEADERS", False, "Use shipped OpenCL headers"),
IncludeDir("CL", []),
@@ -47,7 +48,7 @@ def main():
EXTRA_LIBRARIES = []
if conf["CL_TRACE"]:
- EXTRA_DEFINES["CLPP_TRACE_CL"] = 1
+ EXTRA_DEFINES["PYOPENCL_TRACE"] = 1
INCLUDE_DIRS = ['src/cpp'] + conf["BOOST_INC_DIR"] + conf["CL_INC_DIR"]
@@ -70,6 +71,9 @@ def main():
ext_kwargs = dict()
+ if conf["CL_ENABLE_GL"]:
+ EXTRA_DEFINES["HAVE_GL"] = 1
+
ver_dic = {}
execfile("pyopencl/version.py", ver_dic)
@@ -132,6 +136,7 @@ def main():
install_requires=[
"pytools>=7",
+ "py>=1.0.2"
],
ext_package="pyopencl",
diff --git a/src/wrapper/wrap_cl.cpp b/src/wrapper/wrap_cl.cpp
index a0e05b7ea7aeb4c13f570f9dc75a2eb3c9bc6511..d8bfcb80c1bc5eed7635927263aae5d252b960e2 100644
--- a/src/wrapper/wrap_cl.cpp
+++ b/src/wrapper/wrap_cl.cpp
@@ -65,8 +65,10 @@ namespace
class command_type { };
class command_execution_status { };
class profiling_info { };
-}
+ class gl_object_type { };
+ class gl_texture_info { };
+}
@@ -468,9 +470,14 @@ BOOST_PYTHON_MODULE(_cl)
typedef memory_object cls;
py::class_<cls, boost::noncopyable>("MemoryObject", py::no_init)
.DEF_SIMPLE_METHOD(get_info)
+ .DEF_SIMPLE_METHOD(get_image_info)
.DEF_SIMPLE_METHOD(release)
.def(py::self == py::self)
.def(py::self != py::self)
+#ifdef HAVE_GL
+ .def("get_gl_object_info", get_gl_object_info)
+ .def("get_gl_texture_info", get_gl_texture_info)
+#endif
;
}
@@ -496,6 +503,7 @@ BOOST_PYTHON_MODULE(_cl)
{
typedef cl_image_format cls;
py::class_<cls>("ImageFormat")
+ .def("__init__", py::make_constructor(make_image_format))
.def_readwrite("channel_order", &cls::image_channel_order)
.def_readwrite("channel_data_type", &cls::image_channel_data_type)
;
@@ -503,24 +511,30 @@ BOOST_PYTHON_MODULE(_cl)
DEF_SIMPLE_FUNCTION(get_supported_image_formats);
py::def("create_image_2d", create_image_2D,
- (py::args("context", "flags", "format", "width", "height", "pitch"),
- py::arg("host_buffer")=py::object()),
+ (py::args("context", "flags", "format", "width", "height"),
+ py::arg("pitch")=0,
+ py::arg("host_buffer")=py::object()
+ ),
py::return_value_policy<py::manage_new_object>());
py::def("create_image_3d", create_image_3D,
- (py::args("context", "flags", "format", "width", "height", "depth",
- "row_pitch", "slice_pitch"),
- py::arg("host_buffer")=py::object()),
+ (py::args("context", "flags", "format", "width", "height", "depth"),
+ py::arg("row_pitch")=0,
+ py::arg("slice_pitch")=0,
+ py::arg("host_buffer")=py::object()
+ ),
py::return_value_policy<py::manage_new_object>());
py::def("enqueue_read_image", enqueue_read_image,
- (py::args("queue", "mem", "origin", "region",
- "row_pitch", "slice_pitch", "host_buffer"),
+ (py::args("queue", "mem", "origin", "region", "host_buffer"),
+ py::arg("row_pitch")=0,
+ py::arg("slice_pitch")=0,
py::arg("wait_for")=py::object(),
py::arg("is_blocking")=false),
py::return_value_policy<py::manage_new_object>());
py::def("enqueue_write_image", enqueue_write_image,
- (py::args("queue", "mem", "origin", "region",
- "row_pitch", "slice_pitch", "host_buffer"),
+ (py::args("queue", "mem", "origin", "region", "host_buffer"),
+ py::arg("row_pitch")=0,
+ py::arg("slice_pitch")=0,
py::arg("wait_for")=py::object(),
py::arg("is_blocking")=false),
py::return_value_policy<py::manage_new_object>());
@@ -584,6 +598,7 @@ BOOST_PYTHON_MODULE(_cl)
py::return_self<>())
.def(py::self == py::self)
.def(py::self != py::self)
+ .def("all_kernels", create_kernels_in_program)
;
}
@@ -608,7 +623,6 @@ BOOST_PYTHON_MODULE(_cl)
;
}
- DEF_SIMPLE_FUNCTION(create_kernels_in_program);
py::def("enqueue_nd_range_kernel", enqueue_nd_range_kernel,
(py::args("queue", "kernel"),
py::arg("global_work_size"),
@@ -624,4 +638,44 @@ BOOST_PYTHON_MODULE(_cl)
py::return_value_policy<py::manage_new_object>());
// TODO: clEnqueueNativeKernel
+
+ // GL interop ---------------------------------------------------------------
+ DEF_SIMPLE_FUNCTION(have_gl);
+
+#ifdef HAVE_GL
+ {
+ py::class_<gl_object_type> cls("gl_object_type", py::no_init);
+ ADD_ATTR(GL_OBJECT_, BUFFER);
+ ADD_ATTR(GL_OBJECT_, TEXTURE2D);
+ ADD_ATTR(GL_OBJECT_, TEXTURE3D);
+ ADD_ATTR(GL_OBJECT_, RENDERBUFFER);
+ }
+
+ {
+ py::class_<gl_texture_info> cls("gl_texture_info", py::no_init);
+ ADD_ATTR(GL_, TEXTURE_TARGET);
+ ADD_ATTR(GL_, MIPMAP_LEVEL);
+ }
+
+ py::def("create_from_gl_buffer", create_from_gl_buffer,
+ py::return_value_policy<py::manage_new_object>());
+ py::def("create_from_gl_texture_2d", create_from_gl_texture_2d,
+ py::return_value_policy<py::manage_new_object>());
+ py::def("create_from_gl_texture_3d", create_from_gl_texture_3d,
+ py::return_value_policy<py::manage_new_object>());
+ py::def("create_from_gl_renderbuffer", create_from_gl_renderbuffer,
+ py::return_value_policy<py::manage_new_object>());
+
+ py::def("enqueue_acquire_gl_objects", enqueue_acquire_gl_objects,
+ (py::args("queue", "mem_objects"),
+ py::arg("wait_for")=py::object()
+ ),
+ py::return_value_policy<py::manage_new_object>());
+ py::def("enqueue_release_gl_objects", enqueue_release_gl_objects,
+ (py::args("queue", "mem_objects"),
+ py::arg("wait_for")=py::object()
+ ),
+ py::return_value_policy<py::manage_new_object>());
+
+#endif
}
diff --git a/src/wrapper/wrap_cl.hpp b/src/wrapper/wrap_cl.hpp
index a8d94ef24f8805d6f7feb8c2faa2b68f4bb43d62..19f77a5e5ffda56da26cc2ef56cf778702a23a4d 100644
--- a/src/wrapper/wrap_cl.hpp
+++ b/src/wrapper/wrap_cl.hpp
@@ -10,16 +10,31 @@
#ifdef __APPLE__
+
+// Mac ------------------------------------------------------------------------
#include <OpenCL/opencl.h>
+#ifdef HAVE_GL
+#include <OpenCL/opencl_gl.h>
+#endif
+
#else
+
+// elsewhere ------------------------------------------------------------------
#include <CL/cl.h>
+#ifdef HAVE_GL
+#include <GL/gl.h>
+#include <CL/cl_gl.h>
+#endif
+
#endif
#include <stdexcept>
#include <iostream>
#include <vector>
#include <utility>
+#include <numeric>
#include <boost/foreach.hpp>
+#include <boost/scoped_array.hpp>
#include "wrap_helpers.hpp"
#include "numpy_init.hpp"
@@ -200,6 +215,30 @@
+// buffer creators ------------------------------------------------------------
+#define PYOPENCL_WRAP_BUFFER_CREATOR(NAME, CL_NAME, ARGS, CL_ARGS) \
+ memory_object *NAME ARGS \
+ { \
+ cl_int status_code; \
+ PYOPENCL_PRINT_CALL_TRACE(#CL_NAME); \
+ cl_mem mem = CL_NAME CL_ARGS; \
+ \
+ if (status_code != CL_SUCCESS) \
+ throw pyopencl::error(#CL_NAME, status_code); \
+ \
+ try \
+ { \
+ return new memory_object(mem, false); \
+ } \
+ catch (...) \
+ { \
+ PYOPENCL_CALL_GUARDED(clReleaseMemObject, (mem)); \
+ throw; \
+ } \
+ }
+
+
+
namespace pyopencl
@@ -400,13 +439,13 @@ namespace pyopencl
case CL_DEVICE_VENDOR_ID: DEV_GET_INT_INF(cl_uint);
case CL_DEVICE_MAX_COMPUTE_UNITS: DEV_GET_INT_INF(cl_uint);
case CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS: DEV_GET_INT_INF(cl_uint);
- case CL_DEVICE_MAX_WORK_GROUP_SIZE: DEV_GET_INT_INF(cl_uint);
+ case CL_DEVICE_MAX_WORK_GROUP_SIZE: DEV_GET_INT_INF(size_t);
case CL_DEVICE_MAX_WORK_ITEM_SIZES:
{
std::vector<size_t> result;
PYOPENCL_GET_VEC_INFO(Device, m_device, param_name, result);
- return py::list(result);
+ PYOPENCL_RETURN_VECTOR(size_t, result);
}
case CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR: DEV_GET_INT_INF(cl_uint);
@@ -426,7 +465,7 @@ namespace pyopencl
case CL_DEVICE_IMAGE3D_MAX_HEIGHT: DEV_GET_INT_INF(size_t);
case CL_DEVICE_IMAGE3D_MAX_DEPTH: DEV_GET_INT_INF(size_t);
case CL_DEVICE_IMAGE_SUPPORT: DEV_GET_INT_INF(cl_bool);
- case CL_DEVICE_MAX_PARAMETER_SIZE: DEV_GET_INT_INF(cl_bool);
+ case CL_DEVICE_MAX_PARAMETER_SIZE: DEV_GET_INT_INF(size_t);
case CL_DEVICE_MAX_SAMPLERS: DEV_GET_INT_INF(cl_uint);
case CL_DEVICE_MEM_BASE_ADDR_ALIGN: DEV_GET_INT_INF(cl_uint);
case CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE: DEV_GET_INT_INF(cl_uint);
@@ -611,6 +650,10 @@ namespace pyopencl
handle_from_new_ptr(new platform(
reinterpret_cast<cl_platform_id>(result[i+1]))));
}
+
+ case 0:
+ break;
+
default:
throw error("Context.get_info", CL_INVALID_VALUE,
"unkown context_property key encountered");
@@ -988,28 +1031,9 @@ namespace pyopencl
- memory_object *create_buffer(
- context &ctx,
- cl_mem_flags flags,
- size_t size)
- {
- cl_int status_code;
- PYOPENCL_PRINT_CALL_TRACE("clCreateBuffer");
- cl_mem mem = clCreateBuffer(ctx.data(), flags, size, 0, &status_code);
-
- if (status_code != CL_SUCCESS)
- throw pyopencl::error("clCreateBuffer", status_code);
-
- try
- {
- return new memory_object(mem, false);
- }
- catch (...)
- {
- PYOPENCL_CALL_GUARDED(clReleaseMemObject, (mem));
- throw;
- }
- }
+ PYOPENCL_WRAP_BUFFER_CREATOR(create_buffer, clCreateBuffer,
+ (context &ctx, cl_mem_flags flags, size_t size),
+ (ctx.data(), flags, size, 0, &status_code));
@@ -1058,7 +1082,6 @@ namespace pyopencl
-
event *enqueue_read_buffer(
command_queue &cq,
memory_object &mem,
@@ -1122,6 +1145,14 @@ namespace pyopencl
// images -------------------------------------------------------------------
+ cl_image_format *make_image_format(cl_channel_order ord, cl_channel_type tp)
+ {
+ std::auto_ptr<cl_image_format> result(new cl_image_format);
+ result->image_channel_order = ord;
+ result->image_channel_data_type = tp;
+ return result.release();
+ }
+
py::list get_supported_image_formats(
context const &ctx,
cl_mem_flags flags,
@@ -1135,7 +1166,8 @@ namespace pyopencl
PYOPENCL_CALL_GUARDED(clGetSupportedImageFormats, (
ctx.data(), flags, image_type,
num_image_formats, &formats.front(), 0));
- return py::list(formats);
+
+ PYOPENCL_RETURN_VECTOR(cl_image_format, formats);
}
@@ -1215,8 +1247,8 @@ namespace pyopencl
command_queue &cq,
memory_object &mem,
py::object py_origin, py::object py_region,
- size_t row_pitch, size_t slice_pitch,
py::object buffer,
+ size_t row_pitch, size_t slice_pitch,
py::object py_wait_for,
bool is_blocking
)
@@ -1249,8 +1281,8 @@ namespace pyopencl
command_queue &cq,
memory_object &mem,
py::object py_origin, py::object py_region,
- size_t row_pitch, size_t slice_pitch,
py::object buffer,
+ size_t row_pitch, size_t slice_pitch,
py::object py_wait_for,
bool is_blocking
)
@@ -1629,19 +1661,24 @@ namespace pyopencl
{
std::vector<size_t> result;
PYOPENCL_GET_VEC_INFO(Program, m_program, param_name, result);
- return py::list(result);
+ PYOPENCL_RETURN_VECTOR(size_t, result);
}
case CL_PROGRAM_BINARIES:
{
std::vector<size_t> sizes;
- PYOPENCL_GET_VEC_INFO(Program, m_program, param_name, sizes);
+ PYOPENCL_GET_VEC_INFO(Program, m_program, CL_PROGRAM_BINARY_SIZES, sizes);
+
+ size_t total_size = std::accumulate(sizes.begin(), sizes.end(), 0);
- std::vector<std::vector<unsigned char> > result(sizes.size());
- std::vector<unsigned char *> result_ptrs(sizes.size());
+ boost::scoped_array<unsigned char> result(
+ new unsigned char[total_size]);
+ std::vector<unsigned char *> result_ptrs;
+
+ unsigned char *ptr = result.get();
for (unsigned i = 0; i < sizes.size(); ++i)
{
- result[i].resize(sizes[i]);
- result_ptrs.push_back(&result[i].front());
+ result_ptrs.push_back(ptr);
+ ptr += sizes[i];
}
PYOPENCL_CALL_GUARDED(clGetProgramInfo,
@@ -1649,10 +1686,14 @@ namespace pyopencl
&result_ptrs.front(), 0)); \
py::list py_result;
+ ptr = result.get();
for (unsigned i = 0; i < sizes.size(); ++i)
+ {
py_result.append(py::str(
- reinterpret_cast<char *>(&result[i].front()),
+ reinterpret_cast<char *>(ptr),
sizes[i]));
+ ptr += sizes[i];
+ }
return py_result;
}
@@ -2027,6 +2068,95 @@ namespace pyopencl
PYOPENCL_RETURN_NEW_EVENT(evt);
}
+
+
+
+
+ // gl interop ---------------------------------------------------------------
+ bool have_gl()
+ {
+#ifdef HAVE_GL
+ return true;
+#else
+ return false;
+#endif
+ }
+
+
+
+
+#ifdef HAVE_GL
+ PYOPENCL_WRAP_BUFFER_CREATOR(create_from_gl_buffer, clCreateFromGLBuffer,
+ (context &ctx, cl_mem_flags flags, GLuint bufobj),
+ (ctx.data(), flags, bufobj, &status_code));
+ PYOPENCL_WRAP_BUFFER_CREATOR(create_from_gl_texture_2d, clCreateFromGLTexture2D,
+ (context &ctx, cl_mem_flags flags,
+ GLenum texture_target, GLint miplevel, GLuint texture),
+ (ctx.data(), flags, texture_target, miplevel, texture, &status_code));
+ PYOPENCL_WRAP_BUFFER_CREATOR(create_from_gl_texture_3d, clCreateFromGLTexture3D,
+ (context &ctx, cl_mem_flags flags,
+ GLenum texture_target, GLint miplevel, GLuint texture),
+ (ctx.data(), flags, texture_target, miplevel, texture, &status_code));
+ PYOPENCL_WRAP_BUFFER_CREATOR(create_from_gl_renderbuffer, clCreateFromGLRenderbuffer,
+ (context &ctx, cl_mem_flags flags, GLuint renderbuffer),
+ (ctx.data(), flags, renderbuffer, &status_code));
+
+
+
+
+ py::tuple get_gl_object_info(memory_object const &mem)
+ {
+ cl_gl_object_type otype;
+ GLuint gl_name;
+ PYOPENCL_CALL_GUARDED(clGetGLObjectInfo, (mem.data(), &otype, &gl_name));
+ return py::make_tuple(otype, gl_name);
+ }
+
+
+
+
+ py::object get_gl_texture_info(memory_object const &mem, cl_gl_texture_info param_name)
+ {
+ switch (param_name)
+ {
+ case CL_GL_TEXTURE_TARGET:
+ PYOPENCL_GET_INTEGRAL_INFO(GLTexture, mem.data(), param_name, GLenum);
+ case CL_GL_MIPMAP_LEVEL:
+ PYOPENCL_GET_INTEGRAL_INFO(GLTexture, mem.data(), param_name, GLint);
+
+ default:
+ throw error("MemoryObject.get_gl_texture_info", CL_INVALID_VALUE);
+ }
+ }
+
+
+
+
+#define WRAP_GL_ENQUEUE(what, What) \
+ event *enqueue_##what##_gl_objects( \
+ command_queue &cq, \
+ py::object py_mem_objects, \
+ py::object py_wait_for) \
+ { \
+ PYOPENCL_PARSE_WAIT_FOR; \
+ \
+ std::vector<cl_mem> mem_objects; \
+ PYTHON_FOREACH(mo, py_mem_objects) \
+ mem_objects.push_back(py::extract<memory_object &>(mo)().data()); \
+ \
+ cl_event evt; \
+ PYOPENCL_CALL_GUARDED(clEnqueue##What##GLObjects, ( \
+ cq.data(), \
+ mem_objects.size(), &mem_objects.front(), \
+ num_events_in_wait_list, &event_wait_list.front(), &evt \
+ )); \
+ \
+ PYOPENCL_RETURN_NEW_EVENT(evt); \
+ }
+
+ WRAP_GL_ENQUEUE(acquire, Acquire);
+ WRAP_GL_ENQUEUE(release, Release);
+#endif
}
diff --git a/src/wrapper/wrap_helpers.hpp b/src/wrapper/wrap_helpers.hpp
index 3561235b825505b52daf6d3ae0b66c12878fd7e8..903cfd3f8724789c80a56bebad57eb43f3b544f5 100644
--- a/src/wrapper/wrap_helpers.hpp
+++ b/src/wrapper/wrap_helpers.hpp
@@ -108,7 +108,13 @@ namespace py = boost::python;
else \
throw std::runtime_error("unrecognized order specifier"); \
-
+#define PYOPENCL_RETURN_VECTOR(ITEMTYPE, NAME) \
+ { \
+ py::list pyopencl_result; \
+ BOOST_FOREACH(ITEMTYPE item, NAME) \
+ pyopencl_result.append(item); \
+ return pyopencl_result; \
+ }
namespace
{
diff --git a/test/test_wrapper.py b/test/test_wrapper.py
new file mode 100644
index 0000000000000000000000000000000000000000..b2d39f7d2bca66beb4d917c4c94559d2f2e638b3
--- /dev/null
+++ b/test/test_wrapper.py
@@ -0,0 +1,163 @@
+from __future__ import division
+import numpy
+import numpy.linalg as la
+
+
+
+
+def have_cl():
+ try:
+ import pyopencl
+ return True
+ except:
+ return False
+
+
+if have_cl():
+ import pyopencl as cl
+
+
+
+
+class TestCL:
+ disabled = not have_cl()
+
+ def test_get_info(self):
+ had_failures = [False]
+
+ QUIRKS = [
+ ("NVIDIA", [
+ (cl.Device, cl.device_info.PLATFORM),
+ ]),
+ ]
+
+ def find_quirk(quirk_list, cl_obj, info):
+ for quirk_plat_name, quirks in quirk_list:
+ if quirk_plat_name in platform.name:
+ for quirk_cls, quirk_info in quirks:
+ if (isinstance(cl_obj, quirk_cls)
+ and quirk_info == info):
+ return True
+
+ return False
+
+ def do_test(cl_obj, info_cls, func=None):
+ if func is None:
+ def func(info):
+ cl_obj.get_info(info)
+
+ for info_name in dir(info_cls):
+ if not info_name.startswith("_") and info_name != "to_string":
+ info = getattr(info_cls, info_name)
+
+ try:
+ func(info)
+ except:
+ print "failed get_info", type(cl_obj), info_name
+
+ if find_quirk(QUIRKS, cl_obj, info):
+ print "(known quirk for %s)" % platform.name
+ else:
+ had_failures[0] = True
+ raise
+
+ for platform in cl.get_platforms():
+ do_test(platform, cl.platform_info)
+
+ for device in platform.get_devices():
+ do_test(device, cl.device_info)
+
+ ctx = cl.Context([device])
+ do_test(ctx, cl.context_info)
+
+ props = 0
+ if (device.queue_properties
+ & cl.command_queue_properties.PROFILING_ENABLE):
+ profiling = True
+ props = cl.command_queue_properties.PROFILING_ENABLE
+ queue = cl.CommandQueue(ctx,
+ properties=props)
+ do_test(queue, cl.command_queue_info)
+
+ prg = cl.create_program_with_source(ctx, """
+ __kernel void sum(__global float *a)
+ { a[get_global_id(0)] *= 2; }
+ """).build()
+ do_test(prg, cl.program_info)
+ do_test(prg, cl.program_build_info,
+ lambda info: prg.get_build_info(device, info))
+
+ cl.unload_compiler() # just for the heck of it
+
+ mf = cl.mem_flags
+ n = 2000
+ a_buf = cl.create_buffer(ctx, 0, n*4)
+
+ do_test(a_buf, cl.mem_info)
+
+ kernel = prg.sum
+ do_test(kernel, cl.kernel_info)
+
+ evt = kernel(queue, (n,), a_buf)
+ do_test(evt, cl.event_info)
+
+ if profiling:
+ evt.wait()
+ do_test(evt, cl.profiling_info,
+ lambda info: evt.get_profiling_info(info))
+
+ if device.image_support:
+ if "NVIDIA" not in platform.name:
+ # Samplers are crashy in Nvidia's "conformant" CL release
+ smp = cl.Sampler(ctx, True,
+ cl.addressing_mode.CLAMP,
+ cl.filter_mode.NEAREST)
+ do_test(smp, cl.sampler_info)
+
+ img_format = cl.get_supported_image_formats(
+ ctx, cl.mem_flags.READ_ONLY, cl.mem_object_type.IMAGE2D)[0]
+
+ img = cl.create_image_2d(ctx, cl.mem_flags.READ_ONLY, img_format,
+ 128, 128, 0)
+ do_test(img, cl.image_info,
+ lambda info: img.get_image_info(info))
+ img.image.depth
+
+ if had_failures[0]:
+ raise RuntimeError("get_info testing had errors")
+
+ def test_invalid_kernel_names_cause_failures(self):
+ for platform in cl.get_platforms():
+ for device in platform.get_devices():
+ ctx = cl.Context([device])
+ prg = cl.create_program_with_source(ctx, """
+ __kernel void sum(__global float *a)
+ { a[get_global_id(0)] *= 2; }
+ """).build()
+
+ try:
+ prg.sam
+ raise RuntimeError("invalid kernel name did not cause error")
+ except AttributeError:
+ pass
+
+ def test_image_format_constructor(self):
+ iform = cl.ImageFormat(cl.channel_order.RGBA, cl.channel_type.FLOAT)
+
+ assert iform.channel_order == cl.channel_order.RGBA
+ assert iform.channel_data_type == cl.channel_type.FLOAT
+ assert not iform.__dict__
+
+
+
+
+
+
+
+
+if __name__ == "__main__":
+ # make sure that import failures get reported, instead of skipping the tests.
+ import pyopencl
+
+ from py.test.cmdline import main
+ main([__file__])