Multi-dimensional arrays on the Compute Device

Vector Types

class pyopencl.array.vec

All of OpenCL’s supported vector types, such as float3 and long4 are available as numpy data types within this class. These numpy.dtype instances have field names of x, y, z, and w just like their OpenCL counterparts. They will work both for parameter passing to kernels as well as for passing data back and forth between kernels and Python code. For each type, a make_type function is also provided (e.g. make_float3(x,y,z)).

The Array Class

class pyopencl.array.DefaultAllocator(context, flags=pyopencl.mem_flags.READ_WRITE)

An alias for pyopencl.tools.CLAllocator.

class pyopencl.array.Array(cqa, shape, dtype, order="C", *, allocator=None, base=None, data=None, queue=None)

A numpy.ndarray work-alike that stores its data and performs its computations on the compute device. shape and dtype work exactly as in numpy. Arithmetic methods in Array support the broadcasting of scalars. (e.g. array+5)

cqa can be a pyopencl.Context, pyopencl.CommandQueue or an allocator, as described below. If it is either of the latter two, the queue or allocator arguments may not be passed.

queue (or cqa, as the case may be) specifies the queue in which the array carries out its computations by default.

allocator is a callable that, upon being called with an argument of the number of bytes to be allocated, returns an object that can be cast to an int representing the address of the newly allocated memory. (See DefaultAllocator.)

Changed in version 2011.1: Renamed context to cqa, made it general-purpose.

All arguments beyond order should be considered keyword-only.

data

The pyopencl.MemoryObject instance created for the memory that backs this Array.

shape

The tuple of lengths of each dimension in the array.

dtype

The numpy.dtype of the items in the GPU array.

size

The number of meaningful entries in the array. Can also be computed by multiplying up the numbers in shape.

mem_size

The total number of entries, including padding, that are present in the array.

nbytes

The size of the entire array in bytes. Computed as size times dtype.itemsize.

strides

Tuple of bytes to step in each dimension when traversing an array.

flags

Return an object with attributes c_contiguous, f_contiguous and forc, which may be used to query contiguity properties in analogy to numpy.ndarray.flags.

__len__()

Returns the size of the leading dimension of self.

reshape(shape)

Returns an array containing the same data with a new shape.

ravel()

Returns flattened array containing the same data.

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.

set(ary, queue=None, async=False)

Transfer the contents the numpy.ndarray object ary onto the device.

ary must have the same dtype and size (not necessarily shape) as self.

get(queue=None, ary=None, async=False)

Transfer the contents of self into ary or a newly allocated numpy.ndarray. If ary is given, it must have the right size (not necessarily shape) and dtype.

__str__()
__repr__()
mul_add(self, selffac, other, otherfac, queue=None):

Return selffac*self + otherfac*other.

__add__(other)
__sub__(other)
__iadd__(other)
__isub__(other)
__neg__(other)
__mul__(other)
__div__(other)
__rdiv__(other)
__pow__(other)
__abs__()

Return a Array containing the absolute value of each element of self.

fill(scalar, queue=None)

Fill the array with scalar.

astype(dtype, queue=None)

Return self, cast to dtype.

Constructing Array Instances

pyopencl.array.to_device(queue, ary, allocator=None, async=False)

Return a Array that is an exact copy of the numpy.ndarray instance ary.

See Array for the meaning of allocator.

Changed in version 2011.1: context argument was deprecated.

pyopencl.array.empty(queue, shape, dtype, order="C", allocator=None, base=None, data=None)

A synonym for the Array constructor.

pyopencl.array.zeros(queue, shape, dtype, order="C", allocator=None)

Same as empty(), but the Array is zero-initialized before being returned.

Changed in version 2011.1: context argument was deprecated.

pyopencl.array.empty_like(other_ary)

Make a new, uninitialized Array having the same properties as other_ary.

pyopencl.array.zeros_like(other_ary)

Make a new, zero-initialized Array having the same properties as other_ary.

pyopencl.array.arange(queue, start, stop, step, dtype=None, allocator=None)

Create a Array filled with numbers spaced step apart, starting from start and ending at stop.

For floating point arguments, the length of the result is ceil((stop - start)/step). This rule may result in the last element of the result being greater than stop.

