diff --git a/doc/source/runtime.rst b/doc/source/runtime.rst index 7fa2b130ed861818322825228a7893a025b979bf..2912f4e5ac056e977c68a5f9f53ba89a3720eeb2 100644 --- a/doc/source/runtime.rst +++ b/doc/source/runtime.rst @@ -11,7 +11,7 @@ Version Queries .. data:: VERSION - Gives the numeric version of PyOpenCL as a variable-length tuple + Gives the numeric version of PyOpenCL as a variable-length tuple of integers. Enables easy version checks such as *VERSION >= (0, 93)*. @@ -55,12 +55,12 @@ Constants Platforms, Devices and Contexts ------------------------------- -.. |comparable| replace:: Two instances of this class may be compared +.. |comparable| replace:: Two instances of this class may be compared using *"=="* and *"!="*. -.. |buf-iface| replace:: must implement the Python buffer interface. +.. |buf-iface| replace:: must implement the Python buffer interface. (e.g. by being an :class:`numpy.ndarray`) -.. |explain-waitfor| replace:: *wait_for* - may either be *None* or a list of :class:`Event` instances 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| @@ -108,7 +108,7 @@ Platforms, Devices and Contexts At most one of *devices* and *dev_type* may be not `None`, where *devices* is a list of :class:`Device` instances, and *dev_type* is one of the :class:`device_type` constants. - If neither is specified, a context with a *dev_type* of + If neither is specified, a context with a *dev_type* of :attr:`device_type.DEFAULT` is created. .. note:: @@ -121,7 +121,7 @@ Platforms, Devices and Contexts .. note:: - For + For :attr:`context_properties.CL_GL_CONTEXT_KHR`, :attr:`context_properties.CL_EGL_DISPLAY_KHR`, :attr:`context_properties.CL_GLX_DISPLAY_KHR`, @@ -149,7 +149,7 @@ Platforms, Devices and Contexts Create a :class:`Context` 'somehow'. - If multiple choices for platform and/or device exist, *interactive* + If multiple choices for platform and/or device exist, *interactive* is True, and *sys.stdin.isatty()* is also True, then the user is queried about which device should be chosen. Otherwise, a device is chosen in an implementation-defined manner. @@ -198,7 +198,7 @@ Command Queues and Events .. attribute:: profile.info Lower case versions of the :class:`profiling_info` constants - may be used as attributes on the attribute `profile` of this + may be used as attributes on the attribute `profile` of this class to directly query profiling info. For example, you may use *evt.profile.end* instead of @@ -396,7 +396,7 @@ Images *shape* is a 2- or 3-tuple. - If *hostbuf* is given and *shape* is `None`, then *hostbuf.shape* is + 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`. @@ -405,7 +405,7 @@ Images .. attribute:: info - Lower case versions of the :class:`mem_info` + 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. @@ -464,7 +464,7 @@ Mapping Memory into Host Address Space :return: a tuple *(array, event)*. *array* is a :class:`numpy.ndarray` representing the host side - of the map. Its *.base* member contains a + 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) @@ -476,7 +476,7 @@ Mapping Memory into Host Address Space :return: a tuple *(array, event)*. *array* is a :class:`numpy.ndarray` representing the host side - of the map. Its *.base* member contains a + of the map. Its *.base* member contains a :class:`MemoryMap`. @@ -527,12 +527,12 @@ Programs and Kernels .. method:: build(options="", devices=None) - *options* is a string of compiler flags. + *options* is a string of compiler flags. Returns *self*. .. attribute:: kernel_name - :class:`Kernel` objects can be produced from a built + :class:`Kernel` objects can be produced from a built (see :meth:`build`) program simply by attribute lookup. .. note:: @@ -569,18 +569,18 @@ Programs and Kernels *arg* may be - * `None`: This may be passed for `__global` memory references + * `None`: This may be passed for `__global` memory references to pass a NULL pointer to the kernel. * Anything that satisfies the Python buffer interface, in particular :class:`numpy.ndarray`, :class:`str`, or :mod:`numpy`'s sized scalars, such as :class:`numpy.int32` - or :class:`numpy.float64`. + or :class:`numpy.float64`. - .. note:: + .. note:: Note that Python's own :class:`int` - or :class:`float` objects will not work as-is, but - :mod:`struct` can be used to convert them to binary + or :class:`float` objects will not work as-is, but + :mod:`struct` can be used to convert them to binary data in a :class:`str`, which will work. * An instance of :class:`MemoryObject`. (e.g. :class:`Buffer`, @@ -595,7 +595,7 @@ Programs and Kernels .. method:: __call__(queue, global_size, local_size, *args, global_offset=None, wait_for=None) Use :func:`enqueue_nd_range_kernel` to enqueue a kernel execution, after using - :meth:`set_args` to set each argument in turn. See the documentation for + :meth:`set_args` to set each argument in turn. See the documentation for :meth:`set_arg` to see what argument types are allowed. |std-enqueue-blurb| @@ -604,7 +604,7 @@ Programs and Kernels .. versionchanged:: 0.92 *local_size* was promoted to third positional argument from being a keyword argument. The old keyword argument usage will continue to - be accepted with a warning throughout the 0.92 release cycle. + be accepted with a warning throughout the 0.92 release cycle. This is a backward-compatible change (just barely!) because *local_size* as third positional argument can only be a :class:`tuple` or *None*. :class:`tuple` instances are never valid @@ -637,7 +637,7 @@ Programs and Kernels GL Interoperability ------------------- -Functionality in this section is only available when PyOpenCL is compiled +Functionality in this section is only available when PyOpenCL is compiled with GL support. See :func:`have_gl`. .. versionadded:: 0.91 @@ -667,7 +667,7 @@ 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`. + See :class:`gl_texture_info` for values of *param*. Only available when PyOpenCL is compiled with GL support. See :func:`have_gl`. .. function:: enqueue_acquire_gl_objects(queue, mem_objects, wait_for=None) @@ -683,5 +683,5 @@ with GL support. See :func:`have_gl`. Get information on which CL device corresponds to a given GL/EGL/WGL/CGL device. - See the :class:`Context` constructor for the meaning of + See the :class:`Context` constructor for the meaning of *properties* and :class:`gl_context_info` for *param_name*. diff --git a/src/wrapper/wrap_cl.cpp b/src/wrapper/wrap_cl.cpp index 1e2d5b5d899b827ac244eb7129d3cae4fedfbd95..55d855dce78f055e4179b28a741ec292e9d97abc 100644 --- a/src/wrapper/wrap_cl.cpp +++ b/src/wrapper/wrap_cl.cpp @@ -652,7 +652,7 @@ BOOST_PYTHON_MODULE(_cl) #ifdef CL_VERSION_1_1 py::def("enqueue_read_buffer_rect", enqueue_read_buffer_rect, - (py::args("queue", "mem", "hostbuf", + (py::args("queue", "mem", "hostbuf", "buffer_origin", "host_origin", "region"), py::arg("buffer_pitches")=py::object(), py::arg("host_pitches")=py::object(), @@ -661,7 +661,7 @@ BOOST_PYTHON_MODULE(_cl) ), py::return_value_policy<py::manage_new_object>()); py::def("enqueue_write_buffer_rect", enqueue_write_buffer_rect, - (py::args("queue", "mem", "hostbuf", + (py::args("queue", "mem", "hostbuf", "buffer_origin", "host_origin", "region"), py::arg("buffer_pitches")=py::object(), py::arg("host_pitches")=py::object(), @@ -670,7 +670,7 @@ BOOST_PYTHON_MODULE(_cl) ), py::return_value_policy<py::manage_new_object>()); py::def("enqueue_copy_buffer_rect", enqueue_copy_buffer_rect, - (py::args("queue", "src", "dst", + (py::args("queue", "src", "dst", "src_origin", "dst_origin", "region"), py::arg("src_pitches")=py::object(), py::arg("dst_pitches")=py::object(), diff --git a/src/wrapper/wrap_cl.hpp b/src/wrapper/wrap_cl.hpp index 92c787b8b3ec5ffcac78d2b7b1f89bbddefe2063..5567c02e5775c8160a2de70b90bcec91d072012a 100644 --- a/src/wrapper/wrap_cl.hpp +++ b/src/wrapper/wrap_cl.hpp @@ -4,7 +4,7 @@ -// {{{ includes +// {{{ includes #ifdef __APPLE__ // Mac ------------------------------------------------------------------------ @@ -730,7 +730,7 @@ namespace pyopencl if (py_devices.ptr() != Py_None) { if (py_dev_type.ptr() != Py_None) - throw error("Context", CL_INVALID_VALUE, + throw error("Context", CL_INVALID_VALUE, "one of 'devices' or 'dev_type' must be None"); std::vector<cl_device_id> devices; @@ -1142,7 +1142,7 @@ namespace pyopencl { cl_buffer_region region = { origin, size}; cl_int status_code; - cl_mem mem = clCreateSubBuffer(data(), flags, + cl_mem mem = clCreateSubBuffer(data(), flags, CL_BUFFER_CREATE_TYPE_REGION, ®ion, &status_code); PYOPENCL_PRINT_CALL_TRACE("clCreateSubBuffer"); @@ -1195,7 +1195,7 @@ namespace pyopencl py::object py_hostbuf ) { - if (py_hostbuf.ptr() != Py_None && + if (py_hostbuf.ptr() != Py_None && !(flags & (CL_MEM_USE_HOST_PTR | CL_MEM_COPY_HOST_PTR))) PyErr_Warn(PyExc_UserWarning, "'hostbuf' was passed, " "but no memory flags to make use of it."); @@ -1353,8 +1353,8 @@ namespace pyopencl src.data(), dst.data(), src_offset, dst_offset, byte_count, - num_events_in_wait_list, - event_wait_list.empty( ) ? NULL : &event_wait_list.front(), + num_events_in_wait_list, + event_wait_list.empty( ) ? NULL : &event_wait_list.front(), &evt )); @@ -1479,8 +1479,8 @@ namespace pyopencl src_origin, dst_origin, region, src_pitches[0], src_pitches[1], dst_pitches[0], dst_pitches[1], - num_events_in_wait_list, - event_wait_list.empty( ) ? NULL : &event_wait_list.front(), + num_events_in_wait_list, + event_wait_list.empty( ) ? NULL : &event_wait_list.front(), &evt )); @@ -1574,8 +1574,8 @@ namespace pyopencl case CL_INTENSITY: return 1; case CL_LUMINANCE: return 1; default: - throw pyopencl::error("ImageFormat.channel_dtype_size", - CL_INVALID_VALUE, + throw pyopencl::error("ImageFormat.channel_dtype_size", + CL_INVALID_VALUE, "unrecognized channel order"); } } @@ -1601,8 +1601,8 @@ namespace pyopencl case CL_HALF_FLOAT: return 2; case CL_FLOAT: return 4; default: - throw pyopencl::error("ImageFormat.channel_dtype_size", - CL_INVALID_VALUE, + throw pyopencl::error("ImageFormat.channel_dtype_size", + CL_INVALID_VALUE, "unrecognized channel data type"); } } @@ -1618,7 +1618,7 @@ namespace pyopencl // {{{ image creation - inline + inline image *create_image( context const &ctx, cl_mem_flags flags, @@ -1635,7 +1635,7 @@ namespace pyopencl buffer = host_buffer_deprecated; } - if (buffer.ptr() != Py_None && + if (buffer.ptr() != Py_None && !(flags & (CL_MEM_USE_HOST_PTR | CL_MEM_COPY_HOST_PTR))) PyErr_Warn(PyExc_UserWarning, "'hostbuf' was passed, " "but no memory flags to make use of it."); @@ -1643,7 +1643,7 @@ namespace pyopencl if (shape.ptr() == Py_None) { if (buffer.ptr() == Py_None) - throw pyopencl::error("Image", CL_INVALID_VALUE, + throw pyopencl::error("Image", CL_INVALID_VALUE, "'shape' must be passed if 'hostbuf' is not given"); shape = buffer.attr("shape"); @@ -1683,7 +1683,7 @@ namespace pyopencl if (pitches.ptr() != Py_None) { if (py::len(pitches) != 1) - throw pyopencl::error("Image", CL_INVALID_VALUE, + throw pyopencl::error("Image", CL_INVALID_VALUE, "invalid length of pitch tuple"); pitch = py::extract<size_t>(pitches[0]); } @@ -1691,7 +1691,7 @@ namespace pyopencl // check buffer size cl_int itemsize = get_image_format_item_size(fmt); if (buf && std::max(pitch, width*itemsize)*height > cl_uint(len)) - throw pyopencl::error("Image", CL_INVALID_VALUE, + throw pyopencl::error("Image", CL_INVALID_VALUE, "buffer too small"); mem = clCreateImage2D(ctx.data(), flags, &fmt, @@ -1713,7 +1713,7 @@ namespace pyopencl if (pitches.ptr() != Py_None) { if (py::len(pitches) != 2) - throw pyopencl::error("Image", CL_INVALID_VALUE, + throw pyopencl::error("Image", CL_INVALID_VALUE, "invalid length of pitch tuple"); pitch_x = py::extract<size_t>(pitches[0]); @@ -1722,11 +1722,11 @@ namespace pyopencl // check buffer size cl_int itemsize = get_image_format_item_size(fmt); - if (buf && + if (buf && std::max(pitch_x, width*itemsize) - * std::max(height, pitch_y) + * std::max(height, pitch_y) * depth > cl_uint(len)) - throw pyopencl::error("Image", CL_INVALID_VALUE, + throw pyopencl::error("Image", CL_INVALID_VALUE, "buffer too small"); mem = clCreateImage3D(ctx.data(), flags, &fmt, @@ -1737,7 +1737,7 @@ namespace pyopencl throw pyopencl::error("clCreateImage3D", status_code); } else - throw pyopencl::error("Image", CL_INVALID_VALUE, + throw pyopencl::error("Image", CL_INVALID_VALUE, "invalid dimension"); try @@ -2450,7 +2450,7 @@ namespace pyopencl void set_arg_null(cl_uint arg_index) { cl_mem m = 0; - PYOPENCL_CALL_GUARDED(clSetKernelArg, (m_kernel, arg_index, + PYOPENCL_CALL_GUARDED(clSetKernelArg, (m_kernel, arg_index, sizeof(cl_mem), &m)); } @@ -2786,14 +2786,14 @@ namespace pyopencl (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(gl_renderbuffer, + 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, + context &ctx, cl_mem_flags flags, + GLenum texture_target, GLint miplevel, GLuint texture, unsigned dims) { if (dims == 2) @@ -2801,7 +2801,7 @@ namespace pyopencl else if (dims == 3) return create_from_gl_texture_3d(ctx, flags, texture_target, miplevel, texture); else - throw pyopencl::error("Image", CL_INVALID_VALUE, + throw pyopencl::error("Image", CL_INVALID_VALUE, "invalid dimension"); } @@ -2878,7 +2878,7 @@ namespace pyopencl switch (param_name) { - case CL_CURRENT_DEVICE_FOR_GL_CONTEXT_KHR: + case CL_CURRENT_DEVICE_FOR_GL_CONTEXT_KHR: { cl_device_id param_value; PYOPENCL_CALL_GUARDED(func_ptr, diff --git a/test/test_wrapper.py b/test/test_wrapper.py index 7bd9e26801f6f40016e4632a98df79087f7f3309..56b79369fa159d6ecf5c21507966b823a83a22c2 100644 --- a/test/test_wrapper.py +++ b/test/test_wrapper.py @@ -29,7 +29,7 @@ class TestCL: failure_count = [0] CRASH_QUIRKS = [ - (("NVIDIA Corporation", "NVIDIA CUDA", + (("NVIDIA Corporation", "NVIDIA CUDA", "OpenCL 1.0 CUDA 3.0.1"), [ (cl.Event, cl.event_info.COMMAND_QUEUE), @@ -227,15 +227,15 @@ class TestCL: prg = cl.Program(context, """ __kernel void copy_image( - __global float4 *dest, - __read_only image2d_t src, + __global float4 *dest, + __read_only image2d_t src, sampler_t samp, int width) - { + { int x = get_global_id(0); int y = get_global_id(1); /* - const sampler_t samp = + const sampler_t samp = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST; @@ -293,7 +293,7 @@ def pytest_generate_tests(metafunc): arg_dict["context"] = cl.Context([device]) metafunc.addcall(funcargs=arg_dict.copy(), - id=", ".join("%s=%s" % (arg, value) + id=", ".join("%s=%s" % (arg, value) for arg, value in arg_dict.iteritems())) elif "platform" in metafunc.funcargnames: