Public API

API discovery

class grunnur.API

A generalized GPGPU API.

classmethod all_available() → list[API]

Returns a list of API objects for which backends are available.

classmethod all_by_shortcut(shortcut: str | None = None) → list[API]

If shortcut is a string, returns a list of one API object whose id attribute has its shortcut attribute equal to it (or raises an error if it was not found, or its backend is not available).

If shortcut is None, returns a list of all available API objects.

Parameters:

shortcut – an API shortcut to match.

static any() → API

Returns an API object for some available backend.

Raises RuntimeError if no backends are available.

static cuda() → API

Returns a CUDA API object, if CUDA backend (that is, pycuda package) is available.

classmethod from_api_id(api_id: APIID) → API

Creates an API object out of an identifier.

Parameters:

api_id – API identifier.

static opencl() → API

Returns an OpenCL API object, if OpenCL backend (that is, pyopencl package) is available.

id: APIID

This API’s ID.

property platforms: list[Platform]

A list of this API’s Platform objects.

shortcut: str

A shortcut for this API (to use in all_by_shortcut(), usually coming from some kind of a CLI). Equal to id.shortcut.

class grunnur._adapter_base.APIID

An ID of an API object.

shortcut: str

This API’s shortcut.

Platforms

A platform is an OpenCL term, but we use it for CUDA API as well for the sake of uniformity. Naturally, there will always be a single (dummy) platform in CUDA.

class grunnur.Platform

A generalized GPGPU platform.

classmethod all(api: API) → list[Platform]

Returns a list of platforms available for the given API.

Parameters:

api – the API to search in.

classmethod all_filtered(api: API, filter_: PlatformFilter | None = None) → list[Platform]

Returns a list of all platforms satisfying the given criteria in the given API. If filter is not provided, returns all the platforms.

classmethod from_backend_platform(obj: Any) → Platform

Wraps a backend platform object into a Grunnur platform object.

classmethod from_index(api: API, platform_idx: int) → Platform

Creates a platform based on its index in the list returned by the API.

Parameters:

• api – the API to search in.

• platform_idx – the target platform’s index.

api: API

The API object this platform belongs to.

property devices: list[Device]

A list of this device’s Device objects.

name: str

The platform’s name.

vendor: str

The platform’s vendor.

version: str

The platform’s version.

class grunnur.PlatformFilter

Bases: NamedTuple

A set of filters for platform discovery.

Create new instance of PlatformFilter(include_masks, exclude_masks)

exclude_masks: list[str] | None

A list of strings (treated as regexes), neither of which must match the platform name.

include_masks: list[str] | None

A list of strings (treated as regexes), one of which must match the platform name.

Devices

class grunnur.Device

A generalized GPGPU device.

classmethod all(platform: Platform) → list[Device]

Returns a list of devices available for the given platform.

Parameters:

platform – the platform to search in.

classmethod all_filtered(platform: Platform, filter_: DeviceFilter | None = None) → list[Device]

Returns a list of all devices satisfying the given criteria in the given platform. If filter is not provided, returns all the devices.

classmethod from_backend_device(obj: Any) → Device

Wraps a backend device object into a Grunnur device object.

classmethod from_index(platform: Platform, device_idx: int) → Device

Creates a device based on its index in the list returned by the API.

Parameters:

• platform – the API to search in.

• device_idx – the target device’s index.

name: str

This device’s name.

property params: DeviceParameters

Returns a DeviceParameters object associated with this device.

platform: Platform

The Platform object this device belongs to.

class grunnur.DeviceFilter

A set of filters for device discovery.

Create new instance of DeviceFilter(include_masks, exclude_masks, unique_only, exclude_pure_parallel)

exclude_masks: list[str] | None

A list of strings (treated as regexes), neither of which must match the device name.

exclude_pure_parallel: bool

If True, exclude devices with max_total_local_size equal to 1 (that is the ones that do not support workgroup synchronization).

include_masks: list[str] | None

A list of strings (treated as regexes), one of which must match the device name.

unique_only: bool

If True, only return devices with unique names.

class grunnur.DeviceParameters

An object containing device’s specifications.

abstract align_words(word_size: int) → int

The distance (in word_size units) between global memory base addresses that allow accesses of word-size word_size bytes to get coalesced.

abstract property compute_units: int

The number of multiprocessors (CUDA)/compute units (OpenCL) for the device.

abstract property local_mem_banks: int

The number of independent channels for shared (CUDA)/local (OpenCL) memory, which can be used from one warp without request serialization.

abstract property local_mem_size: int

The size of shared (CUDA)/local (OpenCL) memory (in bytes).

abstract property max_local_sizes: tuple[int, ...]

The maximum number of threads in one block (CUDA), or work items in one work group (OpenCL) for each of the available dimensions.

abstract property max_num_groups: tuple[int, ...]

The maximum number of blocks (CUDA)/work groups (OpenCL) for each of the available dimensions.

abstract property max_total_local_size: int

