Welcome!¶

Bohrium provides automatic acceleration of array operations in Python/NumPy, C, and C++ targeting multi-core CPUs and GP-GPUs. Forget handcrafting CUDA/OpenCL to utilize your GPU and forget threading, mutexes and locks to utilize your multi-core CPU, just use Bohrium!

Features¶

	Architecture Support		Frontends
	Multi-Core CPU	Many-Core GPU	Python2/NumPy	Python3/NumPy	C	C++
Linux	✓	✓	✓	✓	✓	✓
Mac OS	✓	✓	✓		✓	✓

Lazy Evaluation, Bohrium will lazy evaluate all Python/NumPy operations until it encounters a “Python Read” such a printing an array or having a if-statement testing the value of an array.
Views Bohrium supports NumPy views fully thus operating on array slices does not involve data copying.
Loop Fusion, Bohrium uses a fusion algorithm that fuses (or merges) array operations into the same computation kernel that are then JIT-compiled and executed. However, Bohrium can only fuse operations that have some common sized dimension and no horizontal data conflicts.
Lazy CPU/GPU Communication, Bohrium only moves data between the host and the GPU when the data is accessed directly by Python or a Python C-extension.
python -m bohrium, automatically makes import numpy use Bohrium.
Jupyter Support, you can use the magic command %%bohrium to automatically use Bohrium as NumPy.
Zero-copy Interoperability with:
- NumPy
- Cython
- PyOpenCL
- PyCUDA

Please note:

Bohrium is a 64-bit project exclusively.
Source code is available here: https://github.com/bh107/bohrium

Get Started!¶

Installation¶

Bohrium supports Linux and Mac OS.

Linux¶

PyPI Package¶

If you use Bohrium through Python, we strongly recommend to install Bohrium through pypi, which will include BLAS, LAPACK, OpenCV, and OpenCL support:

pip install --user bohrium

Note

On linux, Bohrium requires gcc in $PATH. E.g. on Ubuntu install the build-essential package: sudo apt install build-essential.

Note

On linux, Python development files must be available. E.g. on Ubuntu install python-dev and/or python3-dev.

Anaconda¶

To use Anaconda, simply install the Bohrium PyPI package in an environment:

# Activate the environment where you want to install Bohrium:
source activate my_env
# Install Bohrium using pip
pip install bohrium

Note

Bohrium requires gcc in $PATH. E.g. on Ubuntu install the build-essential package: sudo apt install build-essential.

Install From Source Package¶

Visit Bohrium on github.com and download the latest release: https://github.com/bh107/bohrium/releases/latest. Then build and install Bohrium as described in the following subsections.

Install dependencies, which on Ubuntu is:

sudo apt install build-essential python-pip python-virtualenv cmake git unzip libboost-filesystem-dev libboost-serialization-dev libboost-regex-dev zlib1g-dev libsigsegv-dev

And some additional packages for visualization:

sudo apt-get install freeglut3 freeglut3-dev libxmu-dev libxi-dev

Build and install:

wget https://github.com/bh107/bohrium/archive/master.zip
unzip master.zip
cd bohrium-master
mkdir build
cd build
cmake .. -DCMAKE_INSTALL_PREFIX=<path to install directory>
make
make install

Note

The default install directory is ~/.local

Note

To compile to a custom Python (with valgrind debug support for example), set -DPYTHON_EXECUTABLE=<custom python binary>.

Finally, you need to set the LD_LIBRARY_PATH environment variables and if you didn’t install Bohrium in $HOME/.local/lib your need to set PYTHONPATH as well.

The LD_LIBRARY_PATH should include the path to the installation directory:

export LD_LIBRARY_PATH="<install dir>:$LD_LIBRARY_PATH"

The PYTHONPATH should include the path to the newly installed Bohrium Python module:

export PYTHONPATH="<install dir>/lib/python<python version>/site-packages:$PYTHONPATH"

Check Your Installation¶

Check installation by printing the current runtime stack:

python -m bohrium --info

Mac OS¶

The following explains how to get going on Mac OS.

You need to install the Xcode Developer Tools package, which is found in the App Store.

Note

You might have to manually install some extra header files by running `sudo installer -pkg /Library/Developer/CommandLineTools/Package/macOS_SDK_headers_for_macOS_10.14.pkg -target /` where `10.14` is your current version (more info).

PyPI Package¶

If you use Bohrium through Python, we strongly recommend to install Bohrium through pypi, which will include BLAS, LAPACK, OpenCV, and OpenCL support:

python -m pip install --user bohrium

Note

If you get an error message saying that no package match your criteria it is properly because you are using a Python version for which `no package exist https://pypi.org/project/bohrium-api/#files`_ . Please contact us and we will build a package using your specific Python version.

Install From Source Package¶

Start by installing Homebrew as explained on their website

/usr/bin/ruby -e "$(curl -fsSL https://raw.githubusercontent.com/Homebrew/install/master/install)"

Install dependencies:

brew install python
brew install cmake
brew install boost --with-icu4c
brew install libsigsegv
python3 -m pip install --user numpy cython twine gcc7

Visit Bohrium on github.com, download the latest release: https://github.com/bh107/bohrium/releases/latest or download master, and then build it:

wget https://github.com/bh107/bohrium/archive/master.zip
unzip master.zip
cd bohrium-master
mkdir build
cd build
export PATH="$(brew --prefix)/bin:/usr/local/opt/llvm/bin:/usr/local/opt/opencv3/bin:$PATH"
export CC="clang"
export CXX="clang++"
export C_INCLUDE_PATH=$(llvm-config --includedir)
export CPLUS_INCLUDE_PATH=$(llvm-config --includedir)
export LIBRARY_PATH=$(llvm-config --libdir):$LIBRARY_PATH
cmake .. -DCMAKE_INSTALL_PREFIX=<path to install directory>
make
make install

Note

The default install directory is ~/.local

Note

To compile to a custom Python (with valgrind debug support for example), set -DPYTHON_EXECUTABLE=<custom python binary>.

Finally, you need to set the DYLD_LIBRARY_PATH and LIBRARY_PATH environment variables and if you didn’t install Bohrium in $HOME/.local/lib your need to set PYTHONPATH as well.

The DYLD_LIBRARY_PATH and LIBRARY_PATH should include the path to the installation directory:

export DYLD_LIBRARY_PATH="<install dir>:$DYLD_LIBRARY_PATH"
export LIBRARY_PATH="<install dir>:$LIBRARY_PATH"

The PYTHONPATH should include the path to the newly installed Bohrium Python module:

export PYTHONPATH="<install dir>/lib/python<python version>/site-packages:$PYTHONPATH"

Check Your Installation¶

Check installation by printing the current runtime stack:

python -m bohrium --info

Installation using Spack¶

This guide will install Bohrium using the Spack package manager.

Why use Spack?¶

Spack is a package management tool tailored specifically for supercomputers with a rather dated software stack. It allows to install and maintain packages, starting only from very few dependencies: Pretty much just python2.6, git, curl and some c++ compiler are all that’s needed for the bootstrap.

Needless to say that the request for installing a particular package automatically yields the installation of all dependencies with exactly the right version and configurations. If this causes multiple versions/configurations of the same package to be required, this is no problem and gets resolved automatically, too. As a bonus on top, using an installed package later is super easy as well due to an automatic generation of module files, which set the required environment up.

Installation overview¶

First step is to clone and setup Spack:

export SPACK_ROOT="$PWD/spack"
git clone https://github.com/llnl/spack.git
. $SPACK_ROOT/share/spack/setup-env.sh

Afterwards the installation of Bohrium is instructed:

spack install bohrium

This step will take a while, since Spack will download the sources of all dependencies, unpack, configure and compile them. But since everything happens in the right order automatically, you could easily do this over night.

