Комплексное сокращение числа с помощью numba cuda

Этот ответ полностью верен. Но я также хотел бы указать на другие случаи, когда вы могли бы подумать, что безопасность не имеет значения. Например. команда, которую вы запускаете, жестко запрограммирована или у вас есть 100% контроль или доверие к тому, что ему предоставляется. Даже в этом случае os.system() неверно . Фактически:

  • Вы должны полагаться на внешние инструменты, которые могут отсутствовать или, что еще хуже, у вас может быть команда a с этим именем, но она не " делайте то, что вы ожидаете от этого. Может быть, потому, что у него другая версия или, возможно, потому, что это другая реализация этой команды (например, GNUtar! = BSDtar). Зависимые зависимости python будут намного проще и надежнее.
  • С ошибками справиться сложнее. У вас есть код возврата, который не всегда достаточно, чтобы понять, что происходит. И я надеюсь, что ваше решение этой проблемы заключается не в анализе вывода команды.
  • Переменные среды могут изменить способ работы программы неожиданно. Многие программы используют переменные среды в качестве альтернативы командной строке или параметрам конфигурации. Если ваш скрипт python полагается на определенное поведение из команды, неожиданная переменная в среде пользователя может ее сломать.
  • Если в какой-то момент в будущем вам нужно будет немного настроить поведение вашего скрипта вам нужно будет переписать его с нуля без os.system() или у вас могут быть проблемы с безопасностью.

0
задан Ali Jooya 19 January 2019 в 00:16
поделиться

1 ответ

С вашим кодом было множество проблем. Я не могу охватить все из них, поэтому, пожалуйста, сравните мой код с вашим.

  1. методы numpy array (например, np.sum()) нельзя использовать в ядрах numba CUDA .

  2. Скалярные величины, передаваемые ядру numba cuda (как npart, npts), не нуждаются и не должны обрабатываться массивом, как .to_device(). Просто используйте их как есть. Это причина ошибки Python, которую вы показываете:

    Invalid use of Function(<built-in function lt>) with argument(s) of type(s): (int32, array(int64, 1d, A))
    
  3. Ваше ядро ​​было излишне сложным. Вы в основном выполняете суммы столбцов матрицы, которая была переставлена ​​в соответствии с шаблоном индекса, умноженного на массив коэффициентов. Мы можем выполнить это с помощью одного цикла на поток.

  4. Для приведенной выше реализации нам не нужна двумерная сетка потоков.

  5. У вас были проблемы с отступами в опубликованном вами коде.

Для демонстрации я уменьшил размер вашего набора данных с 1000 столбцов до 15 столбцов. Вот пример, который обращается к вышеупомянутым пунктам:

$ cat t31.py
import numba as nb
import numpy as np
from numba import cuda
import time
import math

def random_choice_noreplace(m,n, axis=-1):
 # m, n are the number of rows, cols of output
 return np.random.rand(m,n).argsort(axis=axis)

@cuda.jit
def cuda_kernel (npart, npts, d_data, d_data_index, d_coef, d_datasum):
 col = cuda.grid(1)
 if col < npts:
     temp = 0
     for i in range(npart):
         temp += d_data[d_data_index[i, col]] * d_coef[i, col]
     d_datasum[col] = temp


def calculate_cuda (data, data_index, coef):

 npart, npts = data_index.shape
 # arrays to copy to GPU memory =====================================
 d_data = cuda.to_device(data)
 d_data_imag = cuda.to_device(data_imag)
 d_data_index = cuda.to_device(data_index)
 d_coef = cuda.to_device(coef)

 d_datasum = cuda.device_array(npts, np.complex64)

 threadsperblock = 64
 blockspergrid = int(math.ceil(npts / threadsperblock))+1
 cuda_kernel[blockspergrid, threadsperblock](npart, npts, d_data, d_data_index, d_coef, d_datasum)
 # Copy data from GPU to CPU ========================================
 final_data_sum = d_datasum.copy_to_host()
 return final_data_sum



def calculate_python (data, data_index, coef):
 npart, npts = data_index.shape
 data_sum = np.zeros(npts, dtype=np.complex64)
 tmp = np.zeros(npts, dtype=np.complex64)
 print(" Calling python function...")
 for i in range(npart):
  tmp[:] = data[data_index[i]]
  data_sum += tmp * coef[i]
 return data_sum

if __name__ == '__main__':

 rows = 31
 cols = 15
 data_size = rows * cols

 data_real = np.random.randn(data_size).astype(np.float32)
 data_imag = np.random.randn(data_size).astype(np.float32)

 data = data_real + 1j * data_imag
 coef = np.random.randn(rows, cols)
 data_index = random_choice_noreplace(rows, cols)

 start_time = time.time()
 gpu_data_sum_python = calculate_python (data, data_index, coef)
 python_time = time.time() - start_time #print("gpu c : ", c_gpu)
 print("---- %s seconds for python ----:" % (python_time))
 print(gpu_data_sum_python)

 start_time = time.time()
 gpu_data_sum = calculate_cuda (data, data_index, coef)
 gpu_time = time.time() - start_time
 print("---- %s seconds for gpu ----:" % (gpu_time))
 print(gpu_data_sum)
$ python t31.py
 Calling python function...
---- 0.000281095504761 seconds for python ----:
[-1.10292518+0.90700358j  2.67771578+2.47935939j -5.22553015-2.22675705j
 -3.55810285+2.39755774j  4.11441088-3.89396238j -2.70894790-0.75690132j
  3.24859619+0.65993834j  1.05531025+2.3959775j  -4.27368307+1.6297332j
  0.17896785-7.0437355j  -6.31506491+6.22674656j -1.85534143-6.08459902j
  0.40037563+6.33309507j -1.71916604-0.55055946j  0.49263301+1.08690035j]
