From 6fe66162d257ef10397c5339f1df0b71213a81f1 Mon Sep 17 00:00:00 2001 From: Andreas Kloeckner Date: Sun, 23 Aug 2009 21:19:16 -0400 Subject: [PATCH] Complete wrapper, beginning docs. --- doc/Makefile | 73 +++ doc/make_constants.py | 41 ++ doc/source/conf.py | 173 ++++++ doc/source/constants.inc | 52 ++ doc/source/index.rst | 48 ++ doc/source/misc.rst | 70 +++ doc/source/reference.rst | 45 ++ doc/upload-docs.sh | 3 + pyopencl/__init__.py | 5 + src/wrapper/numpy_init.hpp | 23 + src/wrapper/wrap_cl.cpp | 160 ++++- src/wrapper/wrap_cl.hpp | 1107 ++++++++++++++++++++++++++++++++-- src/wrapper/wrap_helpers.hpp | 49 ++ 13 files changed, 1773 insertions(+), 76 deletions(-) create mode 100644 doc/Makefile create mode 100644 doc/make_constants.py create mode 100644 doc/source/conf.py create mode 100644 doc/source/constants.inc create mode 100644 doc/source/index.rst create mode 100644 doc/source/misc.rst create mode 100644 doc/source/reference.rst create mode 100755 doc/upload-docs.sh create mode 100644 src/wrapper/numpy_init.hpp diff --git a/doc/Makefile b/doc/Makefile new file mode 100644 index 00000000..2b787d23 --- /dev/null +++ b/doc/Makefile @@ -0,0 +1,73 @@ +# Makefile for Sphinx documentation +# + +# You can set these variables from the command line. +SPHINXOPTS = +SPHINXBUILD = python `which sphinx-build` +PAPER = + +# Internal variables. +PAPEROPT_a4 = -D latex_paper_size=a4 +PAPEROPT_letter = -D latex_paper_size=letter +ALLSPHINXOPTS = -d build/doctrees $(PAPEROPT_$(PAPER)) $(SPHINXOPTS) source + +.PHONY: help clean html web pickle htmlhelp latex changes linkcheck + +help: + @echo "Please use \`make ' where is one of" + @echo " html to make standalone HTML files" + @echo " pickle to make pickle files (usable by e.g. sphinx-web)" + @echo " htmlhelp to make HTML files and a HTML help project" + @echo " latex to make LaTeX files, you can set PAPER=a4 or PAPER=letter" + @echo " changes to make an overview over all changed/added/deprecated items" + @echo " linkcheck to check all external links for integrity" + +constants: + python make_constants.py > source/constants.inc + +clean: + -rm -rf build/* + +html: constants + mkdir -p build/html build/doctrees source/.static + $(SPHINXBUILD) -b html $(ALLSPHINXOPTS) build/html + @echo + @echo "Build finished. The HTML pages are in build/html." + +pickle: constants + mkdir -p build/pickle build/doctrees + $(SPHINXBUILD) -b pickle $(ALLSPHINXOPTS) build/pickle + @echo + @echo "Build finished; now you can process the pickle files or run" + @echo " sphinx-web build/pickle" + @echo "to start the sphinx-web server." + +web: pickle + +htmlhelp: constants + mkdir -p build/htmlhelp build/doctrees + $(SPHINXBUILD) -b htmlhelp $(ALLSPHINXOPTS) build/htmlhelp + @echo + @echo "Build finished; now you can run HTML Help Workshop with the" \ + ".hhp project file in build/htmlhelp." + +latex: constants + mkdir -p build/latex build/doctrees + $(SPHINXBUILD) -b latex $(ALLSPHINXOPTS) build/latex + @echo + @echo "Build finished; the LaTeX files are in build/latex." + @echo "Run \`make all-pdf' or \`make all-ps' in that directory to" \ + "run these through (pdf)latex." + +changes: constants + mkdir -p build/changes build/doctrees + $(SPHINXBUILD) -b changes $(ALLSPHINXOPTS) build/changes + @echo + @echo "The overview file is in build/changes." + +linkcheck: constants + mkdir -p build/linkcheck build/doctrees + $(SPHINXBUILD) -b linkcheck $(ALLSPHINXOPTS) build/linkcheck + @echo + @echo "Link check complete; look for any errors in the above output " \ + "or in build/linkcheck/output.txt." diff --git a/doc/make_constants.py b/doc/make_constants.py new file mode 100644 index 00000000..d0ebe118 --- /dev/null +++ b/doc/make_constants.py @@ -0,0 +1,41 @@ +import pyopencl as cl + +def doc_class(cls): + print ".. class :: %s" % cls.__name__ + print + for i in sorted(dir(cls)): + if not i.startswith("_"): + print " .. attribute :: %s" % i + print + + +print ".. This is an automatically generated file. DO NOT EDIT" +for cls in [ + cl.platform_info, + cl.device_type, + cl.device_info, + cl.device_fp_config, + cl.device_mem_cache_type, + cl.device_local_mem_type, + cl.command_queue_properties, + cl.context_info, + cl.context_properties, + cl.command_queue_info, + cl.mem_flags, + cl.channel_order, + cl.mem_object_type, + cl.mem_info, + cl.image_info, + cl.addressing_mode, + cl.filter_mode, + cl.sampler_info, + cl.map_flags, + cl.program_info, + cl.program_build_info, + cl.kernel_info, + cl.kernel_work_group_info, + cl.event_info, + cl.command_execution_status, + cl.profiling_info, + ]: + doc_class(cls) diff --git a/doc/source/conf.py b/doc/source/conf.py new file mode 100644 index 00000000..552dd8a8 --- /dev/null +++ b/doc/source/conf.py @@ -0,0 +1,173 @@ +# -*- coding: utf-8 -*- +# +# PyOpenCL documentation build configuration file, created by +# sphinx-quickstart on Fri Jun 13 00:51:19 2008. +# +# This file is execfile()d with the current directory set to its containing dir. +# +# The contents of this file are pickled, so don't put values in the namespace +# that aren't pickleable (module imports are okay, they're removed automatically). +# +# All configuration values have a default value; values that are commented out +# serve to show the default value. + +import sys, os + +# If your extensions are in another directory, add it here. If the directory +# is relative to the documentation root, use os.path.abspath to make it +# absolute, like shown here. +#sys.path.append(os.path.abspath('some/directory')) + +# General configuration +# --------------------- + +# Add any Sphinx extension module names here, as strings. They can be extensions +# coming with Sphinx (named 'sphinx.ext.*') or your custom ones. +extensions = [ + 'sphinx.ext.intersphinx', + ] + +# Add any paths that contain templates here, relative to this directory. +templates_path = ['.templates'] + +# The suffix of source filenames. +source_suffix = '.rst' + +# The master toctree document. +master_doc = 'index' + +# General substitutions. +project = 'PyOpenCL' +copyright = '2009, Andreas Kloeckner' + +# The default replacements for |version| and |release|, also used in various +# other places throughout the built documents. +# +# The short X.Y version. +ver_dic = {} +execfile("../../pyopencl/__init__.py", ver_dic) +version = ".".join(str(x) for x in ver_dic["VERSION"]) +# The full version, including alpha/beta/rc tags. +release = ver_dic["VERSION_TEXT"] + +# There are two options for replacing |today|: either, you set today to some +# non-false value, then it is used: +#today = '' +# Else, today_fmt is used as the format for a strftime call. +today_fmt = '%B %d, %Y' + +# List of documents that shouldn't be included in the build. +#unused_docs = [] + +# List of directories, relative to source directories, that shouldn't be searched +# for source files. +#exclude_dirs = [] + +# If true, '()' will be appended to :func: etc. cross-reference text. +#add_function_parentheses = True + +# If true, the current module name will be prepended to all description +# unit titles (such as .. function::). +#add_module_names = True + +# If true, sectionauthor and moduleauthor directives will be shown in the +# output. They are ignored by default. +#show_authors = False + +# The name of the Pygments (syntax highlighting) style to use. +pygments_style = 'sphinx' + + +# Options for HTML output +# ----------------------- + +# The style sheet to use for HTML and HTML Help pages. A file of that name +# must exist either in Sphinx' static/ path, or in one of the custom paths +# given in html_static_path. +html_style = 'default.css' + +# The name for this set of Sphinx documents. If None, it defaults to +# " v documentation". +#html_title = None + +# The name of an image file (within the static path) to place at the top of +# the sidebar. +#html_logo = None + +# Add any paths that contain custom static files (such as style sheets) here, +# relative to this directory. They are copied after the builtin static files, +# so a file named "default.css" will overwrite the builtin "default.css". +html_static_path = ['.static'] + +# If not '', a 'Last updated on:' timestamp is inserted at every page bottom, +# using the given strftime format. +html_last_updated_fmt = '%b %d, %Y' + +# If true, SmartyPants will be used to convert quotes and dashes to +# typographically correct entities. +#html_use_smartypants = True + +# Custom sidebar templates, maps document names to template names. +#html_sidebars = {} + +# Additional templates that should be rendered to pages, maps page names to +# template names. +#html_additional_pages = {} + +# If false, no module index is generated. +#html_use_modindex = True + +# If true, the reST sources are included in the HTML build as _sources/. +#html_copy_source = True + +# If true, an OpenSearch description file will be output, and all pages will +# contain a tag referring to it. The value of this option must be the +# base URL from which the finished HTML is served. +#html_use_opensearch = '' + +# If nonempty, this is the file name suffix for HTML files (e.g. ".xhtml"). +#html_file_suffix = '' + +# Output file base name for HTML help builder. +htmlhelp_basename = 'PyCudadoc' + + +# Options for LaTeX output +# ------------------------ + +# The paper size ('letter' or 'a4'). +#latex_paper_size = 'letter' + +# The font size ('10pt', '11pt' or '12pt'). +#latex_font_size = '10pt' + +# Grouping the document tree into LaTeX files. List of tuples +# (source start file, target name, title, author, document class [howto/manual]). +latex_documents = [ + ('index', 'pyopencl.tex', 'PyOpenCL Documentation', 'Andreas Kloeckner', 'manual'), +] + +# The name of an image file (relative to this directory) to place at the top of +# the title page. +#latex_logo = None + +# For "manual" documents, if this is true, then toplevel headings are parts, +# not chapters. +#latex_use_parts = False + +# Additional stuff for the LaTeX preamble. +#latex_preamble = '' + +# Documents to append as an appendix to all manuals. +#latex_appendices = [] + +# If false, no module index is generated. +#latex_use_modindex = True + +intersphinx_mapping = { + 'http://docs.python.org/dev': None, + 'http://docs.scipy.org/doc/numpy/': None, + 'http://documen.tician.de/codepy/': None, + 'http://documen.tician.de/boostmpi/': None, + } + diff --git a/doc/source/constants.inc b/doc/source/constants.inc new file mode 100644 index 00000000..9595f681 --- /dev/null +++ b/doc/source/constants.inc @@ -0,0 +1,52 @@ +.. This is an automatically generated file. DO NOT EDIT +.. class :: device_info + + .. attribute :: ADDRESS_BITS + .. attribute :: AVAILABLE + .. attribute :: COMPILER_AVAILABLE + .. attribute :: ENDIAN_LITTLE + .. attribute :: ERROR_CORRECTION_SUPPORT + .. attribute :: EXECUTION_CAPABILITIES + .. attribute :: EXTENSIONS + .. attribute :: GLOBAL_MEM_CACHELINE_SIZE + .. attribute :: GLOBAL_MEM_CACHE_SIZE + .. attribute :: GLOBAL_MEM_CACHE_TYPE + .. attribute :: GLOBAL_MEM_SIZE + .. attribute :: IMAGE2D_MAX_HEIGHT + .. attribute :: IMAGE2D_MAX_WIDTH + .. attribute :: IMAGE3D_MAX_DEPTH + .. attribute :: IMAGE3D_MAX_HEIGHT + .. attribute :: IMAGE3D_MAX_WIDTH + .. attribute :: IMAGE_SUPPORT + .. attribute :: LOCAL_MEM_SIZE + .. attribute :: LOCAL_MEM_TYPE + .. attribute :: MAX_CLOCK_FREQUENCY + .. attribute :: MAX_COMPUTE_UNITS + .. attribute :: MAX_CONSTANT_ARGS + .. attribute :: MAX_CONSTANT_BUFFER_SIZE + .. attribute :: MAX_MEM_ALLOC_SIZE + .. attribute :: MAX_PARAMETER_SIZE + .. attribute :: MAX_READ_IMAGE_ARGS + .. attribute :: MAX_SAMPLERS + .. attribute :: MAX_WORK_GROUP_SIZE + .. attribute :: MAX_WORK_ITEM_DIMENSIONS + .. attribute :: MAX_WORK_ITEM_SIZES + .. attribute :: MAX_WRITE_IMAGE_ARGS + .. attribute :: MEM_BASE_ADDR_ALIGN + .. attribute :: MIN_DATA_TYPE_ALIGN_SIZE + .. attribute :: NAME + .. attribute :: PREFERRED_VECTOR_WIDTH_CHAR + .. attribute :: PREFERRED_VECTOR_WIDTH_DOUBLE + .. attribute :: PREFERRED_VECTOR_WIDTH_FLOAT + .. attribute :: PREFERRED_VECTOR_WIDTH_INT + .. attribute :: PREFERRED_VECTOR_WIDTH_LONG + .. attribute :: PREFERRED_VECTOR_WIDTH_SHORT + .. attribute :: PROFILE + .. attribute :: PROFILING_TIMER_RESOLUTION + .. attribute :: QUEUE_PROPERTIES + .. attribute :: SINGLE_FP_CONFIG + .. attribute :: TYPE + .. attribute :: VENDOR + .. attribute :: VENDOR_ID + .. attribute :: VERSION + diff --git a/doc/source/index.rst b/doc/source/index.rst new file mode 100644 index 00000000..cc1d6b18 --- /dev/null +++ b/doc/source/index.rst @@ -0,0 +1,48 @@ +Welcome to PyOpenCL's documentation! +==================================== + +PyOpenCL gives you easy, Pythonic access to the `OpenCL +`_ parallel computation API. +What makes PyOpenCL special? + +* Object cleanup tied to lifetime of objects. This idiom, + often called + `RAII `_ + in C++, makes it much easier to write correct, leak- and + crash-free code. PyOpenCL knows about dependencies, too, so (for example) + it won't detach from a context before all memory allocated in it is also + freed. + +* Completeness. PyOpenCL puts the full power of OpenCL's API at your + disposal, if you wish. + +* Automatic Error Checking. All errors are automatically translated + into Python exceptions. + +* Speed. PyOpenCL's base layer is written in C++, so all the niceties above + are virtually free. + +* Helpful Documentation. You're looking at it. ;) + +Contents +======== + +.. toctree:: + :maxdepth: 2 + + misc + reference + +Note that this guide does not explain OpenCL programming and technology. Please +refer to generic OpenCL tutorials for that. + +PyOpenCL also has its own `web site `_, +where you can find updates, new versions, documentation, and support. + +Indices and tables +================== + +* :ref:`genindex` +* :ref:`modindex` +* :ref:`search` + diff --git a/doc/source/misc.rst b/doc/source/misc.rst new file mode 100644 index 00000000..2dcb914f --- /dev/null +++ b/doc/source/misc.rst @@ -0,0 +1,70 @@ +Installation +============ + +Installation information is maintained collaboratively on the +`PyOpenCL Wiki `_. + +User-visible Changes +==================== + +Version 0.90 +------------ + +* Initial release. + +Acknowledgments +=============== + +* Gert Wohlgemuth ported PyCUDA to MacOS X and contributed large parts of + :class:`pycuda.gpuarray.GPUArray`. +* Alexander Mordvintsev contributed fixes for Windows XP. +* Cosmin Stejerean provided multiple patches for PyCUDA's build system. +* Tom Annau contributed an alternative SourceModule compiler cache as well + as Windows build insight. +* Nicholas Tung improved PyCUDA's documentation. +* Jozef Vesely contributed a massively improved random number generator derived from + the RSA Data Security, Inc. MD5 Message Digest Algorithm. +* Chris Heuser provided a test cases for multi-threaded PyCUDA. +* The reduction templating is based on code by Mark Harris at Nvidia. +* Andrew Wagner provided a test case and contributed the port of the + convolution example. The original convolution code is based on an + example provided by Nvidia. +* Hendrik Riedmann contributed the matrix transpose and list selection + examples. +* Peter Berrington contributed a working example for CUDA-OpenGL + interoperability. + +Licensing +========= + +PyOpenCL is licensed to you under the MIT/X Consortium license: + +Copyright (c) 2009 Andreas Klöckner and Contributors. + +Permission is hereby granted, free of charge, to any person +obtaining a copy of this software and associated documentation +files (the "Software"), to deal in the Software without +restriction, including without limitation the rights to use, +copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the +Software is furnished to do so, subject to the following +conditions: + +The above copyright notice and this permission notice shall be +included in all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, +EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES +OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND +NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT +HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY, +WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING +FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR +OTHER DEALINGS IN THE SOFTWARE. + +Frequently Asked Questions +========================== + +The FAQ is maintained collaboratively on the +`Wiki FAQ page `_. + diff --git a/doc/source/reference.rst b/doc/source/reference.rst new file mode 100644 index 00000000..7e1049fe --- /dev/null +++ b/doc/source/reference.rst @@ -0,0 +1,45 @@ +.. _reference-doc: + +Reference Documentation +======================= + +Version Queries +--------------- + +.. module:: pyopencl +.. moduleauthor:: Andreas Kloeckner + +.. data:: VERSION + + Gives the numeric version of PyCUDA as a variable-length tuple + of integers. Enables easy version checks such as + *VERSION >= (0, 93)*. + +.. data:: VERSION_STATUS + + A text string such as `"rc4"` or `"beta"` qualifying the status + of the release. + +.. data:: VERSION_TEXT + + The full release name (such as `"0.93rc4"`) in string form. + +.. _errors: + +Error Reporting +--------------- + +.. class:: Error + + Base class for all PyOpenCL exceptions. + +.. class:: MemoryError + +.. class:: LogicError + +.. class:: RuntimeError + +Constants +--------- + +.. include:: constants.inc diff --git a/doc/upload-docs.sh b/doc/upload-docs.sh new file mode 100755 index 00000000..9f51d44c --- /dev/null +++ b/doc/upload-docs.sh @@ -0,0 +1,3 @@ +#! /bin/sh + +rsync --progress --verbose --archive --delete build/html/* tikernet@tiker.net:public_html/doc/pyopencl diff --git a/pyopencl/__init__.py b/pyopencl/__init__.py index e69de29b..1d55cca5 100644 --- a/pyopencl/__init__.py +++ b/pyopencl/__init__.py @@ -0,0 +1,5 @@ +VERSION = (0, 90) +VERSION_STATUS = "alpha" +VERSION_TEXT = ".".join(str(x) for x in VERSION) + VERSION_STATUS + +from pyopencl._cl import * diff --git a/src/wrapper/numpy_init.hpp b/src/wrapper/numpy_init.hpp new file mode 100644 index 00000000..03e182c4 --- /dev/null +++ b/src/wrapper/numpy_init.hpp @@ -0,0 +1,23 @@ +#ifndef _FAYHVVAAA_PYCUDA_HEADER_SEEN_NUMPY_INIT_HPP + + + + +#include + + + + +namespace +{ + static struct array_importer + { + array_importer() + { import_array(); } + } _array_importer; +} + + + + +#endif diff --git a/src/wrapper/wrap_cl.cpp b/src/wrapper/wrap_cl.cpp index 5304ec25..42764b47 100644 --- a/src/wrapper/wrap_cl.cpp +++ b/src/wrapper/wrap_cl.cpp @@ -65,25 +65,6 @@ namespace class command_type { }; class command_execution_status { }; class profiling_info { }; - - - - py::list get_platforms(cl_platform_info param_name) - { - cl_uint num_platforms = 0; - PYOPENCL_CALL_GUARDED(clGetPlatformIDs, (0, 0, &num_platforms)); - - std::vector platforms(num_platforms); - PYOPENCL_CALL_GUARDED(clGetPlatformIDs, - (num_platforms, platforms.data(), &num_platforms)); - - py::list result; - BOOST_FOREACH(cl_platform_id pid, platforms) - result.append(handle_from_new_ptr( - new platform(pid))); - - return result; - } } @@ -423,6 +404,8 @@ BOOST_PYTHON_MODULE(_cl) py::class_("Platform", py::no_init) .DEF_SIMPLE_METHOD(get_info) .DEF_SIMPLE_METHOD(get_devices) + .def(py::self == py::self) + .def(py::self != py::self) ; } @@ -430,6 +413,8 @@ BOOST_PYTHON_MODULE(_cl) typedef device cls; py::class_("Device", py::no_init) .DEF_SIMPLE_METHOD(get_info) + .def(py::self == py::self) + .def(py::self != py::self) ; } @@ -438,6 +423,8 @@ BOOST_PYTHON_MODULE(_cl) py::class_("Context", py::init()) .DEF_SIMPLE_METHOD(get_info) + .def(py::self == py::self) + .def(py::self != py::self) ; } @@ -450,12 +437,16 @@ BOOST_PYTHON_MODULE(_cl) .DEF_SIMPLE_METHOD(set_property) .DEF_SIMPLE_METHOD(flush) .DEF_SIMPLE_METHOD(finish) + .def(py::self == py::self) + .def(py::self != py::self) ; } { typedef event cls; py::class_("Event", py::no_init) .DEF_SIMPLE_METHOD(get_info) + .def(py::self == py::self) + .def(py::self != py::self) ; } @@ -468,6 +459,9 @@ BOOST_PYTHON_MODULE(_cl) typedef memory_object cls; py::class_("MemoryObject", py::no_init) .DEF_SIMPLE_METHOD(get_info) + .DEF_SIMPLE_METHOD(release) + .def(py::self == py::self) + .def(py::self != py::self) ; } @@ -489,6 +483,134 @@ BOOST_PYTHON_MODULE(_cl) py::arg("is_blocking")=false), py::return_value_policy()); + // image -------------------------------------------------------------------- + { + typedef cl_image_format cls; + py::class_("ImageFormat") + .def_readwrite("channel_order", &cls::image_channel_order) + .def_readwrite("channel_data_type", &cls::image_channel_data_type) + ; + } + + DEF_SIMPLE_FUNCTION(get_supported_image_formats); + py::def("create_image_2d", create_image_2D, + (py::args("ctx", "flags", "format", "width", "height", "pitch"), + py::arg("buffer")=py::object()), + py::return_value_policy()); + py::def("create_image_3d", create_image_3D, + (py::args("ctx", "flags", "format", "width", "height", "depth", + "row_pitch", "slice_pitch"), + py::arg("buffer")=py::object()), + py::return_value_policy()); + + py::def("enqueue_read_image", enqueue_read_image, + (py::args("cq", "mem", "origin", "region", + "row_pitch", "slice_pitch", "buffer"), + py::arg("wait_for")=py::object(), + py::arg("is_blocking")=false), + py::return_value_policy()); + py::def("enqueue_write_image", enqueue_write_image, + (py::args("cq", "mem", "origin", "region", + "row_pitch", "slice_pitch", "buffer"), + py::arg("wait_for")=py::object(), + py::arg("is_blocking")=false), + py::return_value_policy()); + + py::def("enqueue_copy_image", enqueue_copy_image, + (py::args("cq", "src", "dest", "src_origin", "dest_origin", "region"), + py::arg("wait_for")=py::object()), + py::return_value_policy()); + py::def("enqueue_copy_image_to_buffer", enqueue_copy_image_to_buffer, + (py::args("cq", "src", "dest", "origin", "region", "offset"), + py::arg("wait_for")=py::object()), + py::return_value_policy()); + py::def("enqueue_copy_buffer_to_image", enqueue_copy_image_to_buffer, + (py::args("cq", "src", "dest", "offset", "origin", "region"), + py::arg("wait_for")=py::object()), + py::return_value_policy()); + + // memory_map --------------------------------------------------------------- + { + typedef memory_map cls; + py::class_("MemoryMap", py::no_init) + .def("release", &cls::release, + (py::arg("cq")=0, py::arg("wait_for")=py::object()), + py::return_value_policy()) + ; + } + + py::def("enqueue_map_buffer", enqueue_map_buffer, + (py::args("cq", "buf", "flags", + "offset", + "shape", "dtype", "order"), + py::arg("wait_for")=py::object(), + py::arg("is_blocking")=false)); + py::def("enqueue_map_image", enqueue_map_image, + (py::args("cq", "buf", "flags", + "origin", "region", + "shape", "dtype", "order"), + py::arg("wait_for")=py::object(), + py::arg("is_blocking")=false)); + + + // sampler ------------------------------------------------------------------ + { + typedef sampler cls; + py::class_("Sampler", + py::init()) + .DEF_SIMPLE_METHOD(get_info) + .def(py::self == py::self) + .def(py::self != py::self) + ; + } + + // program ------------------------------------------------------------------ + { + typedef program cls; + py::class_("Program", py::no_init) + .DEF_SIMPLE_METHOD(get_info) + .DEF_SIMPLE_METHOD(get_build_info) + .DEF_SIMPLE_METHOD(build) + .def(py::self == py::self) + .def(py::self != py::self) + ; + } + + py::def("unload_compiler", clUnloadCompiler); + + py::def("create_program_with_source", create_program_with_source, + py::args("ctx", "src"), + py::return_value_policy()); + py::def("create_program_with_binary", create_program_with_binary, + py::args("ctx", "devices", "binaries"), + py::return_value_policy()); + + { + typedef kernel cls; + py::class_("Kernel", + py::init()) + .DEF_SIMPLE_METHOD(get_info) + .DEF_SIMPLE_METHOD(get_work_group_info) + .DEF_SIMPLE_METHOD(set_arg) + .def(py::self == py::self) + .def(py::self != py::self) + ; + } + DEF_SIMPLE_FUNCTION(create_kernels_in_program); + py::def("enqueue_nd_range_kernel", enqueue_nd_range_kernel, + (py::args("cmd_queue", "kernel"), + py::arg("global_work_size"), + py::arg("local_work_size"), + py::arg("global_work_offset")=py::object(), + py::arg("wait_for")=py::object() + ), + py::return_value_policy()); + py::def("enqueue_task", enqueue_task, + (py::args("cmd_queue", "kernel"), + py::arg("wait_for")=py::object() + ), + py::return_value_policy()); + // TODO: clEnqueueNativeKernel } diff --git a/src/wrapper/wrap_cl.hpp b/src/wrapper/wrap_cl.hpp index 40b14f57..7d0de09d 100644 --- a/src/wrapper/wrap_cl.hpp +++ b/src/wrapper/wrap_cl.hpp @@ -4,9 +4,7 @@ -// TODO: Images and samplers // TODO: Memory mapping -// TODO: Profiling // TODO: GL Interop @@ -19,6 +17,7 @@ #include #include #include "wrap_helpers.hpp" +#include "numpy_init.hpp" @@ -188,6 +187,16 @@ +// equality testing ----------------------------------------------------------- +#define PYOPENCL_EQUALITY_TESTS(cls) \ + bool operator==(cls const &other) const \ + { return data() == other.data(); } \ + bool operator!=(cls const &other) const \ + { return data() != other.data(); } + + + + namespace pyopencl @@ -302,6 +311,13 @@ namespace pyopencl : m_platform(pid) { } + cl_platform_id data() const + { + return m_platform; + } + + PYOPENCL_EQUALITY_TESTS(platform); + py::object get_info(cl_platform_info param_name) const { switch (param_name) @@ -319,16 +335,31 @@ namespace pyopencl } py::list get_devices(cl_device_type devtype); - - cl_platform_id data() const - { - return m_platform; - } }; + py::list get_platforms(cl_platform_info param_name) + { + cl_uint num_platforms = 0; + PYOPENCL_CALL_GUARDED(clGetPlatformIDs, (0, 0, &num_platforms)); + + std::vector platforms(num_platforms); + PYOPENCL_CALL_GUARDED(clGetPlatformIDs, + (num_platforms, platforms.data(), &num_platforms)); + + py::list result; + BOOST_FOREACH(cl_platform_id pid, platforms) + result.append(handle_from_new_ptr( + new platform(pid))); + + return result; + } + + + + // device ------------------------------------------------------------------- class device : boost::noncopyable { @@ -344,13 +375,19 @@ namespace pyopencl : m_device(did) { } + cl_device_id data() const + { + return m_device; + } + + PYOPENCL_EQUALITY_TESTS(device); py::object get_info(cl_device_info param_name) const { #define DEV_GET_INT_INF(TYPE) \ PYOPENCL_GET_INTEGRAL_INFO(Device, m_device, param_name, TYPE); - switch (param_name) + switch (param_name) { case CL_DEVICE_TYPE: DEV_GET_INT_INF(cl_device_type); case CL_DEVICE_VENDOR_ID: DEV_GET_INT_INF(cl_uint); @@ -358,7 +395,7 @@ namespace pyopencl case CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS: DEV_GET_INT_INF(cl_uint); case CL_DEVICE_MAX_WORK_GROUP_SIZE: DEV_GET_INT_INF(cl_uint); - case CL_DEVICE_MAX_WORK_ITEM_SIZES: + case CL_DEVICE_MAX_WORK_ITEM_SIZES: { std::vector result; PYOPENCL_GET_VEC_INFO(Device, m_device, param_name, result); @@ -420,11 +457,6 @@ namespace pyopencl throw error("Platform.get_info", CL_INVALID_VALUE); } } - - cl_device_id data() const - { - return m_device; - } }; @@ -461,7 +493,7 @@ namespace pyopencl public: context(cl_context ctx, bool retain) : m_context(ctx) - { + { if (retain) PYOPENCL_CALL_GUARDED(clRetainContext, (ctx)); } @@ -476,7 +508,7 @@ namespace pyopencl { if (len(prop_tuple) != 2) throw error("Context", CL_INVALID_VALUE, "property tuple must have length 2"); - cl_context_properties prop = + cl_context_properties prop = py::extract(prop_tuple[0]); props.push_back(prop); @@ -505,25 +537,33 @@ namespace pyopencl devices.data(), 0, 0, &status_code); + PYOPENCL_PRINT_CALL_TRACE("clCreateContext"); if (status_code != CL_SUCCESS) throw pyopencl::error("Context", status_code); } ~context() { - PYOPENCL_CALL_GUARDED_CLEANUP(clReleaseContext, + PYOPENCL_CALL_GUARDED_CLEANUP(clReleaseContext, (m_context)); } + cl_context data() const + { + return m_context; + } + + PYOPENCL_EQUALITY_TESTS(context); + py::object get_info(cl_context_info param_name) const { - switch (param_name) + switch (param_name) { case CL_CONTEXT_REFERENCE_COUNT: PYOPENCL_GET_INTEGRAL_INFO( Context, m_context, param_name, cl_uint); - case CL_CONTEXT_DEVICES: + case CL_CONTEXT_DEVICES: { std::vector result; PYOPENCL_GET_VEC_INFO(Context, m_context, param_name, result); @@ -541,11 +581,6 @@ namespace pyopencl throw error("Context.get_info", CL_INVALID_VALUE); } } - - cl_context data() const - { - return m_context; - } }; @@ -565,6 +600,12 @@ namespace pyopencl PYOPENCL_CALL_GUARDED(clRetainCommandQueue, (q)); } + command_queue(command_queue const &src) + : m_queue(src.m_queue) + { + PYOPENCL_CALL_GUARDED(clRetainCommandQueue, (m_queue)); + } + command_queue( const context &ctx, const device &dev, @@ -577,19 +618,25 @@ namespace pyopencl props, &status_code); + PYOPENCL_PRINT_CALL_TRACE("clCreateCommandQueue"); if (status_code != CL_SUCCESS) throw pyopencl::error("CommandQueue", status_code); } ~command_queue() { - PYOPENCL_CALL_GUARDED_CLEANUP(clReleaseCommandQueue, + PYOPENCL_CALL_GUARDED_CLEANUP(clReleaseCommandQueue, (m_queue)); } + const cl_command_queue data() const + { return m_queue; } + + PYOPENCL_EQUALITY_TESTS(command_queue); + py::object get_info(cl_command_queue_info param_name) const { - switch (param_name) + switch (param_name) { case CL_QUEUE_CONTEXT: PYOPENCL_GET_OPAQUE_INFO(CommandQueue, m_queue, param_name, @@ -598,10 +645,10 @@ namespace pyopencl PYOPENCL_GET_OPAQUE_INFO(CommandQueue, m_queue, param_name, cl_device_id, device); case CL_QUEUE_REFERENCE_COUNT: - PYOPENCL_GET_INTEGRAL_INFO(CommandQueue, m_queue, param_name, + PYOPENCL_GET_INTEGRAL_INFO(CommandQueue, m_queue, param_name, cl_uint); case CL_QUEUE_PROPERTIES: - PYOPENCL_GET_INTEGRAL_INFO(CommandQueue, m_queue, param_name, + PYOPENCL_GET_INTEGRAL_INFO(CommandQueue, m_queue, param_name, cl_command_queue_properties); default: @@ -614,7 +661,7 @@ namespace pyopencl bool enable) { cl_command_queue_properties old_prop; - PYOPENCL_CALL_GUARDED(clSetCommandQueueProperty, + PYOPENCL_CALL_GUARDED(clSetCommandQueueProperty, (m_queue, prop, PYOPENCL_CAST_BOOL(enable), &old_prop)); return old_prop; } @@ -623,9 +670,6 @@ namespace pyopencl { PYOPENCL_CALL_GUARDED(clFlush, (m_queue)); } void finish() { PYOPENCL_CALL_GUARDED_THREADED(clFlush, (m_queue)); } - - const cl_command_queue data() const - { return m_queue; } }; @@ -640,32 +684,41 @@ namespace pyopencl public: event(cl_event event, bool retain) : m_event(event) - { + { if (retain) PYOPENCL_CALL_GUARDED(clRetainEvent, (event)); } + event(event const &src) + : m_event(src.m_event) + { PYOPENCL_CALL_GUARDED(clRetainEvent, (m_event)); } + ~event() { - PYOPENCL_CALL_GUARDED_CLEANUP(clReleaseEvent, + PYOPENCL_CALL_GUARDED_CLEANUP(clReleaseEvent, (m_event)); } - py::object get_info(cl_command_queue_info param_name) const + const cl_event data() const + { return m_event; } + + PYOPENCL_EQUALITY_TESTS(event); + + py::object get_info(cl_event_info param_name) const { - switch (param_name) + switch (param_name) { case CL_EVENT_COMMAND_QUEUE: PYOPENCL_GET_OPAQUE_INFO(Event, m_event, param_name, cl_command_queue, command_queue); case CL_EVENT_COMMAND_TYPE: - PYOPENCL_GET_INTEGRAL_INFO(Event, m_event, param_name, + PYOPENCL_GET_INTEGRAL_INFO(Event, m_event, param_name, cl_command_type); case CL_EVENT_COMMAND_EXECUTION_STATUS: - PYOPENCL_GET_INTEGRAL_INFO(Event, m_event, param_name, + PYOPENCL_GET_INTEGRAL_INFO(Event, m_event, param_name, cl_int); case CL_EVENT_REFERENCE_COUNT: - PYOPENCL_GET_INTEGRAL_INFO(Event, m_event, param_name, + PYOPENCL_GET_INTEGRAL_INFO(Event, m_event, param_name, cl_uint); default: @@ -673,8 +726,20 @@ namespace pyopencl } } - const cl_event data() const - { return m_event; } + py::object get_profiling_info(cl_profiling_info param_name) const + { + switch (param_name) + { + case CL_PROFILING_COMMAND_QUEUED: + case CL_PROFILING_COMMAND_SUBMIT: + case CL_PROFILING_COMMAND_START: + case CL_PROFILING_COMMAND_END: + PYOPENCL_GET_INTEGRAL_INFO(Event, m_event, param_name, + cl_ulong); + default: + throw error("Event.get_profiling_info", CL_INVALID_VALUE); + } + } }; @@ -686,7 +751,7 @@ namespace pyopencl std::vector event_wait_list(len(events)); PYTHON_FOREACH(evt, events) - event_wait_list[num_events_in_wait_list++] = + event_wait_list[num_events_in_wait_list++] = py::extract(evt)().data(); PYOPENCL_CALL_GUARDED_THREADED(clWaitForEvents, ( @@ -737,13 +802,14 @@ namespace pyopencl class memory_object : boost::noncopyable { private: - cl_mem m_mem; + bool m_valid; + cl_mem m_mem; py::object m_hostbuf; public: memory_object(cl_mem mem, bool retain, py::object *hostbuf=0) - : m_mem(mem) - { + : m_valid(true), m_mem(mem) + { if (retain) PYOPENCL_CALL_GUARDED(clRetainMemObject, (mem)); @@ -751,35 +817,61 @@ namespace pyopencl m_hostbuf = *hostbuf; } - ~memory_object() + memory_object(memory_object &src) + : m_valid(true), m_mem(src.m_mem), m_hostbuf(src.m_hostbuf) + { + PYOPENCL_CALL_GUARDED(clRetainMemObject, (m_mem)); + } + + void release() { - PYOPENCL_CALL_GUARDED_CLEANUP(clReleaseMemObject, + if (!m_valid) + throw error("MemoryObject.free", CL_INVALID_VALUE, + "trying to double-unref mem object"); + PYOPENCL_CALL_GUARDED_CLEANUP(clReleaseMemObject, (m_mem)); + m_valid = false; + } + + ~memory_object() + { + if (m_valid) + release(); } const cl_mem data() const { return m_mem; } + PYOPENCL_EQUALITY_TESTS(memory_object); + + size_t size() const + { + size_t param_value; + PYOPENCL_CALL_GUARDED(clGetMemObjectInfo, + (m_mem, CL_MEM_SIZE, sizeof(param_value), ¶m_value, 0)); + return param_value; + } + py::object get_info(cl_mem_info param_name) const { - switch (param_name) + switch (param_name) { case CL_MEM_TYPE: - PYOPENCL_GET_INTEGRAL_INFO(MemObject, m_mem, param_name, + PYOPENCL_GET_INTEGRAL_INFO(MemObject, m_mem, param_name, cl_mem_object_type); case CL_MEM_FLAGS: - PYOPENCL_GET_INTEGRAL_INFO(MemObject, m_mem, param_name, + PYOPENCL_GET_INTEGRAL_INFO(MemObject, m_mem, param_name, cl_mem_flags); case CL_MEM_SIZE: - PYOPENCL_GET_INTEGRAL_INFO(MemObject, m_mem, param_name, + PYOPENCL_GET_INTEGRAL_INFO(MemObject, m_mem, param_name, size_t); case CL_MEM_HOST_PTR: return m_hostbuf; case CL_MEM_MAP_COUNT: - PYOPENCL_GET_INTEGRAL_INFO(MemObject, m_mem, param_name, + PYOPENCL_GET_INTEGRAL_INFO(MemObject, m_mem, param_name, cl_uint); case CL_MEM_REFERENCE_COUNT: - PYOPENCL_GET_INTEGRAL_INFO(MemObject, m_mem, param_name, + PYOPENCL_GET_INTEGRAL_INFO(MemObject, m_mem, param_name, cl_uint); case CL_MEM_CONTEXT: PYOPENCL_GET_OPAQUE_INFO(MemObject, m_mem, param_name, @@ -792,10 +884,10 @@ namespace pyopencl py::object get_image_info(cl_image_info param_name) const { - switch (param_name) + switch (param_name) { case CL_IMAGE_FORMAT: - PYOPENCL_GET_INTEGRAL_INFO(Image, m_mem, param_name, + PYOPENCL_GET_INTEGRAL_INFO(Image, m_mem, param_name, cl_image_format); case CL_IMAGE_ELEMENT_SIZE: case CL_IMAGE_ROW_PITCH: @@ -820,10 +912,11 @@ namespace pyopencl size_t size) { cl_int status_code; + PYOPENCL_PRINT_CALL_TRACE("clCreateBuffer"); cl_mem mem = clCreateBuffer(ctx.data(), flags, size, 0, &status_code); if (status_code != CL_SUCCESS) - throw pyopencl::error("create_buffer", status_code); + throw pyopencl::error("clCreateBuffer", status_code); try { @@ -861,6 +954,7 @@ namespace pyopencl cl_int status_code; cl_mem mem = clCreateBuffer(ctx.data(), flags, len, buf, &status_code); + PYOPENCL_PRINT_CALL_TRACE("clCreateBuffer"); if (status_code != CL_SUCCESS) throw pyopencl::error("create_host_buffer", status_code); @@ -941,10 +1035,909 @@ namespace pyopencl )); PYOPENCL_RETURN_NEW_EVENT(evt); } -} + // images ------------------------------------------------------------------- + py::list get_supported_image_formats( + context const &ctx, + cl_mem_flags flags, + cl_mem_object_type image_type) + { + cl_uint num_image_formats; + PYOPENCL_CALL_GUARDED(clGetSupportedImageFormats, ( + 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, + num_image_formats, formats.data(), 0)); + return py::list(formats); + } + + + + +#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; \ + } \ + } + + + + +#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 + +#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 + + + + + event *enqueue_read_image( + command_queue &cq, + memory_object &mem, + py::object py_origin, py::object py_region, + size_t row_pitch, size_t slice_pitch, + py::object buffer, + py::object py_wait_for, + bool is_blocking + ) + { + PYOPENCL_PARSE_WAIT_FOR; + COPY_PY_COORD_TRIPLE(origin); + COPY_PY_REGION_TRIPLE(region); + + void *buf; + PYOPENCL_BUFFER_SIZE_T len; + + if (PyObject_AsWriteBuffer(buffer.ptr(), &buf, &len)) + throw py::error_already_set(); + + cl_event evt; + PYOPENCL_CALL_GUARDED(clEnqueueReadImage, ( + cq.data(), + mem.data(), + PYOPENCL_CAST_BOOL(is_blocking), + origin, region, row_pitch, slice_pitch, buf, + num_events_in_wait_list, event_wait_list.data(), &evt + )); + PYOPENCL_RETURN_NEW_EVENT(evt); + } + + + + + event *enqueue_write_image( + command_queue &cq, + memory_object &mem, + py::object py_origin, py::object py_region, + size_t row_pitch, size_t slice_pitch, + py::object buffer, + py::object py_wait_for, + bool is_blocking + ) + { + PYOPENCL_PARSE_WAIT_FOR; + COPY_PY_COORD_TRIPLE(origin); + COPY_PY_REGION_TRIPLE(region); + + const void *buf; + PYOPENCL_BUFFER_SIZE_T len; + + if (PyObject_AsReadBuffer(buffer.ptr(), &buf, &len)) + throw py::error_already_set(); + + cl_event evt; + PYOPENCL_CALL_GUARDED(clEnqueueWriteImage, ( + cq.data(), + mem.data(), + PYOPENCL_CAST_BOOL(is_blocking), + origin, region, row_pitch, slice_pitch, buf, + num_events_in_wait_list, event_wait_list.data(), &evt + )); + PYOPENCL_RETURN_NEW_EVENT(evt); + } + + + + + event *enqueue_copy_image( + command_queue &cq, + memory_object &src, + memory_object &dest, + py::object py_src_origin, + py::object py_dest_origin, + py::object py_region, + py::object py_wait_for + ) + { + PYOPENCL_PARSE_WAIT_FOR; + COPY_PY_COORD_TRIPLE(src_origin); + COPY_PY_COORD_TRIPLE(dest_origin); + COPY_PY_REGION_TRIPLE(region); + + cl_event evt; + PYOPENCL_CALL_GUARDED(clEnqueueCopyImage, ( + cq.data(), src.data(), dest.data(), + src_origin, dest_origin, region, + num_events_in_wait_list, event_wait_list.data(), &evt + )); + PYOPENCL_RETURN_NEW_EVENT(evt); + } + + + + + event *enqueue_copy_image_to_buffer( + command_queue &cq, + memory_object &src, + memory_object &dest, + py::object py_origin, + py::object py_region, + size_t offset, + py::object py_wait_for + ) + { + PYOPENCL_PARSE_WAIT_FOR; + COPY_PY_COORD_TRIPLE(origin); + COPY_PY_REGION_TRIPLE(region); + + cl_event evt; + PYOPENCL_CALL_GUARDED(clEnqueueCopyImageToBuffer, ( + cq.data(), src.data(), dest.data(), + origin, region, offset, + num_events_in_wait_list, event_wait_list.data(), &evt + )); + PYOPENCL_RETURN_NEW_EVENT(evt); + } + + + + + event *enqueue_copy_buffer_to_image( + command_queue &cq, + memory_object &src, + memory_object &dest, + size_t offset, + py::object py_origin, + py::object py_region, + py::object py_wait_for + ) + { + PYOPENCL_PARSE_WAIT_FOR; + COPY_PY_COORD_TRIPLE(origin); + COPY_PY_REGION_TRIPLE(region); + + cl_event evt; + PYOPENCL_CALL_GUARDED(clEnqueueCopyBufferToImage, ( + cq.data(), src.data(), dest.data(), + offset, origin, region, + num_events_in_wait_list, event_wait_list.data(), &evt + )); + PYOPENCL_RETURN_NEW_EVENT(evt); + } + + + + + // maps --------------------------------------------------------------------- + class memory_map + { + private: + bool m_valid; + command_queue m_queue; + memory_object m_mem; + void *m_ptr; + + public: + memory_map(command_queue &cq, memory_object &mem, void *ptr) + : m_valid(true), m_queue(cq), m_mem(mem), m_ptr(ptr) + { + } + + ~memory_map() + { + if (m_valid) + delete release(0, py::object()); + } + + event *release(command_queue *cq, py::object py_wait_for) + { + PYOPENCL_PARSE_WAIT_FOR; + + if (cq == 0) + cq = &m_queue; + + cl_event evt; + PYOPENCL_CALL_GUARDED(clEnqueueUnmapMemObject, ( + cq->data(), m_mem.data(), m_ptr, + num_events_in_wait_list, event_wait_list.data(), &evt + )); + + m_valid = false; + + PYOPENCL_RETURN_NEW_EVENT(evt); + } + }; + + + + + py::object enqueue_map_buffer( + command_queue &cq, + memory_object &buf, + cl_map_flags flags, + cl_map_flags offset, + py::object py_shape, py::object dtype, py::object order_py, + py::object py_wait_for, + bool is_blocking + ) + { + PYOPENCL_PARSE_WAIT_FOR; + PYOPENCL_PARSE_NUMPY_ARRAY_SPEC; + + py::handle<> result = py::handle<>(PyArray_NewFromDescr( + &PyArray_Type, tp_descr, + shape.size(), &shape.front(), /*strides*/ NULL, + 0, ary_flags, /*obj*/NULL)); + + cl_event evt; + cl_int status_code; + PYOPENCL_PRINT_CALL_TRACE("clEnqueueMapBuffer"); + void *mapped = clEnqueueMapBuffer( + cq.data(), buf.data(), + PYOPENCL_CAST_BOOL(is_blocking), flags, + offset, PyArray_NBYTES(result.get()), + num_events_in_wait_list, event_wait_list.data(), &evt, + &status_code); + if (status_code != CL_SUCCESS) + throw pyopencl::error("clEnqueueMapBuffer", status_code); + + event evt_handle(evt, false); + + std::auto_ptr map; + try + { + map = std::auto_ptr(new memory_map(cq, buf, mapped)); + } + catch (...) + { + PYOPENCL_CALL_GUARDED_CLEANUP(clEnqueueUnmapMemObject, ( + cq.data(), buf.data(), mapped, 0, 0, 0)); + throw; + } + + PyArray_BYTES(result.get()) = reinterpret_cast(mapped); + + py::handle<> array_py(handle_from_new_ptr(map.release())); + PyArray_BASE(result.get()) = array_py.get(); + Py_INCREF(array_py.get()); + + return py::make_tuple( + array_py, + handle_from_new_ptr(new event(evt_handle))); + } + + + + + py::object enqueue_map_image( + command_queue &cq, + memory_object &img, + cl_map_flags flags, + py::object py_origin, + py::object py_region, + py::object py_shape, py::object dtype, py::object order_py, + py::object py_wait_for, + bool is_blocking + ) + { + PYOPENCL_PARSE_WAIT_FOR; + PYOPENCL_PARSE_NUMPY_ARRAY_SPEC; + COPY_PY_COORD_TRIPLE(origin); + COPY_PY_REGION_TRIPLE(region); + + cl_event evt; + cl_int status_code; + PYOPENCL_PRINT_CALL_TRACE("clEnqueueMapImage"); + size_t row_pitch, slice_pitch; + void *mapped = clEnqueueMapImage( + cq.data(), img.data(), + PYOPENCL_CAST_BOOL(is_blocking), flags, + origin, region, &row_pitch, &slice_pitch, + num_events_in_wait_list, event_wait_list.data(), &evt, + &status_code); + if (status_code != CL_SUCCESS) + throw pyopencl::error("clEnqueueMapImage", status_code); + + event evt_handle(evt, false); + + std::auto_ptr map; + try + { + map = std::auto_ptr(new memory_map(cq, img, mapped)); + } + catch (...) + { + PYOPENCL_CALL_GUARDED_CLEANUP(clEnqueueUnmapMemObject, ( + cq.data(), img.data(), mapped, 0, 0, 0)); + throw; + } + + py::handle<> result = py::handle<>(PyArray_NewFromDescr( + &PyArray_Type, tp_descr, + shape.size(), &shape.front(), /*strides*/ NULL, + mapped, ary_flags, /*obj*/NULL)); + + py::handle<> array_py(handle_from_new_ptr(map.release())); + PyArray_BASE(result.get()) = array_py.get(); + Py_INCREF(array_py.get()); + + return py::make_tuple( + array_py, + handle_from_new_ptr(new event(evt_handle)), + row_pitch, slice_pitch); + } + + + + + // sampler ------------------------------------------------------------------ + class sampler : boost::noncopyable + { + private: + cl_sampler m_sampler; + + public: + sampler(context const &ctx, bool normalized_coordinates, + cl_addressing_mode am, cl_filter_mode fm) + { + } + + ~sampler() + { + PYOPENCL_CALL_GUARDED_CLEANUP(clReleaseSampler, (m_sampler)); + } + + cl_sampler data() const + { + return m_sampler; + } + + PYOPENCL_EQUALITY_TESTS(sampler); + + py::object get_info(cl_sampler_info param_name) const + { + switch (param_name) + { + case CL_SAMPLER_REFERENCE_COUNT: + PYOPENCL_GET_INTEGRAL_INFO(Sampler, m_sampler, param_name, + cl_uint); + case CL_SAMPLER_CONTEXT: + PYOPENCL_GET_OPAQUE_INFO(Sampler, m_sampler, param_name, + cl_context, context); + case CL_SAMPLER_ADDRESSING_MODE: + PYOPENCL_GET_INTEGRAL_INFO(Sampler, m_sampler, param_name, + cl_addressing_mode); + case CL_SAMPLER_FILTER_MODE: + PYOPENCL_GET_INTEGRAL_INFO(Sampler, m_sampler, param_name, + cl_filter_mode); + case CL_SAMPLER_NORMALIZED_COORDS: + PYOPENCL_GET_INTEGRAL_INFO(Sampler, m_sampler, param_name, + cl_bool); + + default: + throw error("Sampler.get_info", CL_INVALID_VALUE); + } + } + }; + + + + + // program ------------------------------------------------------------------ + class program : boost::noncopyable + { + private: + cl_program m_program; + + public: + program(cl_program prog, bool retain) + : m_program(prog) + { + if (retain) + PYOPENCL_CALL_GUARDED(clRetainProgram, (prog)); + } + + ~program() + { + PYOPENCL_CALL_GUARDED_CLEANUP(clReleaseProgram, (m_program)); + } + + cl_program data() const + { + return m_program; + } + + PYOPENCL_EQUALITY_TESTS(program); + + py::object get_info(cl_program_info param_name) const + { + switch (param_name) + { + case CL_PROGRAM_REFERENCE_COUNT: + PYOPENCL_GET_INTEGRAL_INFO(Program, m_program, param_name, + cl_uint); + case CL_PROGRAM_CONTEXT: + PYOPENCL_GET_OPAQUE_INFO(Program, m_program, param_name, + cl_context, context); + case CL_PROGRAM_NUM_DEVICES: + PYOPENCL_GET_INTEGRAL_INFO(Program, m_program, param_name, + cl_uint); + case CL_PROGRAM_DEVICES: + { + std::vector result; + PYOPENCL_GET_VEC_INFO(Program, m_program, param_name, result); + + py::list py_result; + BOOST_FOREACH(cl_device_id did, result) + py_result.append(handle_from_new_ptr( + new pyopencl::device(did))); + return py_result; + } + case CL_PROGRAM_SOURCE: + PYOPENCL_GET_STR_INFO(Program, m_program, param_name); + case CL_PROGRAM_BINARY_SIZES: + { + std::vector result; + PYOPENCL_GET_VEC_INFO(Program, m_program, param_name, result); + return py::list(result); + } + case CL_PROGRAM_BINARIES: + { + std::vector sizes; + PYOPENCL_GET_VEC_INFO(Program, m_program, param_name, sizes); + + std::vector > result(sizes.size()); + std::vector result_ptrs(sizes.size()); + for (unsigned i = 0; i < sizes.size(); ++i) + { + result[i].resize(sizes[i]); + result_ptrs.push_back(result[i].data()); + } + + PYOPENCL_CALL_GUARDED(clGetProgramInfo, + (m_program, param_name, sizes.size()*sizeof(unsigned char *), + result_ptrs.data(), 0)); \ + + py::list py_result; + for (unsigned i = 0; i < sizes.size(); ++i) + py_result.append(py::str( + reinterpret_cast(result[i].data()), + sizes[i])); + return py_result; + } + + default: + throw error("Program.get_info", CL_INVALID_VALUE); + } + } + + py::object get_build_info( + device const &dev, + cl_program_build_info param_name) const + { + switch (param_name) + { +#define PYOPENCL_FIRST_ARG m_program, dev.data() // hackety hack + case CL_PROGRAM_BUILD_STATUS: + PYOPENCL_GET_INTEGRAL_INFO(ProgramBuild, + PYOPENCL_FIRST_ARG, param_name, + cl_build_status); + case CL_PROGRAM_BUILD_OPTIONS: + case CL_PROGRAM_BUILD_LOG: + PYOPENCL_GET_STR_INFO(ProgramBuild, + PYOPENCL_FIRST_ARG, param_name); +#undef PYOPENCL_FIRST_ARG + + default: + throw error("Program.get_build_info", CL_INVALID_VALUE); + } + } + + void build(std::string options, py::object py_devices) + { + if (py_devices.ptr() == Py_None) + { + PYOPENCL_CALL_GUARDED(clBuildProgram, + (m_program, 0, 0, options.c_str(), 0 ,0)); + } + else + { + std::vector devices; + PYTHON_FOREACH(py_dev, py_devices) + devices.push_back( + py::extract(py_dev)().data()); + PYOPENCL_CALL_GUARDED(clBuildProgram, + (m_program, devices.size(), devices.data(), + options.c_str(), 0 ,0)); + } + + } + }; + + + + + inline program *create_program_with_source( + context &ctx, + std::string const &src) + { + const char *string = src.c_str(); + size_t length = src.size(); + + cl_int status_code; + cl_program result = clCreateProgramWithSource( + ctx.data(), 1, &string, &length, &status_code); + PYOPENCL_PRINT_CALL_TRACE("clCreateProrgramWithSource"); + if (status_code != CL_SUCCESS) + throw pyopencl::error("clCreateProrgramWithSource", status_code); + + try + { + return new program(result, false); + } + catch (...) + { + clReleaseProgram(result); + throw; + } + } + + + + + + inline program *create_program_with_binary( + context &ctx, + py::object py_devices, + py::object py_binaries) + { + std::vector devices; + std::vector binaries; + std::vector sizes; + + size_t num_devices = len(py_devices); + if (len(py_binaries) != num_devices) + throw error("create_program_with_binary", CL_INVALID_VALUE, + "device and binary counts don't match"); + + for (size_t i = 0; i < num_devices; ++i) + { + devices.push_back( + py::extract(py_devices[i])().data()); + const void *buf; + PYOPENCL_BUFFER_SIZE_T len; + + if (PyObject_AsReadBuffer( + py::object(py_binaries[i]).ptr(), &buf, &len)) + throw py::error_already_set(); + + binaries.push_back(reinterpret_cast(buf)); + sizes.push_back(len); + } + + cl_int status_code; + cl_program result = clCreateProgramWithBinary( + ctx.data(), num_devices, + devices.data(), sizes.data(), binaries.data(), + /*binary_status*/ 0, &status_code); + PYOPENCL_PRINT_CALL_TRACE("clCreateProrgramWithBinary"); + if (status_code != CL_SUCCESS) + throw pyopencl::error("clCreateProrgramWithBinary", status_code); + + try + { + return new program(result, false); + } + catch (...) + { + clReleaseProgram(result); + throw; + } + } + + + + + // kernel ------------------------------------------------------------------- + class kernel : boost::noncopyable + { + private: + cl_kernel m_kernel; + + public: + kernel(cl_kernel knl, bool retain) + : m_kernel(knl) + { + if (retain) + PYOPENCL_CALL_GUARDED(clRetainKernel, (knl)); + } + + kernel(program const &prg, std::string const &kernel_name) + { + cl_int status_code; + + m_kernel = clCreateKernel(prg.data(), kernel_name.c_str(), + &status_code); + PYOPENCL_PRINT_CALL_TRACE("clCreateKernel"); + if (status_code != CL_SUCCESS) + throw pyopencl::error("clCreateKernel", status_code); + } + + ~kernel() + { + PYOPENCL_CALL_GUARDED_CLEANUP(clReleaseKernel, (m_kernel)); + } + + cl_kernel data() const + { + return m_kernel; + } + + PYOPENCL_EQUALITY_TESTS(kernel); + + void set_arg_mem(cl_uint arg_index, memory_object &mo) + { + PYOPENCL_CALL_GUARDED(clSetKernelArg, + (m_kernel, arg_index, sizeof(cl_mem), mo.data())); + } + + void set_arg_sampler(cl_uint arg_index, sampler &smp) + { + PYOPENCL_CALL_GUARDED(clSetKernelArg, + (m_kernel, arg_index, sizeof(cl_sampler), smp.data())); + } + + void set_arg_buf(cl_uint arg_index, py::object py_buffer) + { + const void *buf; + PYOPENCL_BUFFER_SIZE_T len; + + if (PyObject_AsReadBuffer(py_buffer.ptr(), &buf, &len)) + PYOPENCL_CALL_GUARDED(clSetKernelArg, + (m_kernel, arg_index, len, buf)); + } + + void set_arg(cl_uint arg_index, py::object arg) + { + py::extract ex_mo(arg); + if (ex_mo.check()) + { + set_arg_mem(arg_index, ex_mo()); + return; + } + + py::extract ex_smp(arg); + if (ex_smp.check()) + { + set_arg_sampler(arg_index, ex_smp()); + return; + } + + set_arg_buf(arg_index, arg); + } + + py::object get_info(cl_kernel_info param_name) const + { + switch (param_name) + { + case CL_KERNEL_FUNCTION_NAME: + PYOPENCL_GET_STR_INFO(Kernel, m_kernel, param_name); + case CL_KERNEL_NUM_ARGS: + case CL_KERNEL_REFERENCE_COUNT: + PYOPENCL_GET_INTEGRAL_INFO(Kernel, m_kernel, param_name, + cl_uint); + case CL_KERNEL_CONTEXT: + PYOPENCL_GET_OPAQUE_INFO(Kernel, m_kernel, param_name, + cl_context, context); + case CL_KERNEL_PROGRAM: + PYOPENCL_GET_OPAQUE_INFO(Kernel, m_kernel, param_name, + cl_program, program); + default: + throw error("Kernel.get_info", CL_INVALID_VALUE); + } + } + + py::object get_work_group_info( + cl_kernel_work_group_info param_name, + device const &dev + ) const + { + switch (param_name) + { +#define PYOPENCL_FIRST_ARG m_kernel, dev.data() // hackety hack + case CL_KERNEL_WORK_GROUP_SIZE: + PYOPENCL_GET_INTEGRAL_INFO(KernelWorkGroup, + PYOPENCL_FIRST_ARG, param_name, + size_t); + case CL_KERNEL_COMPILE_WORK_GROUP_SIZE: + { + std::vector result; + PYOPENCL_GET_VEC_INFO(KernelWorkGroup, + PYOPENCL_FIRST_ARG, param_name, result); + return py::list(result); + } + case CL_KERNEL_LOCAL_MEM_SIZE: + PYOPENCL_GET_INTEGRAL_INFO(KernelWorkGroup, + PYOPENCL_FIRST_ARG, param_name, + cl_ulong); + default: + throw error("Kernel.get_work_group_info", CL_INVALID_VALUE); +#undef PYOPENCL_FIRST_ARG + } + } + }; + + + + + py::list create_kernels_in_program(program &pgm) + { + py::list result; + cl_uint num_kernels; + PYOPENCL_CALL_GUARDED(clCreateKernelsInProgram, ( + pgm.data(), 0, 0, &num_kernels)); + + std::vector kernels; + PYOPENCL_CALL_GUARDED(clCreateKernelsInProgram, ( + pgm.data(), num_kernels, kernels.data(), &num_kernels)); + + BOOST_FOREACH(cl_kernel knl, kernels) + result.append(handle_from_new_ptr(new kernel(knl, true))); + + return result; + } + + + + + + event *enqueue_nd_range_kernel( + command_queue &cq, + kernel &knl, + py::object py_global_work_size, + py::object py_local_work_size, + py::object py_global_work_offset, + py::object py_wait_for) + { + PYOPENCL_PARSE_WAIT_FOR; + + cl_uint work_dim = len(py_global_work_size); + + std::vector global_work_size; + COPY_PY_LIST(size_t, global_work_size); + + size_t *local_work_size_ptr = 0; + std::vector local_work_size; + if (py_local_work_size.ptr() != Py_None) + { + if (work_dim != len(py_local_work_size)) + throw error("enqueue_nd_range_kernel", CL_INVALID_VALUE, + "global/work work sizes have differing dimensions"); + + COPY_PY_LIST(size_t, local_work_size); + + local_work_size_ptr = local_work_size.data(); + } + + size_t *global_work_offset_ptr = 0; + std::vector global_work_offset; + if (py_global_work_offset.ptr() != Py_None) + { + if (work_dim != len(py_local_work_size)) + throw error("enqueue_nd_range_kernel", CL_INVALID_VALUE, + "global work size and offset have differing dimensions"); + + COPY_PY_LIST(size_t, global_work_offset); + + global_work_offset_ptr = global_work_offset.data(); + } + + cl_event evt; + PYOPENCL_CALL_GUARDED(clEnqueueNDRangeKernel, ( + cq.data(), + knl.data(), + work_dim, + global_work_offset_ptr, + global_work_size.data(), + local_work_size_ptr, + num_events_in_wait_list, event_wait_list.data(), &evt + )); + + PYOPENCL_RETURN_NEW_EVENT(evt); + } + + + + + + + event *enqueue_task( + command_queue &cq, + kernel &knl, + py::object py_wait_for) + { + PYOPENCL_PARSE_WAIT_FOR; + + cl_event evt; + PYOPENCL_CALL_GUARDED(clEnqueueTask, ( + cq.data(), + knl.data(), + num_events_in_wait_list, event_wait_list.data(), &evt + )); + + PYOPENCL_RETURN_NEW_EVENT(evt); + } +} + + + #endif diff --git a/src/wrapper/wrap_helpers.hpp b/src/wrapper/wrap_helpers.hpp index 0f502737..f8fb324c 100644 --- a/src/wrapper/wrap_helpers.hpp +++ b/src/wrapper/wrap_helpers.hpp @@ -58,6 +58,55 @@ namespace py = boost::python; boost::python::stl_input_iterator(ITERABLE), \ boost::python::stl_input_iterator())) +#define COPY_PY_LIST(TYPE, NAME) \ + std::copy( \ + boost::python::stl_input_iterator(py_##NAME), \ + boost::python::stl_input_iterator(), \ + std::back_inserter(NAME)); + +#define COPY_PY_COORD_TRIPLE(NAME) \ + size_t NAME[3] = {0, 0, 0}; \ + { \ + size_t my_len = len(py_##NAME); \ + if (my_len > 3) \ + throw error("image copy", CL_INVALID_VALUE, #NAME "has too many components"); \ + for (int i = 0; i < std::min(size_t(3), my_len); ++i) \ + NAME[i] = py::extract(py_##NAME[i])(); \ + } + +#define COPY_PY_REGION_TRIPLE(NAME) \ + size_t NAME[3] = {1, 1, 1}; \ + { \ + size_t my_len = len(py_##NAME); \ + if (my_len > 3) \ + throw error("image copy", CL_INVALID_VALUE, #NAME "has too many components"); \ + for (int i = 0; i < std::min(size_t(3), my_len); ++i) \ + NAME[i] = py::extract(py_##NAME[i])(); \ + } + +#define PYOPENCL_PARSE_NUMPY_ARRAY_SPEC \ + PyArray_Descr *tp_descr; \ + if (PyArray_DescrConverter(dtype.ptr(), &tp_descr) != NPY_SUCCEED) \ + throw py::error_already_set(); \ + \ + py::extract shape_as_int(py_shape); \ + std::vector shape; \ + \ + if (shape_as_int.check()) \ + shape.push_back(shape_as_int()); \ + else \ + COPY_PY_LIST(npy_intp, shape); \ + \ + NPY_ORDER order = PyArray_CORDER; \ + PyArray_OrderConverter(order_py.ptr(), &order); \ + \ + int ary_flags = 0; \ + if (order == PyArray_FORTRANORDER) \ + ary_flags |= NPY_FARRAY; \ + else if (order == PyArray_CORDER) \ + ary_flags |= NPY_CARRAY; \ + else \ + throw std::runtime_error("unrecognized order specifier"); \ -- GitLab