Давайте установим базовый уровень производительности.Добавив определение (1.0) для earthdiam_nm и запустив ваш код под nvprof, мы получим:
$ nvprof python t38.py
1000000 total evaluations in 0.581 seconds
(...)
==1973== Profiling result:
Type Time(%) Time Calls Avg Min Max Name
GPU activities: 55.58% 11.418ms 4 2.8544ms 2.6974ms 3.3044ms [CUDA memcpy HtoD]
28.59% 5.8727ms 1 5.8727ms 5.8727ms 5.8727ms cudapy::__main__::__vectorized_a_cuda$242(Array<double, int=1, A, mutable, aligned>, Array<double, int=1, A, mutable, aligned>, Array<double, int=1, A, mutable, aligned>, Array<double, int=1, A, mutable, aligned>, Array<double, int=1, A, mutable, aligned>)
15.83% 3.2521ms 1 3.2521ms 3.2521ms 3.2521ms [CUDA memcpy DtoH]
(...)
Так что на моей конкретной установке «ядро» само работает за ~ 5,8 мс (маленький, медленный) QuadroK2000 GPU, и время копирования данных составляет в общей сложности 11,4 мс для 4 копий с хоста на устройство и 3,2 мс для передачи результатов обратно на хост.Основное внимание уделяется 4 копиям с хоста на устройство.
Давайте сначала пойдем за низко висящими фруктами.Эта строка кода:
X2,Y2 = np.float64(np.array([30]*nx*ny)),np.float64(np.array([101]*nx*ny))
на самом деле ничего не делает, кроме передачи значений 30 и 101 каждому «работнику».Я использую «работник» здесь, чтобы обратиться к идее конкретного скалярного вычисления в процессе numba «трансляции» функции vectorize по большому набору данных.Процесс векторизации / широковещания numba не требует, чтобы каждый вход был набором данных одинакового размера, просто чтобы передаваемые данные были «широковещательными».Таким образом, можно создать vectorize ufunc, который работает, например, с массивом и скаляром.Это означает, что каждый работник будет использовать свой элемент массива плюс скаляр для выполнения своих вычислений.
Поэтому низко висящий плод состоит в том, чтобы просто удалить эти два массива и передать значения (30, 101) какскаляры к уфунку a_cuda.В то время как мы собираемся использовать «низко висящий фрукт», давайте включим ваше вычисление arctan2 (вместо math.atan2) и ваше окончательное масштабирование на earthdiam_nm в код векторизации, поэтому нам не нужно делать это на хосте.в python / numpy:
$ cat t39.py
import math
from numba import vectorize, float64
import numpy as np
from time import time
earthdiam_nm = 1.0
@vectorize([float64(float64,float64,float64,float64,float64)],target='cuda')
def a_cuda(lat1, lon1, lat2, lon2, s):
a = (math.sin(0.008726645 * (lat2 - lat1))**2) + \
math.cos(0.01745329*(lat1)) * math.cos(0.01745329*(lat2)) * (math.sin(0.008726645 * (lon2 - lon1))**2)
return math.atan2(a, 1-a)*s
def LLA_distance_numba_cuda(lat1, lon1, lat2, lon2):
return a_cuda(np.ascontiguousarray(lat1), np.ascontiguousarray(lon1),
np.ascontiguousarray(lat2), np.ascontiguousarray(lon2), earthdiam_nm)
# generate a mesh of one million evaluation points
nx, ny = 1000,1000
xv, yv = np.meshgrid(np.linspace(29, 31, nx), np.linspace(99, 101, ny))
X, Y = np.float64(xv.reshape(1,nx*ny).flatten()), np.float64(yv.reshape(1,nx*ny).flatten())
# X2,Y2 = np.float64(np.array([30]*nx*ny)),np.float64(np.array([101]*nx*ny))
start = time()
Z=LLA_distance_numba_cuda(X,Y,30.0,101.0)
print('{:d} total evaluations in {:.3f} seconds'.format(nx*ny,time()-start))
#print(Z)
$ nvprof python t39.py
==2387== NVPROF is profiling process 2387, command: python t39.py
1000000 total evaluations in 0.401 seconds
==2387== Profiling application: python t39.py
==2387== Profiling result:
Type Time(%) Time Calls Avg Min Max Name
GPU activities: 48.12% 8.4679ms 1 8.4679ms 8.4679ms 8.4679ms cudapy::__main__::__vectorized_a_cuda$242(Array<double, int=1, A, mutable, aligned>, Array<double, int=1, A, mutable, aligned>, Array<double, int=1, A, mutable, aligned>, Array<double, int=1, A, mutable, aligned>, Array<double, int=1, A, mutable, aligned>, Array<double, int=1, A, mutable, aligned>)
33.97% 5.9774ms 5 1.1955ms 864ns 3.2535ms [CUDA memcpy HtoD]
17.91% 3.1511ms 4 787.77us 1.1840us 3.1459ms [CUDA memcpy DtoH]
(snip)
Теперь мы видим, что операции копирования HtoD были сокращены с итого 11,4 мс до 5,6 мс.Ядро выросло с ~ 5,8 мс до ~ 8,5 мс, потому что мы выполняем больше работы в ядре, но сообщаемое Python время выполнения функции сократилось с ~ 0,58 с до ~ 0,4 с.
Можем ли мысделать лучше?
Мы можем, но для того, чтобы сделать это (я считаю), нам нужно будет использовать другой метод numba cuda.Метод vectorize удобен для скалярных поэлементных операций, но он не может знать, где в общем наборе данных выполняется операция.Нам нужна эта информация, и мы можем получить ее в коде CUDA, но для этого нам потребуется переключиться на декоратор @cuda.jit.
Следующий код преобразует предыдущую функцию vectorize a_cuda вфункция устройства @cuda.jit (по существу без других изменений), а затем мы создаем ядро CUDA, которое выполняет генерацию сетки в соответствии с предоставленными скалярными параметрами и вычисляет результат:
$ cat t40.py
import math
from numba import vectorize, float64, cuda
import numpy as np
from time import time
earthdiam_nm = 1.0
@cuda.jit(device='true')
def a_cuda(lat1, lon1, lat2, lon2, s):
a = (math.sin(0.008726645 * (lat2 - lat1))**2) + \
math.cos(0.01745329*(lat1)) * math.cos(0.01745329*(lat2)) * (math.sin(0.008726645 * (lon2 - lon1))**2)
return math.atan2(a, 1-a)*s
@cuda.jit
def LLA_distance_numba_cuda(lat2, lon2, xb, xe, yb, ye, s, nx, ny, out):
x,y = cuda.grid(2)
if x < nx and y < ny:
lat1 = (((xe-xb) * x)/(nx-1)) + xb # mesh generation
lon1 = (((ye-yb) * y)/(ny-1)) + yb # mesh generation
out[y][x] = a_cuda(lat1, lon1, lat2, lon2, s)
nx, ny = 1000,1000
Z = cuda.device_array((nx,ny), dtype=np.float64)
threads = (32,32)
blocks = (32,32)
start = time()
LLA_distance_numba_cuda[blocks,threads](30.0,101.0, 29.0, 31.0, 99.0, 101.0, earthdiam_nm, nx, ny, Z)
Zh = Z.copy_to_host()
print('{:d} total evaluations in {:.3f} seconds'.format(nx*ny,time()-start))
#print(Zh)
$ nvprof python t40.py
==2855== NVPROF is profiling process 2855, command: python t40.py
1000000 total evaluations in 0.294 seconds
==2855== Profiling application: python t40.py
==2855== Profiling result:
Type Time(%) Time Calls Avg Min Max Name
GPU activities: 75.60% 10.364ms 1 10.364ms 10.364ms 10.364ms cudapy::__main__::LLA_distance_numba_cuda$241(double, double, double, double, double, double, double, __int64, __int64, Array<double, int=2, A, mutable, aligned>)
24.40% 3.3446ms 1 3.3446ms 3.3446ms 3.3446ms [CUDA memcpy DtoH]
(...)
Теперь мы видимчто:
- Время выполнения ядра еще больше - около 10 мс (потому что мы делаем генерацию сетки)
- Нет явного копирования данных с хоста на устройство
- Общее время выполнения функции уменьшено с ~ 0,4 с до ~ 0,3 с