That’s it. If you want to use Bohrium, setup up Spack as above, then load the required modules:

spack module loads -r bohrium > /tmp/bohrium.modules
. /tmp/bohrium.modules

and you are ready to go as the shell environment now contains all required variables (LD_LIBRARY_PATH, PATH, CPATH, PYTHONPATH, …) to get going.

If you get some errors about the command module not being found, you need to install the Spack package environment-modules beforehand. Again, just a plain:

spack install environment-modules

is enough to achieve this.

Tuning the installation procedure¶

Spack offers countless ways to influence how things are installed and what is installed. See the Documentation and especially the Getting Started section for a good overview.

Most importantly the so-called spec allows to specify features or requirements with respect to versions and dependencies, that should be enabled or disabled when building the package. For example:

spec install bohrium~cuda~opencl

Will install Bohrium without CUDA or OpenCL support, which has a dramatic impact on the install time due to the reduced amount of dependencies to be installed. On the other hand:

spec install bohrium@develop

will install specifically the development version of Bohrium. This the current HEAD of the master branch in the github repository. One may also influence the versions of the dependencies by themselves. For example:

spec install bohrium+python^python@3:

will specifically compile Bohrium with a python version larger than 3.

The current list of features the Bohrium package has to offer can be listed by the command:

spack info bohrium

and the list of dependencies which will be installed by a particlar spec can be easily reviewed by something like:

spack spec bohrium@develop~cuda~opencl

User Guide¶

Python/NumPy¶

Runtime Info
Automatic Parallelization
Acceleration
Convert between Bohrium and NumPy
Accelerate Loops
Sliding Views Between Iterations
UserKernel
- OpenMP Example
- OpenCL Example
Interoperability

Runtime Info ¶

Print the current Bohrium runtime stack:

python -m bohrium --info

Automatic Parallelization ¶

Bohrium implements a new python module bohrium that introduces a new array class bohrium.ndarray which inherits from numpy.ndarray. The two array classes are fully compatible thus one only has to replace numpy.ndarray with bohrium.ndarray in order to utilize the Bohrium runtime system.

The following example is a heat-equation solver that uses Bohrium. Note that the only difference between Bohrium code and NumPy code is the first line where we import bohrium as np instead of numpy as np:

import bohrium as np
def heat2d(height, width, epsilon=42):
  G = np.zeros((height+2,width+2),dtype=np.float64)
  G[:,0]  = -273.15
  G[:,-1] = -273.15
  G[-1,:] = -273.15
  G[0,:]  = 40.0
  center = G[1:-1,1:-1]
  north  = G[:-2,1:-1]
  south  = G[2:,1:-1]
  east   = G[1:-1,:-2]
  west   = G[1:-1,2:]
  delta  = epsilon+1
  while delta > epsilon:
    tmp = 0.2*(center+north+south+east+west)
    delta = np.sum(np.abs(tmp-center))
    center[:] = tmp
  return center
heat2d(100, 100)

Alternatively, you can import Bohrium as NumPy through the command line argument -m bohrium:

python -m bohrium heat2d.py

In this case, all instances of import numpy is converted to import bohrium seamlessly. If you need to access the real numpy module use import numpy_force.

Acceleration ¶

The approach of Bohrium is to accelerate all element-wise functions in NumPy (aka universal functions) as well as the reductions and accumulations of element-wise functions. This approach makes it possible to accelerate the heat-equation solver on both multi-core CPUs and GPUs.

Beside element-wise functions, Bohrium also accelerates a selection of common NumPy functions such as dot() and solve(). But the number of functions in NumPy and related projects such as SciPy is enormous thus we cannot hope to accelerate every single function in Bohrium. Instead, Bohrium will automatically convert bohrium.ndarray to numpy.ndarray when encountering a function that Bohrium cannot accelerate. When running on the CPU, this conversion is very cheap but when running on the GPU, this conversion requires the array data to be copied from the GPU to the CPU.

Matplotlib’s matshow() function is example of a function Bohrium cannot accelerate. Say we want to visualize the result of the heat-equation solver, we could use matshow():

from matplotlib import pyplot as plt

res = heat2d(100, 100)
plt.matshow(res, cmap='hot')
plt.show()

Beside producing the image (after approx. 1 min), the execution will raise a Python warning informing you that matplotlib function is handled like a regular NumPy:

/usr/lib/python2.7/site-packages/matplotlib/cbook.py:1506: RuntimeWarning:
Encountering an operation not supported by Bohrium. It will be handled by the original NumPy.
x = np.array(x, subok=True, copy=copy)

Note

Increasing the problem size will improve the performance of Bohrium significantly!

Convert between Bohrium and NumPy ¶

It is possible to convert between Bohrium and NumPy explicitly and thus avoid Python warnings. Let’s walk through an example:

Create a new NumPy array with ones:

np_ary = numpy.ones(42)

Convert any type of array to Bohrium:

bh_ary = bohrium.array(np_ary)

Copy a bohrium array into a new NumPy array:

npy2 = bh_ary.copy2numpy()

Accelerate Loops ¶

As we all know, having for and while loops in Python is bad for performance but is sometimes necessary. E.g. in the case of the heat2d() code, we have to evaluate delta > epsilon in order to know when to stop iterating. To address this issue, Bohrium introduces the function do_while(), which takes a function and calls it repeatedly until either a maximum number of calls has been reached or until the function return False.

The function signature:

def do_while(func, niters, *args, **kwargs):
    """Repeatedly calls the `func` with the `*args` and `**kwargs` as argument.

    The `func` is called while `func` returns True or None and the maximum number
    of iterations, `niters`, hasn't been reached.

    Parameters
    ----------
    func : function
        The function to run in each iterations. `func` can take any argument and may return
        a boolean `bharray` with one element.
    niters: int or None
        Maximum number of iterations in the loop (number of times `func` is called). If None, there is no maximum.
    *args, **kwargs : list and dict
        The arguments to `func`

    Notes
    -----
    `func` can only use operations supported natively in Bohrium.
    """

An example where the function doesn’t return anything:

>>> def loop_body(a):
...     a += 1
>>> a = bh.zeros(4)
>>> bh.do_while(loop_body, 5, a)
>>> a
array([5, 5, 5, 5])

An example where the function returns a bharray with one element and of type bh.bool:

>>> def loop_body(a):
...     a += 1
...     return bh.sum(a) < 10
>>> a = bh.zeros(4)
>>> bh.do_while(loop_body, None, a)
>>> a
array([3, 3, 3, 3])

Sliding Views Between Iterations ¶

It can be useful to increase/decrease the beginning of certain array views between iterations of a loop. This can be achieved using get_iterator(), which returns a special bohrium iterator. The iterator can be given an optional start value (0 by default). The iterator is increased by one for each iteration, but can be changed increase or decrease by multiplying any constant (see example 2).

Iterators only supports addition, subtraction and multiplication. get_iterator() can only be used within Bohrium loops. Views using iterators cannot change shape between iterations. Therefore, views such as a[i:2*i] are not supported.

Example 1. Using iterators to create a loop-based function for calculating the triangular numbers (from 1 to 10). The loop in numpy looks the following:

>>> a = np.arange(1,11)
>>> for i in range(0,9):
...     a[i+1] += a[i]
>>> a
array([1 3 6 10 15 21 28 36 45 55])

The same can be written in Bohrium as:

>>> def loop_body(a):
...    i = get_iterator()
...    a[i+1] += a[i]
>>> a = bh.arange(1,11)
>>> bh.do_while(loop_body, 9, a)
>>> a
array([1 3 6 10 15 21 28 36 45 55])