The maximum total number of threads in one block (CUDA), or work items in one work group (OpenCL).

abstract property type: DeviceType

Device type.

abstract property warp_size: int

The number of threads (CUDA)/work items (OpenCL) that are executed synchronously (within one multiprocessor/compute unit).

class grunnur._adapter_base.DeviceType

An enum representing a device’s type.

CPU = 1

CPU type

GPU = 2

GPU type

Device discovery

grunnur.platforms_and_devices_by_mask(api: API, quantity: int | None = 1, device_filter: DeviceFilter | None = None, platform_filter: PlatformFilter | None = None) → list[tuple[Platform, list[Device]]]

Returns all tuples (platform, list of devices) where the platform name and device names satisfy the given criteria, and there are at least quantity devices in the list.

grunnur.select_devices(api: API, *, interactive: bool = False, quantity: int | None = 1, device_filter: DeviceFilter | None = None, platform_filter: PlatformFilter | None = None) → list[Device]

Using the results from platforms_and_devices_by_mask(), either lets the user select the devices (from the ones matching the criteria) interactively, or takes the first matching list of quantity devices.

Parameters:

• interactive – if True, shows a dialog to select the devices. If False, selects the first matching ones.

• quantity – passed to platforms_and_devices_by_mask().

• device_filters – passed to platforms_and_devices_by_mask().

Contexts

class grunnur.Context

GPGPU context.

deactivate() → None

For CUDA API: deactivates this context, popping all the CUDA context objects from the stack. Other APIs: no effect.

Only call it if you need to manage CUDA contexts manually, and created this object with take_ownership = False. If take_ownership = True contexts will be deactivated automatically in the destructor.

classmethod from_backend_contexts(backend_contexts: Sequence[Any], *, take_ownership: bool = False) → Context

Creates a context from a single or several backend device contexts. If take_ownership is True, this object will be responsible for the lifetime of backend context objects (only important for the CUDA backend).

classmethod from_backend_devices(backend_devices: Sequence[Any]) → Context

Creates a context from a single or several backend device objects.

classmethod from_criteria(api: API, *, interactive: bool = False, devices_num: int | None = 1, device_filter: DeviceFilter | None = None, platform_filter: PlatformFilter | None = None) → Context

Finds devices matching the given criteria and creates a Context object out of them.

Parameters:

• interactive – passed to select_devices().

• devices_num – passed to select_devices() as quantity.

• device_filters – passed to select_devices().

classmethod from_devices(devices: Sequence[Device]) → Context

Creates a context from a device or an iterable of devices.

Parameters:

devices – one or several devices to use.

api: API

The API this context is based on.

property devices: BoundMultiDevice

Returns the BoundMultiDevice encompassing all the devices in this context.

platform: Platform

The platform this context is based on.

class grunnur.BoundDevice

A Device object in a Context.

context: Context

The context this device belongs to.

class grunnur._context.BoundMultiDevice

Bases: Sequence[BoundDevice]

A sequence of bound devices belonging to the same context.

__getitem__(idx: int) → BoundDevice

__getitem__(idx: slice | Iterable[int]) → BoundMultiDevice

Given a single index, returns a single BoundDevice. Given a sequence of indices, returns a BoundMultiDevice object containing respective devices.

The indices correspond to the list of devices used to create this context.

classmethod from_bound_devices(devices: Sequence[BoundDevice]) → BoundMultiDevice

Creates this object from a sequence of bound devices (note that a BoundMultiDevice object itself can serve as such a sequence).

context: Context

The context these devices belong to.

Queues

class grunnur.Queue(device: BoundDevice)

A queue on a single device.

Parameters:

device – a device on which to create a queue.

synchronize() → None

Blocks until sub-queues on all devices are empty.

device: BoundDevice

Device on which this queue operates.

class grunnur.MultiQueue(queues: Sequence[Queue])

A queue on multiple devices.

Parameters:

queues – single-device queues (must belong to distinct devices and the same context).

classmethod on_devices(devices: Iterable[BoundDevice]) → MultiQueue

Creates a queue from provided devices (belonging to the same context).

synchronize() → None

Blocks until queues on all devices are empty.

devices: BoundMultiDevice

Multi-device on which this queue operates.

queues: dict[BoundDevice, Queue]

Single-device queues associated with device indices.

Buffers and arrays

class grunnur.Buffer

A memory buffer on device.

classmethod allocate(device: BoundDevice, size: int) → Buffer

Allocate a buffer of size bytes.

Parameters:

• device – the device on which this buffer will be allocated.

• size – the buffer’s size in bytes.

get(queue: Queue, host_array: NDArray[Any], *, async_: bool = False) → None

Copy the contents of the buffer to the host array.

Parameters:

• queue – the queue to use for the transfer.

• host_array – the destination array.

• async – if True, the transfer is performed asynchronously.

get_sub_region(origin: int, size: int) → Buffer

Return a buffer object describing a subregion of this buffer.

Parameters:

• origin – the offset of the subregion.

