Создание индикатора выполнения в Python с помощью Numba и Cuda

Вы можете использовать numba cuda mapped_array , чтобы помочь с этой задачей. Под капотом это говорит numba создать закрепленное выделение и сделать его доступным для использования на устройстве, что информирует numba не копировать его на устройство, даже несмотря на то, что pinned_array обычно выглядит как массив хостов numba.

В сочетании с этим нам нужно будет убедиться, что numba не пытается копировать массивы, поскольку это приведет к синхронизации в «автоматическом» случае, чего мы не хотим.

Я действительно не знаю, как измерить прогресс этого алгоритма. Например, цикл while в poisson_sample, кажется, повторяется 4 раза для элемента, thread_id которого равен нулю, но я сомневаюсь, что это верно для массива out. (У меня есть лучшее представление о том, как отслеживать прогресс других алгоритмов.)

Если мы знаем, сколько времени должен занять алгоритм в зависимости от прогресса, мы можем просто отслеживать значение, сообщаемое ядром. Когда он достигает 100% (или почти), мы прекращаем мониторинг и приступаем к остальной работе.

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

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

Во всяком случае, в Linux у меня работает следующее, как приблизительный набросок. Это демонстрация с использованием событий, хотя, если вы знаете ход работы алгоритма, события на самом деле не нужны. Вот версия с событиями:

$ cat t1.py
from __future__ import print_function, absolute_import

from numba import cuda
from numba.cuda.random import create_xoroshiro128p_states, xoroshiro128p_uniform_float32
import numpy as np
from math import gamma, exp, ceil


# This function is just an example of what i'd like to put a progress bar on
@cuda.jit
def generate_samples(rng_states, out, rate, progress):
    thread_id = cuda.grid(1)

    def poission_sample(rate, random_number, progress):
        probability_sum = 0
        index = -1
        while probability_sum < random_number:
            index += 1
            probability_sum += ((rate**index)/gamma(index+1)) * exp(-rate)
        cuda.atomic.add(progress, 0, 1)
        return index

    out[thread_id] = poission_sample(rate, xoroshiro128p_uniform_float32(rng_states, thread_id), progress)

number_of_samples = 10000000
progress = cuda.mapped_array(1, dtype=np.int64)
progress[0] = 0;
last_pct = 0
my_e = cuda.event()
threads_per_block = 512
blocks = ceil(number_of_samples/threads_per_block)
my_divisor = (threads_per_block * blocks) // 100
rng_states = create_xoroshiro128p_states(threads_per_block * blocks, seed=1)
out = np.zeros(threads_per_block * blocks, dtype=np.float32)
out_d = cuda.device_array_like(out)
generate_samples[blocks, threads_per_block](rng_states, out_d, 5, progress)
my_e.record()
print(last_pct)
while my_e.query() == False:
    cur_pct = progress[0]/my_divisor
    if cur_pct > last_pct + 10:
        last_pct = cur_pct
        print(cur_pct)
out = out_d.copy_to_host()

print('Average Sample:', out.mean())
$ python3 t1.py
0
10.00129996100117
20.00291991240263
30.004539863804087
40.00519984400468
50.00713978580642
60.00811975640731
70.00941971740848
80.01039968800936
90.01105966820995
Average Sample: 5.000568
$

Вот версия без событий:

$ cat t2.py
from __future__ import print_function, absolute_import

from numba import cuda
from numba.cuda.random import create_xoroshiro128p_states, xoroshiro128p_uniform_float32
import numpy as np
from math import gamma, exp, ceil


# This function is just an example of what i'd like to put a progress bar on
@cuda.jit
def generate_samples(rng_states, out, rate, progress):
    thread_id = cuda.grid(1)

    def poission_sample(rate, random_number, progress):
        probability_sum = 0
        index = -1
        while probability_sum < random_number:
            index += 1
            probability_sum += ((rate**index)/gamma(index+1)) * exp(-rate)
        cuda.atomic.add(progress, 0, 1)
        return index

    out[thread_id] = poission_sample(rate, xoroshiro128p_uniform_float32(rng_states, thread_id), progress)

number_of_samples = 10000000
progress = cuda.mapped_array(1, dtype=np.int64)
progress[0] = 0;
last_pct = 0
threads_per_block = 512
blocks = ceil(number_of_samples/threads_per_block)
my_divisor = (threads_per_block * blocks) // 100
rng_states = create_xoroshiro128p_states(threads_per_block * blocks, seed=1)
out = np.zeros(threads_per_block * blocks, dtype=np.float32)
out_d = cuda.device_array_like(out)
generate_samples[blocks, threads_per_block](rng_states, out_d, 5, progress)
print(last_pct)
while last_pct < 90:
    cur_pct = progress[0]/my_divisor
    if cur_pct > last_pct + 10:
        last_pct = cur_pct
        print(cur_pct)
out = out_d.copy_to_host()

print('Average Sample:', out.mean())
$ python3 t2.py
0
10.000019999400019
20.000039998800037
30.000059998200054
40.000079997600075
50.00009999700009
60.00011999640011
70.00013999580013
80.00015999520015
90.00017999460016
Average Sample: 5.000568
$

Я запускал оба из них на Linux. Версия без использования событий может работать лучше в Windows или, возможно, наоборот (запрос события может ускорить отправку работы). Если вы используете графический процессор дисплея в Windows (т. е. графический процессор не в режиме TCC), то пакетная обработка/планирование работы WDDM может вызвать проблемы. Вы можете попробовать обе настройки для Планирования графического процессора с аппаратным ускорением Windows, чтобы увидеть, работает ли один вариант лучше, чем другой.

Кроме того, это ядро ​​работает на моем графическом процессоре менее чем за секунду (фактически, на моем графическом процессоре GTX 970 продолжительность ядра составляет около 300 мс). Так что это может быть не интересный тестовый пример.