Supported Atomic Operations

Numba-dppy supports some of the atomic operations supported in DPC++. Those that are presently implemented are as follows:

class numba_dppy.ocl.stubs.atomic

atomic namespace

add(ary, idx, val)

Perform atomic ary[idx] += val.

Returns the old value at the index location as if it is loaded atomically.

Note

Supported on int32, int64, float32, float64 operands only.

sub(ary, idx, val)

Perform atomic ary[idx] -= val.

Returns the old value at the index location as if it is loaded atomically.

Note

Supported on int32, int64, float32, float64 operands only.

Example

Here’s an example of how to use atomics add in DPPY:

def main():
    """
    The example demonstrates the use of numba_dppy's ``atomic_add`` intrinsic
    function on a SYCL GPU device. The ``dpctl.select_gpu_device`` is
    equivalent to ``sycl::gpu_selector`` and returns a sycl::device of type GPU.

    If we want to generate native floating point atomics for spported
    SYCL devices we need to set two environment variables:
    NUMBA_DPPY_ACTIVATE_ATOMICS_FP_NATIVE=1
    NUMBA_DPPY_LLVM_SPIRV_ROOT=/path/to/dpcpp/provided/llvm_spirv

    To run this example:
    NUMBA_DPPY_ACTIVATE_ATOMICS_FP_NATIVE=1 NUMBA_DPPY_LLVM_SPIRV_ROOT=/path/to/dpcpp/provided/llvm_spirv python atomic_op.py

    Without these two environment variables Numba_dppy will use other
    implementation for floating point atomics.
    """

    @dppy.kernel
    def atomic_add(a):
        dppy.atomic.add(a, 0, 1)

    global_size = 100
    a = np.array([0], dtype=np.float32)

    # Use the environment variable SYCL_DEVICE_FILTER to change the default device.
    # See https://github.com/intel/llvm/blob/sycl/sycl/doc/EnvironmentVariables.md#sycl_device_filter.
    device = dpctl.select_default_device()
    print("Using device ...")
    device.print_device_info()

    with dppy.offload_to_sycl_device(device):
        atomic_add[global_size, dppy.DEFAULT_LOCAL_SIZE](a)

    # Expected 100, because global_size = 100
    print(a)

    print("Done...")

Note

The numba_dppy.atomic.add function is analogous to The numba.cuda.atomic.add provided by the numba.cuda backend.

Generating Native FP Atomics

Numba-dppy supports generating native floating-point atomics. This feature is experimental. Users will need to provide the following environment variables to activate it.

NUMBA_DPPY_ACTIVATE_ATOMICS_FP_NATIVE=1 NUMBA_DPPY_LLVM_SPIRV_ROOT=/path/to/dpcpp/provided/llvm_spirv

Example command:

NUMBA_DPPY_ACTIVATE_ATOMICS_FP_NATIVE=1 \
NUMBA_DPPY_LLVM_SPIRV_ROOT=/path/to/dpcpp/provided/llvm_spirv \
python program.py

Full examples

• numba_dppy/examples/atomic_op.py