В 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).