From e1f7d40ab885dfbf35fde3405ac1931a584bb6a2 Mon Sep 17 00:00:00 2001 From: Andreas Kloeckner Date: Sat, 20 Jul 2013 18:55:06 -0400 Subject: [PATCH] Add int_ptr, from_int_ptr for interoperability --- doc/misc.rst | 20 +++ doc/runtime.rst | 41 ++++- doc/subst.rst | 1 - pyopencl/__init__.py | 296 ++++++++++++++++++--------------- src/wrapper/wrap_cl.hpp | 57 +------ src/wrapper/wrap_cl_part_1.cpp | 41 +++-- src/wrapper/wrap_cl_part_2.cpp | 13 +- src/wrapper/wrap_helpers.hpp | 28 +++- test/test_wrapper.py | 58 ++++++- 9 files changed, 338 insertions(+), 217 deletions(-) diff --git a/doc/misc.rst b/doc/misc.rst index 5af41afb..d49969b3 100644 --- a/doc/misc.rst +++ b/doc/misc.rst @@ -69,6 +69,25 @@ C interface to Python: to guess which one should get a method for the operation. Instead, simply leave that command to be a function. +.. _interoperability: + +Interoperability with other OpenCL software +------------------------------------------- + +Just about every object in :mod:`pyopncl` supports the following +interface (here shown as an example for :class:`pyopencl.MemoryObject`, +from which :class:`pyopencl.Buffer` and :class:`pyopencl.Image` inherit): + +* :meth:`pyopencl.MemoryObject.from_int_ptr` +* :attr:`pyopencl.MemoryObject.int_ptr` + +This allows retrieving the C-level pointer to an OpenCL object as a Python +integer, which may then be passed to other C libraries whose interfaces expose +OpenCL objects. It also allows turning C-level OpenCL objects obtained from +other software to be turned into the corresponding :mod:`pyopencl` objects. + +.. versionadded:: 2013.2 + User-visible Changes ==================== @@ -84,6 +103,7 @@ Version 2013.2 :func:`pyopencl.enqueue_map_image`. * :class:`pyopencl.ImageFormat` was made comparable and hashable. * :mod:`pyopencl.reduction` supports slicing (contributed by Alex Nitz) +* Added :ref:`interoperability` * Bug fixes Version 2013.1 diff --git a/doc/runtime.rst b/doc/runtime.rst index 5eff4034..d6999f0d 100644 --- a/doc/runtime.rst +++ b/doc/runtime.rst @@ -78,6 +78,9 @@ Platforms, Devices and Contexts Return a list of devices matching *device_type*. See :class:`device_type` for values of *device_type*. + .. automethod:: from_int_ptr + .. autoattribute:: int_ptr + |comparable| .. class:: Device @@ -92,6 +95,9 @@ Platforms, Devices and Contexts See :class:`device_info` for values of *param*. + .. automethod:: from_int_ptr + .. autoattribute:: int_ptr + Two instances of this class may be compared using *=="* and *"!="*. .. class:: Context(devices=None, properties=None, dev_type=None) @@ -173,6 +179,9 @@ Platforms, Devices and Contexts .. versionadded:: 2011.1 + .. automethod:: from_int_ptr + .. autoattribute:: int_ptr + |comparable| .. function:: create_some_context(interactive=True) @@ -226,6 +235,9 @@ Command Queues and Events .. method:: flush() .. method:: finish() + .. automethod:: from_int_ptr + .. autoattribute:: int_ptr + |comparable| .. class:: Event @@ -257,6 +269,9 @@ Command Queues and Events .. method:: wait() + .. automethod:: from_int_ptr + .. autoattribute:: int_ptr + |comparable| .. function:: wait_for_events(events) @@ -331,6 +346,9 @@ Memory area as a :class:`numpy.ndarray` of the given *shape*, *dtype* and *order*. + .. automethod:: from_int_ptr + .. autoattribute:: int_ptr + |comparable| .. function:: enqueue_migrate_mem_objects(queue, mem_objects, flags=0, wait_for=None) @@ -360,7 +378,7 @@ Buffers If *hostbuf* is specified, *size* defaults to the size of the specified buffer if it is passed as zero. - :class:`Buffer` is a subclass of :class:`MemoryObject`. + :class:`Buffer` inherits from :class:`MemoryObject`. Note that actual memory allocation in OpenCL may be deferred. Buffers are attached to a :class:`Context` and are only @@ -453,7 +471,7 @@ Images If *hostbuf* is given and *shape* is `None`, then *hostbuf.shape* is used as the *shape* parameter. - :class:`Image` is a subclass of :class:`MemoryObject`. + :class:`Image` inherits from :class:`MemoryObject`. .. note:: @@ -602,6 +620,9 @@ Samplers See :class:`sampler_info` for values of *param*. + .. automethod:: from_int_ptr + .. autoattribute:: int_ptr + |comparable| Programs and Kernels @@ -664,12 +685,15 @@ Programs and Kernels in the same name space and take precedence over :class:`Kernel` names. - |comparable| - .. method:: all_kernels() Returns a list of all :class:`Kernel` objects in the :class:`Program`. + .. automethod:: from_int_ptr + .. autoattribute:: int_ptr + + |comparable| + .. function:: create_program_with_built_in_kernels(context, devices, kernel_names) Only available with CL 1.2. @@ -830,6 +854,9 @@ Programs and Kernels .. versionadded:: 2013.1 + .. automethod:: from_int_ptr + .. autoattribute:: int_ptr + |comparable| .. class:: LocalMemory(size) @@ -889,20 +916,20 @@ with GL support. See :func:`have_gl`. .. class:: GLBuffer(context, flags, bufobj) - :class:`GLBuffer` is a subclass of :class:`MemoryObject`. + :class:`GLBuffer` inherits from :class:`MemoryObject`. .. attribute:: gl_object .. class:: GLRenderBuffer(context, flags, bufobj) - :class:`GLRenderBuffer` is a subclass of :class:`MemoryObject`. + :class:`GLRenderBuffer` inherits from :class:`MemoryObject`. .. attribute:: gl_object .. class:: GLTexture(context, flags, texture_target, miplevel, texture, dims) *dims* is either 2 or 3. - :class:`GLTexture` is a subclass of :class:`Image`. + :class:`GLTexture` inherits from :class:`Image`. .. attribute:: gl_object diff --git a/doc/subst.rst b/doc/subst.rst index 2d7393d8..cb2b9695 100644 --- a/doc/subst.rst +++ b/doc/subst.rst @@ -11,4 +11,3 @@ .. |copy-depr| replace:: **Note:** This function is deprecated as of PyOpenCL 2011.1. Use :func:`enqueue_copy` instead. - diff --git a/pyopencl/__init__.py b/pyopencl/__init__.py index 57566ca7..85791809 100644 --- a/pyopencl/__init__.py +++ b/pyopencl/__init__.py @@ -61,25 +61,6 @@ def compiler_output(text): "to see more.", CompilerWarning) -# {{{ Kernel - -class Kernel(_cl._Kernel): - def __init__(self, prg, name): - if not isinstance(prg, _cl._Program): - prg = prg._get_prg() - - _cl._Kernel.__init__(self, prg, name) - self._source = getattr(prg, "_source", None) - - def capture_call(self, filename, queue, global_size, local_size, - *args, **kwargs): - from pyopencl.capture_call import capture_kernel_call - capture_kernel_call(self, filename, queue, global_size, local_size, - *args, **kwargs) - -# }}} - - # {{{ Program (including caching support) class Program(object): @@ -130,6 +111,15 @@ class Program(object): def all_kernels(self): return self._get_prg().all_kernels() + def int_ptr(self): + return self._get_prg().int_ptr + int_ptr = property(int_ptr, doc=_cl._Program.int_ptr.__doc__) + + def from_int_ptr(int_ptr_value): + return Program(_cl._Program.from_int_ptr(int_ptr_value)) + from_int_ptr.__doc__ = _cl._Program.from_int_ptr.__doc__ + from_int_ptr = staticmethod(from_int_ptr) + def __getattr__(self, attr): try: knl = Kernel(self, attr) @@ -221,6 +211,15 @@ class Program(object): options = " ".join(options) return self._prg().compile(options, devices, headers) + def __eq__(self, other): + return self._get_prg() == other._get_prg() + + def __ne__(self, other): + return self._get_prg() == other._get_prg() + + def __hash__(self): + return hash(self._get_prg()) + def create_program_with_built_in_kernels(context, devices, kernel_names): if not isinstance(kernel_names, str): @@ -237,122 +236,6 @@ def link_program(context, programs, options=[], devices=None): # }}} -# {{{ Image - -class Image(_cl._ImageBase): - def __init__(self, context, flags, format, shape=None, pitches=None, - hostbuf=None, is_array=False, buffer=None): - - if shape is None and hostbuf is None: - raise Error("'shape' must be passed if 'hostbuf' is not given") - - if shape is None and hostbuf is not None: - shape = hostbuf.shape - - if hostbuf is not None and not \ - (flags & (mem_flags.USE_HOST_PTR | mem_flags.COPY_HOST_PTR)): - from warnings import warn - warn("'hostbuf' was passed, but no memory flags to make use of it.") - - if hostbuf is None and pitches is not None: - raise Error("'pitches' may only be given if 'hostbuf' is given") - - if context._get_cl_version() >= (1, 2) and get_cl_header_version() >= (1, 2): - if buffer is not None and is_array: - raise ValueError( - "'buffer' and 'is_array' are mutually exclusive") - - if len(shape) == 3: - if buffer is not None: - raise TypeError( - "'buffer' argument is not supported for 3D arrays") - elif is_array: - image_type = mem_object_type.IMAGE2D_ARRAY - else: - image_type = mem_object_type.IMAGE3D - - elif len(shape) == 2: - if buffer is not None: - raise TypeError( - "'buffer' argument is not supported for 2D arrays") - elif is_array: - image_type = mem_object_type.IMAGE1D_ARRAY - else: - image_type = mem_object_type.IMAGE2D - - elif len(shape) == 1: - if buffer is not None: - image_type = mem_object_type.IMAGE1D_BUFFER - elif is_array: - raise TypeError("array of zero-dimensional images not supported") - else: - image_type = mem_object_type.IMAGE1D - - else: - raise ValueError("images cannot have more than three dimensions") - - desc = ImageDescriptor() - - desc.image_type = image_type - desc.shape = shape # also sets desc.array_size - - if pitches is None: - desc.pitches = (0, 0) - else: - desc.pitches = pitches - - desc.num_mip_levels = 0 # per CL 1.2 spec - desc.num_samples = 0 # per CL 1.2 spec - desc.buffer = buffer - - _cl._ImageBase.__init__(self, context, flags, format, desc, hostbuf) - else: - # legacy init for CL 1.1 and older - if is_array: - raise TypeError("'is_array=True' is not supported for CL < 1.2") - #if num_mip_levels is not None: - #raise TypeError( - # "'num_mip_levels' argument is not supported for CL < 1.2") - #if num_samples is not None: - #raise TypeError( - # "'num_samples' argument is not supported for CL < 1.2") - if buffer is not None: - raise TypeError("'buffer' argument is not supported for CL < 1.2") - - _cl._ImageBase.__init__(self, context, flags, format, shape, - pitches, hostbuf) - - class _ImageInfoGetter: - def __init__(self, event): - from warnings import warn - warn("Image.image.attr is deprecated. " - "Use Image.attr directly, instead.") - - self.event = event - - def __getattr__(self, name): - try: - inf_attr = getattr(_cl.image_info, name.upper()) - except AttributeError: - raise AttributeError("%s has no attribute '%s'" - % (type(self), name)) - else: - return self.event.get_image_info(inf_attr) - - image = property(_ImageInfoGetter) - - @property - def shape(self): - if self.type == mem_object_type.IMAGE2D: - return (self.width, self.height) - elif self.type == mem_object_type.IMAGE3D: - return (self.width, self.height, self.depth) - else: - raise LogicError("only images have shapes") - -# }}} - - def _add_functionality(): cls_to_info_cls = { _cl.Platform: @@ -368,7 +251,7 @@ def _add_functionality(): _cl.MemoryObjectHolder: (MemoryObjectHolder.get_info, _cl.mem_info), Image: - (_cl._ImageBase.get_image_info, _cl.image_info), + (_cl.Image.get_image_info, _cl.image_info), Program: (Program.get_info, _cl.program_info), Kernel: @@ -410,17 +293,19 @@ def _add_functionality(): # }}} # {{{ Platform + def platform_repr(self): - return "" % (self.name, self.obj_ptr) + return "" % (self.name, self.int_ptr) Platform.__repr__ = platform_repr # }}} # {{{ Device + def device_repr(self): return "" % ( - self.name.strip(), self.platform.name.strip(), self.obj_ptr) + self.name.strip(), self.platform.name.strip(), self.int_ptr) Device.__repr__ = device_repr @@ -555,6 +440,16 @@ def _add_functionality(): # }}} # {{{ Kernel + + kernel_old_init = Kernel.__init__ + + def kernel_init(self, prg, name): + if not isinstance(prg, _cl._Program): + prg = prg._get_prg() + + kernel_old_init(self, prg, name) + self._source = getattr(prg, "_source", None) + def kernel_call(self, queue, global_size, local_size, *args, **kwargs): global_offset = kwargs.pop("global_offset", None) g_times_l = kwargs.pop("g_times_l", False) @@ -622,9 +517,17 @@ def _add_functionality(): else: raise + def kernel_capture_call(self, filename, queue, global_size, local_size, + *args, **kwargs): + from pyopencl.capture_call import capture_kernel_call + capture_kernel_call(self, filename, queue, global_size, local_size, + *args, **kwargs) + + Kernel.__init__ = kernel_init Kernel.__call__ = kernel_call Kernel.set_scalar_arg_dtypes = kernel_set_scalar_arg_dtypes Kernel.set_args = kernel_set_args + Kernel.capture_call = kernel_capture_call # }}} @@ -654,6 +557,123 @@ def _add_functionality(): # }}} + # {{{ Image + + image_old_init = Image.__init__ + + def image_init(self, context, flags, format, shape=None, pitches=None, + hostbuf=None, is_array=False, buffer=None): + + if shape is None and hostbuf is None: + raise Error("'shape' must be passed if 'hostbuf' is not given") + + if shape is None and hostbuf is not None: + shape = hostbuf.shape + + if hostbuf is not None and not \ + (flags & (mem_flags.USE_HOST_PTR | mem_flags.COPY_HOST_PTR)): + from warnings import warn + warn("'hostbuf' was passed, but no memory flags to make use of it.") + + if hostbuf is None and pitches is not None: + raise Error("'pitches' may only be given if 'hostbuf' is given") + + if context._get_cl_version() >= (1, 2) and get_cl_header_version() >= (1, 2): + if buffer is not None and is_array: + raise ValueError( + "'buffer' and 'is_array' are mutually exclusive") + + if len(shape) == 3: + if buffer is not None: + raise TypeError( + "'buffer' argument is not supported for 3D arrays") + elif is_array: + image_type = mem_object_type.IMAGE2D_ARRAY + else: + image_type = mem_object_type.IMAGE3D + + elif len(shape) == 2: + if buffer is not None: + raise TypeError( + "'buffer' argument is not supported for 2D arrays") + elif is_array: + image_type = mem_object_type.IMAGE1D_ARRAY + else: + image_type = mem_object_type.IMAGE2D + + elif len(shape) == 1: + if buffer is not None: + image_type = mem_object_type.IMAGE1D_BUFFER + elif is_array: + raise TypeError("array of zero-dimensional images not supported") + else: + image_type = mem_object_type.IMAGE1D + + else: + raise ValueError("images cannot have more than three dimensions") + + desc = ImageDescriptor() + + desc.image_type = image_type + desc.shape = shape # also sets desc.array_size + + if pitches is None: + desc.pitches = (0, 0) + else: + desc.pitches = pitches + + desc.num_mip_levels = 0 # per CL 1.2 spec + desc.num_samples = 0 # per CL 1.2 spec + desc.buffer = buffer + + image_old_init(self, context, flags, format, desc, hostbuf) + else: + # legacy init for CL 1.1 and older + if is_array: + raise TypeError("'is_array=True' is not supported for CL < 1.2") + #if num_mip_levels is not None: + #raise TypeError( + # "'num_mip_levels' argument is not supported for CL < 1.2") + #if num_samples is not None: + #raise TypeError( + # "'num_samples' argument is not supported for CL < 1.2") + if buffer is not None: + raise TypeError("'buffer' argument is not supported for CL < 1.2") + + image_old_init(self, context, flags, format, shape, + pitches, hostbuf) + + class _ImageInfoGetter: + def __init__(self, event): + from warnings import warn + warn("Image.image.attr is deprecated. " + "Use Image.attr directly, instead.") + + self.event = event + + def __getattr__(self, name): + try: + inf_attr = getattr(_cl.image_info, name.upper()) + except AttributeError: + raise AttributeError("%s has no attribute '%s'" + % (type(self), name)) + else: + return self.event.get_image_info(inf_attr) + + def image_shape(self): + if self.type == mem_object_type.IMAGE2D: + return (self.width, self.height) + elif self.type == mem_object_type.IMAGE3D: + return (self.width, self.height, self.depth) + else: + raise LogicError("only images have shapes") + + Image.__init__ = image_init + Image.image = property(_ImageInfoGetter) + Image.shape = property(image_shape) + + # }}} + # {{{ Error def error_str(self): diff --git a/src/wrapper/wrap_cl.hpp b/src/wrapper/wrap_cl.hpp index d4057b6a..dbdced1a 100644 --- a/src/wrapper/wrap_cl.hpp +++ b/src/wrapper/wrap_cl.hpp @@ -432,11 +432,6 @@ namespace pyopencl return m_platform; } - npy_intp obj_ptr() const - { - return (npy_intp) data(); - } - PYOPENCL_EQUALITY_TESTS(platform); py::object get_info(cl_platform_info param_name) const @@ -570,11 +565,6 @@ namespace pyopencl return m_device; } - npy_intp obj_ptr() const - { - return (npy_intp) data(); - } - PYOPENCL_EQUALITY_TESTS(device); py::object get_info(cl_device_info param_name) const @@ -902,11 +892,6 @@ namespace pyopencl return m_context; } - npy_intp obj_ptr() const - { - return (npy_intp) data(); - } - PYOPENCL_EQUALITY_TESTS(context); py::object get_info(cl_context_info param_name) const @@ -1186,11 +1171,6 @@ namespace pyopencl const cl_command_queue data() const { return m_queue; } - npy_intp obj_ptr() const - { - return (npy_intp) data(); - } - PYOPENCL_EQUALITY_TESTS(command_queue); py::object get_info(cl_command_queue_info param_name) const @@ -1271,11 +1251,6 @@ namespace pyopencl const cl_event data() const { return m_event; } - npy_intp obj_ptr() const - { - return (npy_intp) data(); - } - PYOPENCL_EQUALITY_TESTS(event); py::object get_info(cl_event_info param_name) const @@ -1560,11 +1535,6 @@ namespace pyopencl const cl_mem data() const { return m_mem; } - npy_intp obj_ptr() const - { - return (npy_intp) data(); - } - }; #if PYOPENCL_CL_VERSION >= 0x1020 @@ -2825,6 +2795,13 @@ namespace pyopencl throw pyopencl::error("Sampler", status_code); } + sampler(cl_sampler samp, bool retain) + : m_sampler(samp) + { + if (retain) + PYOPENCL_CALL_GUARDED(clRetainSampler, (samp)); + } + ~sampler() { PYOPENCL_CALL_GUARDED_CLEANUP(clReleaseSampler, (m_sampler)); @@ -2835,11 +2812,6 @@ namespace pyopencl return m_sampler; } - npy_intp obj_ptr() const - { - return (npy_intp) data(); - } - PYOPENCL_EQUALITY_TESTS(sampler); py::object get_info(cl_sampler_info param_name) const @@ -2904,11 +2876,6 @@ namespace pyopencl return m_program_kind; } - npy_intp obj_ptr() const - { - return (npy_intp) data(); - } - PYOPENCL_EQUALITY_TESTS(program); py::object get_info(cl_program_info param_name) const @@ -3309,11 +3276,6 @@ namespace pyopencl return m_kernel; } - npy_intp obj_ptr() const - { - return (npy_intp) data(); - } - PYOPENCL_EQUALITY_TESTS(kernel); void set_arg_null(cl_uint arg_index) @@ -3883,11 +3845,6 @@ namespace pyopencl // {{{ deferred implementation bits - inline py::object create_event_wrapper_from_int(intptr_t cl_event_as_int) - { - return py::object(handle_from_new_ptr(new event((cl_event)cl_event_as_int, true))); - } - inline py::object create_mem_object_wrapper(cl_mem mem) { cl_mem_object_type mem_obj_type; diff --git a/src/wrapper/wrap_cl_part_1.cpp b/src/wrapper/wrap_cl_part_1.cpp index dba0dec7..f3448aca 100644 --- a/src/wrapper/wrap_cl_part_1.cpp +++ b/src/wrapper/wrap_cl_part_1.cpp @@ -10,6 +10,9 @@ using namespace pyopencl; void pyopencl_expose_part_1() { + py::docstring_options doc_op; + doc_op.disable_cpp_signatures(); + py::def("get_cl_header_version", get_cl_header_version); // {{{ platform @@ -21,10 +24,10 @@ void pyopencl_expose_part_1() .DEF_SIMPLE_METHOD(get_info) .def("get_devices", &cls::get_devices, py::arg("device_type")=CL_DEVICE_TYPE_ALL) - .add_property("obj_ptr", &cls::obj_ptr) .def(py::self == py::self) .def(py::self != py::self) .def("__hash__", &cls::hash) + PYOPENCL_EXPOSE_TO_FROM_INT_PTR(cl_platform_id) ; } @@ -35,7 +38,6 @@ void pyopencl_expose_part_1() typedef device cls; py::class_("Device", py::no_init) .DEF_SIMPLE_METHOD(get_info) - .add_property("obj_ptr", &cls::obj_ptr) .def(py::self == py::self) .def(py::self != py::self) .def("__hash__", &cls::hash) @@ -45,6 +47,7 @@ void pyopencl_expose_part_1() #if PYOPENCL_CL_VERSION >= 0x1020 .DEF_SIMPLE_METHOD(create_sub_devices) #endif + PYOPENCL_EXPOSE_TO_FROM_INT_PTR(cl_device_id) ; } @@ -54,7 +57,7 @@ void pyopencl_expose_part_1() { typedef context cls; - py::class_ >("Context", py::no_init) .def("__init__", make_constructor(create_context, py::default_call_policies(), @@ -63,10 +66,10 @@ void pyopencl_expose_part_1() py::arg("dev_type")=py::object() ))) .DEF_SIMPLE_METHOD(get_info) - .add_property("obj_ptr", &cls::obj_ptr) .def(py::self == py::self) .def(py::self != py::self) .def("__hash__", &cls::hash) + PYOPENCL_EXPOSE_TO_FROM_INT_PTR(cl_context) ; } @@ -85,10 +88,10 @@ void pyopencl_expose_part_1() #endif .DEF_SIMPLE_METHOD(flush) .DEF_SIMPLE_METHOD(finish) - .add_property("obj_ptr", &cls::obj_ptr) .def(py::self == py::self) .def(py::self != py::self) .def("__hash__", &cls::hash) + PYOPENCL_EXPOSE_TO_FROM_INT_PTR(cl_command_queue) ; } @@ -101,12 +104,14 @@ void pyopencl_expose_part_1() .DEF_SIMPLE_METHOD(get_info) .DEF_SIMPLE_METHOD(get_profiling_info) .DEF_SIMPLE_METHOD(wait) - .add_property("obj_ptr", &cls::obj_ptr) .def(py::self == py::self) .def(py::self != py::self) .def("__hash__", &cls::hash) - .def("from_cl_event_as_int", create_event_wrapper_from_int, - py::args("cl_event_as_int")) + PYOPENCL_EXPOSE_TO_FROM_INT_PTR(cl_event) + + // deprecated, remove in 2015.x. + .def("from_cl_event_as_int", from_int_ptr, + py::return_value_policy()) .staticmethod("from_cl_event_as_int") ; } @@ -162,6 +167,12 @@ void pyopencl_expose_part_1() .def(py::self == py::self) .def(py::self != py::self) .def("__hash__", &cls::hash) + + .add_property("int_ptr", to_int_ptr, + "Return an integer corresponding to the pointer value " + "of the underlying :c:type:`cl_mem`. " + "Use :meth:`from_int_ptr` to turn back into a Python object." + "\n\n.. versionadded:: 2013.2\n") ; } { @@ -169,10 +180,18 @@ void pyopencl_expose_part_1() py::class_ >( "MemoryObject", py::no_init) .DEF_SIMPLE_METHOD(release) - .add_property("obj_ptr", &cls::obj_ptr) .add_property("hostbuf", &cls::hostbuf) - .def("from_cl_mem_as_int", memory_object_from_int, - py::args("cl_mem_as_int")) + + .def("from_int_ptr", memory_object_from_int, + "(static method) Return a new Python object referencing the C-level " \ + ":c:type:`cl_mem` object at the location pointed to " \ + "by *int_ptr_value*. The relevant :c:func:`clRetain*` function " \ + "will be called." \ + "\n\n.. versionadded:: 2013.2\n") \ + .staticmethod("from_int_ptr") + + // deprecated, remove in 2015.x + .def("from_cl_mem_as_int", memory_object_from_int) .staticmethod("from_cl_mem_as_int") ; } diff --git a/src/wrapper/wrap_cl_part_2.cpp b/src/wrapper/wrap_cl_part_2.cpp index 87b60800..4d010796 100644 --- a/src/wrapper/wrap_cl_part_2.cpp +++ b/src/wrapper/wrap_cl_part_2.cpp @@ -47,6 +47,9 @@ using namespace pyopencl; void pyopencl_expose_part_2() { + py::docstring_options doc_op; + doc_op.disable_cpp_signatures(); + // {{{ image #if PYOPENCL_CL_VERSION >= 0x1020 @@ -67,7 +70,7 @@ void pyopencl_expose_part_2() { typedef image cls; py::class_, boost::noncopyable>( - "_ImageBase", py::no_init) + "Image", py::no_init) .def("__init__", make_constructor(create_image, py::default_call_policies(), (py::args("context", "flags", "format"), @@ -173,10 +176,10 @@ void pyopencl_expose_part_2() py::class_("Sampler", py::init()) .DEF_SIMPLE_METHOD(get_info) - .add_property("obj_ptr", &cls::obj_ptr) .def(py::self == py::self) .def(py::self != py::self) .def("__hash__", &cls::hash) + PYOPENCL_EXPOSE_TO_FROM_INT_PTR(cl_sampler) ; } @@ -225,11 +228,11 @@ void pyopencl_expose_part_2() py::return_value_policy()) .staticmethod("link") #endif - .add_property("obj_ptr", &cls::obj_ptr) .def(py::self == py::self) .def(py::self != py::self) .def("__hash__", &cls::hash) .def("all_kernels", create_kernels_in_program) + PYOPENCL_EXPOSE_TO_FROM_INT_PTR(cl_program) ; } @@ -243,7 +246,7 @@ void pyopencl_expose_part_2() { typedef kernel cls; - py::class_("_Kernel", + py::class_("Kernel", py::init()) .DEF_SIMPLE_METHOD(get_info) .DEF_SIMPLE_METHOD(get_work_group_info) @@ -251,10 +254,10 @@ void pyopencl_expose_part_2() #if PYOPENCL_CL_VERSION >= 0x1020 .DEF_SIMPLE_METHOD(get_arg_info) #endif - .add_property("obj_ptr", &cls::obj_ptr) .def(py::self == py::self) .def(py::self != py::self) .def("__hash__", &cls::hash) + PYOPENCL_EXPOSE_TO_FROM_INT_PTR(cl_kernel) ; } diff --git a/src/wrapper/wrap_helpers.hpp b/src/wrapper/wrap_helpers.hpp index 69e666f4..dac179c7 100644 --- a/src/wrapper/wrap_helpers.hpp +++ b/src/wrapper/wrap_helpers.hpp @@ -141,9 +141,35 @@ namespace return boost::python::handle<>( typename boost::python::manage_new_object::apply::type()(ptr)); } -} + template + inline T *from_int_ptr(intptr_t obj_ref) + { + ClType clobj = (ClType) obj_ref; + return new T(clobj, /* retain */ true); + } + template + inline intptr_t to_int_ptr(T const &obj) + { + return (intptr_t) obj.data(); + } +} +#define PYOPENCL_EXPOSE_TO_FROM_INT_PTR(CL_TYPENAME) \ + .def("from_int_ptr", from_int_ptr, \ + py::return_value_policy(), \ + py::arg("int_ptr_value"), \ + "(static method) Return a new Python object referencing the C-level " \ + ":c:type:`" #CL_TYPENAME "` object at the location pointed to " \ + "by *int_ptr_value*. The relevant :c:func:`clRetain*` function " \ + "will be called." \ + "\n\n.. versionadded:: 2013.2\n") \ + .staticmethod("from_int_ptr") \ + .add_property("int_ptr", to_int_ptr, \ + "Return an integer corresponding to the pointer value " \ + "of the underlying :c:type:`" #CL_TYPENAME "`. " \ + "Use :meth:`from_int_ptr` to turn back into a Python object." \ + "\n\n.. versionadded:: 2013.2\n") \ #endif diff --git a/test/test_wrapper.py b/test/test_wrapper.py index d92eb844..7b5a65e2 100644 --- a/test/test_wrapper.py +++ b/test/test_wrapper.py @@ -42,7 +42,11 @@ else: @pytools.test.mark_test.opencl -def test_get_info(platform, device): +def test_get_info(ctx_factory): + ctx = ctx_factory() + device, = ctx.devices + platform = device.platform + failure_count = [0] pocl_quirks = [ @@ -125,10 +129,7 @@ def test_get_info(platform, device): failure_count[0] += 1 do_test(platform, cl.platform_info) - do_test(device, cl.device_info) - - ctx = cl.Context([device]) do_test(ctx, cl.context_info) props = 0 @@ -188,6 +189,55 @@ def test_get_info(platform, device): lambda info: img.get_image_info(info)) +@pytools.test.mark_test.opencl +def test_int_ptr(ctx_factory): + def do_test(obj): + new_obj = type(obj).from_int_ptr(obj.int_ptr) + assert obj == new_obj + assert type(obj) is type(new_obj) + + ctx = ctx_factory() + device, = ctx.devices + platform = device.platform + do_test(device) + do_test(platform) + do_test(ctx) + + queue = cl.CommandQueue(ctx) + do_test(queue) + + evt = cl.enqueue_marker(queue) + do_test(evt) + + prg = cl.Program(ctx, """ + __kernel void sum(__global float *a) + { a[get_global_id(0)] *= 2; } + """).build() + + do_test(prg) + do_test(prg.sum) + + n = 2000 + a_buf = cl.Buffer(ctx, 0, n*4) + do_test(a_buf) + + # crashes on intel... + if device.image_support and platform.vendor not in [ + "Intel(R) Corporation", + "The pocl project", + ]: + smp = cl.Sampler(ctx, False, + cl.addressing_mode.CLAMP, + cl.filter_mode.NEAREST) + do_test(smp) + + img_format = cl.get_supported_image_formats( + ctx, cl.mem_flags.READ_ONLY, cl.mem_object_type.IMAGE2D)[0] + + img = cl.Image(ctx, cl.mem_flags.READ_ONLY, img_format, (128, 256)) + do_test(img) + + @pytools.test.mark_test.opencl def test_invalid_kernel_names_cause_failures(ctx_factory): ctx = ctx_factory() -- GitLab