Example 2. Increasing every second element by one, starting at both ends, in the same loop. As it can be seen: i is increased by 2, while j is descreased by 2 for each iteration:

>>> def loop_body(a):
...   i = get_iterator(1)
...   a[2*i] += a[2*(i-1)]
...   j = i+1
...   a[1-2*j] += a[1-2*(j-1)]
>>> a = bh.ones(10)
>>> bh.for_loop(loop_body, 4, a)
>>> a
array([1 5 2 4 3 3 4 2 5 1])

Nested loops is also available in do_while by using grids. A grid is a set of iterators that depend on each other, just as with nested loops. A grid can have arbitrary size and is available via. the function get_grid(), which is only usable within a do_while loop body. The function takes an amount of integers as parameters, corresponding to the range of the loops (from outer to inner). It returns the same amount of iterators, which functions as a grid. An example of this can be seen in Example 3 below. Example 3. Creating a range in an array with multiple dimensions. In Numpy it can be written as:

>>> a = bh.zeros((3,3))
>>> counter = bh.zeros(1)
>>> for i in range(3):
...    for j in range(3):
...        counter += 1
...        a[i,j] += counter
>>> a
[[1. 2. 3.]
 [4. 5. 6.]
 [7. 8. 9.]]

The same can done within a do_while loop by using a grid:

>>> def kernel(a, counter):
...    i, j = get_grid(3,3)
...    counter += 1
...    a[i,j] += counter
>>> a = bh.zeros((3,3))
>>> counter = bh.zeros(1)
>>> bh.do_while(kernel, 3*3, a, counter)
>>> a
[[1. 2. 3.]
 [4. 5. 6.]
 [7. 8. 9.]]

UserKernel ¶

Bohrium supports user kernel, which makes it possible to implement a specialized handwritten kernel. The idea is that if you encounter a problem that you cannot implement using array programming and Bohrium cannot accelerate, you can write a kernel in C99 that calls other libraries or do the calculation itself.

OpenMP Example ¶

In order to write and run your own kernel use bh.user_kernel.execute():

import bohrium as bh

def fftn(ary):
    # Making sure that `ary` is complex, contiguous, and uses no offset
    ary = bh.user_kernel.make_behaving(ary, dtype=bh.complex128)
    res = bh.empty_like(a)

    # Indicates the direction of the transform you are interested in;
    # technically, it is the sign of the exponent in the transform.
    sign = ["FFTW_FORWARD", "FFTW_BACKWARD"]

    kernel = """
    #include <stdint.h>
    #include <stdlib.h>
    #include <complex.h>
    #include <fftw3.h>

    #if defined(_OPENMP)
        #include <omp.h>
    #else
        static inline int omp_get_max_threads() { return 1; }
        static inline int omp_get_thread_num()  { return 0; }
        static inline int omp_get_num_threads() { return 1; }
    #endif

    void execute(double complex *in, double complex *out) {
        const int ndim = %(ndim)d;
        const int shape[] = {%(shape)s};
        const int sign = %(sign)s;

        fftw_init_threads();
        fftw_plan_with_nthreads(omp_get_max_threads());

        fftw_plan p = fftw_plan_dft(ndim, shape, in, out, sign, FFTW_ESTIMATE);
        if(p == NULL) {
            printf("fftw plan fail!\\n");
            exit(-1);
        }
        fftw_execute(p);
        fftw_destroy_plan(p);
        fftw_cleanup_threads();
    }
    """ % {'ndim': a.ndim, 'shape': str(a.shape)[1:-1], 'sign': sign[0]}

    # Adding some extra link options to the compiler command
    cmd = bh.user_kernel.get_default_compiler_command() + " -lfftw3 -lfftw3_threads"
    bh.user_kernel.execute(kernel, [ary, res], compiler_command=cmd)
    return res

OpenCL Example ¶

In order to use the OpenCL backend, use the tag and param of bh.user_kernel.execute():

import bohrium as bh

kernel = """
#pragma OPENCL EXTENSION cl_khr_fp64 : enable

kernel void execute(global double *a, global double *b) {
    int i0 = get_global_id(0);
    int i1 = get_global_id(1);
    int gid = i0 * 5 + i1;
    b[gid] = a[gid] + gid;
}
"""
a = bh.ones(10*5, bh.double).reshape(10,5)
res = bh.empty_like(a)
# Notice, the OpenCL backend requires global_work_size and local_work_size
bh.user_kernel.execute(kernel, [a, res],
                       tag="opencl",
                       param={"global_work_size": [10, 5], "local_work_size": [1, 1]})
print(res)

Note

Remember to use the OpenCL backend by setting BH_STACK=opencl.

Interoperability ¶

Bohrium is interoperable with other popular Python projects such as Cython and PyOpenCL. The idea is that if you encounter a problem that you cannot implement using array programming and Bohrium cannot accelerate, you can manually accelerate that problem using Cython or PyOpenCL.

NumPy ¶

One example of such a problem is bincount() from NumPy. bincount() computes a histogram of an array, which isn’t possible to implement efficiently through array programming. One approach is simply to use the implementation of NumPy:

import numpy
import bohrium

def bincount_numpy(ary):
    # Make a NumPy copy of the Bohrium array
    np_ary = ary.copy2numpy()
    # Let NumPy handle the calculation
    result = numpy.bincount(np_ary)
    # Copy the result back into a new Bohrium array
    return bohrium.array(result)

In this case, we use bohrium.copy2numpy() and bohrium.array() to copy the Bohrium to NumPy and back again.

Cython ¶

In order to parallelize bincount() for a multi-core CPU, one can use Cython:

import numpy as np
import bohrium
import cython
from cython.parallel import prange, parallel
from libc.stdlib cimport abort, malloc, free
cimport numpy as cnp
cimport openmp
ctypedef cnp.uint64_t uint64

@cython.boundscheck(False) # turn off bounds-checking
@cython.cdivision(True) # turn off division-by-zero checking
cdef _count(uint64[:] x, uint64[:] out):
    cdef int num_threads, thds_id
    cdef uint64 i, start, end
    cdef uint64* local_histo

    with nogil, parallel():
        num_threads = openmp.omp_get_num_threads()
        thds_id = openmp.omp_get_thread_num()
        start = (x.shape[0] / num_threads) * thds_id
        if thds_id == num_threads-1:
            end = x.shape[0]
        else:
            end = start + (x.shape[0] / num_threads)

        if not(thds_id < num_threads-1 and x.shape[0] < num_threads):
            local_histo = <uint64 *> malloc(sizeof(uint64) * out.shape[0])
            if local_histo == NULL:
                abort()
            for i in range(out.shape[0]):
                local_histo[i] = 0

            for i in range(start, end):
                local_histo[x[i]] += 1

            with gil:
                for i in range(out.shape[0]):
                    out[i] += local_histo[i]
            free(local_histo)


def bincount_cython(x, minlength=None):
    # The output `ret` has the size of the max element plus one
    ret = bohrium.zeros(x.max()+1, dtype=x.dtype)

    # To reduce overhead, we use `interop_numpy.get_array()` instead of `copy2numpy()`
    # This approach means that `x_buf` and `ret_buf` points to the same memory as `x` and `ret`.
    # Therefore, only change or deallocate `x` and `ret` when you are finished using `x_buf` and `ret_buf`.
    x_buf = bohrium.interop_numpy.get_array(x)
    ret_buf = bohrium.interop_numpy.get_array(ret))

    # Now, we can run the Cython function
    _count(x_buf, ret_buf))

    # Since `ret_buf` points to the memory of `ret`, we can simply return `ret`.
    return ret

The function _count() is a regular Cython function that performs the histogram calculation. The function bincount_cython() uses bohrium.interop_numpy.get_array() to retrieve data pointers from the Bohrium arrays without any data copying.

