Supported Atomic Operations

Several atomic operations supported in DPC++ are also supported in the extension. Those that are presently implemented are as follows:

class numba_dpex.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_dpex'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_DPEX_ACTIVATE_ATOMICS_FP_NATIVE=1
    NUMBA_DPEX_LLVM_SPIRV_ROOT=/path/to/dpcpp/provided/llvm_spirv

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

    Without these two environment variables numba-dpex 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 dpctl.device_context(device):
        atomic_add[global_size, dppy.DEFAULT_LOCAL_SIZE](a)

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

    print("Done...")

Note

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

Generating Native FP Atomics

Generating native floating-point atomics is supported. This feature is experimental. Users will need to provide the following environment variables to activate it.

NUMBA_DPEX_ACTIVATE_ATOMICS_FP_NATIVE=1 NUMBA_DPEX_LLVM_SPIRV_ROOT=/path/to/dpcpp/provided/llvm_spirv

Example command:

NUMBA_DPEX_ACTIVATE_ATOMICS_FP_NATIVE=1 \
NUMBA_DPEX_LLVM_SPIRV_ROOT=/path/to/dpcpp/provided/llvm_spirv \
python program.py

Full examples

• numba_dpex/examples/atomic_op.py