• size – the size of the subregion.

set(queue: Queue, buf: NDArray[Any] | Buffer, *, sync: bool = False) → None

Copy the contents of the host array or another buffer to this buffer.

Parameters:

• queue – the queue to use for the transfer.

• buf – the source - numpy array or a Buffer object.

• sync – if True, the transfer blocks until completion.

device: BoundDevice

Device on which this buffer is allocated.

property offset: int

Offset of this buffer (in bytes) from the beginning of the physical allocation it resides in.

property size: int

The buffer’s size (in bytes).

class grunnur.ArrayMetadata(shape: Iterable[int] | int, dtype: DTypeLike, *, strides: Iterable[int] | None = None, first_element_offset: int | None = None, buffer_size: int | None = None)

Bases: AsArrayMetadata

A helper object for array-like classes that handles shape/strides/buffer size checks without actual data attached to it.

__getitem__(slices: slice | tuple[slice, ...]) → ArrayMetadata

Returns the view of this metadata with the given ranges, with the offsets and buffer size corresponding to the original buffer.

as_array_metadata() → ArrayMetadata

Returns array metadata representing this object.

get_sub_region(origin: int, size: int) → ArrayMetadata

Returns the same metadata shape-wise, but for the given subregion of the original buffer.

with_(dtype: DTypeLike | None = None) → ArrayMetadata

Replaces a property of the metadata and returns a new metadata object.

buffer_size: int

The size of the buffer this array resides in.

dtype: numpy.dtype[Any]

Array item data type.

first_element_offset: int

The offset of the first element (that is, the one with the all indices equal to 0).

is_contiguous: bool

If True, means that array’s data forms a continuous chunk of memory.

min_offset: int

The minimum offset of an array element described by this metadata.

shape: tuple[int, ...]

Array shape.

span: int

The minimum size of the buffer that fits all the elements described by this metadata.

strides: tuple[int, ...]

Array strides.

class grunnur.AsArrayMetadata

An abstract class for any object allowing conversion to ArrayMetadata.

abstract as_array_metadata() → ArrayMetadata

Returns array metadata representing this object.

class grunnur.ArrayLike

Bases: Protocol

A protocol for an array-like object supporting views via __getitem__(). numpy.ndarray or Array follow this protocol.

__getitem__(slices: slice | tuple[slice, ...]) → _ArrayLike

Returns a view of this array.

property dtype: dtype[Any]

The type of an array element.

property shape: tuple[int, ...]

Array shape.

class grunnur._array._ArrayLike

Any type that follows the ArrayLike protocol.

alias of TypeVar(‘_ArrayLike’, bound=ArrayLike)

class grunnur.Array

Bases: AsArrayMetadata

Array on a single device.

__getitem__(slices: slice | tuple[slice, ...]) → Array

Returns a view of this array.

as_array_metadata() → ArrayMetadata

Returns array metadata representing this object.

classmethod empty(device: BoundDevice, shape: Iterable[int] | int, dtype: DTypeLike, strides: Iterable[int] | None = None, first_element_offset: int = 0, allocator: Callable[[BoundDevice, int], Buffer] | None = None) → Array

Creates an empty array.

classmethod empty_like(device: BoundDevice, array_like: AsArrayMetadata | NDArray[Any], allocator: Callable[[BoundDevice, int], Buffer] | None = None) → Array

Creates an empty array with the same metadata as array_like.

In case of a numpy array, uses its shape, dtype, and strides.

classmethod empty_with_metadata(device: BoundDevice, metadata: ArrayMetadata, allocator: Callable[[BoundDevice, int], Buffer] | None = None) → Array

Creates an empty array with the given metadata.

classmethod from_host(queue_or_device: Queue | BoundDevice, host_arr: NDArray[Any]) → Array

Creates an array object from a host array.

Parameters:

• queue_or_device – the queue to use for the transfer, or the target device. In the latter case an operation will be performed synchronously.

• host_arr – the source array.

get(queue: Queue, dest: NDArray[Any] | None = None, *, async_: bool = False) → NDArray[Any]

Copies the contents of the array to the host array and returns it.

Parameters:

• queue – the queue to use for the transfer.

• dest – the destination array. If None, the target array is created.

• async – if True, the transfer is performed asynchronously.

minimum_subregion() → Array

Returns a new array with the same metadata and the buffer substituted with the minimum-sized subregion of the original buffer, such that all the elements described by the metadata still fit in it.

set(queue: Queue, array: NDArray[Any] | Array, *, sync: bool = False) → None

Copies the contents of the host array to the array.

Parameters:

• queue – the queue to use for the transfer.

• array – the source array.

• sync – if True, the transfer blocks until completion.

device: BoundDevice

Device this array is allocated on.

dtype: dtype[Any]

Array item data type.

metadata: ArrayMetadata

Array metadata object.

shape: tuple[int, ...]

Array shape.

strides: tuple[int, ...]

Array strides.

class grunnur._array.BaseSplay