PyOpenCL ¶

In order to parallelize bincount() for a GPGPU, one can use PyOpenCL:

import bohrium
import pyopencl as cl

def bincount_pyopencl(x):
    # Check that PyOpenCL is installed and that the Bohrium runtime uses the OpenCL backend
    if not interop_pyopencl.available():
        raise NotImplementedError("OpenCL not available")

    # Get the OpenCL context from Bohrium
    ctx = bohrium.interop_pyopencl.get_context()
    queue = cl.CommandQueue(ctx)

    x_max = int(x.max())

    # Check that the size of histogram doesn't exceeds the memory capacity of the GPU
    if x_max >= interop_pyopencl.max_local_memory(queue.device) // x.itemsize:
        raise NotImplementedError("OpenCL: max element is too large for the GPU")

    # Let's create the output array and retrieve the in-/output OpenCL buffers
    # NB: we always return uint32 array
    ret = bohrium.empty((x_max+1, ), dtype=np.uint32)
    x_buf = bohrium.interop_pyopencl.get_buffer(x)
    ret_buf = bohrium.interop_pyopencl.get_buffer(ret)

    # The OpenCL kernel is based on the book "OpenCL Programming Guide" by Aaftab Munshi at al.
    source = """
    kernel void histogram_partial(
        global DTYPE *input,
        global uint *partial_histo,
        uint input_size
    ){
        int local_size = (int)get_local_size(0);
        int group_indx = get_group_id(0) * HISTO_SIZE;
        int gid = get_global_id(0);
        int tid = get_local_id(0);

        local uint tmp_histogram[HISTO_SIZE];

        int j = HISTO_SIZE;
        int indx = 0;

        // clear the local buffer that will generate the partial histogram
        do {
            if (tid < j)
                tmp_histogram[indx+tid] = 0;
            j -= local_size;
            indx += local_size;
        } while (j > 0);

        barrier(CLK_LOCAL_MEM_FENCE);

        if (gid < input_size) {
            atomic_inc(&tmp_histogram[input[gid]]);
        }

        barrier(CLK_LOCAL_MEM_FENCE);

        // copy the partial histogram to appropriate location in
        // histogram given by group_indx
        if (local_size >= HISTO_SIZE){
            if (tid < HISTO_SIZE)
                partial_histo[group_indx + tid] = tmp_histogram[tid];
        }else{
            j = HISTO_SIZE;
            indx = 0;
            do {
                if (tid < j)
                    partial_histo[group_indx + indx + tid] = tmp_histogram[indx + tid];

                j -= local_size;
                indx += local_size;
            } while (j > 0);
        }
    }

    kernel void histogram_sum_partial_results(
        global uint *partial_histogram,
        int num_groups,
        global uint *histogram
    ){
        int gid = (int)get_global_id(0);
        int group_indx;
        int n = num_groups;
        local uint tmp_histogram[HISTO_SIZE];

        tmp_histogram[gid] = partial_histogram[gid];
        group_indx = HISTO_SIZE;
        while (--n > 0) {
            tmp_histogram[gid] += partial_histogram[group_indx + gid];
            group_indx += HISTO_SIZE;
        }
        histogram[gid] = tmp_histogram[gid];
    }
    """
    source = source.replace("HISTO_SIZE", "%d" % ret.shape[0])
    source = source.replace("DTYPE", interop_pyopencl.type_np2opencl_str(x.dtype))
    prg = cl.Program(ctx, source).build()

    # Calculate sizes for the kernel execution
    local_size = interop_pyopencl.kernel_info(prg.histogram_partial, queue)[0]  # Max work-group size
    num_groups = int(math.ceil(x.shape[0] / float(local_size)))
    global_size = local_size * num_groups

    # First we compute the partial histograms
    partial_res_g = cl.Buffer(ctx, cl.mem_flags.WRITE_ONLY, num_groups * ret.nbytes)
    prg.histogram_partial(queue, (global_size,), (local_size,), x_buf, partial_res_g, np.uint32(x.shape[0]))

    # Then we sum the partial histograms into the final histogram
    prg.histogram_sum_partial_results(queue, ret.shape, None, partial_res_g, np.uint32(num_groups), ret_buf)
    return ret

The implementation is regular PyOpenCL and the OpenCL kernel is based on the book “OpenCL Programming Guide” by Aaftab Munshi et al. However, notice that we use bohrium.interop_pyopencl.get_context() to get the PyOpenCL context rather than pyopencl.create_some_context(). In order to avoid copying data between host and device memory, we use bohrium.interop_pyopencl.get_buffer() to create a OpenCL buffer that points to the device memory of the Bohrium arrays.

PyCUDA ¶

The PyCUDA implementation is very similar to the PyOpenCL. Besides some minor difference in the kernel source code, we use interop_pycuda.init() to initiate PyCUDA and use interop_pycuda.get_gpuarray() to get the CUDA buffers from the Bohrium arrays:

def bincount_pycuda(x, minlength=None):
    """PyCUDA implementation of `bincount()`"""

    if not interop_pycuda.available():
        raise NotImplementedError("CUDA not available")

    import pycuda
    from pycuda.compiler import SourceModule

    interop_pycuda.init()

    x_max = int(x.max())
    if x_max < 0:
        raise RuntimeError("bincount(): first argument must be a 1 dimensional, non-negative int array")
    if x_max > np.iinfo(np.uint32).max:
        raise NotImplementedError("CUDA: the elements in the first argument must fit in a 32bit integer")
    if minlength is not None:
        x_max = max(x_max, minlength)

    # TODO: handle large max element by running multiple bincount() on a range
    if x_max >= interop_pycuda.max_local_memory() // x.itemsize:
        raise NotImplementedError("CUDA: max element is too large for the GPU")

    # Let's create the output array and retrieve the in-/output CUDA buffers
    # NB: we always return uint32 array
    ret = array_create.ones((x_max+1, ), dtype=np.uint32)
    x_buf = interop_pycuda.get_gpuarray(x)
    ret_buf = interop_pycuda.get_gpuarray(ret)

    # CUDA kernel is based on the book "OpenCL Programming Guide" by Aaftab Munshi et al.
    source = """
    __global__ void histogram_partial(
        DTYPE *input,
        uint *partial_histo,
        uint input_size
    ){
        int local_size = blockDim.x;
        int group_indx = blockIdx.x * HISTO_SIZE;
        int gid = (blockIdx.x * blockDim.x + threadIdx.x);
        int tid = threadIdx.x;

        __shared__ uint tmp_histogram[HISTO_SIZE];

        int j = HISTO_SIZE;
        int indx = 0;

        // clear the local buffer that will generate the partial histogram
        do {
            if (tid < j)
                tmp_histogram[indx+tid] = 0;
            j -= local_size;
            indx += local_size;
        } while (j > 0);

        __syncthreads();

        if (gid < input_size) {
            atomicAdd(&tmp_histogram[input[gid]], 1);
        }

        __syncthreads();

        // copy the partial histogram to appropriate location in
        // histogram given by group_indx
        if (local_size >= HISTO_SIZE){
            if (tid < HISTO_SIZE)
                partial_histo[group_indx + tid] = tmp_histogram[tid];
        }else{
            j = HISTO_SIZE;
            indx = 0;
            do {
                if (tid < j)
                    partial_histo[group_indx + indx + tid] = tmp_histogram[indx + tid];

                j -= local_size;
                indx += local_size;
            } while (j > 0);
        }
    }

    __global__ void histogram_sum_partial_results(
        uint *partial_histogram,
        int num_groups,
        uint *histogram
    ){
        int gid = (blockIdx.x * blockDim.x + threadIdx.x);
        int group_indx;
        int n = num_groups;
        __shared__ uint tmp_histogram[HISTO_SIZE];

        tmp_histogram[gid] = partial_histogram[gid];
        group_indx = HISTO_SIZE;
        while (--n > 0) {
            tmp_histogram[gid] += partial_histogram[group_indx + gid];
            group_indx += HISTO_SIZE;
        }
        histogram[gid] = tmp_histogram[gid];
    }
    """
    source = source.replace("HISTO_SIZE", "%d" % ret.shape[0])
    source = source.replace("DTYPE", interop_pycuda.type_np2cuda_str(x.dtype))
    prg = SourceModule(source)

    # Calculate sizes for the kernel execution
    kernel = prg.get_function("histogram_partial")
    local_size = kernel.get_attribute(pycuda.driver.function_attribute.MAX_THREADS_PER_BLOCK)  # Max work-group size
    num_groups = int(math.ceil(x.shape[0] / float(local_size)))
    global_size = local_size * num_groups

    # First we compute the partial histograms
    partial_res_g = pycuda.driver.mem_alloc(num_groups * ret.nbytes)
    kernel(x_buf, partial_res_g, np.uint32(x.shape[0]), block=(local_size, 1, 1), grid=(num_groups, 1))

    # Then we sum the partial histograms into the final histogram
    kernel = prg.get_function("histogram_sum_partial_results")
    kernel(partial_res_g, np.uint32(num_groups), ret_buf, block=(1, 1, 1), grid=(ret.shape[0], 1))
    return ret

Performance Comparison ¶

Finally, let’s compare the performance of the difference approaches. We run on a Intel(R) Core(TM) i5-6600K CPU @ 3.50GHz with 4 CPU-cores and a GeForce GTX Titan X (maxwell). The timing is wall-clock time including everything, in particular the host/device communication overhead.

(Source code, png, hires.png, pdf)

The timing code:

import numpy as np
import time

SIZE = 500000000
ITER = 100

t1 = time.time()
a = np.minimum(np.arange(SIZE, dtype=np.int64), 64)
for _ in range(ITER):
    b = np.bincount(a)
t2 = time.time()
s = b.sum()
print ("Sum: %d, time: %f sec" % (s, t2 - t1))

Conclusion ¶

Interoperability makes it possible to accelerate code that Bohrium doesn’t accelerate automatically. The Bohrium team constantly works on improving the performance and increase the number of NumPy operations automatically accelerated but in some cases we simply have to give the user full control.

C++ library¶

The C++ interface of Bohrium is similar to NumPy but is still very basic.

Indexing / Slicing¶

Bohrium C++ only support single index indexing:

// Create a new empty array (4 by 5)
bhxx::BhArray<double> A = bhxx::empty<double>({4, 5});
// Create view of the third row of A
bhxx::BhArray<double> B = A[2];

If you need more flexible slicing, you can set the shape and stride manually:

// Create a new array (4 by 5) of ones
bhxx::BhArray<double> A = bhxx::ones<double>({4, 5});
// Create view of the complete A.
bhxx::BhArray<double> B = A;
// B is now a 2 by 5 view with a step of two in the first dimension.
// In NumPy, this corresponds to: `B = A[::2, :]`
B.setShapeAndStride({2, 5}, {10, 1});

Code Snippets¶

You can find some examples in the source tree and some code snippets here:

#include<bhxx/bhxx.hpp>

/** Return a new empty array */
bhxx::BhArray<double> A = bhxx::empty<double>({4, 5});

/** Return the rank (number of dimensions) of the array */
int rank = A.rank();

/** Return the offset of the array */
uint64_t offset = A.offset();

/** Return the shape of the array */
Shape shape = A.shape();

/** Return the stride of the array */
Stride stride = A.stride();

/** Return the total number of elements of the array */
uint64_t size = A.size();

/** Return a pointer to the base of the array */
std::shared_ptr<BhBase> base = A.base();

/** Return whether the view is contiguous and row-major */
bool is_contig = A.isContiguous();

/** Return a new copy of the array */
bhxx::BhArray<double> copy = A.copy();

/** Return a copy of the array as a standard vector */
std::vector<double> vec = A.vec();

/** Print the content of A */
std::cout << A << "\n";

// Return a new transposed view
bhxx::BhArray<double> A_T = A.transpose();

// Return a new reshaped view (the array must be contiguous)
bhxx::BhArray<double> A_reshaped = A.reshape(Shape shape);

/** Return a new view with a "new axis" inserted.
 *
 *  The "new axis" is inserted just before `axis`.
 *  If negative, the count is backwards
 */
bhxx::BhArray<double> A_new_axis = A.newAxis(1);

// Return a new empty array
auto A = bhxx::empty<float>({3,4});

// Return a new empty array that has the same shape as `ary`
auto B = bhxx::empty_like<float>(A);

// Return a new array filled with zeros
auto A = bhxx::zeros<float>({3,4});

// Return evenly spaced values within a given interval.
auto A = bhxx::arange(1, 3, 2); // start, stop, step
auto A = bhxx::arange(1, 3); // start, stop, step=1
auto A = bhxx::arange(3); // start=0, stop, step=1

// Random array, interval [0.0, 1.0)
auto A = bhxx::random.randn<double>({3, 4});

// Element-wise `static_cast`.
bhxx::BhArray<int> B = bhxx::cast<int>(A);

// Alias, A and B points to the same underlying data.
bhxx::empty<float> A = bhxx::empty<float>({3,4});
bhxx::empty<float> B = A;

// a is an alias
void add_inplace(bhxx::BhArray<double> a,
                 bhxx::BhArray<double> b) {
    a += b;
}
add_inplace(A, B);

// Create the data of A into a new array B.
bhxx::empty<float> A = bhxx::empty<float>({3,4});
bhxx::empty<float> B = A.copy();

// Copy the data of B into the existing array A.
A = B;

// Copying and converting the data of A into C.
bhxx::empty<double> C = bhxx::cast<double>(A);

// Alias, A and B points to the same underlying data.
bhxx::empty<float> A = bhxx::empty<float>({3,4});
bhxx::empty<float> B = bhxx::empty<float>({4});
B.reset(A);

// Evaluation triggers:
bhxx::flush();
std::cout << A << "\n";
A.vec();
A.data();

// Operator overloads
A + B - C * E / G;

// Standard functions
bhxx::sin(A) + bhxx::cos(B) + bhxx::sqrt(C) + ...

// Reductions (sum, product, maximum, etc.)
bhxx::add_reduce(A, 0); // Sum of axis 0
bhxx::multiply_reduce(B, 1); // Product of axis 1
bhxx::maximum_reduce(C, 2); // Maximum of axis 2

The API¶

The following is the complete API as defined in the header file:

template <typename T> class BhArray : public bhxx::BhArrayUnTypedCore ¶

#include <BhArray.hpp>

Representation of a multidimensional array that point to a BhBase array.

Template Parameters

T: The data type of the array and the underlying base array

Public Types

typedef T scalar_type¶: The data type of each array element.

Public Functions

BhArray()¶: Default constructor that leave the instance completely uninitialized.

BhArray(Shape shape, Stride stride)¶

Create a new array. Shape and Stride must have the same length.

Parameters

shape: Shape of the new array
stride: Stride of the new array

BhArray(Shape shape)¶: Create a new array (contiguous stride, row-major)

BhArray(std::shared_ptr<BhBase> base, Shape shape, Stride stride, uint64_t offset = 0)¶

Create a array that points to the given base

Note: The caller should make sure that the shared pointer uses the RuntimeDeleter as its deleter, since this is implicitly assumed throughout, i.e. if one wants to construct a BhBase object, use the make_base_ptr helper function.

BhArray(std::shared_ptr<BhBase> base, Shape shape)¶

Create a view that points to the given base (contiguous stride, row-major)

Note: The caller should make sure that the shared pointer uses the RuntimeDeleter as its deleter, since this is implicitly assumed throughout, i.e. if one wants to construct a BhBase object, use the make_base_ptr helper function.

template <typename InType, typename std::enable_if< type_traits::is_safe_numeric_cast< scalar_type, InType >::value, int >::type = 0> BhArray(const BhArray<InType> &ary)¶

Create a copy of ary using a Bohrium identity operation, which copies the underlying array data.

Note: This function implements implicit type conversion for all widening type casts

BhArray(const BhArray&)¶: Copy constructor that only copies meta data. The underlying array data is untouched

BhArray(BhArray&&)¶: Move constructor that only moves meta data. The underlying array data is untouched

BhArray<T> &operator=(const BhArray<T> &other)¶: Copy the data of other into the array using a Bohrium identity operation

BhArray<T> &operator=(BhArray<T> &&other)¶

Copy the data of other into the array using a Bohrium identity operation

Note: A move assignment is the same as a copy assignment.

template <typename InType, typename std::enable_if< type_traits::is_arithmetic< InType >::value, int >::type = 0> BhArray<T> &operator=(const InType &scalar_value)¶: Copy the scalar of scalar_value into the array using a Bohrium identity operation

BhArray<T> copy() const¶: Return a new copy of the array using a Bohrium identity operation

void reset(BhArray<T> ary)¶: Reset the array to ary

void reset()¶: Reset the array by cleaning all meta data and leave the array uninitialized.

int rank() const¶: Return the rank (number of dimensions) of the array

uint64_t size() const¶: Return the total number of elements of the array

bool isContiguous() const¶: Return whether the view is contiguous and row-major

bool isDataInitialised() const¶: Is the data referenced by this view’s base array already allocated, i.e. initialised

const T *data(bool flush = true) const¶

Obtain the data pointer of the array, not taking ownership of any kind.

Return

The data pointer that might be a nullptr if the data in the base data is not initialised.

Parameters

flush: Should we flush the runtime system before retrieving the data pointer

T *data(bool flush = true)¶: The non-const version of .data()

std::vector<T> vec() const¶

Return a copy of the array as a standard vector

Note: The array must be contiguous

void pprint(std::ostream &os, int current_nesting_level, int max_nesting_level) const¶

Pretty printing the content of the array

Parameters

os: The output stream to write to.
current_nesting_level: The nesting level to print at (typically 0).
max_nesting_level: The maximum nesting level to print at (typically rank()-1).

BhArray<T> operator[](int64_t idx) const¶: Returns a new view of the idx dimension. Negative index counts from the back.

BhArray<T> transpose() const¶: Return a new transposed view.

BhArray<T> reshape(Shape shape) const¶: Return a new reshaped view (the array must be contiguous)

BhArray<T> newAxis(int axis) const¶

Return a new view with a “new axis” inserted.

Return

The new array

Parameters

axis: The “new axis” is inserted just before axis. If negative, the count is backwards (e.g -1 insert a “new axis” at the end of the array)

class BhArrayUnTypedCore¶

#include <BhArray.hpp>

Core class that represent the core attributes of a view that isn’t typed by its dtype

Subclassed by bhxx::BhArray< T >

Public Functions

BhArrayUnTypedCore()¶: Default constructor that leave the instance completely uninitialized

BhArrayUnTypedCore(uint64_t offset, Shape shape, Stride stride, std::shared_ptr<BhBase> base)¶: Constructor to initiate all but the _slides attribute

bh_view getBhView() const¶: Return a bh_view of the array

uint64_t offset() const¶: Return the offset of the array

const Shape &shape() const¶: Return the shape of the array

const Stride &stride() const¶: Return the stride of the array

const std::shared_ptr<BhBase> &base() const¶: Return the base of the array

std::shared_ptr<BhBase> &base()¶: Return the base of the array

void setShapeAndStride(Shape shape, Stride stride)¶: Set the shape and stride of the array (both must have the same length)

const bh_slide &slides() const¶: Return the slides object of the array

bh_slide &slides()¶: Return the slides object of the array

Protected Attributes

uint64_t _offset = 0¶: The array offset (from the start of the base in number of elements)

Shape _shape¶: The array shape (size of each dimension in number of elements)

Stride _stride¶: The array stride (the absolute stride of each dimension in number of elements)

std::shared_ptr<BhBase> _base¶: Pointer to the base of this array.

bh_slide _slides¶: Metadata to support sliding views.

Friends

void swap(BhArrayUnTypedCore &a, BhArrayUnTypedCore &b)¶: Swapping a and b

class BhBase : public bh_base¶

#include <BhBase.hpp>

The base underlying (multiple) arrays

Public Functions

bool ownMemory()¶

Is the memory managed referenced by bh_base’s data pointer managed by Bohrium or is it owned externally

Note: If this flag is false, the class will make sure that the memory is not deleted when going out of scope.

template <typename T> BhBase(size_t nelem, T *memory)¶

Construct a base array with nelem elements using externally managed storage.

The class will make sure, that the storage is not deleted when going out of scope. Needless to say that the memory should be large enough to incorporate nelem_ elements.

Template Parameters

T: The type of each element

Parameters

nelem: Number of elements
memory: Pointer to the external memory

template <typename InputIterator, typename T = typename std::iterator_traits<InputIterator>::value_type> BhBase(InputIterator begin, InputIterator end)¶

Construct a base array and initialise it with the elements provided by an iterator range.

The values are copied into the Bohrium storage. If you want to provide external storage to Bohrium use the constructor BhBase(size_t nelem, T* memory) instead.

template <typename T> BhBase(T dummy, size_t nelem)¶

Construct a base array with nelem elements

Note

The use of this particular constructor is discouraged. It is only needed from BhArray to construct base objects which are uninitialised and do not yet hold any deta. If you wish to construct an uninitialised BhBase object, do this via the BhArray interface and not using this constructor.

Parameters

dummy: Dummy argument to fix the type of elements used. It may only have ever have the value 0 in the appropriate type.
nelem: Number of elements

~BhBase()¶: Destructor

BhBase(const BhBase&)¶: Deleted copy constructor

BhBase &operator=(const BhBase&)¶: Deleted copy assignment

BhBase &operator=(BhBase &&other)¶: Delete move assignment

BhBase(BhBase &&other)¶: Move another BhBase object here

Private Members

bool m_own_memory¶

class Random¶

#include <random.hpp>

Random class that maintain the state of the random number generation

Public Functions

Random(uint64_t seed = std::random_device{}())¶

Create a new random instance

Parameters

seed: T he seed of the random number generation. If not set, std::random_device is used.

BhArray<uint64_t> random123(uint64_t size)¶

New 1D random array using the Random123 algorithm https://www.deshawresearch.com/resources_random123.html

Return

The new random array

Parameters

size: Size of the new 1D random array

void reset(uint64_t seed = std::random_device{}())¶

Reset the random instance

Parameters

seed: The seed of the random number generation. If not set, std::random_device is used.

template <typename T> BhArray<T> randn(Shape shape)¶

Return random floats in the half-open interval [0.0, 1.0) using Random123

Return

Real array

Parameters

shape: The shape of the returned array

Private Members

uint64_t _seed¶

uint64_t _count = 0¶

namespace bhxx¶

Typedefs

typedef BhStaticVector<uint64_t> Shape¶: Static allocated shape that is interchangeable with standard C++ vector as long as the vector is smaller than BH_MAXDIM.

typedef BhStaticVector<int64_t> Stride¶: Static allocated stride that is interchangeable with standard C++ vector as long as the vector is smaller than BH_MAXDIM.

Functions

template <typename T> BhArray<T> arange(int64_t start, int64_t stop, int64_t step)¶

Return evenly spaced values within a given interval.

Return

New 1D array

Template Parameters

T: Data type of the returned array

Parameters

start: Start of interval. The interval includes this value.
stop: End of interval. The interval does not include this value.
step: Spacing between values. For any output out, this is the distance between two adjacent values, out[i+1] - out[i].

void flush()¶: Force the execution of all lazy evaluated array operations

template <typename T> BhArray<T> empty(Shape shape)¶

Return a new empty array

Return

The new array

Template Parameters

T: The data type of the new array

Parameters

shape: The shape of the new array

template <typename OutType, typename InType> BhArray<OutType> empty_like(const bhxx::BhArray<InType> &ary)¶

Return a new empty array that has the same shape as ary

Return

The new array

Template Parameters

OutType: The data type of the returned new array
InType: The data type of the input array

Parameters

ary: The array to take the shape from

template <typename T> BhArray<T> full(Shape shape, T value)¶

Return a new array filled with value

Return

The new array

Template Parameters

T: The data type of the new array

Parameters

shape: The shape of the new array
value: The value to fill the new array with

template <typename T> BhArray<T> zeros(Shape shape)¶

Return a new array filled with zeros

Return

The new array

Template Parameters

T: The data type of the new array

Parameters

shape: The shape of the new array

template <typename T> BhArray<T> ones(Shape shape)¶

Return a new array filled with ones

Return

The new array

Template Parameters

T: The data type of the new array

Parameters

shape: The shape of the new array

template <typename T> BhArray<T> arange(int64_t start, int64_t stop)¶

Return evenly spaced values within a given interval using steps of 1.

Return

New 1D array

Template Parameters

T: Data type of the returned array

Parameters

start: Start of interval. The interval includes this value.
stop: End of interval. The interval does not include this value.

template <typename T> BhArray<T> arange(int64_t stop)¶

Return evenly spaced values from 0 to stop using steps of 1.

Return

New 1D array

Template Parameters

T: Data type of the returned array

Parameters

stop: End of interval. The interval does not include this value.

template <typename OutType, typename InType> BhArray<OutType> cast(const bhxx::BhArray<InType> &ary)¶

Element-wise static_cast.

Return

New array

Template Parameters

OutType: The data type of the returned array
InType: The data type of the input array

Parameters

ary: Input array to cast

Stride contiguous_stride(const Shape &shape)¶: Return a contiguous stride (row-major) based on shape

template <typename T> std::ostream &operator<<(std::ostream &os, const BhArray<T> &ary)¶

Pretty printing the data of an array to a stream Example:

auto A = bhxx::arange<double>(3);
std::cout << A << std::endl;

Return

A reference to os

Template Parameters

T: The data of ary

Parameters

os: The output stream to write to
ary: The array to print

template <typename T> BhArray<T> as_contiguous(BhArray<T> ary)¶

Create an contiguous view or a copy of an array. The array is only copied if it isn’t already contiguous.

Return

Either a view of ary or a new copy of ary.

Template Parameters

T: The data type of ary.

Parameters

ary: The array to make contiguous.

template <int N> Shape broadcasted_shape(std::array<Shape, N> shapes)¶

Return the result of broadcasting shapes against each other

Return

Broadcasted shape

Parameters

shapes: Array of shapes

template <typename T> BhArray<T> broadcast_to(BhArray<T> ary, const Shape &shape)¶

Return a new view of ary that is broadcasted to shape We use the term broadcast as defined by NumPy. Let ret be the broadcasted view of ary: 1) One-sized dimensions are prepended to ret.shape() until it has the same number of dimension as ary. 2) The stride of each one-sized dimension in ret is set to zero. 3) The shape of ary is set to shape

Note

See: https://docs.scipy.org/doc/numpy-1.15.0/user/basics.broadcasting.html

Return

The broadcasted array

Parameters

ary: Input array
shape: The new shape

template <typename T1, typename T2> bool is_same_array(const BhArray<T1> &a, const BhArray<T2> &b)¶

Check whether a and b are the same view pointing to the same base

Return

The boolean answer.

Template Parameters

T1: The data type of a.
T2: The data type of b.

Parameters

a: The first array to compare.
b: The second array to compare.

template <typename T1, typename T2> bool may_share_memory(const BhArray<T1> &a, const BhArray<T2> &b)¶

Check whether a and b can share memory

Note

A return of True does not necessarily mean that the two arrays share any element. It just means that they might.

Return

The boolean answer.

Template Parameters

T1: The data type of a.
T2: The data type of b.

Parameters

a: The first array to compare.
b: The second array to compare.