dtype, if not specified, is taken as the largest common type of start, stop and step.

Changed in version 2011.1: context argument was deprecated.

Changed in version 2011.2: allocator keyword argument was added.

pyopencl.array.take(a, indices, out=None, queue=None)

Return the Array [a[indices[0]], ..., a[indices[n]]]. For the moment, a must be a type that can be bound to a texture.

Conditionals

pyopencl.array.if_positive(criterion, then_, else_, out=None, queue=None)

Return an array like then_, which, for the element at index i, contains then_[i] if criterion[i]>0, else else_[i]. (added in 0.94)

pyopencl.array.maximum(a, b, out=None, queue=None)

Return the elementwise maximum of a and b. (added in 0.94)

pyopencl.array.minimum(a, b, out=None, queue=None)

Return the elementwise minimum of a and b. (added in 0.94)

Reductions

pyopencl.array.sum(a, dtype=None, queue=None)
pyopencl.array.dot(a, b, dtype=None, queue=None)
pyopencl.array.subset_dot(subset, a, b, dtype=None, queue=None)
pyopencl.array.max(a, queue=None)
pyopencl.array.min(a, queue=None)
pyopencl.array.subset_max(subset, a, queue=None)
pyopencl.array.subset_min(subset, a, queue=None)

See also Custom Reductions.

Elementwise Functions on Arrray Instances

The pyopencl.clmath module contains exposes array versions of the C functions available in the OpenCL standard. (See table 6.8 in the spec.)

pyopencl.clmath.acos(array, queue=None)
pyopencl.clmath.acosh(array, queue=None)
pyopencl.clmath.acospi(array, queue=None)
pyopencl.clmath.asin(array, queue=None)
pyopencl.clmath.asinh(array, queue=None)
pyopencl.clmath.asinpi(array, queue=None)
pyopencl.clmath.atan(array, queue=None)
pyopencl.clmath.atanh(array, queue=None)
pyopencl.clmath.atanpi(array, queue=None)
pyopencl.clmath.cbrt(array, queue=None)
pyopencl.clmath.ceil(array, queue=None)
pyopencl.clmath.cos(array, queue=None)
pyopencl.clmath.cosh(array, queue=None)
pyopencl.clmath.cospi(array, queue=None)
pyopencl.clmath.erfc(array, queue=None)
pyopencl.clmath.erf(array, queue=None)
pyopencl.clmath.exp(array, queue=None)
pyopencl.clmath.exp2(array, queue=None)
pyopencl.clmath.exp10(array, queue=None)
pyopencl.clmath.expm1(array, queue=None)
pyopencl.clmath.fabs(array, queue=None)
pyopencl.clmath.floor(array, queue=None)
pyopencl.clmath.fmod(arg, mod, queue=None)

Return the floating point remainder of the division arg/mod, for each element in arg and mod.

pyopencl.clmath.frexp(arg, queue=None)

Return a tuple (significands, exponents) such that arg == significand * 2**exponent.

pyopencl.clmath.ilogb(array, queue=None)
pyopencl.clmath.ldexp(significand, exponent, queue=None)

Return a new array of floating point values composed from the entries of significand and exponent, paired together as result = significand * 2**exponent.

pyopencl.clmath.lgamma(array, queue=None)
pyopencl.clmath.log(array, queue=None)
pyopencl.clmath.log2(array, queue=None)
pyopencl.clmath.log10(array, queue=None)
pyopencl.clmath.log1p(array, queue=None)
pyopencl.clmath.logb(array, queue=None)
pyopencl.clmath.modf(arg, queue=None)

Return a tuple (fracpart, intpart) of arrays containing the integer and fractional parts of arg.

pyopencl.clmath.nan(array, queue=None)
pyopencl.clmath.rint(array, queue=None)
pyopencl.clmath.round(array, queue=None)
pyopencl.clmath.sin(array, queue=None)
pyopencl.clmath.sinh(array, queue=None)
pyopencl.clmath.sinpi(array, queue=None)
pyopencl.clmath.sqrt(array, queue=None)
pyopencl.clmath.tan(array, queue=None)
pyopencl.clmath.tanh(array, queue=None)
pyopencl.clmath.tanpi(array, queue=None)
pyopencl.clmath.tgamma(array, queue=None)
pyopencl.clmath.trunc(array, queue=None)