Base class for splay strategies for MultiArray.

abstract __call__(arr: _ArrayLike, devices: Sequence[BoundDevice]) → dict[BoundDevice, _ArrayLike]

Creates a dictionary of views of an array-like object for each of the given devices.

Parameters:

• arr – an array-like object.

• devices – a multi-device object.

class grunnur.MultiArray

An array on multiple devices.

class CloneSplay

Copies the given array to each device.

class EqualSplay

Splays the given array equally between the devices using the outermost dimension. The outermost dimension should be larger or equal to the number of devices.

classmethod empty(devices: BoundMultiDevice, shape: Iterable[int], dtype: DTypeLike, allocator: Callable[[BoundDevice, int], Buffer] | None = None, splay: BaseSplay | None = None) → MultiArray

Creates an empty array.

Parameters:

• devices – devices on which the sub-arrays will be allocated.

• shape – array shape.

• dtype – array data type.

• allocator – an optional callable taking two integer arguments (the device to allocate it on and the buffer size in bytes) and returning a Buffer object. If None, will use Buffer.allocate().

• splay – the splay strategy (if None, an EqualSplay object is used).

classmethod from_host(mqueue: MultiQueue, host_arr: NDArray[Any], splay: BaseSplay | None = None) → MultiArray

Creates an array object from a host array.

Parameters:

• mqueue – the queue to use for the transfer.

• host_arr – the source array.

• splay – the splay strategy (if None, an EqualSplay object is used).

get(mqueue: MultiQueue, dest: NDArray[Any] | None = None, *, async_: bool = False) → NDArray[Any]

Copies the contents of the array to the host array and returns it.

Parameters:

• mqueue – the queue to use for the transfer.

• dest – the destination array. If None, the target array is created.

• async – if True, the transfer is performed asynchronously.

set(mqueue: MultiQueue, array: NDArray[Any] | MultiArray, *, sync: bool = False) → None

Copies the contents of the host array to the array.

Parameters:

• mqueue – the queue to use for the transfer.

• array – the source array.

• sync – if True, the transfer blocks until completion.

devices: BoundMultiDevice

Devices on which the sub-arrays are allocated

dtype: dtype[Any]

Array item data type.

shape: tuple[int, ...]

Array shape.

shapes: dict[BoundDevice, tuple[int, ...]]

Sub-array shapes matched to device indices.

Programs and kernels

class grunnur.Program(devices: Sequence[BoundDevice], template_src: str | Callable[..., str] | DefTemplate | Snippet, *, no_prelude: bool = False, fast_math: bool = False, render_args: Sequence[Any] = (), render_globals: Mapping[str, Any] = {}, compiler_options: Sequence[str] = [], keep: bool = False, constant_arrays: Mapping[str, AsArrayMetadata] = {})

A compiled program on device(s).

Parameters:

• devices – a single- or a multi-device object on which to compile this program.

• template_src – a string with the source code, or a Mako template source to render.

• no_prelude – do not add prelude to the rendered source.

• fast_math – compile using fast (but less accurate) math functions.

• render_args – a list of positional args to pass to the template.

• render_globals – a dictionary of globals to pass to the template.

• compiler_options – a list of options to pass to the backend compiler.

• keep – keep the intermediate files in a temporary directory.

• constant_arrays – (CUDA only) a dictionary name: (size, dtype) of global constant arrays to be declared in the program.

set_constant_array(queue: Queue, name: str, arr: Array | Buffer | NDArray[Any]) → None

Uploads a constant array to the context’s devices (CUDA only).

Parameters:

• queue – the queue to use for the transfer.

• name – the name of the constant array symbol in the code.

• arr – either a device or a host array.

devices: BoundMultiDevice

The devices on which this program was compiled.

kernel: KernelHub

An object whose attributes are Kernel objects with the corresponding names.

sources: dict[BoundDevice, str]

Source files used for each device.

class grunnur._program.KernelHub

An object providing access to the host program’s kernels.

__getattr__(kernel_name: str) → Kernel

Returns a Kernel object for a function (CUDA)/kernel (OpenCL) with the name kernel_name.

class grunnur._program.Kernel

A kernel compiled for multiple devices.

__call__(queue: Queue | MultiQueue, global_size: Sequence[int] | Mapping[BoundDevice, Sequence[int]], local_size: Sequence[int] | None | Mapping[BoundDevice, Sequence[int] | None] = None, *args: Mapping[BoundDevice, Array | Buffer | generic] | MultiArray | Array | Buffer | generic, cu_dynamic_local_mem: int = 0) → Any

A shortcut for Kernel.prepare() and subsequent PreparedKernel.__call__(). See their doc entries for details.

prepare(global_size: Sequence[int] | Mapping[BoundDevice, Sequence[int]], local_size: Sequence[int] | None | Mapping[BoundDevice, Sequence[int] | None] = None) → PreparedKernel

Prepares the kernel for execution.

If local_size or global_size are integer, they will be treated as 1-tuples.

One can pass specific global and local sizes for each device using dictionaries keyed with device indices. This achieves another purpose: the kernel will only be prepared for those devices, and not for all devices available in the context.

Parameters:

• global_size – the total number of threads (CUDA)/work items (OpenCL) in each dimension (column-major). Note that there may be a maximum size in each dimension as well as the maximum number of dimensions. See DeviceParameters for details.

• local_size – the number of threads in a block (CUDA)/work items in a work group (OpenCL) in each dimension (column-major). If None, it will be chosen automatically.

property max_total_local_sizes: dict[BoundDevice, int]

The maximum possible number of threads in a block (CUDA)/work items in a work group (OpenCL) for this kernel.

class grunnur._program.PreparedKernel

A kernel specialized for execution on a set of devices with all possible preparations and checks performed.

__call__(queue: Queue | MultiQueue, *args: Mapping[BoundDevice, Array | Buffer | generic] | MultiArray | Array | Buffer | generic, cu_dynamic_local_mem: int = 0) → Any

Enqueues the kernel on the devices in the given queue. The kernel must have been prepared for all of these devices.

If an argument is a Array or Buffer object, it must belong to the device on which the kernel is being executed (so queue must only have one device).

If an argument is a MultiArray, it should have subarrays on all the devices from the given queue.

If an argument is a numpy scalar, it will be passed to the kernel directly.

If an argument is a integer-keyed dict, its values corresponding to the device indices the kernel is executed on will be passed as kernel arguments.

Parameters:

• cu_dynamic_local_mem – CUDA only. The size of dynamically allocated local (shared in CUDA terms) memory, in bytes. That is, the size of extern __shared__ arrays in CUDA kernels.

• args – kernel arguments.

Returns:

a list of Event objects for enqueued kernels in case of PyOpenCL.

Static kernels

class grunnur.StaticKernel(devices: Sequence[BoundDevice], template_src: str | Callable[..., str] | DefTemplate | Snippet, name: str, global_size: Sequence[int] | Mapping[BoundDevice, Sequence[int]], *, local_size: Sequence[int] | None | Mapping[BoundDevice, Sequence[int] | None] = None, render_args: Sequence[Any] = (), render_globals: Mapping[str, Any] = {}, constant_arrays: Mapping[str, AsArrayMetadata] = {}, keep: bool = False, fast_math: bool = False, compiler_options: Iterable[str] = [])

An object containing a GPU kernel with fixed call sizes.

The globals for the source template will contain an object with the name static of the type VsizeModules containing the id/size functions to be used instead of regular ones.

Parameters:

• devices – a single- or a multi-device object on which to compile this program.

• template_src – a string with the source code, or a Mako template source to render.

• name – the kernel’s name.

• global_size – see prepare().

• local_size – see prepare().

• render_globals – a dictionary of globals to pass to the template.

• constant_arrays – (CUDA only) a dictionary name: (size, dtype) of global constant arrays to be declared in the program.

__call__(queue: Queue | MultiQueue, *args: Mapping[BoundDevice, Array | Buffer | numpy.generic] | MultiArray | Array | Buffer | numpy.generic) → Any

Execute the kernel. In case of the OpenCL backend, returns a pyopencl.Event object.

Parameters:

• queue – the multi-device queue to use.

• args – kernel arguments. See grunnur._program.PreparedKernel.__call__().

set_constant_array(queue: Queue, name: str, arr: Array | Buffer | NDArray[Any]) → None

Uploads a constant array to the context’s devices (CUDA only).

Parameters:

• queue – the queue to use for the transfer.

• name – the name of the constant array symbol in the code.

• arr – either a device or a host array.

devices: BoundMultiDevice

Devices on which this kernel was compiled.

queue: Queue

The queue this static kernel was compiled and prepared for.

sources: dict[BoundDevice, str]

Source files used for each device.

class grunnur._vsize.VsizeModules(local_id: Module, local_size: Module, group_id: Module, num_groups: Module, global_id: Module, global_size: Module, global_flat_id: Module, global_flat_size: Module, skip: Module)

A collection of modules passed to grunnur.StaticKernel. Should be used instead of regular group/thread id functions.

Create new instance of VsizeModules(local_id, local_size, group_id, num_groups, global_id, global_size, global_flat_id, global_flat_size, skip)

global_flat_id: Module

Provides the function VSIZE_T ${global_flat_id}() returning the global id of the current thread with all dimensions flattened.

global_flat_size: Module

Provides the function VSIZE_T ${global_flat_size}(). returning the global size of with all dimensions flattened.

global_id: Module

Provides the function VSIZE_T ${global_id}(int dim) returning the global id of the current thread.

global_size: Module

Provides the function VSIZE_T ${global_size}(int dim) returning the global size along dimension dim.

group_id: Module

Provides the function VSIZE_T ${group_id}(int dim) returning the group id of the current thread.

local_id: Module

Provides the function VSIZE_T ${local_id}(int dim) returning the local id of the current thread.

