Параллелизированная переменная счетчика ядра Cuda

Question

Есть ли способ иметь целочисленную переменную-счетчик, которую можно увеличивать / уменьшать во всех потоках в распараллеленном ядре cuda? Приведенный ниже код выводит «[1]», поскольку изменения в массиве счетчиков из одного потока не применяются в других.

import numpy as np
from numba import cuda


@cuda.jit('void(int32[:])')
def func(counter):
    counter[0] = counter[0] + 1


counter = cuda.to_device(np.zeros(1, dtype=np.int32))
threadsperblock = 64
blockspergrid = 18
func[blockspergrid, threadsperblock](counter)
print(counter.copy_to_host())

Answer 1

Один из подходов заключается в использовании numba cuda atomics :

$ cat t18.py
import numpy as np
from numba import cuda


@cuda.jit('void(int32[:])')
def func(counter):
    cuda.atomic.add(counter, 0, 1)


counter = cuda.to_device(np.zeros(1, dtype=np.int32))
threadsperblock = 64
blockspergrid = 18
print blockspergrid * threadsperblock
func[blockspergrid, threadsperblock](counter)
print(counter.copy_to_host())
$ python t18.py
1152
[1152]
$

атомарная операция выполняет неделимую операцию чтения-изменения-записи для цели, поэтому потоки не мешают друг другу при обновлении целевой переменной.

Конечно, возможны и другие методы, в зависимости от ваших реальных потребностей, такие как классическое параллельное сокращение . Numba обеспечивает также уменьшение сахара .