3.5. Supported Atomic Operations¶
Numba provides access to some of the atomic operations supported in CUDA, in the
Those that are presently implemented are as follows:
Namespace for atomic operations
add(ary, idx, val)
Perform atomic ary[idx] += val. Supported on int32, float32, and float64 operands only.
max(ary, idx, val)
Perform atomic ary[idx] = max(ary[idx], val). NaN is treated as a missing value, so max(NaN, n) == max(n, NaN) == n. Note that this differs from Python and Numpy behaviour, where max(a, b) is always a when either a or b is a NaN.
Supported on float64 operands only.
The following code demonstrates the use of
find the maximum value in an array. Note that this is not the most efficient way
of finding a maximum in this case, but that it serves as an example:
from numba import cuda import numpy as np @cuda.jit def max_example(result, values): """Find the maximum value in values and store in result""" tid = cuda.threadIdx.x bid = cuda.blockIdx.x bdim = cuda.blockDim.x i = (bid * bdim) + tid cuda.atomic.max(result, 0, values[i]) arr = np.random.rand(16384) result = np.zeros(1, dtype=np.float64) max_example[256,64](result, arr) print(result) # Found using cuda.atomic.max print(max(arr)) # Print max(arr) for comparision (should be equal!)