local_size: Module

Provides the function VSIZE_T ${local_size}(int dim) returning the size of the current group.

num_groups: Module

Provides the function VSIZE_T ${num_groups}(int dim) returning the number of groups in dimension dim.

skip: Module

Provides the function bool ${skip}() that should be used at the start of a static kernel function to see if the current thread/work item is inside the padding area and needs to be skipped. Usually one would write if (${skip}()) return;.

Utilities

class grunnur.Template(mako_template: Template)

A wrapper for mako Template objects.

classmethod from_associated_file(filename: str) → Template

Returns a Template object created from the file which has the same name as filename and the extension .mako. Typically used in computation modules as Template.from_associated_file(__file__).

classmethod from_string(template_source: str) → Template

Returns a Template object created from source.

get_def(name: str) → DefTemplate

Returns the template def with the name name.

class grunnur.DefTemplate(name: str, mako_def_template: DefTemplate, source: str)

A wrapper for Mako DefTemplate objects.

classmethod from_callable(name: str, callable_obj: Callable[..., str]) → DefTemplate

Creates a template def from a callable returning a string. The parameter list of the callable is used to create the pararameter list of the resulting template def; the callable should return the body of a Mako template def regardless of the arguments it receives.

classmethod from_string(name: str, argnames: Iterable[str], source: str) → DefTemplate

Creates a template def from a string with its body and a list of argument names.

render(*args: Any, **globals_: Any) → str

Renders the template def with given arguments and globals.

class grunnur.RenderError(exception: Exception, args: Sequence[Any], globals_: Mapping[str, Any], source: str, mako_traceback: str)

A custom wrapper for Mako template render errors, to facilitate debugging.

exception: Exception

The original exception thrown by Mako’s render().

globals: dict[str, Any]

The globals used to render the template.

source: str

The source of the template.

class grunnur.Snippet(template: DefTemplate, render_globals: Mapping[str, Any] = {})

Contains a source snippet - a template function that will be rendered in place, with possible context that can include other Snippet or Module objects.

Creates a snippet out of a prepared template.

classmethod from_callable(callable_obj: Callable[..., str], name: str = '_snippet', render_globals: Mapping[str, Any] = {}) → Snippet

Creates a snippet from a callable returning a string. The parameter list of the callable is used to create the pararameter list of the resulting template def; the callable should return the body of a Mako template def regardless of the arguments it receives.

Parameters:

• callable_obj – a callable returning the template source.

• name – the snippet’s name (will simplify debugging)

• render_globals – a dictionary of “globals” to be used when rendering the template.

classmethod from_string(argnames: Iterable[str], source: str, name: str = '_snippet', render_globals: Mapping[str, Any] = {}) → Snippet

Creates a snippet from a template source and a list of arguments.

Parameters:

• argnames – names of the arguments for the created template.

• source – a string with the template source.

• name – the snippet’s name (will simplify debugging)

• render_globals – a dictionary of “globals” to be used when rendering the template.

class grunnur.Module(template: DefTemplate, render_globals: Mapping[str, Any] = {})

Contains a source module - a template function that will be rendered at root level, and the place where it was called will receive its unique identifier (prefix), which is used to prefix all module’s functions, types and macros in the global namespace.

Creates a module out of a prepared template.

Parameters:

• template

• render_globals

classmethod from_callable(callable_obj: Callable[..., str], name: str = '_module', render_globals: Mapping[str, Any] = {}) → Module

Creates a module from a callable returning a string. The parameter list of the callable is used to create the pararameter list of the resulting template def; the callable should return the body of a Mako template def regardless of the arguments it receives.

The prefix will be passed as the first argument to the template def on render.

Parameters:

• callable_obj – a callable returning the template source.

• name – the module’s name (will simplify debugging)

• render_globals – a dictionary of “globals” to be used when rendering the template.

classmethod from_string(source: str, name: str = '_module', render_globals: Mapping[str, Any] = {}) → Module

Creates a module from a template source, treated as a body of a template def with a single argument (prefix).

Parameters:

• source – a string with the template source.

• name – the module’s name (will simplify debugging)

• render_globals – a dictionary of “globals” to be used when rendering the template.

Data type utilities

class numpy.typing.DTypeLike

intersphinx fails to pick this up. See numpy.typing.DTypeLike for the actual documentation.

class numpy.typing.NDArray

intersphinx fails to pick this up. See numpy.typing.NDArray for the actual documentation.

C interop

grunnur.dtypes.ctype(dtype: DTypeLike) → str | Module

Returns an object that can be passed as a global to Program() and used to render a C equivalent of the given numpy dtype. If there is a built-in C equivalent, the object is just a string with the type name; otherwise it is a Module object containing the corresponding struct declaration.

The given dtype cannot be an array type (that is, its shape must be ()).

Modules are cached, and the function returns a single module instance for equal dtype’s. Therefore inside a kernel it will be rendered with the same prefix everywhere it is used. This results in a behavior characteristic for a structural type system, same as for the basic dtype-ctype conversion.