Generating Arrays of Random Numbers

class pyopencl.clrandom.RanluxGenerator(self, queue, num_work_items, max_work_items, luxury=2, seed=None)
Parameters:
  • queuepyopencl.CommandQueue, only used for initialization
  • luxury – the “luxury value” of the generator, and should be 0-4, where 0 is fastest and 4 produces the best numbers. It can also be >=24, in which case it directly sets the p-value of RANLUXCL.
  • num_work_items – is the number of generators to initialize, usually corresponding to the number of work-items in the NDRange RANLUXCL will be used with.
  • max_work_items – should reflect the maximum number of work-items that will be used on any parallel instance of RANLUXCL. So for instance if we are launching 5120 work-items on GPU1 and 10240 work-items on GPU2, GPU1’s RANLUXCLTab would be generated by calling ranluxcl_intialization with numWorkitems = 5120 while GPU2’s RANLUXCLTab would use numWorkitems = 10240. However maxWorkitems must be at least 10240 for both GPU1 and GPU2, and it must be set to the same value for both.

New in version 2011.2.

state

A pyopencl.array.Array containing the state of the generator.

nskip

nskip is an integer which can (optionally) be defined in the kernel code as RANLUXCL_NSKIP. If this is done the generator will be faster for luxury setting 0 and 1, or when the p-value is manually set to a multiple of 24.

fill_uniform(ary, a=0, b=1, queue=None)

Fill ary with uniformly distributed random numbers in the interval (a, b), endpoints excluded.

uniform(queue, shape, dtype, order="C", allocator=None, base=None, data=None, a=0, b=1)

Make a new empty array, apply fill_uniform() to it.

fill_normal(ary, mu=0, sigma=1, queue=None):

Fill ary with normally distributed numbers with mean mu and standard deviation sigma.

normal(queue, shape, dtype, order="C", allocator=None, base=None, data=None, mu=0, sigma=1)

Make a new empty array, apply fill_normal() to it.

synchronize()

The generator gets inefficient when different work items invoke the generator a differing number of times. This function ensures efficiency.

pyopencl.clrandom.rand(queue, shape, dtype, a=0, b=1)

Return an array of shape filled with random values of dtype in the range [a,b).

pyopencl.clrandom.fill_rand(result, queue=None, a=0, b=1)

Fill result with random values of dtype in the range [0,1).

PyOpenCL now includes and uses the RANLUXCL random number generator by Ivar Ursin Nikolaisen. In addition to being usable through the convenience functions above, it is available in any piece of code compiled through PyOpenCL by:

#include <pyopencl-ranluxcl.cl>

The RANLUX generator is described in the following two articles. If you use the generator for scientific purposes, please consider citing them:

Single-pass Custom Expression Evaluation

Evaluating involved expressions on pyopencl.array.Array instances can be somewhat inefficient, because a new temporary is created for each intermediate result. The functionality in the module pyopencl.elementwise contains tools to help generate kernels that evaluate multi-stage expressions on one or several operands in a single pass.

class pyopencl.elementwise.ElementwiseKernel(context, arguments, operation, name="kernel", preamble="", options=[])

Generate a kernel that takes a number of scalar or vector arguments and performs the scalar operation on each entry of its arguments, if that argument is a vector.

arguments is specified as a string formatted as a C argument list. operation is specified as a C assignment statement, without a semicolon. Vectors in operation should be indexed by the variable i.

name specifies the name as which the kernel is compiled, and options are passed unmodified to pyopencl.Program.build().

preamble is a piece of C source code that gets inserted outside of the function context in the elementwise operation’s kernel source code.

__call__(*args)

Invoke the generated scalar kernel. The arguments may either be scalars or GPUArray instances.

Here’s a usage example:

import pyopencl as cl
import pyopencl.array as cl_array
import numpy

ctx = cl.create_some_context()
queue = cl.CommandQueue(ctx)

n = 10
a_gpu = cl_array.to_device(
        ctx, queue, numpy.random.randn(n).astype(numpy.float32))
