From 3c9eff945313057e281f94cc220274370c2cf7bc Mon Sep 17 00:00:00 2001 From: Andreas Kloeckner Date: Mon, 24 Aug 2009 13:12:37 -0400 Subject: [PATCH] First successful test again AMD Stream API. --- MANIFEST.in | 14 ++ README | 9 ++ doc/source/constants.inc | 273 ++++++++++++++++++++++----------------- doc/source/index.rst | 38 ++++++ doc/source/reference.rst | 87 ++++++++++++- examples/demo.py | 32 +++++ pyopencl/__init__.py | 52 +++++++- setup.py | 5 +- src/wrapper/wrap_cl.cpp | 15 ++- src/wrapper/wrap_cl.hpp | 144 +++++++++++++++------ 10 files changed, 504 insertions(+), 165 deletions(-) create mode 100644 MANIFEST.in create mode 100644 README create mode 100644 examples/demo.py diff --git a/MANIFEST.in b/MANIFEST.in new file mode 100644 index 00000000..f2d2e47b --- /dev/null +++ b/MANIFEST.in @@ -0,0 +1,14 @@ +include src/cl/*.hpp +include src/wrapper/*.hpp +include src/wrapper/*.cpp +include test/*.py +include examples/*.py +include doc/source/*.rst +include doc/Makefile +include doc/source/conf.py + +include ez_setup.py +include configure.py +include Makefile.in +include aksetup_helper.py +include README_SETUP.txt diff --git a/README b/README new file mode 100644 index 00000000..947a93c9 --- /dev/null +++ b/README @@ -0,0 +1,9 @@ +Hi there, welcome to PyOpenCL! +---------------------------- + +You can find installation instructions and documentation at + + http://mathema.tician.de/software/pyopencl + +Have fun, +Andreas diff --git a/doc/source/constants.inc b/doc/source/constants.inc index 63b42119..6052535e 100644 --- a/doc/source/constants.inc +++ b/doc/source/constants.inc @@ -1,19 +1,103 @@ .. This is an automatically generated file. DO NOT EDIT -.. class :: platform_info +.. class :: addressing_mode - .. attribute :: EXTENSIONS - .. attribute :: NAME - .. attribute :: PROFILE - .. attribute :: VENDOR - .. attribute :: VERSION + .. attribute :: CLAMP + .. attribute :: CLAMP_TO_EDGE + .. attribute :: NONE + .. attribute :: REPEAT -.. class :: device_type +.. class :: channel_order - .. attribute :: ACCELERATOR - .. attribute :: ALL - .. attribute :: CPU - .. attribute :: DEFAULT - .. attribute :: GPU + .. attribute :: A + .. attribute :: BGRA + .. attribute :: INTENSITY + .. attribute :: LUMINANCE + .. attribute :: R + .. attribute :: RA + .. attribute :: RG + .. attribute :: RGB + .. attribute :: RGBA + +.. class :: channel_type + + .. attribute :: FLOAT + .. attribute :: HALF_FLOAT + .. attribute :: SIGNED_INT16 + .. attribute :: SIGNED_INT32 + .. attribute :: SIGNED_INT8 + .. attribute :: SNORM_INT16 + .. attribute :: SNORM_INT8 + .. attribute :: UNORM_INT16 + .. attribute :: UNORM_INT8 + .. attribute :: UNORM_INT_101010 + .. attribute :: UNORM_SHORT_555 + .. attribute :: UNORM_SHORT_565 + .. attribute :: UNSIGNED_INT16 + .. attribute :: UNSIGNED_INT32 + .. attribute :: UNSIGNED_INT8 + +.. class :: command_execution_status + + .. attribute :: COMPLETE + .. attribute :: QUEUED + .. attribute :: RUNNING + .. attribute :: SUBMITTED + +.. class :: command_queue_info + + .. attribute :: CONTEXT + .. attribute :: DEVICE + .. attribute :: PROPERTIES + .. attribute :: REFERENCE_COUNT + +.. class :: command_queue_properties + + .. attribute :: OUT_OF_ORDER_EXEC_MODE_ENABLE + .. attribute :: PROFILING_ENABLE + +.. class :: command_type + + .. attribute :: ACQUIRE_GL_OBJECTS + .. attribute :: COPY_BUFFER + .. attribute :: COPY_BUFFER_TO_IMAGE + .. attribute :: COPY_IMAGE + .. attribute :: COPY_IMAGE_TO_BUFFER + .. attribute :: MAP_BUFFER + .. attribute :: MAP_IMAGE + .. attribute :: MARKER + .. attribute :: NATIVE_KERNEL + .. attribute :: NDRANGE_KERNEL + .. attribute :: READ_BUFFER + .. attribute :: READ_IMAGE + .. attribute :: RELEASE_GL_OBJECTS + .. attribute :: TASK + .. attribute :: UNMAP_MEM_OBJECT + .. attribute :: WRITE_BUFFER + .. attribute :: WRITE_IMAGE + +.. class :: context_info + + .. attribute :: DEVICES + .. attribute :: PROPERTIES + .. attribute :: REFERENCE_COUNT + +.. class :: context_properties + + .. attribute :: PLATFORM + +.. class :: device_exec_capabilities + + .. attribute :: KERNEL + .. attribute :: NATIVE_KERNEL + +.. class :: device_fp_config + + .. attribute :: DENORM + .. attribute :: FMA + .. attribute :: INF_NAN + .. attribute :: ROUND_TO_INF + .. attribute :: ROUND_TO_NEAREST + .. attribute :: ROUND_TO_ZERO .. class :: device_info @@ -67,14 +151,10 @@ .. attribute :: VENDOR_ID .. attribute :: VERSION -.. class :: device_fp_config +.. class :: device_local_mem_type - .. attribute :: DENORM - .. attribute :: FMA - .. attribute :: INF_NAN - .. attribute :: ROUND_TO_INF - .. attribute :: ROUND_TO_NEAREST - .. attribute :: ROUND_TO_ZERO + .. attribute :: GLOBAL + .. attribute :: LOCAL .. class :: device_mem_cache_type @@ -82,33 +162,54 @@ .. attribute :: READ_ONLY_CACHE .. attribute :: READ_WRITE_CACHE -.. class :: device_local_mem_type +.. class :: device_type - .. attribute :: GLOBAL - .. attribute :: LOCAL + .. attribute :: ACCELERATOR + .. attribute :: ALL + .. attribute :: CPU + .. attribute :: DEFAULT + .. attribute :: GPU -.. class :: command_queue_properties +.. class :: event_info - .. attribute :: OUT_OF_ORDER_EXEC_MODE_ENABLE - .. attribute :: PROFILING_ENABLE + .. attribute :: COMMAND_EXECUTION_STATUS + .. attribute :: COMMAND_QUEUE + .. attribute :: COMMAND_TYPE + .. attribute :: REFERENCE_COUNT -.. class :: context_info +.. class :: filter_mode - .. attribute :: DEVICES - .. attribute :: PROPERTIES - .. attribute :: REFERENCE_COUNT + .. attribute :: LINEAR + .. attribute :: NEAREST -.. class :: context_properties +.. class :: image_info - .. attribute :: PLATFORM + .. attribute :: DEPTH + .. attribute :: ELEMENT_SIZE + .. attribute :: FORMAT + .. attribute :: HEIGHT + .. attribute :: ROW_PITCH + .. attribute :: SLICE_PITCH + .. attribute :: WIDTH -.. class :: command_queue_info +.. class :: kernel_info .. attribute :: CONTEXT - .. attribute :: DEVICE - .. attribute :: PROPERTIES + .. attribute :: FUNCTION_NAME + .. attribute :: NUM_ARGS + .. attribute :: PROGRAM .. attribute :: REFERENCE_COUNT +.. class :: kernel_work_group_info + + .. attribute :: COMPILE_WORK_GROUP_SIZE + .. attribute :: WORK_GROUP_SIZE + +.. class :: map_flags + + .. attribute :: READ + .. attribute :: WRITE + .. class :: mem_flags .. attribute :: ALLOC_HOST_PTR @@ -118,24 +219,6 @@ .. attribute :: USE_HOST_PTR .. attribute :: WRITE_ONLY -.. class :: channel_order - - .. attribute :: A - .. attribute :: BGRA - .. attribute :: INTENSITY - .. attribute :: LUMINANCE - .. attribute :: R - .. attribute :: RA - .. attribute :: RG - .. attribute :: RGB - .. attribute :: RGBA - -.. class :: mem_object_type - - .. attribute :: BUFFER - .. attribute :: IMAGE2D - .. attribute :: IMAGE3D - .. class :: mem_info .. attribute :: CONTEXT @@ -146,40 +229,32 @@ .. attribute :: SIZE .. attribute :: TYPE -.. class :: image_info - - .. attribute :: DEPTH - .. attribute :: ELEMENT_SIZE - .. attribute :: FORMAT - .. attribute :: HEIGHT - .. attribute :: ROW_PITCH - .. attribute :: SLICE_PITCH - .. attribute :: WIDTH - -.. class :: addressing_mode +.. class :: mem_object_type - .. attribute :: CLAMP - .. attribute :: CLAMP_TO_EDGE - .. attribute :: NONE - .. attribute :: REPEAT + .. attribute :: BUFFER + .. attribute :: IMAGE2D + .. attribute :: IMAGE3D -.. class :: filter_mode +.. class :: platform_info - .. attribute :: LINEAR - .. attribute :: NEAREST + .. attribute :: EXTENSIONS + .. attribute :: NAME + .. attribute :: PROFILE + .. attribute :: VENDOR + .. attribute :: VERSION -.. class :: sampler_info +.. class :: profiling_info - .. attribute :: ADDRESSING_MODE - .. attribute :: CONTEXT - .. attribute :: FILTER_MODE - .. attribute :: NORMALIZED_COORDS - .. attribute :: REFERENCE_COUNT + .. attribute :: END + .. attribute :: QUEUED + .. attribute :: START + .. attribute :: SUBMIT -.. class :: map_flags +.. class :: program_build_info - .. attribute :: READ - .. attribute :: WRITE + .. attribute :: LOG + .. attribute :: OPTIONS + .. attribute :: STATUS .. class :: program_info @@ -191,43 +266,11 @@ .. attribute :: REFERENCE_COUNT .. attribute :: SOURCE -.. class :: program_build_info - - .. attribute :: LOG - .. attribute :: OPTIONS - .. attribute :: STATUS - -.. class :: kernel_info +.. class :: sampler_info + .. attribute :: ADDRESSING_MODE .. attribute :: CONTEXT - .. attribute :: FUNCTION_NAME - .. attribute :: NUM_ARGS - .. attribute :: PROGRAM - .. attribute :: REFERENCE_COUNT - -.. class :: kernel_work_group_info - - .. attribute :: COMPILE_WORK_GROUP_SIZE - .. attribute :: WORK_GROUP_SIZE - -.. class :: event_info - - .. attribute :: COMMAND_EXECUTION_STATUS - .. attribute :: COMMAND_QUEUE - .. attribute :: COMMAND_TYPE + .. attribute :: FILTER_MODE + .. attribute :: NORMALIZED_COORDS .. attribute :: REFERENCE_COUNT -.. class :: command_execution_status - - .. attribute :: COMPLETE - .. attribute :: QUEUED - .. attribute :: RUNNING - .. attribute :: SUBMITTED - -.. class :: profiling_info - - .. attribute :: END - .. attribute :: QUEUED - .. attribute :: START - .. attribute :: SUBMIT - diff --git a/doc/source/index.rst b/doc/source/index.rst index cc1d6b18..75281c54 100644 --- a/doc/source/index.rst +++ b/doc/source/index.rst @@ -24,6 +24,44 @@ What makes PyOpenCL special? * Helpful Documentation. You're looking at it. ;) +Here's an example, to given you an impression:: + + import pyopencl as cl + import numpy + import numpy.linalg as la + + a = numpy.random.rand(50000).astype(numpy.float32) + b = numpy.random.rand(50000).astype(numpy.float32) + + 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) + + prg = cl.create_program_with_source(ctx, """ + __kernel void sum(__global const float *a, + __global const float *b, __global float *c) + { + int gid = get_global_id(0); + c[gid] = a[gid] + b[gid]; + } + """).build() + + prg.sum(queue, a.shape, a_buf, b_buf, dest_buf) + + a_plus_b = numpy.empty_like(a) + cl.enqueue_read_buffer(queue, dest_buf, a_plus_b).wait() + + print la.norm(a_plus_b - (a+b)) + +(You can find this example as :file:`examples/demo.py` in the PyOpenCL +source distribution.) + Contents ======== diff --git a/doc/source/reference.rst b/doc/source/reference.rst index 6edd5064..dcdd024c 100644 --- a/doc/source/reference.rst +++ b/doc/source/reference.rst @@ -63,6 +63,12 @@ Platforms, Devices and Contexts 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_devices(device_type) Return a list of devices matching *device_type*. @@ -76,6 +82,12 @@ Platforms, Devices and Contexts 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. + Two instances of this class may be compared using *=="* and *"!="*. .. class:: Context(devices, properties=[]) @@ -87,20 +99,37 @@ Platforms, Devices and Contexts 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. + |comparable| +.. function:: create_context_from_type(dev_type, properties=[]) + Command Queues and Events ------------------------- -.. class:: CommandQueue(context, device, properties=[]) +.. class:: CommandQueue(context, device=None, properties=[]) Create a new command queue. *properties* is a list of key-value tuples, where each key must be one of :class:`command_queue_properties`. + 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:: set_property(prop, enable) See :class:`command_queue_properties` for possible values of *prop*. @@ -116,10 +145,18 @@ Command Queues and Events See :class:`event_info` for values of *param*. + .. attribute:: info + + Lower case versions of the :class:`event_info` constants + may be used as attributes on instances of this class + to directly query info attributes. + .. method:: get_profiling_info(param) See :class:`profiling_info` for values of *param*. + .. method:: wait() + |comparable| .. function:: wait_for_events(events) @@ -140,6 +177,12 @@ Memory See :class:`mem_info` for values of *param*. + .. attribute:: info + + Lower case versions of the :class:`mem_info` constants + may be used as attributes on instances of this class + to directly query info attributes. + .. method:: get_image_info(param) See :class:`image_info` for values of *param*. @@ -254,6 +297,12 @@ Samplers 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. + |comparable| Programs and Kernels @@ -265,13 +314,31 @@ Programs and Kernels See :class:`program_info` for values of *param*. + .. attribute:: info + + Lower case versions of the :class:`program_info` constants + may be used as attributes on instances of this class + to directly query info attributes. + .. method:: get_build_info(param, device) See :class:`program_build_info` for values of *param*. - .. method:: build(options, devices=None) + .. method:: build(options="", devices=None) + + *options* is a string of compiler flags. + Returns *self*. + + .. attribute:: kernel_name + + :class:`Kernel` objects can be produced from a built + (see :meth:`build`) program simply by attribute lookup. - *options* is a string of compiler flags. + .. note:: + + The :class:`program_info` attributes live + in the same name space and take precedence over + :class:`Kernel` names. |comparable| @@ -287,16 +354,28 @@ Programs and Kernels 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_work_group_info(param, device) See :class:`kernel_work_group_info` for values of *param*. + .. method:: __call__(queue, global_size, *args, global_offset=None, local_size=None, wait_for=None) + + |enqueue-waitfor| + + |comparable| .. function:: enqueue_nd_range_kernel(queue, kernel, global_work_size, local_work_size, global_work_offset=None, wait_for=None) |enqueue-waitfor| -.. function:: enqueue_task(queue, kernel, global_work_offset, wait_for=None) +.. function:: enqueue_task(queue, kernel, wait_for=None) |enqueue-waitfor| diff --git a/examples/demo.py b/examples/demo.py new file mode 100644 index 00000000..f1c2d46f --- /dev/null +++ b/examples/demo.py @@ -0,0 +1,32 @@ +import pyopencl as cl +import numpy +import numpy.linalg as la + +a = numpy.random.rand(50000).astype(numpy.float32) +b = numpy.random.rand(50000).astype(numpy.float32) + +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) + +prg = cl.create_program_with_source(ctx, """ + __kernel void sum(__global const float *a, + __global const float *b, __global float *c) + { + int gid = get_global_id(0); + c[gid] = a[gid] + b[gid]; + } + """).build() + +prg.sum(queue, a.shape, a_buf, b_buf, dest_buf) + +a_plus_b = numpy.empty_like(a) +cl.enqueue_read_buffer(queue, dest_buf, a_plus_b).wait() + +print la.norm(a_plus_b - (a+b)) diff --git a/pyopencl/__init__.py b/pyopencl/__init__.py index 589f1ae7..efb1da56 100644 --- a/pyopencl/__init__.py +++ b/pyopencl/__init__.py @@ -1,5 +1,5 @@ VERSION = (0, 90) -VERSION_STATUS = "alpha" +VERSION_STATUS = "" VERSION_TEXT = ".".join(str(x) for x in VERSION) + VERSION_STATUS import pyopencl._cl as _cl @@ -19,17 +19,63 @@ def _add_functionality(): _cl.CommandQueue: _cl.command_queue_info, _cl.Event: _cl.event_info, _cl.MemoryObject: _cl.mem_info, - _cl.Program: _cl.program_info, _cl.Kernel: _cl.kernel_info, } def make_getattr(info_cls): def result(self, name): - return self.get_info(getattr(info_cls, name.upper())) + try: + inf_attr = getattr(info_cls, name.upper()) + except AttributeError: + raise AttributeError("%s has no attribute '%s'" + % (type(self), name)) + else: + return self.get_info(inf_attr) return result for cls, info_cls in cls_to_info_cls.iteritems(): cls.__getattr__ = make_getattr(info_cls) + def program_getattr(self, attr): + try: + pi_attr = getattr(program_info, attr.upper()) + except AttributeError: + try: + return Kernel(self, attr) + except LogicError: + raise AttributeError("'%s' was not found as a program info attribute or as a kernel name" + % attr) + else: + return self.get_info(pi_attr) + + Program.__getattr__ = program_getattr + + def kernel_call(self, queue, global_size, *args, **kwargs): + for i, arg in enumerate(args): + self.set_arg(i, arg) + + global_offset = kwargs.pop("global_offset", None) + local_size = kwargs.pop("local_size", None) + wait_for = kwargs.pop("wait_for", None) + + if kwargs: + raise TypeError( + "Kernel.__call__ recived unexpected keyword arguments: %s" + % ", ".join(kwargs.keys())) + + return enqueue_nd_range_kernel(queue, self, global_size, local_size, + global_offset, wait_for) + + Kernel.__call__ = kernel_call + + def event_wait(self): + wait_for_events([self]) + return self + + Event.wait = event_wait + + + + _add_functionality() diff --git a/setup.py b/setup.py index 3fc148ba..dc890948 100644 --- a/setup.py +++ b/setup.py @@ -17,7 +17,7 @@ def get_config_schema(): IncludeDir("CL", []), LibraryDir("CL", []), - Libraries("CL", []), + Libraries("CL", ["OpenCL"]), StringListOption("CXXFLAGS", [], help="Any extra C++ compiler options to include"), @@ -88,7 +88,8 @@ def main(): * Speed. PyOpenCL's base layer is written in C++, so all the niceties above are virtually free. - * Helpful `Documentation `_. + * Helpful and complete `Documentation `_ + as well as a `Wiki `_. For, now the best next stop is the `source tree `_. """, diff --git a/src/wrapper/wrap_cl.cpp b/src/wrapper/wrap_cl.cpp index 57e1313b..95ee8562 100644 --- a/src/wrapper/wrap_cl.cpp +++ b/src/wrapper/wrap_cl.cpp @@ -74,7 +74,7 @@ namespace BOOST_PYTHON_MODULE(_cl) { #define DECLARE_EXC(NAME, BASE) \ - CL##NAME = py::handle<>(PyErr_NewException("pyopencl._cl." #NAME, BASE, NULL)); \ + CL##NAME = py::handle<>(PyErr_NewException("pyopencl." #NAME, BASE, NULL)); \ py::scope().attr(#NAME) = CL##NAME; { @@ -428,11 +428,15 @@ BOOST_PYTHON_MODULE(_cl) ; } + py::def("create_context_from_type", create_context_from_type, + (py::arg("dev_type"), py::arg("properties")=py::list()), + py::return_value_policy()); + { typedef command_queue cls; py::class_("CommandQueue", - py::init >()) + py::init >()) .DEF_SIMPLE_METHOD(get_info) .DEF_SIMPLE_METHOD(set_property) .DEF_SIMPLE_METHOD(flush) @@ -468,7 +472,7 @@ BOOST_PYTHON_MODULE(_cl) py::def("create_buffer", create_buffer, py::return_value_policy()); - py::def("create_host_buffer", create_buffer, + py::def("create_host_buffer", create_host_buffer, py::return_value_policy()); py::def("enqueue_read_buffer", enqueue_read_buffer, @@ -572,7 +576,8 @@ BOOST_PYTHON_MODULE(_cl) .DEF_SIMPLE_METHOD(get_info) .DEF_SIMPLE_METHOD(get_build_info) .def("build", &cls::build, - (py::arg("options"), py::arg("devices")=py::object())) + (py::arg("options")="", py::arg("devices")=py::object()), + py::return_self<>()) .def(py::self == py::self) .def(py::self != py::self) ; diff --git a/src/wrapper/wrap_cl.hpp b/src/wrapper/wrap_cl.hpp index 39b3f87e..81005b03 100644 --- a/src/wrapper/wrap_cl.hpp +++ b/src/wrapper/wrap_cl.hpp @@ -161,15 +161,14 @@ // event helpers -------------------------------------------------------------- #define PYOPENCL_PARSE_WAIT_FOR \ cl_uint num_events_in_wait_list = 0; \ - std::vector event_wait_list(len(py_wait_for)); \ + std::vector event_wait_list; \ \ + if (py_wait_for.ptr() != Py_None) \ { \ - if (py_wait_for.ptr() != Py_None) \ - { \ - PYTHON_FOREACH(evt, py_wait_for) \ - event_wait_list[num_events_in_wait_list++] = \ - py::extract(evt)().data(); \ - } \ + event_wait_list.resize(len(py_wait_for)); \ + PYTHON_FOREACH(evt, py_wait_for) \ + event_wait_list[num_events_in_wait_list++] = \ + py::extract(evt)().data(); \ } #define PYOPENCL_RETURN_NEW_EVENT(EVT) \ @@ -484,6 +483,31 @@ namespace pyopencl // context ------------------------------------------------------------------ +#define PYOPENCL_PARSE_CONTEXT_PROPERTIES \ + std::vector props; \ + \ + PYTHON_FOREACH(prop_tuple, py_properties) \ + { \ + if (len(prop_tuple) != 2) \ + throw error("Context", CL_INVALID_VALUE, "property tuple must have length 2"); \ + cl_context_properties prop = \ + py::extract(prop_tuple[0]); \ + props.push_back(prop); \ + \ + if (prop == CL_CONTEXT_PLATFORM) \ + { \ + py::extract value(prop_tuple[1]); \ + props.push_back( \ + reinterpret_cast(value().data())); \ + } \ + else \ + throw error("Context", CL_INVALID_VALUE, "invalid context property"); \ + } \ + props.push_back(0); \ + + + + class context : public boost::noncopyable { private: @@ -501,26 +525,7 @@ namespace pyopencl py::list py_devices, py::list py_properties=py::list()) { - std::vector props; - - PYTHON_FOREACH(prop_tuple, py_properties) - { - if (len(prop_tuple) != 2) - throw error("Context", CL_INVALID_VALUE, "property tuple must have length 2"); - cl_context_properties prop = - py::extract(prop_tuple[0]); - props.push_back(prop); - - if (prop == CL_CONTEXT_PLATFORM) - { - py::extract value(prop_tuple[1]); - props.push_back( - reinterpret_cast(value().data())); - } - else - throw error("Context", CL_INVALID_VALUE, "invalid context property"); - } - props.push_back(0); + PYOPENCL_PARSE_CONTEXT_PROPERTIES; std::vector devices; PYTHON_FOREACH(py_dev, py_devices) @@ -574,7 +579,33 @@ namespace pyopencl return py_result; } - case CL_CONTEXT_PROPERTIES: // FIXME: complicated + case CL_CONTEXT_PROPERTIES: + { + std::vector result; + PYOPENCL_GET_VEC_INFO(Context, m_context, param_name, result); + + py::list py_result; + for (size_t i = 0; i < result.size(); i+=2) + { + cl_context_properties key = result[i]; + py::object value; + switch (key) + { + case CL_CONTEXT_PLATFORM: + { + value = py::object( + handle_from_new_ptr(new platform( + reinterpret_cast(result[i+1])))); + } + default: + throw error("Context.get_info", CL_INVALID_VALUE, + "unkown context_property key encountered"); + } + + py_result.append(py::make_tuple(result[i], value)); + } + return py_result; + } default: throw error("Context.get_info", CL_INVALID_VALUE); @@ -584,6 +615,34 @@ namespace pyopencl + context *create_context_from_type( + cl_device_type dev_type, + py::list py_properties) + { + PYOPENCL_PARSE_CONTEXT_PROPERTIES; + + cl_int status_code; + cl_context ctx = clCreateContextFromType( + &props.front(), + dev_type, + 0, 0, &status_code); + + PYOPENCL_PRINT_CALL_TRACE("clCreateContextFromType"); + if (status_code != CL_SUCCESS) + throw pyopencl::error("Context", status_code); + + try + { + return new context(ctx, false); + } + catch (...) + { + PYOPENCL_CALL_GUARDED_CLEANUP(clReleaseContext, (ctx)); + throw; + } + } + + // command_queue ------------------------------------------------------------ class command_queue @@ -607,15 +666,25 @@ namespace pyopencl command_queue( const context &ctx, - const device &dev, + const device *py_dev=0, cl_command_queue_properties props=0) { + cl_device_id dev; + if (py_dev) + dev = py_dev->data(); + else + { + std::vector devs; + PYOPENCL_GET_VEC_INFO(Context, ctx.data(), CL_CONTEXT_DEVICES, devs); + if (devs.size() == 0) + throw pyopencl::error("CommandQueue", CL_INVALID_VALUE, + "context doesn't have any devices? -- don't know which one to default to"); + dev = devs[0]; + } + cl_int status_code; m_queue = clCreateCommandQueue( - ctx.data(), - dev.data(), - props, - &status_code); + ctx.data(), dev, props, &status_code); PYOPENCL_PRINT_CALL_TRACE("clCreateCommandQueue"); if (status_code != CL_SUCCESS) @@ -1600,7 +1669,7 @@ namespace pyopencl } } - void build(std::string options, py::object py_devices) + program &build(std::string options, py::object py_devices) { if (py_devices.ptr() == Py_None) { @@ -1618,6 +1687,7 @@ namespace pyopencl options.c_str(), 0 ,0)); } + return *this; } }; @@ -1744,14 +1814,16 @@ namespace pyopencl void set_arg_mem(cl_uint arg_index, memory_object &mo) { + cl_mem m = mo.data(); PYOPENCL_CALL_GUARDED(clSetKernelArg, - (m_kernel, arg_index, sizeof(cl_mem), mo.data())); + (m_kernel, arg_index, sizeof(cl_mem), &m)); } void set_arg_sampler(cl_uint arg_index, sampler &smp) { + cl_sampler s = smp.data(); PYOPENCL_CALL_GUARDED(clSetKernelArg, - (m_kernel, arg_index, sizeof(cl_sampler), smp.data())); + (m_kernel, arg_index, sizeof(cl_sampler), &s)); } void set_arg_buf(cl_uint arg_index, py::object py_buffer) -- GitLab