Почему numba cuda работает медленно после повторного вызова? - PullRequest
Новая
       89

Почему numba cuda работает медленно после повторного вызова?

0 голосов
/ 10 сентября 2018

Я экспериментирую, как использовать cuda внутри numba.Однако я столкнулся с чем-то отличным от моих ожиданий.Вот мой код 

from numba import cuda
@cuda.jit
def matmul(A, B, C):
"""Perform square matrix multiplication of C = A * B
"""
d=cuda.local.array((3,3),dtype=numba.float64)
i, j = cuda.grid(2)
if i < C.shape[0] and j < C.shape[1]:
    tmp = 0.
    for k in range(A.shape[1]):
        tmp += A[i, k] * B[k, j]
    C[i, j] = tmp

Это матричная функция, которую я сам определил для тестирования с использованием numba.cuda.Перед запуском тестов я также загрузил массивы в следующем коде: 

import numpy as np
a=np.random.rand(2000,2000)
b=np.random.rand(2000,2000)
c=np.empty((2000,2000))
a1=cuda.to_device(a)
b1=cuda.to_device(b)
c1=cuda.to_device(c)

Затем я использовал следующий код для эксперимента: 

from time import time
count =0
start=time()
for i in range(2000):
  matmul[(256,256),(16,16)](a1,b1,c1)
  count +=1
  print(count)

Циклы for выполнялись очень быстро дляпервые 1028 пробежек.Однако после 1028-го он работал очень медленно. Что именно вызвало это и как мне это исправить.Кстати, я работаю на win10.

Вот моя информация о cuda, вызванная из numba.cuda 

from numba import cuda
gpu = cuda.get_current_device()
print("name = %s" % gpu.name)
print("maxThreadsPerBlock = %s" % str(gpu.MAX_THREADS_PER_BLOCK))
print("maxBlockDimX = %s" % str(gpu.MAX_BLOCK_DIM_X))
print("maxBlockDimY = %s" % str(gpu.MAX_BLOCK_DIM_Y))
print("maxBlockDimZ = %s" % str(gpu.MAX_BLOCK_DIM_Z))
print("maxGridDimX = %s" % str(gpu.MAX_GRID_DIM_X))
print("maxGridDimY = %s" % str(gpu.MAX_GRID_DIM_Y))
print("maxGridDimZ = %s" % str(gpu.MAX_GRID_DIM_Z))
print("maxSharedMemoryPerBlock = %s" % 
str(gpu.MAX_SHARED_MEMORY_PER_BLOCK))
print("asyncEngineCount = %s" % str(gpu.ASYNC_ENGINE_COUNT))
print("canMapHostMemory = %s" % str(gpu.CAN_MAP_HOST_MEMORY))
print("multiProcessorCount = %s" % str(gpu.MULTIPROCESSOR_COUNT))
print("warpSize = %s" % str(gpu.WARP_SIZE))
print("unifiedAddressing = %s" % str(gpu.UNIFIED_ADDRESSING))
print("pciBusID = %s" % str(gpu.PCI_BUS_ID))
print("pciDeviceID = %s" % str(gpu.PCI_DEVICE_ID))

, и вывод:

name = b'GeForceGTX 1050 Ti '

maxThreadsPerBlock = 1024

maxBlockDimX = 1024

maxBlockDimY = 1024

maxBlockDimZ = 64

maxGridDimX = 21474836471027 *

maxGridDimY = 65535

maxGridDimZ = 65535

maxSharedMemoryPerBlock = 49152

asyncEngineCount = 2

canMapHostMemory = 1 * 1037multiProcessorCount = 6

warpSize = 32

unifiedAddressing = 1

pciBusID = 3

pciDeviceID = 0

1 Ответ

0 голосов
/ 10 сентября 2018

Это вызвано асинхронной очередью запуска, связанной с запусками ядра GPU.

Когда вы указываете numba передать ядро ​​GPU: 

matmul[(256,256),(16,16)](a1,b1,c1)

Этот запрос переходит в очередь, ипоток ЦП (т. е. python), который выдал этот вызов ядра, может продолжаться, даже если ядро ​​GPU еще не завершено или даже не запущено.

Среда выполнения CUDA ставит эти запросы в очередь и выдает их, так как GPU готов кбольше работы.

То, что вы наблюдаете изначально во время очень быстрого увеличения цикла for, - это очередь, заполняемая рабочими запросами.Это не отражает фактического времени, которое требуется графическому процессору для выполнения работы.

В конечном итоге очередь заполняется, и среда выполнения CUDA останавливает поток ЦП (т. Е. Python) в точке запуска ядра до появления очереди.слот открывается.В этот момент цикл for может продолжаться еще одну итерацию.Именно в этот момент (возможно, около 1028 итераций) вы начинаете видеть «замедление».После этого цикл for продолжается на приблизительно скорости, с которой ядра GPU выполняются и удаляются из очереди обработки.

Здесь нечего исправлять;это ожидаемое поведение.

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

Например, numba предоставляет numba.cuda.synchronize () Так что, если вы измените свой цикл for следующим образом: 

for i in range(2000):
  matmul[(256,256),(16,16)](a1,b1,c1)
  cuda.synchronize()
  count +=1
  print(count)

Вы увидите, что цикл for продолжится с фактическойскорость завершения работы графического процессора, а не скорость «заполнения очереди».

...