Прежде всего, основная операция, которую вы показали, состоит в том, чтобы взять две матрицы, перенести их в графический процессор, выполнить несколько поэлементных умножений для получения третьего массива и передать этот третий массив обратно на хост.
Может быть возможно сделать реализацию numba / cuda guvectorize (или ядро cuda.jit), которая может работать быстрее, чем наивная последовательная реализация Python, но я сомневаюсь, что было бы возможно превысить производительность хорошо написанного хоста код (например, используя некоторый метод распараллеливания, например, guvectorize), чтобы сделать то же самое. Это связано с тем, что арифметическая интенсивность на байт, передаваемая между хостом и устройством, слишком мала. Эта операция слишком проста.
Во-вторых, я считаю, что важно начать с понимания того, для чего предназначены numba vectorize и guvectorize. Основной принцип состоит в том, чтобы написать определение ufunc с точки зрения «что будет делать работник?» и затем позвольте numba раскрутить несколько рабочих из этого. Способ, которым вы указываете Numba ускорять работу нескольких рабочих, состоит в том, чтобы передать набор данных, который больше, чем подписи, которые вы дали . Следует отметить, что numba не знает, как распараллелить цикл for внутри определения ufunc . Он получает параллельную «силу», принимая ваше определение ufunc и выполняя его среди параллельных работников, где каждый работник обрабатывает «срез» данных, но запускает все ваше определение ufunc на этом срезе. Как дополнительное чтение, я рассмотрел некоторые из этих оснований здесь также.
Итак, проблема, с которой мы сталкиваемся в вашей реализации, заключается в том, что вы написали подпись (и ufunc), которая отображает весь входной набор данных на одного работника. Как показала @talonmies, ваше базовое ядро набрало в общей сложности 64 потока / рабочих (что очень мало, чтобы быть интересным для графического процессора, даже если не считать вышеприведенных утверждений об арифметической интенсивности), но на самом деле я подозреваю, что На самом деле 64 - это просто минимальный размер потокового блока numba, и фактически только один поток в этом потоковом блоке выполняет какую-либо полезную вычислительную работу. Этот поток выполняет весь ваш ufunc, включая все циклы for, последовательно.
Это явно не то, что кто-то хотел бы для рационального использования vectorize или guvectorize.
Итак, давайте вернемся к тому, что вы пытаетесь сделать. В конечном итоге ваш ufunc хочет умножить входное значение из одного массива на входное значение из другого массива и сохранить результат в 3-м массиве. Мы хотим повторить этот процесс много раз. Если бы все 3 размера массива были одинаковыми, мы могли бы реализовать это с помощью vectorize и даже не пришлось бы прибегать к более сложному guvectorize. Давайте сравним этот подход с вашим оригиналом, сосредоточив внимание на выполнении ядра CUDA. Вот рабочий пример, где t14.py - ваш исходный код, запускаемый с профилировщиком, а t15.py - его версия vectorize, признавая, что мы изменили размер вашего multBy массива, чтобы он соответствовал cv и discount:
$ nvprof --print-gpu-trace python t14.py
==4145== NVPROF is profiling process 4145, command: python t14.py
Function: discount factor cumVest duration (seconds):1.24354910851
==4145== Profiling application: python t14.py
==4145== Profiling result:
Start Duration Grid Size Block Size Regs* SSMem* DSMem* Size Throughput SrcMemType DstMemType Device Context Stream Name
312.36ms 1.2160us - - - - - 8B 6.2742MB/s Pageable Device Quadro K2000 (0 1 7 [CUDA memcpy HtoD]
312.81ms 27.392us - - - - - 156.25KB 5.4400GB/s Pageable Device Quadro K2000 (0 1 7 [CUDA memcpy HtoD]
313.52ms 5.8696ms - - - - - 15.259MB 2.5387GB/s Pageable Device Quadro K2000 (0 1 7 [CUDA memcpy HtoD]
319.74ms 1.0880us - - - - - 8B 7.0123MB/s Pageable Device Quadro K2000 (0 1 7 [CUDA memcpy HtoD]
319.93ms 896ns - - - - - 8B 8.5149MB/s Pageable Device Quadro K2000 (0 1 7 [CUDA memcpy HtoD]
321.40ms 1.22538s (1 1 1) (64 1 1) 63 0B 0B - - - - Quadro K2000 (0 1 7 cudapy::__main__::__gufunc_cVestDiscount$242(Array<__int64, int=1, A, mutable, aligned>, Array<double, int=3, A, mutable, aligned>, Array<double, int=4, A, mutable, aligned>, Array<__int64, int=1, A, mutable, aligned>, Array<__int64, int=1, A, mutable, aligned>, Array<double, int=4, A, mutable, aligned>) [37]
1.54678s 7.1816ms - - - - - 15.259MB 2.0749GB/s Device Pageable Quadro K2000 (0 1 7 [CUDA memcpy DtoH]
Regs: Number of registers used per CUDA thread. This number includes registers used internally by the CUDA driver and/or tools and can be more than what the compiler shows.
SSMem: Static shared memory allocated per CUDA block.
DSMem: Dynamic shared memory allocated per CUDA block.
SrcMemType: The type of source memory accessed by memory operation/copy
DstMemType: The type of destination memory accessed by memory operation/copy
$ cat t15.py
import numpy as np
from numba import guvectorize,vectorize
import time
from timeit import default_timer as timer
@vectorize(['float64(float64, float64)'], target='cuda')
def cVestDiscount (a, b):
return a * b
discount = np.float64(np.arange(2000000).reshape(100,4000,5))
multBy = np.full_like(discount, 1)
cv = np.empty_like(discount)
func_start = timer()
cv = cVestDiscount(multBy, discount)
timing=timer()-func_start
print("Function: discount factor cumVest duration (seconds):" + str(timing))
$ nvprof --print-gpu-trace python t15.py
==4167== NVPROF is profiling process 4167, command: python t15.py
Function: discount factor cumVest duration (seconds):0.37507891655
==4167== Profiling application: python t15.py
==4167== Profiling result:
Start Duration Grid Size Block Size Regs* SSMem* DSMem* Size Throughput SrcMemType DstMemType Device Context Stream Name
193.92ms 6.2729ms - - - - - 15.259MB 2.3755GB/s Pageable Device Quadro K2000 (0 1 7 [CUDA memcpy HtoD]
201.09ms 5.7101ms - - - - - 15.259MB 2.6096GB/s Pageable Device Quadro K2000 (0 1 7 [CUDA memcpy HtoD]
364.92ms 842.49us (15625 1 1) (128 1 1) 13 0B 0B - - - - Quadro K2000 (0 1 7 cudapy::__main__::__vectorized_cVestDiscount$242(Array<double, int=1, A, mutable, aligned>, Array<double, int=1, A, mutable, aligned>, Array<double, int=1, A, mutable, aligned>) [31]
365.77ms 7.1528ms - - - - - 15.259MB 2.0833GB/s Device Pageable Quadro K2000 (0 1 7 [CUDA memcpy DtoH]
Regs: Number of registers used per CUDA thread. This number includes registers used internally by the CUDA driver and/or tools and can be more than what the compiler shows.
SSMem: Static shared memory allocated per CUDA block.
DSMem: Dynamic shared memory allocated per CUDA block.
SrcMemType: The type of source memory accessed by memory operation/copy
DstMemType: The type of destination memory accessed by memory operation/copy
$
Мы видим, что ваше приложение сообщило о времени выполнения около 1,244 секунды, тогда как версия векторизации сообщает о времени выполнения около 0,375 секунды. Но в обоих этих числах есть накладные расходы на питона. Если мы посмотрим на сгенерированную длительность ядра CUDA в профилировщике, разница станет еще более очевидной. Мы видим, что исходное ядро занимало около 1,225 секунды, тогда как ядро векторизации выполняется примерно за 842 микросекунды (то есть менее 1 миллисекунды). Мы также отмечаем, что время ядра вычислений теперь намного, намного меньше, чем время, необходимое для передачи трех массивов в / из графического процессора (что занимает около 20 миллисекунд), и мы отмечаем, что измерения ядра теперь составляют 15625 блоков из 128 каждый поток для общего числа потоков / рабочих 2000000, точно совпадающих с общим числом операций умножения, которые должны быть выполнены, и значительно больше, чем ничтожные 64 потока (и, возможно, действительно только один поток) в действии с вашим исходным кодом.
Учитывая простоту описанного выше подхода vectorize, если вы действительно хотите сделать это поэлементное умножение, то вы можете подумать о простой репликации multBy, чтобы он по размеру соответствовал двум другим массивам, и сделано с этим.
Но остается вопрос: как обрабатывать входные массивы разного размера, как в исходной задаче? Для этого, я думаю, нам нужно перейти на guvectorize (или, как указывалось @talonmies, написать собственное ядро @cuda.jit, что, вероятно, лучший совет, несмотря на возможность того, что ни один из этих подходов не сможет преодолеть накладные расходы на передачу данных). в / из устройства, как уже упоминалось).
Чтобы справиться с этим с помощью guvectorize, нам нужно более тщательно подумать об уже упомянутой концепции «нарезки». Давайте перепишем ваше ядро guvectorize так, чтобы оно работало только с «срезом» общих данных, а затем позволили функции запуска guvectorize ускорить работу нескольких рабочих, по одному рабочему на срез.
В CUDA нам нравится иметь много рабочих; ты действительно не можешь иметь слишком много. Так что это повлияет на то, как мы «нарежем» наши массивы, чтобы дать возможность нескольким работникам действовать. Если бы мы нарезали вдоль 3-го (последнего, n) измерения, у нас было бы только 5 срезов для работы, то есть максимум 5 рабочих. Аналогично, если мы нарежем по первому или countRow измерению, у нас будет 100 срезов, то есть максимум 100 рабочих. В идеале, мы бы нарезали вдоль 2-го или countCol измерения. Однако для простоты я нарежу по первому или countRow размеру. Это явно неоптимально, но посмотрите ниже на работающий пример того, как вы можете подойти к проблеме нарезки по второму измерению. Нарезка по первому измерению означает, что мы удалим первый цикл for из нашего ядра guvectorize и позволим системе ufunc распараллеливать это измерение (основываясь на размерах передаваемых нами массивов). Код может выглядеть примерно так:
$ cat t16.py
import numpy as np
from numba import guvectorize
import time
from timeit import default_timer as timer
@guvectorize(['void(float64[:,:], float64[:,:], int64, int64, float64[:,:])'], '(m,o),(m,o),(),() -> (m,o)', target='cuda', nopython=True)
def cVestDiscount (multBy, discount, n, countCol, cv):
for ID in range(0,countCol):
for num in range(0,n):
cv[ID][num] = multBy[ID][num] * discount[ID][num]
multBy = np.float64(np.arange(20000).reshape(4000,5))
discount = np.float64(np.arange(2000000).reshape(100,4000,5))
n = np.int64(5)
countCol = np.int64(4000)
cv = np.zeros(shape=(100,4000,5), dtype=np.float64)
func_start = timer()
cv = cVestDiscount(multBy, discount, n, countCol, cv)
timing=timer()-func_start
print("Function: discount factor cumVest duration (seconds):" + str(timing))
$ nvprof --print-gpu-trace python t16.py
==4275== NVPROF is profiling process 4275, command: python t16.py
Function: discount factor cumVest duration (seconds):0.0670170783997
==4275== Profiling application: python t16.py
==4275== Profiling result:
Start Duration Grid Size Block Size Regs* SSMem* DSMem* Size Throughput SrcMemType DstMemType Device Context Stream Name
307.05ms 27.392us - - - - - 156.25KB 5.4400GB/s Pageable Device Quadro K2000 (0 1 7 [CUDA memcpy HtoD]
307.79ms 5.9293ms - - - - - 15.259MB 2.5131GB/s Pageable Device Quadro K2000 (0 1 7 [CUDA memcpy HtoD]
314.34ms 1.3440us - - - - - 8B 5.6766MB/s Pageable Device Quadro K2000 (0 1 7 [CUDA memcpy HtoD]
314.54ms 896ns - - - - - 8B 8.5149MB/s Pageable Device Quadro K2000 (0 1 7 [CUDA memcpy HtoD]
317.27ms 47.398ms (2 1 1) (64 1 1) 63 0B 0B - - - - Quadro K2000 (0 1 7 cudapy::__main__::__gufunc_cVestDiscount$242(Array<double, int=3, A, mutable, aligned>, Array<double, int=3, A, mutable, aligned>, Array<__int64, int=1, A, mutable, aligned>, Array<__int64, int=1, A, mutable, aligned>, Array<double, int=3, A, mutable, aligned>) [35]
364.67ms 7.3799ms - - - - - 15.259MB 2.0192GB/s Device Pageable Quadro K2000 (0 1 7 [CUDA memcpy DtoH]
Regs: Number of registers used per CUDA thread. This number includes registers used internally by the CUDA driver and/or tools and can be more than what the compiler shows.
SSMem: Static shared memory allocated per CUDA block.
DSMem: Dynamic shared memory allocated per CUDA block.
SrcMemType: The type of source memory accessed by memory operation/copy
DstMemType: The type of destination memory accessed by memory operation/copy
$
Замечания:
Изменения кода были связаны с удалением параметра countCol, удалением первого цикла for из ядра guvectorize и внесением соответствующих изменений в сигнатуру функции, чтобы отразить это. Мы также изменили наши трехмерные функции в сигнатуре на двумерные. В конце концов, мы берем двумерный «кусочек» трехмерных данных и позволяем каждому работнику работать со срезом.
Размеры ядра, о которых сообщает профилировщик, теперь составляют 2 блока вместо 1. Это имеет смысл, потому что в исходной реализации действительно был представлен только 1 «срез», и поэтому требовался 1 рабочий, и, следовательно, 1 поток (но numba раскрутил 1 поток из 64 потоков). В этой реализации есть 100 срезов, и numba решила увеличить количество потоков до 64 рабочих / потоков, чтобы обеспечить необходимые 100 рабочих / потоков.
Производительность ядра, сообщаемая профилировщиком в 47,4 мс, теперь находится где-то между оригинальной (~ 1,224 с) и массивно параллельной версией vectorize (при ~ 0,001 с). Таким образом, переход от 1 до 100 рабочих значительно ускорился, но возможен прирост производительности. Если вы поймете, как нарезать размер countCol, вы, вероятно, сможете приблизиться к версии vectorize с точки зрения производительности (см. Ниже). Обратите внимание, что разницы между тем, где мы находимся здесь (~ 47 мс), и версией векторизации (~ 1 мс) более чем достаточно, чтобы компенсировать дополнительную стоимость переноса (~ 5 мс или менее) передачи немного большей матрицы multBy. к устройству, чтобы облегчить простоту vectorize.
Некоторые дополнительные комментарии по времени синхронизации Python: я считаю, что точное поведение Python по компиляции необходимых ядер для исходных, векторизованных и guvectorize улучшенных версий отличается. Если мы изменим код t15.py для запуска прогона «разогрева», то, по крайней мере, время Python будет соответствовать по тренду общему времени стены и времени только для ядра:
$ cat t15.py
import numpy as np
from numba import guvectorize,vectorize
import time
from timeit import default_timer as timer
@vectorize(['float64(float64, float64)'], target='cuda')
def cVestDiscount (a, b):
return a * b
multBy = np.float64(np.arange(20000).reshape(4000,5))
discount = np.float64(np.arange(2000000).reshape(100,4000,5))
multBy = np.full_like(discount, 1)
cv = np.empty_like(discount)
#warm-up run
cv = cVestDiscount(multBy, discount)
func_start = timer()
cv = cVestDiscount(multBy, discount)
timing=timer()-func_start
print("Function: discount factor cumVest duration (seconds):" + str(timing))
[bob@cluster2 python]$ time python t14.py
Function: discount factor cumVest duration (seconds):1.24376320839
real 0m2.522s
user 0m1.572s
sys 0m0.809s
$ time python t15.py
Function: discount factor cumVest duration (seconds):0.0228319168091
real 0m1.050s
user 0m0.473s
sys 0m0.445s
$ time python t16.py
Function: discount factor cumVest duration (seconds):0.0665760040283
real 0m1.252s
user 0m0.680s
sys 0m0.441s
$
Теперь, эффективно отвечая на вопрос в комментариях: «Как бы я переделал проблему, чтобы разрезать ее вдоль 4000 (countCol, или« среднего ») измерения?»
Мы можем руководствоваться тем, что работало, чтобы разрезать вдоль первого измерения. Один из возможных подходов состоит в том, чтобы изменить форму массивов таким образом, чтобы измерение 4000 было первым измерением, а затем удалить его, аналогично тому, что мы делали в предыдущей обработке guvectorize. Вот рабочий пример:
$ cat t17.py
import numpy as np
from numba import guvectorize
import time
from timeit import default_timer as timer
@guvectorize(['void(int64, float64[:], float64[:,:], int64, float64[:,:])'], '(),(o),(m,o),() -> (m,o)', target='cuda', nopython=True)
def cVestDiscount (countCol, multBy, discount, n, cv):
for ID in range(0,countCol):
for num in range(0,n):
cv[ID][num] = multBy[num] * discount[ID][num]
countRow = np.int64(100)
multBy = np.float64(np.arange(20000).reshape(4000,5))
discount = np.float64(np.arange(2000000).reshape(4000,100,5))
n = np.int64(5)
countCol = np.int64(4000)
cv = np.zeros(shape=(4000,100,5), dtype=np.float64)
func_start = timer()
cv = cVestDiscount(countRow, multBy, discount, n, cv)
timing=timer()-func_start
print("Function: discount factor cumVest duration (seconds):" + str(timing))
[bob@cluster2 python]$ python t17.py
Function: discount factor cumVest duration (seconds):0.0266749858856
$ nvprof --print-gpu-trace python t17.py
==8544== NVPROF is profiling process 8544, command: python t17.py
Function: discount factor cumVest duration (seconds):0.0268459320068
==8544== Profiling application: python t17.py
==8544== Profiling result:
Start Duration Grid Size Block Size Regs* SSMem* DSMem* Size Throughput SrcMemType DstMemType Device Context Stream Name
304.92ms 1.1840us - - - - - 8B 6.4437MB/s Pageable Device Quadro K2000 (0 1 7 [CUDA memcpy HtoD]
305.36ms 27.392us - - - - - 156.25KB 5.4400GB/s Pageable Device Quadro K2000 (0 1 7 [CUDA memcpy HtoD]
306.08ms 6.0208ms - - - - - 15.259MB 2.4749GB/s Pageable Device Quadro K2000 (0 1 7 [CUDA memcpy HtoD]
312.44ms 1.0880us - - - - - 8B 7.0123MB/s Pageable Device Quadro K2000 (0 1 7 [CUDA memcpy HtoD]
313.59ms 8.9961ms (63 1 1) (64 1 1) 63 0B 0B - - - - Quadro K2000 (0 1 7 cudapy::__main__::__gufunc_cVestDiscount$242(Array<__int64, int=1, A, mutable, aligned>, Array<double, int=2, A, mutable, aligned>, Array<double, int=3, A, mutable, aligned>, Array<__int64, int=1, A, mutable, aligned>, Array<double, int=3, A, mutable, aligned>) [35]
322.59ms 7.2772ms - - - - - 15.259MB 2.0476GB/s Device Pageable Quadro K2000 (0 1 7 [CUDA memcpy DtoH]
Regs: Number of registers used per CUDA thread. This number includes registers used internally by the CUDA driver and/or tools and can be more than what the compiler shows.
SSMem: Static shared memory allocated per CUDA block.
DSMem: Dynamic shared memory allocated per CUDA block.
SrcMemType: The type of source memory accessed by memory operation/copy
DstMemType: The type of destination memory accessed by memory operation/copy
$
Несколько предсказуемо, мы наблюдаем, что время выполнения сократилось с ~ 47 мс, когда мы нарезали на 100 рабочих, до ~ 9 мс, когда мы нарезали на 4000 рабочих. Точно так же мы наблюдаем, что numba выбирает раскрутить 63 блока по 64 потока каждый на общую сумму 4032 потока, чтобы обработать 4000 рабочих, необходимых для этой «нарезки».
По-прежнему не так быстро, как ядро ~ 1ms vectorize (в котором имеется много доступных параллельных «кусочков» для рабочих), но немного быстрее, чем ядро ~ 1.2s, предложенное в первоначальном вопросе. И общее время ожидания кода Python примерно в 2 раза быстрее, даже при всех накладных расходах Python.
В качестве заключительного замечания давайте вернемся к заявлению, которое я сделал ранее (и похоже на утверждения, сделанные в комментарии и в другом ответе):
"Я сомневаюсь, что было бы возможно превысить производительность хорошо написанного основного кода (например, используя какой-либо метод распараллеливания, такой как guvectorize), чтобы сделать то же самое."
Теперь у нас есть удобные тестовые случаи в t16.py или t17.py, с которыми мы можем работать, чтобы проверить это. Для простоты я выберу t16.py. Мы можем «преобразовать это обратно в код процессора», просто удалив целевое обозначение из guvectorize ufunc:
$ cat t16a.py
import numpy as np
from numba import guvectorize
import time
from timeit import default_timer as timer
@guvectorize(['void(float64[:,:], float64[:,:], int64, int64, float64[:,:])'], '(m,o),(m,o),(),() -> (m,o)')
def cVestDiscount (multBy, discount, n, countCol, cv):
for ID in range(0,countCol):
for num in range(0,n):
cv[ID][num] = multBy[ID][num] * discount[ID][num]
multBy = np.float64(np.arange(20000).reshape(4000,5))
discount = np.float64(np.arange(2000000).reshape(100,4000,5))
n = np.int64(5)
countCol = np.int64(4000)
cv = np.zeros(shape=(100,4000,5), dtype=np.float64)
func_start = timer()
cv = cVestDiscount(multBy, discount, n, countCol, cv)
timing=timer()-func_start
print("Function: discount factor cumVest duration (seconds):" + str(timing))
$ time python t16a.py
Function: discount factor cumVest duration (seconds):0.00657796859741
real 0m0.528s
user 0m0.474s
sys 0m0.047s
$
Итак, мы видим, что эта версия только для ЦПУ выполняет функцию примерно за 6 миллисекунд, и у нее нет никаких «накладных расходов» GPU, таких как инициализация CUDA и копирование данных в / из GPU. Общее время простоя также является нашим лучшим измерением: около 0,5 с по сравнению с 1,0 с для нашего лучшего графического процессора. Так что эта конкретная проблема, из-за ее низкой арифметической интенсивности на байт передачи данных, вероятно, не очень подходит для вычислений на GPU.