diff --git a/doc/source/index.rst b/doc/source/index.rst index 9bd36e1e873a34f4fa1cd93286afc63a326c1b32..a695b9dddb06f1c7012613e346bf9cedf3e5c6a6 100644 --- a/doc/source/index.rst +++ b/doc/source/index.rst @@ -40,13 +40,11 @@ Here's an example, to give you an impression:: queue = cl.CommandQueue(ctx) mf = cl.mem_flags - a_buf = cl.create_host_buffer( - ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, a) - b_buf = cl.create_host_buffer( - ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, b) - dest_buf = cl.create_buffer(ctx, mf.WRITE_ONLY, b.nbytes) + a_buf = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=a) + b_buf = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=b) + dest_buf = cl.Buffer(ctx, mf.WRITE_ONLY, b.nbytes) - prg = cl.create_program_with_source(ctx, """ + prg = cl.Program(ctx, """ __kernel void sum(__global const float *a, __global const float *b, __global float *c) { diff --git a/doc/source/misc.rst b/doc/source/misc.rst index 1343d8222d9c689f3e6cd2e98884636e69bf89aa..138b3923de62d619abf0f76da4e990dea5237f23 100644 --- a/doc/source/misc.rst +++ b/doc/source/misc.rst @@ -9,6 +9,52 @@ Acknowledgments * James Snyder provided patches to make PyOpenCL work on OS X 10.6. * Roger Pau Monné supplied the example :file:`examples/benchmark-all.py`. +* David Garcia contributed significantly to PyOpenCL's API design + and reported many bugs. + +Guidelines +========== + +.. _api-compatibility: + +API Stability +------------- + +I consider PyOpenCL's API "stable". That doesn't mean it can't +change. But if it does, your code will generally continue to run. It +may however start spewing warnings about things you need to change to +stay compatible with future versions. + +Deprecation warnings will be around for a whole release cycle, as +identified by the second number in the release name. (the "90" in +"0.90") Further, the stability promise applies for any code that's +part of a released version. It doesn't apply to undocumented bits of +the API, and it doesn't apply to unreleased code downloaded from git. + +.. _versus-c: + +Relation with OpenCL's C Bindings +--------------------------------- + +We've tried to follow these guidelines when binding the OpenCL's +C interface to Python: + +* Remove the `cl_`, `CL_` and `cl` prefix from data types, macros and + function names. +* Follow :pep:`8`, i.e. + + * Make function names lowercase. + * If a data type or function name is composed of more than one word, + separate the words with a single underscore. + +* `get_info` functions become attributes. +* Object creation is done by constructors, to the extent possible. + (i.e. minimize use of "factory functions") + +* If an operation involves two or more "complex" objects (like e.g. a + kernel enqueue involves a kernel and a queue), refuse the temptation + to guess which one should get a method for the operation. + Instead, simply leave that command to be a function. User-visible Changes ==================== @@ -28,10 +74,23 @@ Version 0.91 :func:`pyopencl.enqueue_write_image` are now defaulted to zero. The argument order of `enqueue_{read,write}_image` has changed for this reason. +* Deprecate + :func:`pyopencl.create_image_2d`, + :func:`pyopencl.create_image_3d` + in favor of the :class:`pyopencl.Image` constructor. +* Deprecate + :func:`pyopencl.create_program_with_source`, + :func:`pyopencl.create_program_with_binary` + in favor of the :class:`pyopencl.Program` constructor. +* Deprecate + :func:`pyopencl.create_buffer`, + :func:`pyopencl.create_host_buffer` + in favor of the :class:`pyopencl.Buffer` constructor. * :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 -------------- diff --git a/doc/source/reference.rst b/doc/source/reference.rst index bb1976cdc153a75860d2ff4dd0a928355cd9c5bc..13c92a63460fcdee8f62aa707366885ec6e18b89 100644 --- a/doc/source/reference.rst +++ b/doc/source/reference.rst @@ -51,9 +51,10 @@ 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`. *wait_for* +.. |explain-waitfor| replace:: *wait_for* may either be *None* or a list of :class:`Event` instances for whose completion this command waits before starting exeuction. +.. |std-enqueue-blurb| replace:: Returns a new :class:`Event`. |explain-waitfor| .. function:: get_platforms() @@ -61,16 +62,16 @@ Platforms, Devices and Contexts .. class:: Platform - .. method:: get_info(param) - - See :class:`platform_info` for values of *param*. - .. attribute:: info Lower case versions of the :class:`platform_info` constants may be used as attributes on instances of this class to directly query info attributes. + .. method:: get_info(param) + + See :class:`platform_info` for values of *param*. + .. method:: get_devices(device_type=device_type.ALL) Return a list of devices matching *device_type*. @@ -80,16 +81,16 @@ Platforms, Devices and Contexts .. class:: Device - .. method:: get_info(param) - - See :class:`device_info` for values of *param*. - .. attribute:: info Lower case versions of the :class:`device_info` constants may be used as attributes on instances of this class to directly query info attributes. + .. method:: get_info(param) + + See :class:`device_info` for values of *param*. + Two instances of this class may be compared using *=="* and *"!="*. .. class:: Context(devices, properties=None) @@ -97,16 +98,16 @@ Platforms, Devices and Contexts Create a new context. *properties* is a list of key-value tuples, where each key must be one of :class:`context_properties`. - .. method:: get_info(param) - - See :class:`context_info` for values of *param*. - .. attribute:: info Lower case versions of the :class:`context_info` constants may be used as attributes on instances of this class to directly query info attributes. + .. method:: get_info(param) + + See :class:`context_info` for values of *param*. + |comparable| .. function:: create_context_from_type(dev_type, properties=[]) @@ -122,19 +123,20 @@ Command Queues and Events if *device* is None, one of the devices in *context* is chosen in an implementation-defined manner. - .. method:: get_info(param) - - See :class:`command_queue_info` for values of *param*. - .. attribute:: info Lower case versions of the :class:`command_queue_info` constants may be used as attributes on instances of this class to directly query info attributes. + .. method:: get_info(param) + + See :class:`command_queue_info` for values of *param*. + .. method:: set_property(prop, enable) See :class:`command_queue_properties` for possible values of *prop*. + *enable* is a :class:`bool`. .. method:: flush() .. method:: finish() @@ -143,9 +145,11 @@ Command Queues and Events .. class:: Event - .. method:: get_info(param) + .. attribute:: info - See :class:`event_info` for values of *param*. + Lower case versions of the :class:`event_info` constants + may be used as attributes on instances of this class + to directly query info attributes. .. attribute:: profile.info @@ -156,11 +160,9 @@ Command Queues and Events For example, you may use *evt.profile.end* instead of *evt.get_profiling_info(pyopencl.profiling_info.END)*. - .. attribute:: info + .. method:: get_info(param) - Lower case versions of the :class:`event_info` constants - may be used as attributes on instances of this class - to directly query info attributes. + See :class:`event_info` for values of *param*. .. method:: get_profiling_info(param) @@ -186,59 +188,31 @@ Memory .. class:: MemoryObject - .. method:: get_info(param) - - 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. - .. attribute:: image.info - - 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. + .. method:: get_info(param) - For example, you may use *img.image.depth* instead of - *img.get_image_info(pyopencl.image_info.DEPTH)*. + See :class:`mem_info` for values of *param*. .. 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 ^^^^^^^ -.. function:: create_buffer(context, flags, size) +.. class:: Buffer(context, flags, size=0, hostbuf=None) + Create a :class:`Buffer`. See :class:`mem_flags` for values of *flags*. - Returns a new buffer-type :class:`MemoryObject`. + If *hostbuf* is specified, *size* defaults to the size of + the specified buffer if it is passed as zero. -.. function:: create_host_buffer(context, flags, buffer) - - Create a buffer :class:`MemoryObject` using host memory. *host_buffer* |buf-iface|. - - See :class:`mem_flags` for values of *flags*. + :class:`Buffer` is a subclass of :class:`MemoryObject`. .. function:: enqueue_read_buffer(queue, mem, host_buffer, device_offset=0, wait_for=None, is_blocking=False) @@ -277,48 +251,52 @@ Image Formats Images ^^^^^^ +.. class:: Image(context, flags, format, shape, pitches=None, host_buffer=None) -.. function:: create_image_2d(context, flags, format, width, height, pitch=0, host_buffer=None) + :class:`Image` is a subclass of :class:`MemoryObject`. - See :class:`mem_flags` for possible values of *flags*. - Returns a new image-type :class:`MemoryObject`. + .. versionadded:: 0.91 - .. versionchanged:: 0.91 - *pitch* argument defaults to zero, moved. + .. attribute:: info -.. function:: create_image_3d(context, flags, format, width, height, depth, row_pitch=0, slice_pitch=0, host_buffer=None) + Lower case versions of the :class:`mem_info` + and :class:`image_info` constants + may be used as attributes on instances of this class + to directly query info attributes. - See :class:`mem_flags` for possible values of *flags*. - Returns a new image-type :class:`MemoryObject`. + .. method:: get_image_info(param) - .. versionchanged:: 0.91 - *pitch* arguments defaults to zero, moved. + See :class:`image_info` for values of *param*. + + .. method:: release() + + |comparable| .. function:: enqueue_read_image(queue, mem, origin, region, host_buffer, row_pitch=0, slice_pitch=0, wait_for=None, is_blocking=False) - |enqueue-waitfor| + |std-enqueue-blurb| .. 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| + |std-enqueue-blurb| .. 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| + |std-enqueue-blurb| .. function:: enqueue_copy_image_to_buffer(queue, src, dest, origin, region, offset, wait_for=None) - |enqueue-waitfor| + |std-enqueue-blurb| .. function:: enqueue_copy_buffer_to_image(queue, src, dest, offset, origin, region, wait_for=None) - |enqueue-waitfor| + |std-enqueue-blurb| Mapping Memory into Host Address Space ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ @@ -329,46 +307,59 @@ Mapping Memory into Host Address Space .. function:: enqueue_map_buffer(queue, buf, flags, offset, shape, dtype, order, wait_for=None, is_blocking=False) - |enqueue-waitfor| + |explain-waitfor| *shape*, *dtype*, and *order* have the same meaning as in :func:`numpy.empty`. See :class:`map_flags` for possible values of *flags*. + :return: a tuple *(array, event)*. *array* is a + :class:`numpy.ndarray` representing the host side + of the map. Its *.base* member contains a + :class:`MemoryMap`. + .. function:: enqueue_map_image(queue, buf, flags, origin, region, shape, dtype, order, wait_for=None, is_blocking=False) - |enqueue-waitfor| + |explain-waitfor| *shape*, *dtype*, and *order* have the same meaning as in :func:`numpy.empty`. See :class:`map_flags` for possible values of *flags*. + :return: a tuple *(array, event)*. *array* is a + :class:`numpy.ndarray` representing the host side + of the map. Its *.base* member contains a + :class:`MemoryMap`. + + Samplers ^^^^^^^^ .. class:: Sampler(context, normalized_coords, addressing_mode, filter_mode) + *normalized_coords* is a :class:`bool` indicating whether + to use coordinates between 0 and 1 (*True*) or the texture's + natural pixel size (*False*). See :class:`addressing_mode` and :class:`filter_mode` for possible argument values. - .. method:: get_info(param) - - See :class:`sampler_info` for values of *param*. - .. attribute:: info Lower case versions of the :class:`sampler_info` constants may be used as attributes on instances of this class to directly query info attributes. + .. method:: get_info(param) + + See :class:`sampler_info` for values of *param*. + |comparable| Programs and Kernels -------------------- -.. class:: Program +.. class:: Program(context, src) + Program(context, devices, binaries) - .. method:: get_info(param) - - See :class:`program_info` for values of *param*. + *binaries* must contain one binary for each entry in *devices*. .. attribute:: info @@ -376,6 +367,10 @@ Programs and Kernels may be used as attributes on instances of this class to directly query info attributes. + .. method:: get_info(param) + + See :class:`program_info` for values of *param*. + .. method:: get_build_info(device, param) See :class:`program_build_info` for values of *param*. @@ -403,23 +398,18 @@ Programs and 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) - - *binaries* must contain one binary for each entry in *devices*. .. class:: Kernel(program, name) - .. method:: get_info(param) - - See :class:`kernel_info` for values of *param*. - .. attribute:: info Lower case versions of the :class:`kernel_info` constants may be used as attributes on instances of this class to directly query info attributes. + .. method:: get_info(param) + + See :class:`kernel_info` for values of *param*. .. method:: get_work_group_info(param, device) @@ -427,18 +417,18 @@ Programs and Kernels .. method:: __call__(queue, global_size, *args, global_offset=None, local_size=None, wait_for=None) - |enqueue-waitfor| + |std-enqueue-blurb| |comparable| .. function:: enqueue_nd_range_kernel(queue, kernel, global_work_size, local_work_size, global_work_offset=None, wait_for=None) - |enqueue-waitfor| + |std-enqueue-blurb| .. function:: enqueue_task(queue, kernel, wait_for=None) - |enqueue-waitfor| + |std-enqueue-blurb| .. _gl-interop: @@ -454,36 +444,34 @@ with GL support. See :func:`have_gl`. Return *True* if PyOpenCL was compiled with OpenGL interoperability, otherwise *False*. -.. function:: create_from_gl_buffer(context, mem_flags, gl_buffer_obj) +.. class:: GLBuffer(context, flags, bufobj) - See :class:`mem_flags` for values of *flags*. - Returns a new :class:`MemoryObject`. + :class:`GLBuffer` is a subclass of :class:`MemoryObject`. -.. function:: create_from_gl_texture_2d(context, mem_flags, texture_target, miplevel, texture) + .. attribute:: gl_object - See :class:`mem_flags` for values of *flags*. - Returns a new :class:`MemoryObject`. +.. class:: GLRenderBuffer(context, flags, bufobj) -.. function:: create_from_gl_texture_3d(context, mem_flags, texture_target, miplevel, texture) + :class:`GLRenderBuffer` is a subclass of :class:`MemoryObject`. - See :class:`mem_flags` for values of *flags*. - Returns a new :class:`MemoryObject`. + .. attribute:: gl_object -.. function:: create_from_gl_renderbuffer(context, mem_flags, gl_renderbuffer) +.. class:: GLTexture(context, flags, texture_target, miplevel, texture, dims) - See :class:`mem_flags` for values of *flags*. - Returns a new :class:`MemoryObject`. + *dims* is either 2 or 3. + :class:`GLTexture` is a subclass of :class:`Image`. -.. function:: enqueue_acquire_gl_objects(queue, mem_objects, wait_for=None) + .. attribute:: gl_object - *mem_objects* is a list of :class:`MemoryObject` instances. - |enqueue-waitfor| + .. method:: get_gl_texture_info(param) -.. function:: enqueue_release_gl_objects(queue, mem_objects, wait_for=None) + See :class:`gl_texture_info` for values of *param*. Only available when PyOpenCL is compiled with GL support. See :func:`have_gl`. - *mem_objects* is a list of :class:`MemoryObject` instances. |enqueue-waitfor| +.. function:: enqueue_acquire_gl_objects(queue, mem_objects, wait_for=None) -.. seealso:: + *mem_objects* is a list of :class:`MemoryObject` instances. + |std-enqueue-blurb| - :meth:`MemoryObject.get_gl_object_info`, :meth:`MemoryObject.get_gl_texture_info`. +.. function:: enqueue_release_gl_objects(queue, mem_objects, wait_for=None) + *mem_objects* is a list of :class:`MemoryObject` instances. |std-enqueue-blurb| diff --git a/examples/benchmark-all.py b/examples/benchmark-all.py index 93685006dc8817307c12025d5813f9c5485f1e6d..09fc407e2007e9f37823f7d51783215947426a8d 100644 --- a/examples/benchmark-all.py +++ b/examples/benchmark-all.py @@ -41,13 +41,11 @@ for platform in cl.get_platforms(): properties=cl.command_queue_properties.PROFILING_ENABLE) mf = cl.mem_flags - a_buf = cl.create_host_buffer( - ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, a) - b_buf = cl.create_host_buffer( - ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, b) - dest_buf = cl.create_buffer(ctx, mf.WRITE_ONLY, b.nbytes) + a_buf = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=a) + b_buf = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=b) + dest_buf = cl.Buffer(ctx, mf.WRITE_ONLY, b.nbytes) - prg = cl.create_program_with_source(ctx, """ + prg = cl.Program(ctx, """ __kernel void sum(__global const float *a, __global const float *b, __global float *c) { diff --git a/examples/demo.py b/examples/demo.py index 0fb13624787b26619b9254b9acf919e1c16cbcb9..863588d3684ac7085ae8151b17a53aed8c2482ea 100644 --- a/examples/demo.py +++ b/examples/demo.py @@ -9,13 +9,11 @@ ctx = cl.create_context_from_type(cl.device_type.ALL) queue = cl.CommandQueue(ctx) mf = cl.mem_flags -a_buf = cl.create_host_buffer( - ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, a) -b_buf = cl.create_host_buffer( - ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, b) -dest_buf = cl.create_buffer(ctx, mf.WRITE_ONLY, b.nbytes) +a_buf = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=a) +b_buf = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=b) +dest_buf = cl.Buffer(ctx, mf.WRITE_ONLY, b.nbytes) -prg = cl.create_program_with_source(ctx, """ +prg = cl.Program(ctx, """ __kernel void sum(__global const float *a, __global const float *b, __global float *c) { diff --git a/examples/narray.py b/examples/narray.py index c90f6cabf130376792e6f05583240a83f3476aa4..3f333fe5d23ef133764920bf707693408a830ce4 100644 --- a/examples/narray.py +++ b/examples/narray.py @@ -7,10 +7,9 @@ ctx = cl.create_context_from_type(cl.device_type.GPU) queue = cl.CommandQueue(ctx) mf = cl.mem_flags -demo_buf = cl.create_buffer(ctx, mf.WRITE_ONLY, demo_r.nbytes) +demo_buf = cl.Buffer(ctx, mf.WRITE_ONLY, demo_r.nbytes) - -prg = cl.create_program_with_source(ctx, +prg = cl.Program(ctx, """ __kernel void demo(__global uint *demo) { diff --git a/pyopencl/__init__.py b/pyopencl/__init__.py index 09d7539822426ef6278b25e13be3f2bf178d554a..7977933ac9bed3c6a33188e5d532ef5eab712e01 100644 --- a/pyopencl/__init__.py +++ b/pyopencl/__init__.py @@ -11,13 +11,31 @@ CONSTANT_CLASSES = [ def _add_functionality(): cls_to_info_cls = { - _cl.Platform: _cl.platform_info, - _cl.Device: _cl.device_info, - _cl.Context: _cl.context_info, - _cl.CommandQueue: _cl.command_queue_info, - _cl.Event: _cl.event_info, - _cl.MemoryObject: _cl.mem_info, - _cl.Kernel: _cl.kernel_info, + _cl.Platform: [ + (_cl.Platform.get_info, _cl.platform_info), + ], + _cl.Device: [ + (_cl.Device.get_info, _cl.device_info) + ], + _cl.Context: [ + (_cl.Context.get_info, _cl.context_info), + ], + _cl.CommandQueue: [ + (_cl.CommandQueue.get_info, _cl.command_queue_info) + ], + _cl.Event: [ + (_cl.Event.get_info, _cl.event_info), + ], + _cl.MemoryObject: [ + (MemoryObject.get_info,_cl.mem_info), + ], + _cl.Image: [ + (Image.get_image_info, _cl.image_info), + (MemoryObject.get_info,_cl.mem_info), + ], + _cl.Kernel: [ + (Kernel.get_info, _cl.kernel_info), + ], } def to_string(cls, value): @@ -31,58 +49,34 @@ def _add_functionality(): for cls in CONSTANT_CLASSES: cls.to_string = classmethod(to_string) - class ProfilingInfoGetter: - def __init__(self, event): - self.event = event - - def __getattr__(self, name): - info_cls = _cl.profiling_info + # get_info attributes ----------------------------------------------------- + def make_getattr(info_classes): + name_to_info = dict( + (intern(info_name.lower()), (info_method, info_value)) + for info_method, info_class in info_classes[::-1] + for info_name, info_value in + info_class.__dict__.iteritems() + if info_name != "to_string" and not info_name.startswith("_") + ) - try: - inf_attr = getattr(info_cls, name.upper()) - except AttributeError: - raise AttributeError("%s has no attribute '%s'" - % (type(self), name)) - else: - return self.event.get_profiling_info(inf_attr) - - _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: - inf_attr = getattr(info_cls, name.upper()) - except AttributeError: + inf_method, inf_attr = name_to_info[name] + except KeyError: raise AttributeError("%s has no attribute '%s'" % (type(self), name)) else: - return self.get_info(inf_attr) + return inf_method(self, inf_attr) return result - for cls, info_cls in cls_to_info_cls.iteritems(): - cls.__getattr__ = make_getattr(info_cls) + for cls, info_classes in cls_to_info_cls.iteritems(): + cls.__getattr__ = make_getattr(info_classes) + # Program ----------------------------------------------------------------- def program_getattr(self, attr): try: - pi_attr = getattr(program_info, attr.upper()) + pi_attr = getattr(_cl.program_info, attr.upper()) except AttributeError: try: knl = Kernel(self, attr) @@ -98,6 +92,25 @@ def _add_functionality(): Program.__getattr__ = program_getattr + # Event ------------------------------------------------------------------- + class ProfilingInfoGetter: + def __init__(self, event): + self.event = event + + def __getattr__(self, name): + info_cls = _cl.profiling_info + + try: + inf_attr = getattr(info_cls, name.upper()) + except AttributeError: + raise AttributeError("%s has no attribute '%s'" + % (type(self), name)) + else: + return self.event.get_profiling_info(inf_attr) + + _cl.Event.profile = property(ProfilingInfoGetter) + + # Kernel ------------------------------------------------------------------ def kernel_call(self, queue, global_size, *args, **kwargs): for i, arg in enumerate(args): self.set_arg(i, arg) @@ -116,6 +129,7 @@ def _add_functionality(): Kernel.__call__ = kernel_call + # ImageFormat ------------------------------------------------------------- def image_format_repr(self): return "ImageFormat(%s, %s)" % ( channel_order.to_string(self.channel_order), @@ -123,13 +137,83 @@ def _add_functionality(): ImageFormat.__repr__ = image_format_repr + # Image ------------------------------------------------------------------- + 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) + + _cl.Image.image = property(ImageInfoGetter) + def event_wait(self): wait_for_events([self]) return self Event.wait = event_wait + if _cl.have_gl(): + def gl_object_get_gl_object(self): + return self.get_gl_object_info()[1] - + GLBuffer.gl_object = property(gl_object_get_gl_object) + GLTexture.gl_object = property(gl_object_get_gl_object) _add_functionality() + + + + +# backward compatibility ------------------------------------------------------ +def create_image_2d(context, flags, format, width, height, pitch=0, host_buffer=None): + from warnings import warn + warn("create_image_2d is deprecated. Use the Image() constructor instead.", + DeprecationWarning) + return Image(context, flags, format, (width, height), (pitch,), host_buffer) + +def create_image_3d(context, flags, format, width, height, depth, + row_pitch=0, slice_pitch=0, host_buffer=None): + from warnings import warn + warn("create_image_3d is deprecated. Use the Image() constructor instead.", + DeprecationWarning) + return Image(context, flags, format, (width, height, depth), + (row_pitch, slice_pitch), host_buffer) + +def create_program_with_source(context, source): + from warnings import warn + warn("create_program_with_source is deprecated. Use the Program() constructor instead.", + DeprecationWarning) + + return Program(context, source) + +def create_program_with_binary(context, devices, binaries): + from warnings import warn + warn("create_program_with_binary is deprecated. Use the Program() constructor instead.", + DeprecationWarning) + + return Program(context, devices, binaries) + +def create_buffer(context, flags, size): + from warnings import warn + warn("create_buffer is deprecated. Use the Buffer() constructor instead.", + DeprecationWarning) + + return Buffer(context, flags, size=size) + +def create_host_buffer(context, flags, hostbuf): + from warnings import warn + warn("create_host_buffer is deprecated. Use the Buffer() constructor instead.", + DeprecationWarning) + + return Buffer(context, flags, hostbuf=hostbuf) diff --git a/src/wrapper/wrap_cl.cpp b/src/wrapper/wrap_cl.cpp index d8bfcb80c1bc5eed7635927263aae5d252b960e2..a89a212477d3c95f142c3260fc0e8c0ae1d358ef 100644 --- a/src/wrapper/wrap_cl.cpp +++ b/src/wrapper/wrap_cl.cpp @@ -466,25 +466,29 @@ BOOST_PYTHON_MODULE(_cl) py::return_value_policy()); DEF_SIMPLE_FUNCTION(enqueue_wait_for_events); + // memory_object ------------------------------------------------------------ { typedef memory_object cls; py::class_("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 ; } - py::def("create_buffer", create_buffer, - py::return_value_policy()); - py::def("create_host_buffer", create_host_buffer, - py::return_value_policy()); + { + typedef buffer cls; + py::class_, boost::noncopyable>( + "Buffer", py::no_init) + .def("__init__", make_constructor(create_buffer, + py::default_call_policies(), + (py::args("context", "flags"), + py::arg("size")=0, + py::arg("hostbuf")=py::object() + ))) + ; + } py::def("enqueue_read_buffer", enqueue_read_buffer, (py::args("queue", "mem", "host_buffer"), @@ -500,6 +504,20 @@ BOOST_PYTHON_MODULE(_cl) py::return_value_policy()); // image -------------------------------------------------------------------- + { + typedef image cls; + py::class_, boost::noncopyable>( + "Image", py::no_init) + .def("__init__", make_constructor(create_image, + py::default_call_policies(), + (py::args("context", "flags", "format", "shape"), + py::arg("pitches")=py::object(), + py::arg("host_buffer")=py::object() + ))) + .DEF_SIMPLE_METHOD(get_image_info) + ; + } + { typedef cl_image_format cls; py::class_("ImageFormat") @@ -510,19 +528,6 @@ 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"), - py::arg("pitch")=0, - py::arg("host_buffer")=py::object() - ), - py::return_value_policy()); - py::def("create_image_3d", create_image_3D, - (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::def("enqueue_read_image", enqueue_read_image, (py::args("queue", "mem", "origin", "region", "host_buffer"), @@ -591,6 +596,14 @@ BOOST_PYTHON_MODULE(_cl) { typedef program cls; py::class_("Program", py::no_init) + .def("__init__", make_constructor( + create_program_with_source, + py::default_call_policies(), + py::args("context", "src"))) + .def("__init__", make_constructor( + create_program_with_binary, + py::default_call_policies(), + py::args("context", "devices", "binaries"))) .DEF_SIMPLE_METHOD(get_info) .DEF_SIMPLE_METHOD(get_build_info) .def("build", &cls::build, @@ -604,13 +617,6 @@ BOOST_PYTHON_MODULE(_cl) py::def("unload_compiler", unload_compiler); - py::def("create_program_with_source", create_program_with_source, - py::args("context", "src"), - py::return_value_policy()); - py::def("create_program_with_binary", create_program_with_binary, - py::args("context", "devices", "binaries"), - py::return_value_policy()); - { typedef kernel cls; py::class_("Kernel", @@ -657,14 +663,41 @@ BOOST_PYTHON_MODULE(_cl) ADD_ATTR(GL_, MIPMAP_LEVEL); } - py::def("create_from_gl_buffer", create_from_gl_buffer, - py::return_value_policy()); - py::def("create_from_gl_texture_2d", create_from_gl_texture_2d, - py::return_value_policy()); - py::def("create_from_gl_texture_3d", create_from_gl_texture_3d, - py::return_value_policy()); - py::def("create_from_gl_renderbuffer", create_from_gl_renderbuffer, - py::return_value_policy()); + { + typedef gl_buffer cls; + py::class_, boost::noncopyable>( + "GLBuffer", py::no_init) + .def("__init__", make_constructor(create_from_gl_buffer, + py::default_call_policies(), + (py::args("context", "flags" "bufobj")))) + .def("get_gl_object_info", get_gl_object_info) + ; + } + + { + typedef gl_renderbuffer cls; + py::class_, boost::noncopyable>( + "GLRenderBuffer", py::no_init) + .def("__init__", make_constructor(create_from_gl_renderbuffer, + py::default_call_policies(), + (py::args("context", "flags" "bufobj")))) + .def("get_gl_object_info", get_gl_object_info) + ; + } + + { + typedef gl_texture cls; + py::class_, boost::noncopyable>( + "GLTexture", py::no_init) + .def("__init__", make_constructor(create_from_gl_texture, + py::default_call_policies(), + (py::args("context", "flags", + "texture_target", "miplevel", + "texture", "dims")))) + .def("get_gl_object_info", get_gl_object_info) + .DEF_SIMPLE_METHOD(get_gl_texture_info) + ; + } py::def("enqueue_acquire_gl_objects", enqueue_acquire_gl_objects, (py::args("queue", "mem_objects"), diff --git a/src/wrapper/wrap_cl.hpp b/src/wrapper/wrap_cl.hpp index 19f77a5e5ffda56da26cc2ef56cf778702a23a4d..4b9df49158f6afdebe15d041dede02f633d0a06f 100644 --- a/src/wrapper/wrap_cl.hpp +++ b/src/wrapper/wrap_cl.hpp @@ -59,6 +59,17 @@ +#define PYOPENCL_DEPRECATED(WHAT, KILL_VERSION, EXTRA_MSG) \ + { \ + PyErr_Warn( \ + PyExc_DeprecationWarning, \ + WHAT " is deprecated and will stop working in PyOpenCL " KILL_VERSION". " \ + EXTRA_MSG); \ + } + + + + // tracing and error reporting ------------------------------------------------ #ifdef PYOPENCL_TRACE #define PYOPENCL_PRINT_CALL_TRACE(NAME) \ @@ -215,32 +226,6 @@ -// 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 { class error : public std::runtime_error @@ -951,7 +936,7 @@ namespace pyopencl void release() { if (!m_valid) - throw error("MemoryObject.free", CL_INVALID_VALUE, + throw error("MemoryObject.free", CL_INVALID_VALUE, "trying to double-unref mem object"); PYOPENCL_CALL_GUARDED_CLEANUP(clReleaseMemObject, (m_mem)); @@ -1006,68 +991,63 @@ namespace pyopencl throw error("MemoryObject.get_info", CL_INVALID_VALUE); } } - - py::object get_image_info(cl_image_info param_name) const - { - switch (param_name) - { - case CL_IMAGE_FORMAT: - PYOPENCL_GET_INTEGRAL_INFO(Image, m_mem, param_name, - cl_image_format); - case CL_IMAGE_ELEMENT_SIZE: - case CL_IMAGE_ROW_PITCH: - case CL_IMAGE_SLICE_PITCH: - case CL_IMAGE_WIDTH: - case CL_IMAGE_HEIGHT: - case CL_IMAGE_DEPTH: - PYOPENCL_GET_INTEGRAL_INFO(Image, m_mem, param_name, size_t); - - default: - throw error("MemoryObject.get_image_info", CL_INVALID_VALUE); - } - } }; - PYOPENCL_WRAP_BUFFER_CREATOR(create_buffer, clCreateBuffer, - (context &ctx, cl_mem_flags flags, size_t size), - (ctx.data(), flags, size, 0, &status_code)); + class buffer : public memory_object + { + public: + buffer(cl_mem mem, bool retain, py::object *hostbuf=0) + : memory_object(mem, retain, hostbuf) + { } + }; - memory_object *create_host_buffer( + memory_object *create_buffer( context &ctx, cl_mem_flags flags, - py::object buffer) + size_t size, + py::object buffer + ) { - void *buf; - PYOPENCL_BUFFER_SIZE_T len; - if (flags & CL_MEM_USE_HOST_PTR) - { - if (PyObject_AsWriteBuffer(buffer.ptr(), &buf, &len)) - throw py::error_already_set(); - } - else + void *buf = 0; + py::object *retained_buf_obj = 0; + if (buffer.ptr() != Py_None) { - if (PyObject_AsReadBuffer( - buffer.ptr(), const_cast(&buf), &len)) - throw py::error_already_set(); + PYOPENCL_BUFFER_SIZE_T len; + if (flags & CL_MEM_USE_HOST_PTR) + { + if (PyObject_AsWriteBuffer(buffer.ptr(), &buf, &len)) + throw py::error_already_set(); + } + else + { + if (PyObject_AsReadBuffer( + buffer.ptr(), const_cast(&buf), &len)) + throw py::error_already_set(); + } + + if (flags & CL_MEM_USE_HOST_PTR) + retained_buf_obj = &buffer; + + if (size > size_t(len)) + throw pyopencl::error("Buffer", CL_INVALID_VALUE, + "specified size is greater than host buffer size"); + if (size == 0) + size = len; } cl_int status_code; - cl_mem mem = clCreateBuffer(ctx.data(), flags, len, buf, &status_code); + cl_mem mem = clCreateBuffer(ctx.data(), flags, size, buf, &status_code); PYOPENCL_PRINT_CALL_TRACE("clCreateBuffer"); if (status_code != CL_SUCCESS) throw pyopencl::error("create_host_buffer", status_code); - py::object *retained_buf_obj = 0; - if (flags & CL_MEM_USE_HOST_PTR) - retained_buf_obj = &buffer; - try { return new memory_object(mem, false, retained_buf_obj); @@ -1145,6 +1125,37 @@ namespace pyopencl // images ------------------------------------------------------------------- + class image : public memory_object + { + public: + image(cl_mem mem, bool retain, py::object *hostbuf=0) + : memory_object(mem, retain, hostbuf) + { } + + py::object get_image_info(cl_image_info param_name) const + { + switch (param_name) + { + case CL_IMAGE_FORMAT: + PYOPENCL_GET_INTEGRAL_INFO(Image, data(), param_name, + cl_image_format); + case CL_IMAGE_ELEMENT_SIZE: + case CL_IMAGE_ROW_PITCH: + case CL_IMAGE_SLICE_PITCH: + case CL_IMAGE_WIDTH: + case CL_IMAGE_HEIGHT: + case CL_IMAGE_DEPTH: + PYOPENCL_GET_INTEGRAL_INFO(Image, data(), param_name, size_t); + + default: + throw error("MemoryObject.get_image_info", CL_INVALID_VALUE); + } + } + }; + + + + cl_image_format *make_image_format(cl_channel_order ord, cl_channel_type tp) { std::auto_ptr result(new cl_image_format); @@ -1160,11 +1171,11 @@ namespace pyopencl { cl_uint num_image_formats; PYOPENCL_CALL_GUARDED(clGetSupportedImageFormats, ( - ctx.data(), flags, image_type, 0, 0, & num_image_formats)); + ctx.data(), flags, image_type, 0, 0, &num_image_formats)); std::vector formats(num_image_formats); PYOPENCL_CALL_GUARDED(clGetSupportedImageFormats, ( - ctx.data(), flags, image_type, + ctx.data(), flags, image_type, num_image_formats, &formats.front(), 0)); PYOPENCL_RETURN_VECTOR(cl_image_format, formats); @@ -1173,79 +1184,107 @@ namespace pyopencl -#define PYOPENCL_MAKE_CREATE_IMAGE(ITYPE, IMG_ARG_DECLS, IMG_ARGS) \ - inline memory_object *create_image_##ITYPE( \ - context const &ctx, \ - cl_mem_flags flags, \ - cl_image_format const &fmt, \ - IMG_ARG_DECLS, \ - py::object buffer) \ - { \ - void *buf = 0; \ - PYOPENCL_BUFFER_SIZE_T len; \ - py::object *retained_buf_obj = 0; \ - \ - if (buffer.ptr() != Py_None) \ - { \ - if (flags & CL_MEM_USE_HOST_PTR) \ - { \ - if (PyObject_AsWriteBuffer(buffer.ptr(), &buf, &len)) \ - throw py::error_already_set(); \ - } \ - else \ - { \ - if (PyObject_AsReadBuffer( \ - buffer.ptr(), const_cast(&buf), &len)) \ - throw py::error_already_set(); \ - } \ - \ - if (flags & CL_MEM_USE_HOST_PTR) \ - retained_buf_obj = &buffer; \ - } \ - \ - cl_int status_code; \ - cl_mem mem = clCreateImage##ITYPE(ctx.data(), flags, &fmt, \ - IMG_ARGS, buf, &status_code); \ - \ - PYOPENCL_PRINT_CALL_TRACE("clCreateImage" #ITYPE); \ - if (status_code != CL_SUCCESS) \ - throw pyopencl::error("create_image_" #ITYPE, status_code); \ - \ - try \ - { \ - return new memory_object(mem, false, retained_buf_obj); \ - } \ - catch (...) \ - { \ - PYOPENCL_CALL_GUARDED(clReleaseMemObject, (mem)); \ - throw; \ - } \ - } + inline image *create_image( + context const &ctx, + cl_mem_flags flags, + cl_image_format const &fmt, + py::object shape, + py::object pitches, + py::object buffer) + { + void *buf = 0; + PYOPENCL_BUFFER_SIZE_T len; + py::object *retained_buf_obj = 0; + if (buffer.ptr() != Py_None) + { + if (flags & CL_MEM_USE_HOST_PTR) + { + if (PyObject_AsWriteBuffer(buffer.ptr(), &buf, &len)) + throw py::error_already_set(); + } + else + { + if (PyObject_AsReadBuffer( + buffer.ptr(), const_cast(&buf), &len)) + throw py::error_already_set(); + } + if (flags & CL_MEM_USE_HOST_PTR) + retained_buf_obj = &buffer; + } + unsigned dims = py::len(shape); + cl_int status_code; + cl_mem mem; + if (dims == 2) + { + size_t width = py::extract(shape[0]); + size_t height = py::extract(shape[1]); + + size_t pitch = 0; + if (pitches.ptr() != Py_None) + { + if (py::len(pitches) != 1) + throw pyopencl::error("Image", CL_INVALID_VALUE, + "invalid length of pitch tuple"); + pitch = py::extract(pitches[0]); + } + + mem = clCreateImage2D(ctx.data(), flags, &fmt, + width, height, pitch, buf, &status_code); + + PYOPENCL_PRINT_CALL_TRACE("clCreateImage2D"); + if (status_code != CL_SUCCESS) + throw pyopencl::error("clCreateImage2D", status_code); + } + else if (dims == 3) + { + size_t width = py::extract(shape[0]); + size_t height = py::extract(shape[1]); + size_t depth = py::extract(shape[2]); -#define PYOPENCL_IMG_ARG_DECLS \ - size_t width, size_t height, size_t pitch -#define PYOPENCL_IMG_ARGS width, height, pitch - PYOPENCL_MAKE_CREATE_IMAGE(2D, PYOPENCL_IMG_ARG_DECLS, PYOPENCL_IMG_ARGS) -#undef PYOPENCL_IMG_ARG_DECLS -#undef PYOPENCL_IMG_ARGS + size_t pitch_x = 0; + size_t pitch_y = 0; -#define PYOPENCL_IMG_ARG_DECLS \ - size_t width, size_t height, size_t depth, \ - size_t row_pitch, size_t slice_pitch -#define PYOPENCL_IMG_ARGS width, height, depth, row_pitch, slice_pitch - PYOPENCL_MAKE_CREATE_IMAGE(3D, PYOPENCL_IMG_ARG_DECLS, PYOPENCL_IMG_ARGS) -#undef PYOPENCL_IMG_ARG_DECLS -#undef PYOPENCL_IMG_ARGS + if (pitches.ptr() != Py_None) + { + if (py::len(pitches) != 2) + throw pyopencl::error("Image", CL_INVALID_VALUE, + "invalid length of pitch tuple"); + + pitch_x = py::extract(pitches[0]); + pitch_y = py::extract(pitches[1]); + } + + mem = clCreateImage3D(ctx.data(), flags, &fmt, + width, height, depth, pitch_x, pitch_y, buf, &status_code); + + PYOPENCL_PRINT_CALL_TRACE("clCreateImage3D"); + if (status_code != CL_SUCCESS) + throw pyopencl::error("clCreateImage3D", status_code); + } + else + throw pyopencl::error("Image", CL_INVALID_VALUE, + "invalid dimension"); + + try + { + return new image(mem, false, retained_buf_obj); + } + catch (...) + { + PYOPENCL_CALL_GUARDED(clReleaseMemObject, (mem)); + throw; + } + } event *enqueue_read_image( command_queue &cq, - memory_object &mem, + image &img, py::object py_origin, py::object py_region, py::object buffer, size_t row_pitch, size_t slice_pitch, @@ -1266,7 +1305,7 @@ namespace pyopencl cl_event evt; PYOPENCL_CALL_GUARDED(clEnqueueReadImage, ( cq.data(), - mem.data(), + img.data(), PYOPENCL_CAST_BOOL(is_blocking), origin, region, row_pitch, slice_pitch, buf, num_events_in_wait_list, &event_wait_list.front(), &evt @@ -1279,7 +1318,7 @@ namespace pyopencl event *enqueue_write_image( command_queue &cq, - memory_object &mem, + image &img, py::object py_origin, py::object py_region, py::object buffer, size_t row_pitch, size_t slice_pitch, @@ -1300,7 +1339,7 @@ namespace pyopencl cl_event evt; PYOPENCL_CALL_GUARDED(clEnqueueWriteImage, ( cq.data(), - mem.data(), + img.data(), PYOPENCL_CAST_BOOL(is_blocking), origin, region, row_pitch, slice_pitch, buf, num_events_in_wait_list, &event_wait_list.front(), &evt @@ -1315,8 +1354,8 @@ namespace pyopencl command_queue &cq, memory_object &src, memory_object &dest, - py::object py_src_origin, - py::object py_dest_origin, + py::object py_src_origin, + py::object py_dest_origin, py::object py_region, py::object py_wait_for ) @@ -1342,7 +1381,7 @@ namespace pyopencl command_queue &cq, memory_object &src, memory_object &dest, - py::object py_origin, + py::object py_origin, py::object py_region, size_t offset, py::object py_wait_for @@ -1369,7 +1408,7 @@ namespace pyopencl memory_object &src, memory_object &dest, size_t offset, - py::object py_origin, + py::object py_origin, py::object py_region, py::object py_wait_for ) @@ -1381,7 +1420,7 @@ namespace pyopencl cl_event evt; PYOPENCL_CALL_GUARDED(clEnqueueCopyBufferToImage, ( cq.data(), src.data(), dest.data(), - offset, origin, region, + offset, origin, region, num_events_in_wait_list, &event_wait_list.front(), &evt )); PYOPENCL_RETURN_NEW_EVENT(evt); @@ -1455,7 +1494,7 @@ namespace pyopencl cl_int status_code; PYOPENCL_PRINT_CALL_TRACE("clEnqueueMapBuffer"); void *mapped = clEnqueueMapBuffer( - cq.data(), buf.data(), + cq.data(), buf.data(), PYOPENCL_CAST_BOOL(is_blocking), flags, offset, PyArray_NBYTES(result.get()), num_events_in_wait_list, &event_wait_list.front(), &evt, @@ -1484,7 +1523,7 @@ namespace pyopencl Py_INCREF(array_py.get()); return py::make_tuple( - array_py, + array_py, handle_from_new_ptr(new event(evt_handle))); } @@ -1495,7 +1534,7 @@ namespace pyopencl command_queue &cq, memory_object &img, cl_map_flags flags, - py::object py_origin, + py::object py_origin, py::object py_region, py::object py_shape, py::object dtype, py::object order_py, py::object py_wait_for, @@ -1512,7 +1551,7 @@ namespace pyopencl PYOPENCL_PRINT_CALL_TRACE("clEnqueueMapImage"); size_t row_pitch, slice_pitch; void *mapped = clEnqueueMapImage( - cq.data(), img.data(), + cq.data(), img.data(), PYOPENCL_CAST_BOOL(is_blocking), flags, origin, region, &row_pitch, &slice_pitch, num_events_in_wait_list, &event_wait_list.front(), &evt, @@ -1544,7 +1583,7 @@ namespace pyopencl Py_INCREF(array_py.get()); return py::make_tuple( - array_py, + array_py, handle_from_new_ptr(new event(evt_handle)), row_pitch, slice_pitch); } @@ -2086,21 +2125,111 @@ namespace pyopencl #ifdef HAVE_GL - PYOPENCL_WRAP_BUFFER_CREATOR(create_from_gl_buffer, clCreateFromGLBuffer, + class gl_buffer : public memory_object + { + public: + gl_buffer(cl_mem mem, bool retain, py::object *hostbuf=0) + : memory_object(mem, retain, hostbuf) + { } + }; + + + + + class gl_renderbuffer : public memory_object + { + public: + gl_renderbuffer(cl_mem mem, bool retain, py::object *hostbuf=0) + : memory_object(mem, retain, hostbuf) + { } + }; + + + + + class gl_texture : public image + { + public: + gl_texture(cl_mem mem, bool retain, py::object *hostbuf=0) + : image(mem, retain, hostbuf) + { } + + py::object get_gl_texture_info(cl_gl_texture_info param_name) + { + switch (param_name) + { + case CL_GL_TEXTURE_TARGET: + PYOPENCL_GET_INTEGRAL_INFO(GLTexture, data(), param_name, GLenum); + case CL_GL_MIPMAP_LEVEL: + PYOPENCL_GET_INTEGRAL_INFO(GLTexture, data(), param_name, GLint); + + default: + throw error("MemoryObject.get_gl_texture_info", CL_INVALID_VALUE); + } + } + }; + + + + +#define PYOPENCL_WRAP_BUFFER_CREATOR(TYPE, NAME, CL_NAME, ARGS, CL_ARGS) \ + TYPE *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 TYPE(mem, false); \ + } \ + catch (...) \ + { \ + PYOPENCL_CALL_GUARDED(clReleaseMemObject, (mem)); \ + throw; \ + } \ + } + + + + + PYOPENCL_WRAP_BUFFER_CREATOR(gl_buffer, + 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, + PYOPENCL_WRAP_BUFFER_CREATOR(gl_texture, + 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, + PYOPENCL_WRAP_BUFFER_CREATOR(gl_texture, + 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, + PYOPENCL_WRAP_BUFFER_CREATOR(gl_renderbuffer, + create_from_gl_renderbuffer, clCreateFromGLRenderbuffer, (context &ctx, cl_mem_flags flags, GLuint renderbuffer), (ctx.data(), flags, renderbuffer, &status_code)); + gl_texture *create_from_gl_texture( + context &ctx, cl_mem_flags flags, + GLenum texture_target, GLint miplevel, + GLuint texture, unsigned dims) + { + if (dims == 2) + return create_from_gl_texture_2d(ctx, flags, texture_target, miplevel, texture); + else if (dims == 3) + return create_from_gl_texture_3d(ctx, flags, texture_target, miplevel, texture); + else + throw pyopencl::error("Image", CL_INVALID_VALUE, + "invalid dimension"); + } + + @@ -2112,26 +2241,6 @@ namespace pyopencl 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, \ diff --git a/test/test_wrapper.py b/test/test_wrapper.py index b2d39f7d2bca66beb4d917c4c94559d2f2e638b3..7ecaef49bf708164332b5a1bd708c8f4a93545d2 100644 --- a/test/test_wrapper.py +++ b/test/test_wrapper.py @@ -22,7 +22,7 @@ if have_cl(): class TestCL: disabled = not have_cl() - def test_get_info(self): + def test_get_info(self, platform, device): had_failures = [False] QUIRKS = [ @@ -41,7 +41,7 @@ class TestCL: return False - def do_test(cl_obj, info_cls, func=None): + def do_test(cl_obj, info_cls, func=None, try_attr_form=True): if func is None: def func(info): cl_obj.get_info(info) @@ -61,67 +61,82 @@ class TestCL: 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 try_attr_form: + try: + getattr(cl_obj, info_name.lower()) + except: + print "failed attr-based 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 + + do_test(platform, cl.platform_info) + + 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.Program(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), + try_attr_form=False) + + cl.unload_compiler() # just for the heck of it + + mf = cl.mem_flags + n = 2000 + a_buf = cl.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), + try_attr_form=False) + + 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) + #img = cl.Image(ctx, cl.mem_flags.READ_ONLY, img_format, + #(128, 128)) + img.depth + img.image.depth + do_test(img, cl.image_info, + lambda info: img.get_image_info(info)) if had_failures[0]: raise RuntimeError("get_info testing had errors") @@ -130,7 +145,7 @@ class TestCL: for platform in cl.get_platforms(): for device in platform.get_devices(): ctx = cl.Context([device]) - prg = cl.create_program_with_source(ctx, """ + prg = cl.Program(ctx, """ __kernel void sum(__global float *a) { a[get_global_id(0)] *= 2; } """).build() @@ -142,17 +157,47 @@ class TestCL: pass def test_image_format_constructor(self): + # doesn't need image support to succeed 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__ + def test_nonempty_supported_image_formats(self, device, context): + if device.image_support: + assert len(cl.get_supported_image_formats( + context, cl.mem_flags.READ_ONLY, cl.mem_object_type.IMAGE2D)) > 0 + else: + from py.test import skip + skip("images not supported on %s" % device.name) + + +def pytest_generate_tests(metafunc): + if ("device" in metafunc.funcargnames + or "context" in metafunc.funcargnames): + arg_dict = {} + for platform in cl.get_platforms(): + if "platform" in metafunc.funcargnames: + arg_dict["platform"] = platform + + for device in platform.get_devices(): + if "device" in metafunc.funcargnames: + arg_dict["device"] = device + if "context" in metafunc.funcargnames: + arg_dict["context"] = cl.Context([device]) + + metafunc.addcall(funcargs=arg_dict.copy()) + + elif "platform" in metafunc.funcargnames: + for platform in cl.get_platforms(): + metafunc.addcall( + funcargs=dict(platform=platform)) if __name__ == "__main__":