diff --git a/doc/source/array.rst b/doc/source/array.rst
index e49f8305e28bd56ee2a980677cbda05a7f508c80..d72444240c53d006114010a9758311cca2ee925e 100644
--- a/doc/source/array.rst
+++ b/doc/source/array.rst
@@ -90,6 +90,19 @@ The :class:`Array` Class
Returns the size of the leading dimension of *self*.
+ .. method :: reshape(shape)
+
+ Returns an array containing the same data with a new shape.
+
+ .. method :: ravel()
+
+ Returns flattened array containing the same data.
+
+ .. metod :: view(dtype=None)
+
+ Returns view of array with the same data. If *dtype* is different from
+ current dtype, the actual bytes of memory will be reinterpreted.
+
.. method :: set(ary, queue=None, async=False)
Transfer the contents the :class:`numpy.ndarray` object *ary*
diff --git a/pyopencl/array.py b/pyopencl/array.py
index 3a7c8d1b5a97fb6ad6a01b60764c0ace597b49e9..c27ea813cacf4232822ee7972efd694f21cef37f 100644
--- a/pyopencl/array.py
+++ b/pyopencl/array.py
@@ -50,8 +50,7 @@ class vec:
def _create_vector_types():
field_names = ["x", "y", "z", "w"]
- name_to_dtype = {}
- dtype_to_name = {}
+ from pyopencl.tools import register_dtype
counts = [2, 3, 4, 8, 16]
for base_name, base_type in [
@@ -78,8 +77,7 @@ def _create_vector_types():
formats=[base_type]*count,
titles=titles))
- name_to_dtype[name] = dtype
- dtype_to_name[dtype] = name
+ register_dtype(dtype, name)
setattr(vec, name, dtype)
@@ -92,9 +90,6 @@ def _create_vector_types():
% (my_field_names_defaulted, my_field_names),
dict(array=np.array, my_dtype=dtype))))
- vec._dtype_to_c_name = dtype_to_name
- vec._c_name_to_dtype = name_to_dtype
-
_create_vector_types()
# }}}
@@ -293,9 +288,6 @@ class Array(object):
self.data = self.allocator(self.size * self.dtype.itemsize)
else:
self.data = None
-
- if base is not None:
- raise ValueError("If data is specified, base must be None.")
else:
self.data = data
@@ -308,6 +300,28 @@ class Array(object):
def flags(self):
return _ArrayFlags(self)
+ def _new_with_changes(self, data, shape=None, dtype=None, strides=None, queue=None,
+ base=None):
+ if shape is None:
+ shape = self.shape
+ if dtype is None:
+ dtype = self.dtype
+ if strides is None:
+ strides = self.strides
+ if queue is None:
+ queue = self.queue
+ if base is None and data is self.data:
+ base = self
+
+ if queue is not None:
+ return Array(queue, shape, dtype,
+ allocator=self.allocator, strides=strides, base=base)
+ elif self.allocator is not None:
+ return Array(self.allocator, shape, dtype, queue=queue,
+ strides=strides, base=base)
+ else:
+ return Array(self.context, shape, dtype, strides=strides, base=base)
+
#@memoize_method FIXME: reenable
def get_sizes(self, queue, kernel_specific_max_wg_size=None):
return splay(queue, self.mem_size,
@@ -661,6 +675,37 @@ class Array(object):
def __gt__(self, other):
raise NotImplementedError
+ # {{{ views
+
+ def reshape(self, *shape):
+ # TODO: add more error-checking, perhaps
+ if isinstance(shape[0], tuple) or isinstance(shape[0], list):
+ shape = tuple(shape[0])
+ size = reduce(lambda x, y: x * y, shape, 1)
+ if size != self.size:
+ raise ValueError("total size of new array must be unchanged")
+
+ return self._new_with_changes(data=self.data, shape=shape)
+
+ def ravel(self):
+ return self.reshape(self.size)
+
+ def view(self, dtype=None):
+ if dtype is None:
+ dtype = self.dtype
+
+ old_itemsize = self.dtype.itemsize
+ itemsize = np.dtype(dtype).itemsize
+
+ if self.shape[-1] * old_itemsize % itemsize != 0:
+ raise ValueError("new type not compatible with array")
+
+ shape = self.shape[:-1] + (self.shape[-1] * old_itemsize // itemsize,)
+
+ return self._new_with_changes(data=self.data, shape=shape, dtype=dtype)
+
+ # }}
+
# }}}
# {{{ creation helpers
@@ -711,15 +756,7 @@ def zeros(*args, **kwargs):
return _zeros(*args, **kwargs)
def empty_like(ary):
- if ary.queue is not None:
- return Array(ary.queue, ary.shape, ary.dtype,
- allocator=ary.allocator, strides=ary.strides)
- elif ary.allocator is not None:
- return Array(ary.allocator, ary.shape, ary.dtype, queue=ary.queue,
- strides=ary.strides)
- else:
- return Array(ary.context, ary.shape, ary.dtype,
- strides=ary.strides)
+ return ary._new_with_changes(data=None)
def zeros_like(ary):
result = empty_like(ary)
diff --git a/pyopencl/compyte b/pyopencl/compyte
index 9bebf5c44e95690139ab0940d39cddcf7fa13593..4118040aa50b5a7ef9b50d4b6b9f20fb1c764781 160000
--- a/pyopencl/compyte
+++ b/pyopencl/compyte
@@ -1 +1 @@
-Subproject commit 9bebf5c44e95690139ab0940d39cddcf7fa13593
+Subproject commit 4118040aa50b5a7ef9b50d4b6b9f20fb1c764781
diff --git a/pyopencl/reduction.py b/pyopencl/reduction.py
index aecc202d34cd06e43adcd5fb23bc7f54414307a9..0a5aaf422217b070cf268df5a12f77b10e2e0d23 100644
--- a/pyopencl/reduction.py
+++ b/pyopencl/reduction.py
@@ -55,11 +55,10 @@ KERNEL = """
#pragma OPENCL EXTENSION cl_amd_fp64: enable
% endif
+ ${preamble}
typedef ${out_type} out_type;
- ${preamble}
-
__kernel void ${name}(
__global out_type *out, ${arguments},
unsigned int seq_count, unsigned int n)
@@ -217,7 +216,7 @@ def get_reduction_source(
-def get_reduction_kernel(
+def get_reduction_kernel(stage,
ctx, out_type, out_type_size,
neutral, reduce_expr, map_expr=None, arguments=None,
name="reduce_kernel", preamble="",
@@ -225,7 +224,7 @@ def get_reduction_kernel(
if map_expr is None:
map_expr = "in[i]"
- if arguments is None:
+ if stage == 2:
arguments = "__global const %s *in" % out_type
inf = get_reduction_source(
@@ -266,7 +265,7 @@ class ReductionKernel:
trip_count = 0
while True:
- self.stage_1_inf = get_reduction_kernel(ctx,
+ self.stage_1_inf = get_reduction_kernel(1, ctx,
dtype_to_ctype(dtype_out), dtype_out.itemsize,
neutral, reduce_expr, map_expr, arguments,
name=name+"_stage1", options=options, preamble=preamble,
@@ -285,9 +284,9 @@ class ReductionKernel:
assert trip_count <= 2
# stage 2 has only one input and no map expression
- self.stage_2_inf = get_reduction_kernel(ctx,
+ self.stage_2_inf = get_reduction_kernel(2, ctx,
dtype_to_ctype(dtype_out), dtype_out.itemsize,
- neutral, reduce_expr,
+ neutral, reduce_expr, arguments=arguments,
name=name+"_stage2", options=options, preamble=preamble,
max_group_size=max_group_size)
diff --git a/pyopencl/tools.py b/pyopencl/tools.py
index 7981a76224d9310e05e06b90ad36d7876289f018..8e1accbbc0e06e6bd1caccb6cdec643719b96af6 100644
--- a/pyopencl/tools.py
+++ b/pyopencl/tools.py
@@ -33,6 +33,11 @@ import numpy as np
from decorator import decorator
import pyopencl as cl
+from pyopencl.compyte.dtypes import (
+ register_dtype, _fill_dtype_registry,
+ dtype_to_ctype)
+
+_fill_dtype_registry(respect_windows=True)
@@ -186,51 +191,8 @@ def pytest_generate_tests_for_pyopencl(metafunc):
-# {{{ C code generation helpers -----------------------------------------------
-def dtype_to_ctype(dtype):
- if dtype is None:
- raise ValueError("dtype may not be None")
-
- dtype = np.dtype(dtype)
- if dtype == np.int64:
- return "long"
- elif dtype == np.uint64:
- return "unsigned long"
- elif dtype == np.int32:
- return "int"
- elif dtype == np.uint32:
- return "unsigned int"
- elif dtype == np.int16:
- return "short int"
- elif dtype == np.uint16:
- return "short unsigned int"
- elif dtype == np.int8:
- return "signed char"
- elif dtype == np.uint8:
- return "unsigned char"
- elif dtype == np.bool:
- return "bool"
- elif dtype == np.float32:
- return "float"
- elif dtype == np.float64:
- return "double"
- elif dtype == np.complex64:
- return "complex float"
- elif dtype == np.complex128:
- return "complex double"
- else:
- import pyopencl.array as cl_array
- try:
- return cl_array.vec._dtype_to_c_name[dtype]
- except KeyError:
- raise ValueError, "unable to map dtype '%s'" % dtype
-
-# }}}
-
-
-
+# {{{ C argument lists
-# {{{ C argument lists --------------------------------------------------------
class Argument:
def __init__(self, dtype, name):
self.dtype = np.dtype(dtype)
@@ -256,59 +218,12 @@ class ScalarArg(Argument):
def parse_c_arg(c_arg):
c_arg = (c_arg
- .replace("const", "")
- .replace("volatile", "")
.replace("__global", "")
.replace("__local", "")
.replace("__constant", ""))
- # process and remove declarator
- import re
- decl_re = re.compile(r"(\**)\s*([_a-zA-Z0-9]+)(\s*\[[ 0-9]*\])*\s*$")
- decl_match = decl_re.search(c_arg)
-
- if decl_match is None:
- raise ValueError("couldn't parse C declarator '%s'" % c_arg)
-
- name = decl_match.group(2)
-
- if decl_match.group(1) or decl_match.group(3) is not None:
- arg_class = VectorArg
- else:
- arg_class = ScalarArg
-
- tp = c_arg[:decl_match.start()]
- tp = " ".join(tp.split())
-
- type_re = re.compile(r"^([a-z0-9 ]+)$")
- type_match = type_re.match(tp)
- if not type_match:
- raise RuntimeError("type '%s' did not match expected shape of type"
- % tp)
-
- tp = type_match.group(1)
-
- if tp == "float": dtype = np.float32
- elif tp == "double": dtype = np.float64
- elif tp in ["int", "signed int"]: dtype = np.int32
- elif tp in ["unsigned", "unsigned int"]: dtype = np.uint32
- elif tp in ["long", "long int"]: dtype = np.int64
- elif tp in ["unsigned long", "unsigned long int", "long unsigned int"]:
- dtype = np.uint64
- elif tp in ["short", "short int"]: dtype = np.int16
- elif tp in ["unsigned short", "unsigned short int", "short unsigned int"]:
- dtype = np.uint16
- elif tp in ["char", "signed char"]: dtype = np.int8
- elif tp in ["unsigned char"]: dtype = np.uint8
- elif tp in ["bool"]: dtype = np.bool
- else:
- import pyopencl.array as cl_array
- try:
- dtype = cl_array.vec._c_name_to_dtype[tp]
- except KeyError:
- raise ValueError("unknown type '%s'" % tp)
-
- return arg_class(dtype, name)
+ from pyopencl.compyte.dtypes import parse_c_arg_backend
+ return parse_c_arg_backend(c_arg, ScalarArg, VectorArg)
# }}}
diff --git a/test/test_array.py b/test/test_array.py
index 120f3132ee19e88f56bd03824e854f7042819057..8dd2fab72231f6e6097e847f2863ee4488ab93b8 100644
--- a/test/test_array.py
+++ b/test/test_array.py
@@ -639,6 +639,90 @@ def test_mem_pool_with_arrays(ctx_factory):
assert b_dev.allocator is mem_pool
assert result.allocator is mem_pool
+@pytools.test.mark_test.opencl
+def test_view(ctx_factory):
+ context = ctx_factory()
+ queue = cl.CommandQueue(context)
+
+ a = np.arange(128).reshape(8, 16).astype(np.float32)
+ a_dev = cl_array.to_device(queue, a)
+
+ # same dtype
+ view = a_dev.view()
+ assert view.shape == a_dev.shape and view.dtype == a_dev.dtype
+
+ # larger dtype
+ view = a_dev.view(np.complex64)
+ assert view.shape == (8, 8) and view.dtype == np.complex64
+
+ # smaller dtype
+ view = a_dev.view(np.int16)
+ assert view.shape == (8, 32) and view.dtype == np.int16
+
+mmc_dtype = np.dtype([("cur_min", np.float32), ("cur_max", np.float32)])
+
+from pyopencl.tools import register_dtype
+register_dtype(mmc_dtype, "minmax_collector")
+
+@pytools.test.mark_test.opencl
+def test_struct_reduce(ctx_factory):
+ context = ctx_factory()
+ queue = cl.CommandQueue(context)
+
+ preamble = """//CL//
+ struct minmax_collector
+ {
+ float cur_min;
+ float cur_max;
+ };
+
+ typedef struct minmax_collector minmax_collector;
+
+ minmax_collector mmc_neutral()
+ {
+ // FIXME: needs infinity literal in real use, ok here
+ minmax_collector result = {10000, -10000};
+ return result;
+ }
+
+ minmax_collector mmc_from_scalar(float x)
+ {
+ minmax_collector result = {x, x};
+ return result;
+ }
+
+ minmax_collector agg_mmc(minmax_collector a, minmax_collector b)
+ {
+ minmax_collector result = {
+ fmin(a.cur_min, b.cur_min),
+ fmax(a.cur_max, b.cur_max),
+ };
+ return result;
+ }
+
+ """
+
+
+ from pyopencl.clrandom import rand as clrand
+ a_gpu = clrand(queue, (20000,), dtype=np.float32)
+ a = a_gpu.get()
+
+
+ from pyopencl.reduction import ReductionKernel
+ red = ReductionKernel(context, mmc_dtype,
+ neutral="mmc_neutral()",
+ reduce_expr="agg_mmc(a, b)", map_expr="mmc_from_scalar(x[i])",
+ arguments="__global float *x", preamble=preamble)
+
+ minmax = red(a_gpu).get()
+ #print minmax["cur_min"], minmax["cur_max"]
+ #print np.min(a), np.max(a)
+
+ assert minmax["cur_min"] == np.min(a)
+ assert minmax["cur_max"] == np.max(a)
+
+
+
if __name__ == "__main__":
# make sure that import failures get reported, instead of skipping the
@@ -651,3 +735,5 @@ if __name__ == "__main__":
else:
from py.test.cmdline import main
main([__file__])
+
+# vim: filetype=pyopencl