Note

If dtype is a struct type, it needs to be aligned (manually or with align()) in order for the field offsets on the host and the devices coincide.

Warning

numpy’s isalignedstruct attribute, or aligned parameter is neither necessary nor sufficient for the memory layout on the host and the device to coincide. Use align() if you are not sure how to align your struct type.

grunnur.dtypes.complex_ctr(dtype: DTypeLike) → str

Returns name of the constructor for the given dtype.

grunnur.dtypes.c_constant(val: complex | numpy.generic | NDArray[Any], dtype: DTypeLike | None = None) → str

Returns a C-style numerical constant. If val has a struct dtype, the generated constant will have the form { ... } and can be used as an initializer for a variable.

grunnur.dtypes.align(dtype: DTypeLike) → dtype[Any]

Returns a new struct dtype with the field offsets changed to the ones a compiler would use (without being given any explicit alignment qualifiers). Ignores all existing explicit itemsizes and offsets.

Struct helpers

grunnur.dtypes.flatten_dtype(dtype: DTypeLike) → list[FieldInfo]

Returns a list of tuples (path, dtype) for each of the basic dtypes in a (possibly nested) dtype. path is a list of field names/array indices leading to the corresponding element.

grunnur.dtypes.extract_field(arr: NDArray[Any], path: list[str | int]) → numpy.generic | NDArray[Any]

Extracts an element from an array of struct dtype. The path is the sequence of field names/array indices returned from flatten_dtype().

class grunnur.dtypes.FieldInfo(path: list[str | int], dtype: numpy.dtype[Any], offset: int)

Metadata of a structure field.

Create new instance of FieldInfo(path, dtype, offset)

Data type checks and conversions

grunnur.dtypes.is_complex(dtype: DTypeLike) → bool

Returns True if dtype is complex.

grunnur.dtypes.is_double(dtype: DTypeLike) → bool

Returns True if dtype is double precision floating point.

grunnur.dtypes.is_integer(dtype: DTypeLike) → bool

Returns True if dtype is an integer.

grunnur.dtypes.is_real(dtype: DTypeLike) → bool

Returns True if dtype is a real number (but not complex).

grunnur.dtypes.result_type(*dtypes: DTypeLike) → dtype[Any]

Wrapper for numpy.result_type() which takes into account types supported by GPUs.

grunnur.dtypes.min_scalar_type(val: complex | number[Any]) → dtype[Any]

Wrapper for numpy.min_scalar_type() which takes into account types supported by GPUs.

grunnur.dtypes.complex_for(dtype: DTypeLike) → dtype[Any]

Returns complex dtype corresponding to given floating point dtype.

grunnur.dtypes.real_for(dtype: DTypeLike) → dtype[Any]

Returns floating point dtype corresponding to given complex dtype.

Function modules

Module factories which are used to compensate for the lack of complex number operations in OpenCL, and the lack of C++ synthax which would allow one to write them.

grunnur.functions.add(*in_dtypes: DTypeLike, out_dtype: DTypeLike | None = None) → Module

Returns a Module with a function of len(in_dtypes) arguments that adds values of types in_dtypes. If out_dtype is given, it will be set as a return type for this function.

This is necessary since on some platforms complex numbers are based on 2-vectors, and therefore the + operator for a complex and a real number works in an unexpected way (returning (a.x + b, a.y + b) instead of (a.x + b, a.y)).

grunnur.functions.cast(in_dtype: DTypeLike, out_dtype: DTypeLike) → Module

Returns a Module with a function of one argument that casts values of in_dtype to out_dtype.

grunnur.functions.conj(dtype: DTypeLike) → Module

Returns a Module with a function of one argument that conjugates the value of type dtype (if it is not a complex data type, the value will not be modified).

grunnur.functions.div(dividend_dtype: DTypeLike, divisor_dtype: DTypeLike, out_dtype: DTypeLike | None = None) → Module

Returns a Module with a function of two arguments that divides a value of type dividend_dtype by a value of type divisor_dtype. If out_dtype is given, it will be set as a return type for this function.

grunnur.functions.exp(dtype: DTypeLike) → Module

Returns a Module with a function of one argument that exponentiates the value of type dtype (must be a real or a complex data type).

grunnur.functions.mul(*in_dtypes: DTypeLike, out_dtype: DTypeLike | None = None) → Module

Returns a Module with a function of len(in_dtypes) arguments that multiplies values of types in_dtypes. If out_dtype is given, it will be set as a return type for this function.

grunnur.functions.norm(dtype: DTypeLike) → Module

Returns a Module with a function of one argument that returns the 2-norm of the value of type dtype (product by the complex conjugate if the value is complex, square otherwise).

grunnur.functions.polar(dtype: DTypeLike) → Module

Returns a Module with a function of two arguments that returns the complex-valued rho * exp(i * theta) for values rho, theta of type dtype (must be a real data type).

grunnur.functions.polar_unit(dtype: DTypeLike) → Module

Returns a Module with a function of one argument that returns a complex number exp(i * theta) == (cos(theta), sin(theta)) for a value theta of type dtype (must be a real data type).

grunnur.functions.pow(base_dtype: DTypeLike, exponent_dtype: DTypeLike | None = None, out_dtype: DTypeLike | None = None) → Module

Returns a Module with a function of two arguments that raises the first argument of type base_dtype to the power of the second argument of type exponent_dtype (an integer or real data type).

If exponent_dtype or out_dtype are not given, they default to base_dtype. If base_dtype is not the same as out_dtype, the input is cast to out_dtype before exponentiation. If exponent_dtype is real, but both base_dtype and out_dtype are integer, a ValueError is raised.

Virtual buffers

Often one needs temporary buffers that are only used in one place in the code, but used many times. Allocating them each time they are used may involve too much overhead; allocating real buffers and storing them increases the program’s memory requirements. A possible middle ground is using virtual allocations, where several of them can use the samy physical allocation. The virtual allocation manager will make sure that two virtual buffers that are used simultaneously (as declared by the user) will not share the same physical space.

class grunnur.VirtualManager(device: BoundDevice)

Base class for a manager of virtual allocations.

Parameters:

context – an instance of Context.

allocator(dependencies: Any | None = None) → VirtualAllocator

Create a callable to use for Array creation.

Parameters:

dependencies – can be a Array instance (the ones containing persistent allocations will be ignored), an iterable with valid values, or an object with the attribute __virtual_allocations__ which is a valid value (the last two will be processed recursively).

pack(queue: Queue) → None

Packs the real allocations possibly reducing total memory usage. This process can be slow and may synchronize the base queue.

statistics() → VirtualAllocationStatistics

Returns allocation statistics.

class grunnur.TrivialManager(device: BoundDevice)

Trivial manager — allocates a separate buffer for each allocation request.

class grunnur.ZeroOffsetManager(device: BoundDevice)

Tries to assign several allocation requests to a single real allocation, if dependencies allow that. All virtual allocations start from the beginning of real allocations.

class grunnur.VirtualAllocator(manager: VirtualManager, dependencies: set[int])

A helper callable object to use as an allocator for Array creation. Encapsulates the dependencies (as identifiers, doesn’t hold references for actual objects).

class grunnur.VirtualAllocationStatistics

Virtual allocation details.

real_num: int

The number of physical allocations.

real_size_total: int

The total size of physical allocations (in bytes).

real_sizes: dict[int, int]

A dictionary size: count with the counts for physical allocations of each size.

virtual_num: int

The number of virtual allocations.

virtual_size_total: int

The total size of virtual allocations (in bytes).

virtual_sizes: dict[int, int]

A dictionary size: count with the counts for virtual allocations of each size.

Kernel toolbox

There is a set of macros attached to any kernel depending on the API it is being compiled for:

GRUNNUR_CUDA_API

If defined, specifies that the kernel is being compiled for CUDA API.

GRUNNUR_OPENCL_API

If defined, specifies that the kernel is being compiled for CUDA API.

GRUNNUR_FAST_MATH

If defined, specifies that the compilation for this kernel was requested with fast_math == True.

LOCAL_BARRIER

Synchronizes threads inside a block.

FUNCTION

Modifier for a device-only function declaration.

KERNEL

Modifier for a kernel function declaration.

GLOBAL_MEM

Modifier for a global memory pointer argument.

LOCAL_MEM_DECL

Modifier for a statically allocated local memory variable.

LOCAL_MEM_DYNAMIC

Modifier for a dynamically allocated local memory variable (CUDA only).

LOCAL_MEM

Modifier for a local memory argument in device-only functions.

CONSTANT_MEM_DECL

Modifier for a statically allocated constant memory variable.

CONSTANT_MEM

Modifier for a constant memory argument in device-only functions.

INLINE

Modifier for inline functions.

SIZE_T

The type of local/global IDs and sizes. Equal to unsigned int for CUDA, and size_t for OpenCL (which can be 32- or 64-bit unsigned integer, depending on the device).

SIZE_T get_local_id(unsigned int dim)

SIZE_T get_group_id(unsigned int dim)

SIZE_T get_global_id(unsigned int dim)

SIZE_T get_local_size(unsigned int dim)

SIZE_T get_num_groups(unsigned int dim)

SIZE_T get_global_size(unsigned int dim)

Local, group and global identifiers and sizes. In case of CUDA mimic the behavior of corresponding OpenCL functions.

VSIZE_T

The type of local/global IDs in the virtual grid. It is separate from SIZE_T because the former is intended to be equivalent to what the backend is using, while VSIZE_T is a separate type and can be made larger than SIZE_T in the future if necessary.

ALIGN(int)

Used to specify an explicit alignment (in bytes) for fields in structures, as

typedef struct {
    char ALIGN(4) a;
    int b;
} MY_STRUCT;