b_gpu = cl_array.to_device(
        ctx, queue, numpy.random.randn(n).astype(numpy.float32))

from pyopencl.elementwise import ElementwiseKernel
lin_comb = ElementwiseKernel(ctx,
        "float a, float *x, "
        "float b, float *y, "
        "float *z",
        "z[i] = a*x[i] + b*y[i]",
        "linear_combination")

c_gpu = cl_array.empty_like(a_gpu)
lin_comb(5, a_gpu, 6, b_gpu, c_gpu)

import numpy.linalg as la
assert la.norm((c_gpu - (5*a_gpu+6*b_gpu)).get()) < 1e-5

(You can find this example as examples/demo_elementwise.py in the PyOpenCL distribution.)

Custom Reductions

class pyopencl.reduction.ReductionKernel(ctx, dtype_out, neutral, reduce_expr, map_expr=None, arguments=None, name="reduce_kernel", options=[], preamble="")

Generate a kernel that takes a number of scalar or vector arguments (at least one vector argument), performs the map_expr on each entry of the vector argument and then the reduce_expr on the outcome of that. neutral serves as an initial value. preamble offers the possibility to add preprocessor directives and other code (such as helper functions) to be added before the actual reduction kernel code.

Vectors in map_expr should be indexed by the variable i. reduce_expr uses the formal values “a” and “b” to indicate two operands of a binary reduction operation. If you do not specify a map_expr, “in[i]” – and therefore the presence of only one input argument – is automatically assumed.

dtype_out specifies the numpy.dtype in which the reduction is performed and in which the result is returned. neutral is specified as float or integer formatted as string. reduce_expr and map_expr are specified as string formatted operations and arguments is specified as a string formatted as a C argument list. name specifies the name as which the kernel is compiled. options are passed unmodified to pyopencl.Program.build(). preamble specifies a string of code that is inserted before the actual kernels.

__call__(*args, queue=None)

Here’s a usage example:

a = pyopencl.array.arange(queue, 400, dtype=numpy.float32)
b = pyopencl.array.arange(queue, 400, dtype=numpy.float32)

krnl = ReductionKernel(ctx, numpy.float32, neutral="0",
        reduce_expr="a+b", map_expr="x[i]*y[i]",
        arguments="__global float *x, __global float *y")

my_dot_prod = krnl(a, b).get()

Parallel Scan / Prefix Sum

class pyopencl.scan.ExclusiveScanKernel(ctx, dtype, scan_expr, neutral, name_prefix="scan", options=[], preamble="", devices=None)

Generates a kernel that can compute a prefix sum using any associative operation given as scan_expr. scan_expr uses the formal values “a” and “b” to indicate two operands of an associative binary operation. neutral is the neutral element of scan_expr, obeying scan_expr(a, neutral) == a.

dtype specifies the type of the arrays being operated on. name_prefix is used for kernel names to ensure recognizability in profiles and logs. options is a list of compiler options to use when building. preamble specifies a string of code that is inserted before the actual kernels. devices may be used to restrict the set of devices on which the kernel is meant to run. (defaults to all devices in the context ctx.

__call__(self, input_ary, output_ary=None, allocator=None, queue=None)
class pyopencl.scan.InclusiveScanKernel(dtype, scan_expr, neutral=None, name_prefix="scan", options=[], preamble="", devices=None)

Works like ExclusiveScanKernel. Unlike the exclusive case, neutral is not required.

Here’s a usage example:

knl = InclusiveScanKernel(context, np.int32, "a+b")

n = 2**20-2**18+5
host_data = np.random.randint(0, 10, n).astype(np.int32)
dev_data = cl_array.to_device(queue, host_data)

knl(dev_data)
assert (dev_data.get() == np.cumsum(host_data, axis=0)).all()

Custom data types in Reduction and Scan

If you would like to use your own (struct/union/whatever) data types in scan and reduction, define those types in the preamble and let PyOpenCL know about them using this function:

pyopencl.tools.register_dtype(dtype, name)

dtype is a numpy.dtype().

Fast Fourier Transforms

Bogdan Opanchuk’s pyfft package offers a variety of GPU-based FFT implementations.