---- 0.593510866165 seconds for gpu ----:
[-1.10292506+0.9070037j   2.67771506+2.47935939j -5.22553062-2.22675681j
 -3.55810285+2.39755821j  4.11441135-3.89396238j -2.70894790-0.75690138j
  3.24859619+0.65993822j  1.05531013+2.39597774j -4.27368307+1.62973344j
  0.17896791-7.0437355j  -6.31506491+6.22674656j -1.85534155-6.08459902j
  0.40037528+6.33309603j -1.71916604-0.55055946j  0.49263287+1.08690035j]
$

Обратите внимание, что существуют небольшие различия в численном выражении между результатами вычислений хоста и устройства, начиная с 6-го знака после запятой, в некоторых случаях. Я связываю это с возможными различиями порядка вычислений между кодом хоста и устройства в сочетании с пределами float32 (или complex64) типа данных numpy.

Поскольку у вас есть встроенный тайминг в ваш код, вы можете быть заинтересованы в производительности. Для numba python я рекомендую типичную практику бенчмаркинга, которая заключается не в измерении первого прогона, а в измерении второго прогона. Это позволяет избежать однократных накладных расходов при измерении. Кроме того, мы хотели бы выбрать гораздо больший размер набора данных, чем 15 столбцов, чтобы дать графическому процессору достаточно большой объем работы для амортизации различных затрат. С этими модификациями, вот пример, показывающий, что версия GPU в этом коде может быть быстрее, чем версия CPU в этом коде:

$ cat t31.py
import numba as nb
import numpy as np
from numba import cuda
import time
import math

def random_choice_noreplace(m,n, axis=-1):
 # m, n are the number of rows, cols of output
 return np.random.rand(m,n).argsort(axis=axis)

@cuda.jit
def cuda_kernel (npart, npts, d_data, d_data_index, d_coef, d_datasum):
 col = cuda.grid(1)
 if col < npts:
     temp = 0
     for i in range(npart):
         temp += d_data[d_data_index[i, col]] * d_coef[i, col]
     d_datasum[col] = temp


def calculate_cuda (data, data_index, coef):

 npart, npts = data_index.shape
 # arrays to copy to GPU memory =====================================
 d_data = cuda.to_device(data)
 d_data_imag = cuda.to_device(data_imag)
 d_data_index = cuda.to_device(data_index)
 d_coef = cuda.to_device(coef)

 d_datasum = cuda.device_array(npts, np.complex64)

 threadsperblock = 64
 blockspergrid = int(math.ceil(npts / threadsperblock))+1
 cuda_kernel[blockspergrid, threadsperblock](npart, npts, d_data, d_data_index, d_coef, d_datasum)
 # Copy data from GPU to CPU ========================================
 final_data_sum = d_datasum.copy_to_host()
 return final_data_sum



def calculate_python (data, data_index, coef):
 npart, npts = data_index.shape
 data_sum = np.zeros(npts, dtype=np.complex64)
 tmp = np.zeros(npts, dtype=np.complex64)
 print(" Calling python function...")
 for i in range(npart):
  tmp[:] = data[data_index[i]]
  data_sum += tmp * coef[i]
 return data_sum

if __name__ == '__main__':

 rows = 31
 cols = 1000000
 data_size = rows * cols

 data_real = np.random.randn(data_size).astype(np.float32)
 data_imag = np.random.randn(data_size).astype(np.float32)

 data = data_real + 1j * data_imag
 coef = np.random.randn(rows, cols)
 data_index = random_choice_noreplace(rows, cols)

 gpu_data_sum_python = calculate_python (data, data_index, coef)
 start_time = time.time()
 gpu_data_sum_python = calculate_python (data, data_index, coef)
 python_time = time.time() - start_time #print("gpu c : ", c_gpu)
 print("---- %s seconds for python ----:" % (python_time))
 print(gpu_data_sum_python)

 gpu_data_sum = calculate_cuda (data, data_index, coef)
 start_time = time.time()
 gpu_data_sum = calculate_cuda (data, data_index, coef)
 gpu_time = time.time() - start_time
 print("---- %s seconds for gpu ----:" % (gpu_time))
 print(gpu_data_sum)
$ python t31.py
 Calling python function...
 Calling python function...
---- 0.806931018829 seconds for python ----:
[  6.56164026-7.95271683j  -7.42586899+3.68758106j   3.48999476+3.10376763j
 ...,  13.12746525+4.61855698j   0.08796659+0.9710322j
  -6.54224586+4.89485168j]
---- 0.417661905289 seconds for gpu ----:
[  6.56164074-7.95271683j  -7.42586851+3.68758035j   3.48999500+3.10376763j
 ...,  13.12746525+4.61855745j   0.08796643+0.97103256j
  -6.54224634+4.89485121j]
$

С этими модификациями код GPU кажется примерно в 2 раза быстрее, чем код процессора.

Это было измерено на CUDA 9.2, Fedora 27, Quadro K2000 (относительно небольшой, медленный GPU). Я бы не стал читать слишком много в этом сравнении, поскольку код процессора почти наверняка также неоптимален, и это все еще относительно небольшой объем работы на каждую точку выходных данных, чтобы ускорение графического процессора было интересным.

0
ответ дан Robert Crovella 19 January 2019 в 00:16
поделиться
Другие вопросы по тегам:

Похожие вопросы: