Atomic operations

In Taichi, augmented assignments (e.g., x[i] += 1) are automatically atomic.

Warning

When modifying global variables in parallel, make sure you use atomic operations. For example, to sum up all the elements in x,

@ti.kernel
def sum():
    for i in x:
        # Approach 1: OK
        total[None] += x[i]

        # Approach 2: OK
        ti.atomic_add(total[None], x[i])

        # Approach 3: Wrong result since the operation is not atomic.
        total[None] = total[None] + x[i]

Note

When atomic operations are applied to local values, the Taichi compiler will try to demote these operations into their non-atomic counterparts.

Apart from the augmented assignments, explicit atomic operations, such as ti.atomic_add, also do read-modify-write atomically. These operations additionally return the old value of the first argument.

Below is a list of all explicit atomic operations:

ti.atomic_add(x, y)
ti.atomic_sub(x, y)

Atomically compute x + y or x - y and store the result in x.

Returns:The old value of x.

For example,

x[i] = 3
y[i] = 4
z[i] = ti.atomic_add(x[i], y[i])
# now x[i] = 7, y[i] = 4, z[i] = 3
ti.atomic_and(x, y)
ti.atomic_or(x, y)
ti.atomic_xor(x, y)

Atomically compute x & y (bitwise and), x | y (bitwise or), or x ^ y (bitwise xor), and store the result in x.

Returns:The old value of x.

Note

Supported atomic operations on each backend:

type CPU/CUDA OpenGL Metal
i32 OK OK OK
f32 OK OK OK
i64 OK EXT N/A
f64 OK EXT N/A

(OK: supported; EXT: require extension; N/A: not available)