BhArray<bool> add(const BhArray<bool> &in1, const BhArray<bool> &in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Array input.

BhArray<bool> add(const BhArray<bool> &in1, bool in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Scalar input.

BhArray<bool> add(bool in1, const BhArray<bool> &in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Scalar input.
in2: Array input.

BhArray<std::complex<double>> add(const BhArray<std::complex<double>> &in1, const BhArray<std::complex<double>> &in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Array input.

BhArray<std::complex<double>> add(const BhArray<std::complex<double>> &in1, std::complex<double> in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Scalar input.

BhArray<std::complex<double>> add(std::complex<double> in1, const BhArray<std::complex<double>> &in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Scalar input.
in2: Array input.

BhArray<std::complex<float>> add(const BhArray<std::complex<float>> &in1, const BhArray<std::complex<float>> &in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Array input.

BhArray<std::complex<float>> add(const BhArray<std::complex<float>> &in1, std::complex<float> in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Scalar input.

BhArray<std::complex<float>> add(std::complex<float> in1, const BhArray<std::complex<float>> &in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Scalar input.
in2: Array input.

BhArray<float> add(const BhArray<float> &in1, const BhArray<float> &in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Array input.

BhArray<float> add(const BhArray<float> &in1, float in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Scalar input.

BhArray<float> add(float in1, const BhArray<float> &in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Scalar input.
in2: Array input.

BhArray<double> add(const BhArray<double> &in1, const BhArray<double> &in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Array input.

BhArray<double> add(const BhArray<double> &in1, double in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Scalar input.

BhArray<double> add(double in1, const BhArray<double> &in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Scalar input.
in2: Array input.

BhArray<int16_t> add(const BhArray<int16_t> &in1, const BhArray<int16_t> &in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Array input.

BhArray<int16_t> add(const BhArray<int16_t> &in1, int16_t in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Scalar input.

BhArray<int16_t> add(int16_t in1, const BhArray<int16_t> &in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Scalar input.
in2: Array input.

BhArray<int32_t> add(const BhArray<int32_t> &in1, const BhArray<int32_t> &in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Array input.

BhArray<int32_t> add(const BhArray<int32_t> &in1, int32_t in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Scalar input.

BhArray<int32_t> add(int32_t in1, const BhArray<int32_t> &in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Scalar input.
in2: Array input.

BhArray<int64_t> add(const BhArray<int64_t> &in1, const BhArray<int64_t> &in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Array input.

BhArray<int64_t> add(const BhArray<int64_t> &in1, int64_t in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Scalar input.

BhArray<int64_t> add(int64_t in1, const BhArray<int64_t> &in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Scalar input.
in2: Array input.

BhArray<int8_t> add(const BhArray<int8_t> &in1, const BhArray<int8_t> &in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Array input.

BhArray<int8_t> add(const BhArray<int8_t> &in1, int8_t in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Scalar input.

BhArray<int8_t> add(int8_t in1, const BhArray<int8_t> &in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Scalar input.
in2: Array input.

BhArray<uint16_t> add(const BhArray<uint16_t> &in1, const BhArray<uint16_t> &in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Array input.

BhArray<uint16_t> add(const BhArray<uint16_t> &in1, uint16_t in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Scalar input.

BhArray<uint16_t> add(uint16_t in1, const BhArray<uint16_t> &in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Scalar input.
in2: Array input.

BhArray<uint32_t> add(const BhArray<uint32_t> &in1, const BhArray<uint32_t> &in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Array input.

BhArray<uint32_t> add(const BhArray<uint32_t> &in1, uint32_t in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Scalar input.

BhArray<uint32_t> add(uint32_t in1, const BhArray<uint32_t> &in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Scalar input.
in2: Array input.

BhArray<uint64_t> add(const BhArray<uint64_t> &in1, const BhArray<uint64_t> &in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Array input.

BhArray<uint64_t> add(const BhArray<uint64_t> &in1, uint64_t in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Scalar input.

BhArray<uint64_t> add(uint64_t in1, const BhArray<uint64_t> &in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Scalar input.
in2: Array input.

BhArray<uint8_t> add(const BhArray<uint8_t> &in1, const BhArray<uint8_t> &in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Array input.

BhArray<uint8_t> add(const BhArray<uint8_t> &in1, uint8_t in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Scalar input.

BhArray<uint8_t> add(uint8_t in1, const BhArray<uint8_t> &in2)¶

Add arguments element-wise.

Return

Output array.

Parameters

in1: Scalar input.
in2: Array input.

BhArray<std::complex<double>> subtract(const BhArray<std::complex<double>> &in1, const BhArray<std::complex<double>> &in2)¶

Subtract arguments, element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Array input.

BhArray<std::complex<double>> subtract(const BhArray<std::complex<double>> &in1, std::complex<double> in2)¶

Subtract arguments, element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Scalar input.

BhArray<std::complex<double>> subtract(std::complex<double> in1, const BhArray<std::complex<double>> &in2)¶

Subtract arguments, element-wise.

Return

Output array.

Parameters

in1: Scalar input.
in2: Array input.

BhArray<std::complex<float>> subtract(const BhArray<std::complex<float>> &in1, const BhArray<std::complex<float>> &in2)¶

Subtract arguments, element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Array input.

BhArray<std::complex<float>> subtract(const BhArray<std::complex<float>> &in1, std::complex<float> in2)¶

Subtract arguments, element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Scalar input.

BhArray<std::complex<float>> subtract(std::complex<float> in1, const BhArray<std::complex<float>> &in2)¶

Subtract arguments, element-wise.

Return

Output array.

Parameters

in1: Scalar input.
in2: Array input.

BhArray<float> subtract(const BhArray<float> &in1, const BhArray<float> &in2)¶

Subtract arguments, element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Array input.

BhArray<float> subtract(const BhArray<float> &in1, float in2)¶

Subtract arguments, element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Scalar input.

BhArray<float> subtract(float in1, const BhArray<float> &in2)¶

Subtract arguments, element-wise.

Return

Output array.

Parameters

in1: Scalar input.
in2: Array input.

BhArray<double> subtract(const BhArray<double> &in1, const BhArray<double> &in2)¶

Subtract arguments, element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Array input.

BhArray<double> subtract(const BhArray<double> &in1, double in2)¶

Subtract arguments, element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Scalar input.

BhArray<double> subtract(double in1, const BhArray<double> &in2)¶

Subtract arguments, element-wise.

Return

Output array.

Parameters

in1: Scalar input.
in2: Array input.

BhArray<int16_t> subtract(const BhArray<int16_t> &in1, const BhArray<int16_t> &in2)¶

Subtract arguments, element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Array input.

BhArray<int16_t> subtract(const BhArray<int16_t> &in1, int16_t in2)¶

Subtract arguments, element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Scalar input.

BhArray<int16_t> subtract(int16_t in1, const BhArray<int16_t> &in2)¶

Subtract arguments, element-wise.

Return

Output array.

Parameters

in1: Scalar input.
in2: Array input.

BhArray<int32_t> subtract(const BhArray<int32_t> &in1, const BhArray<int32_t> &in2)¶

Subtract arguments, element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Array input.

BhArray<int32_t> subtract(const BhArray<int32_t> &in1, int32_t in2)¶

Subtract arguments, element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Scalar input.

BhArray<int32_t> subtract(int32_t in1, const BhArray<int32_t> &in2)¶

Subtract arguments, element-wise.

Return

Output array.

Parameters

in1: Scalar input.
in2: Array input.

BhArray<int64_t> subtract(const BhArray<int64_t> &in1, const BhArray<int64_t> &in2)¶

Subtract arguments, element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Array input.

BhArray<int64_t> subtract(const BhArray<int64_t> &in1, int64_t in2)¶

Subtract arguments, element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Scalar input.

BhArray<int64_t> subtract(int64_t in1, const BhArray<int64_t> &in2)¶

Subtract arguments, element-wise.

Return

Output array.

Parameters

in1: Scalar input.
in2: Array input.

BhArray<int8_t> subtract(const BhArray<int8_t> &in1, const BhArray<int8_t> &in2)¶

Subtract arguments, element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Array input.

BhArray<int8_t> subtract(const BhArray<int8_t> &in1, int8_t in2)¶

Subtract arguments, element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Scalar input.

BhArray<int8_t> subtract(int8_t in1, const BhArray<int8_t> &in2)¶

Subtract arguments, element-wise.

Return

Output array.

Parameters

in1: Scalar input.
in2: Array input.

BhArray<uint16_t> subtract(const BhArray<uint16_t> &in1, const BhArray<uint16_t> &in2)¶

Subtract arguments, element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Array input.

BhArray<uint16_t> subtract(const BhArray<uint16_t> &in1, uint16_t in2)¶

Subtract arguments, element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Scalar input.

BhArray<uint16_t> subtract(uint16_t in1, const BhArray<uint16_t> &in2)¶

Subtract arguments, element-wise.

Return

Output array.

Parameters

in1: Scalar input.
in2: Array input.

BhArray<uint32_t> subtract(const BhArray<uint32_t> &in1, const BhArray<uint32_t> &in2)¶

Subtract arguments, element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Array input.

BhArray<uint32_t> subtract(const BhArray<uint32_t> &in1, uint32_t in2)¶

Subtract arguments, element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Scalar input.

BhArray<uint32_t> subtract(uint32_t in1, const BhArray<uint32_t> &in2)¶

Subtract arguments, element-wise.

Return

Output array.

Parameters

in1: Scalar input.
in2: Array input.

BhArray<uint64_t> subtract(const BhArray<uint64_t> &in1, const BhArray<uint64_t> &in2)¶

Subtract arguments, element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Array input.

BhArray<uint64_t> subtract(const BhArray<uint64_t> &in1, uint64_t in2)¶

Subtract arguments, element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Scalar input.

BhArray<uint64_t> subtract(uint64_t in1, const BhArray<uint64_t> &in2)¶

Subtract arguments, element-wise.

Return

Output array.

Parameters

in1: Scalar input.
in2: Array input.

BhArray<uint8_t> subtract(const BhArray<uint8_t> &in1, const BhArray<uint8_t> &in2)¶

Subtract arguments, element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Array input.

BhArray<uint8_t> subtract(const BhArray<uint8_t> &in1, uint8_t in2)¶

Subtract arguments, element-wise.

Return

Output array.

Parameters

in1: Array input.
in2: Scalar input.

BhArray<uint8_t> subtract(uint8_t in1, const BhArray<uint8_t> &in2)¶

Subtract arguments, element-wise.

Return

Output array.

Parameters

in1: Scalar input.
in2: Array input.

BhArray<bool> multiply(const BhArray<bool> &in1, const BhArray<bool> &in2)¶