Преобразование функции в NumbaPro CUDA

Question

Я сравниваю несколько Python модулей/расширений или методы для достижения следующих действий:

import numpy as np 

def fdtd(input_grid, steps): 
    grid = input_grid.copy() 
    old_grid = np.zeros_like(input_grid) 
    previous_grid = np.zeros_like(input_grid) 

    l_x = grid.shape[0] 
    l_y = grid.shape[1] 

    for i in range(steps): 
     np.copyto(previous_grid, old_grid) 
     np.copyto(old_grid, grid) 

     for x in range(l_x): 
      for y in range(l_y): 
       grid[x,y] = 0.0 
       if 0 < x+1 < l_x: 
        grid[x,y] += old_grid[x+1,y] 
       if 0 < x-1 < l_x: 
        grid[x,y] += old_grid[x-1,y] 
       if 0 < y+1 < l_y: 
        grid[x,y] += old_grid[x,y+1] 
       if 0 < y-1 < l_y: 
        grid[x,y] += old_grid[x,y-1] 

       grid[x,y] /= 2.0 
       grid[x,y] -= previous_grid[x,y] 

    return grid 

Эта функция является очень базовая реализация метода временной области (FDTD) конечных разностей. Я реализовал эту функцию несколькими способами:

• с более NumPy Подпрограммы
• в Cython
• использованием Numba (авто) JIT.

Теперь я хотел бы сравнить производительность с NumbaPro CUDA.

Это первый раз, когда я пишу код для CUDA, и я придумал код ниже.

from numbapro import cuda, float32, int16 
import numpy as np 

@cuda.jit(argtypes=(float32[:,:], float32[:,:], float32[:,:], int16, int16, int16)) 
def kernel(grid, old_grid, previous_grid, steps, l_x, l_y): 

    x,y = cuda.grid(2) 

    for i in range(steps): 
     previous_grid[x,y] = old_grid[x,y] 
     old_grid[x,y] = grid[x,y] 

    for i in range(steps): 

     grid[x,y] = 0.0 

     if 0 < x+1 and x+1 < l_x: 
      grid[x,y] += old_grid[x+1,y] 
     if 0 < x-1 and x-1 < l_x: 
      grid[x,y] += old_grid[x-1,y] 
     if 0 < y+1 and y+1 < l_x: 
      grid[x,y] += old_grid[x,y+1] 
     if 0 < y-1 and y-1 < l_x: 
      grid[x,y] += old_grid[x,y-1] 

     grid[x,y] /= 2.0 
     grid[x,y] -= previous_grid[x,y] 


def fdtd(input_grid, steps): 

    grid = cuda.to_device(input_grid) 
    old_grid = cuda.to_device(np.zeros_like(input_grid)) 
    previous_grid = cuda.to_device(np.zeros_like(input_grid)) 

    l_x = input_grid.shape[0] 
    l_y = input_grid.shape[1] 

    kernel[(16,16),(32,8)](grid, old_grid, previous_grid, steps, l_x, l_y) 

    return grid.copy_to_host() 

К сожалению, я получаю следующее сообщение об ошибке:

File ".../fdtd_numbapro.py", line 98, in fdtd 
    return grid.copy_to_host() 
    File "/opt/anaconda1anaconda2anaconda3/lib/python2.7/site-packages/numbapro/cudadrv/devicearray.py", line 142, in copy_to_host 
    File "/opt/anaconda1anaconda2anaconda3/lib/python2.7/site-packages/numbapro/cudadrv/driver.py", line 1702, in device_to_host 
    File "/opt/anaconda1anaconda2anaconda3/lib/python2.7/site-packages/numbapro/cudadrv/driver.py", line 772, in check_error 
numbapro.cudadrv.error.CudaDriverError: CUDA_ERROR_LAUNCH_FAILED 
Failed to copy memory D->H 

Я использовал grid.to_host(), так и будет работать ни. CUDA определенно работает с NumbaPro в этой системе.

Answer 1

Проблема разрешена пользователем. Я перекрестные ссылками обсуждения списка рассылки Anaconda для этой проблемы: https://groups.google.com/a/continuum.io/forum/#!searchin/anaconda/fdtd/anaconda/VgiN4h37UrA/18tAc60EIkcJ

Answer 2

Я сделал некоторые незначительные изменения в исходный код, чтобы он работает в Parakeet:

1) Сплит сложных сравнений, такие как «0 < x-1 < l_x "в" 0 < x-1 и x-1 < l_x ".

2) Заменено np.copyto с явным индексированным присваиванием (previous_grid [:,:] = old_grid).

После этого я сравниваю Попугай автономной работы для движков C, OpenMP и CUDA против первоначального времени Python и autojit Numba в на 1000x1000 сетке с шагом = 20.

Parakeet (backend = c) cold: fdtd : 0.5590s 
Parakeet (backend = c) warm: fdtd : 0.1260s 

Parakeet (backend = openmp) cold: fdtd : 0.4317s 
Parakeet (backend = openmp) warm: fdtd : 0.1693s 

Parakeet (backend = cuda) cold: fdtd : 2.6357s 
Parakeet (backend = cuda) warm: fdtd : 0.2455s 

Numba (autojit) cold: 672.3666s 
Numba (autojit) warm: 657.8858s 

Python: 203.3907s 

Поскольку мало легкодоступными параллелизм в вашем коде, параллельные бэкэнды действительно хуже, чем последовательные. Во многом это связано с различием в том, что оптимизация циклов запускается Parakeet для каждого бэкэнд вместе с некоторыми дополнительными накладными расходами, связанными с передачами CUDA-памяти и запуском групп потоков OpenMP. Я не уверен, почему Autojit Numba здесь настолько медленный, я уверен, что это будет быстрее с аннотациями типа или с помощью NumbaPro.