Этот ответ полностью верен. Но я также хотел бы указать на другие случаи, когда вы могли бы подумать, что безопасность не имеет значения. Например. команда, которую вы запускаете, жестко запрограммирована или у вас есть 100% контроль или доверие к тому, что ему предоставляется. Даже в этом случае os.system()
неверно . Фактически: os.system()
или у вас могут быть проблемы с безопасностью.
С вашим кодом было множество проблем. Я не могу охватить все из них, поэтому, пожалуйста, сравните мой код с вашим.
методы numpy array (например, np.sum()
) нельзя использовать в ядрах numba CUDA .
Скалярные величины, передаваемые ядру 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))
Ваше ядро было излишне сложным. Вы в основном выполняете суммы столбцов матрицы, которая была переставлена в соответствии с шаблоном индекса, умноженного на массив коэффициентов. Мы можем выполнить это с помощью одного цикла на поток.
Для приведенной выше реализации нам не нужна двумерная сетка потоков.
У вас были проблемы с отступами в опубликованном вами коде.
Для демонстрации я уменьшил размер вашего набора данных с 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). Я бы не стал читать слишком много в этом сравнении, поскольку код процессора почти наверняка также неоптимален, и это все еще относительно небольшой объем работы на каждую точку выходных данных, чтобы ускорение графического процессора было интересным.