Безопасно ли реализовывать cuda gridsync () в Numba следующим образом - PullRequest
0 голосов
/ 09 февраля 2019

В Numba отсутствует команда cuda-C gridsync (), поэтому нет постоянного метода синхронизации по всей сетке.Доступны только синхронизации на уровне блоков.

Если cudaKernal1 - очень быстрое время выполнения, то следующий код будет работать в 1000 раз быстрее

for i in range(10000):
   X = X + cudaKernel1[(100,100),(32,32)] (X)

, если поместить цикл в то же ядро, чтобы избежатьВремя настройки ядра GPU.Но вы не можете этого сделать, потому что вам требуется, чтобы вся сетка завершилась до начала следующей итерации, и в Numba нет команды gridsync ().

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

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

Я не могу понять,что небезопасно или почему возникнет состояние расы, которое станет ловушкой.

Что не так с чем-то вроде следующего.

@numba.cuda.jit('void()')
def gpu_initGridSync():
    if ( cuda.threadIdx.x == 0): 
        Global_u[0] = 0
        Global_u[1] = 0

@numba.cuda.jit('void(int32)'device=True)
def gpu_fakeGridSync(i):
    ###wait till the the entire grid has finished doSomething()
    # in Cuda-C we'd call gridsync()
    # but lack that in Numba so do the following instead.

    #Syncthreads in current block
    numba.cuda.syncthreads()

    #increment global counter, once per block
    if ( cuda.threadIdx.x == 0 ):  numba.atomic.add( Global_u, 0, 1 )

    # idle in a loop
    while ( Global_u[0] < (i+1)*cuda.gridDim.x-1 ) ):  pass   #2

    #regroup the block threads after the slow global memory reads.
    numba.cuda.syncthreads()

    # now, to avoid a race condition of blocks re-entering the above while
    # loop before other blocks have exited we do this global sync a second time

     #increment global counter, once per block
    if ( cuda.threadIdx.x == 0 ):  numba.atomic.add( Global_u,1, 1 )

    # idle in a loop
    while ( Global_u[1] > (i+2)*cuda.gridDim.x ) ):  pass   #2

    #regroup the block threads after the slow global memory reads.
    numba.cuda.syncthreads()

Это затем используется следующим образом:

@numba.cuda.jit('void(float32[:])')):
def ReallyReallyFast(X):
    i = numba.cuda.grid(1)
    for h in range(1,40000,4):
        temp = calculateSomething(X)
        gpu_fakeGridSync(h)
        X[i] = X[i]+temp
        gpu_fakeGridSync(h+2)

gpu_initGridSync[(1,),(1,)]()
ReallyReallyFast[(1000,), (32,) ](X)


@numba.cuda.jit('float32(float32[:])',device=True):
def calculateSomething(X):  # A dummy example of a very fast kernel operation
    i = numba.cuda.grid(1)
    if (i>0):
        return (X[i]-X[i-1])/2.0
    return 0.0

Мне кажется, это логично.Существует один тонкий шаг для инициализации глобального счетчика.Это должно быть сделано в его собственном вызове ядра, чтобы избежать состояния гонки.Но после этого я могу свободно вызывать fakeGridSync без его повторной инициализации.Я должен следить за тем, как в какой итерации цикла я его называю (отсюда и передается параметр in в gridSync).

Признаю, я вижу, что некоторые усилия потрачены впустую, но разве это убийца сделки?Например, в операторе № 2 этот цикл while означает, что все потоки во всех законченных блоках вращают свои колеса с потерянным усилием.Я предполагаю, что это может слегка замедлить блоки сетки, которые все еще пытаются выполнить «doSomething».Однако я не уверен, насколько плохи эти напрасные усилия.Второй провал в утверждении № 2 заключается в том, что все потоки борются за одну и ту же глобальную память, поэтому они будут медленнее обращаться к ней.Это может быть даже хорошо, если планировщик откладывает их выполнение и позволяет более часто выполнять полезные потоки.Можно улучшить этот простой код, если бы в каждом блоке проверялось только значение thread (0).

1 Ответ

0 голосов
/ 10 февраля 2019

Я думаю, что комментарий Роберта Кровеллы указывает на правильный ответ о том, почему этот метод потерпит неудачу.

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

В настоящее время графические процессоры Nvidia не имеют приоритетных многозадачных планировщиков.Работа завершается.

Таким образом, возможно, что как только достаточное количество блоков войдет в цикл while для ожидания, остальные блоки не будут запущены планировщиком.Таким образом, цикл ожидания будет ждать вечно.

Я вижу, что есть исследовательские работы, предлагающие, как Nvidia могла бы сделать свой планировщик приоритетным.https://www.computer.org/csdl/proceedings/snpd/2012/2120/00/06299288.pdf Но, очевидно, сейчас это не так.

Мне просто интересно, как cuda-C удалось выполнить команду gridSync ().Если это может быть сделано в C, должен быть какой-то общий способ обойти эти ограничения. Это загадка, я надеюсь, что кто-то прокомментирует ниже

Это действительно позор, оставить 1000-кратное ускорение натаблица.

...