#performance #cuda #numba
Вопрос:
Есть несколько тем, посвященных схожим темам, но, к сожалению, они кажутся мне слишком сложными, поэтому я хотел бы задать аналогичный вопрос, надеясь, что кто-нибудь специально посмотрит на мой код, чтобы сказать мне, не ошибся ли я в чем-то.
Я изучаю numba cuda прямо сейчас, начиная с простых примеров, которые можно найти в сети. Я начал с этого урока здесь:
https://github.com/ContinuumIO/gtc2017-numba/blob/master/4 — Writing CUDA Kernels.ipynb
который показывает, как параллельно выполнять добавление массивов. Конфигурация системы, которую они использовали для оценки времени, не указана. Для репликации кода я использую Geforce GTX 1080 Ti и процессор Intel Core i7 8700K.
Я в основном скопировал сценарий добавления из учебника, но добавил также последовательный код для сравнения:
from numba import cuda
import numpy as np
import time
import math
@cuda.jit
def addition_kernel(x, y, out):
tx = cuda.threadIdx.x
ty = cuda.blockIdx.x
block_size = cuda.blockDim.x
grid_size = cuda.gridDim.x
start = tx ty * block_size
stride = block_size * grid_size
for i in range(start, x.shape[0], stride):
out[i] = y[i] x[i]
def add(n, x, y):
for i in range(n):
y[i] = y[i] x[i]
if __name__ =="__main__":
print(cuda.gpus[0])
print("")
n = 100000
x = np.arange(n).astype(np.float32)
y = 2 * x
out = np.empty_like(x)
x_device = cuda.to_device(x)
y_device = cuda.to_device(y)
out_device = cuda.device_array_like(x)
# Set the number of threads in a block
threadsperblock = 128
# Calculate the number of thread blocks in the grid
blockspergrid = 30#math.ceil(n[0] / threadsperblock)
# Now start the kernel
start = time.process_time()
cuda.synchronize()
addition_kernel[blockspergrid, threadsperblock](x_device, y_device, out_device)
cuda.synchronize()
end = time.process_time()
out_global_mem = out_device.copy_to_host()
print("parallel time: ", end - start)
start = time.process_time()
add(n,x,y)
end = time.process_time()
print("sequential time: ", end-start)
Параллельное время составляет в среднем около 0,14 секунды, в то время как код без ядра GPU занимает всего 0,02 секунды.
Мне это кажется довольно странным. Есть ли что-нибудь, что я сделал не так? Или эта проблема не является хорошим примером параллелизма? (что я не думаю, так как вы можете запускать цикл for параллельно)
Что странно, так это то, что я почти не замечаю разницы, если не использую to_device()
функции. Насколько я понял, это должно быть важно, так как они избегают связи между процессором и графическим процессором после каждой итерации.
Ответ №1:
addition_kernel
компилируется во время выполнения, когда он вызывается в первый раз, так что в середине вашего измеренного времени! Компиляция ядра-довольно трудоемкая операция. Вы можете ускорить вычисление заранее, предоставив типы в Numba.
Обратите внимание, что массивы слишком малы, поэтому вы можете увидеть значительное улучшение графических процессоров. Более того, сравнение с версией процессора на самом деле несправедливо: вы также должны использовать Numba для реализации процессора или, по крайней мере, Numpy (но не интерпретируемый цикл pure-CPython).
Вот пример:
import numba as nb
@cuda.jit('void(float32[::1], float32[::1], float32[::1])')
def addition_kernel(x, y, out):
tx = cuda.threadIdx.x
ty = cuda.blockIdx.x
block_size = cuda.blockDim.x
grid_size = cuda.gridDim.x
start = tx ty * block_size
stride = block_size * grid_size
for i in range(start, x.shape[0], stride):
out[i] = y[i] x[i]
@nb.njit('void(int64, float32[::1], float32[::1])')
def add(n, x, y):
for i in range(n):
y[i] = y[i] x[i]
Комментарии:
1. У вас есть пример того, как я могу предоставить типы данных в numba? После поиска в Google я прочитал, что AOT-компиляция недоступна для cuda numba, но, возможно, вы имеете в виду что-то другое? Что касается справедливости, я согласен, но код, который я добавил для последовательного варианта, был самым медленным методом, который я мог придумать. Поэтому я был особенно удивлен, когда процессор оказался быстрее графического процессора. Однако я заметил, что графический процессор обгоняет процессор после увеличения размера массива (в 10 раз).
2. Я добавил пример. На моей машине это исправляет проблему с кодом. Обратите внимание, что я должен умножить размер (n) на 1000, чтобы тайминги не округлялись до 0. Реализация GPU на моем Nvidia 1660 Super в 2 раза быстрее по сравнению с последовательной реализацией на моем i5-9600KF. Реализация GPU может быть оптимизирована с помощью развертывания, а реализация CPU также может быть оптимизирована с использованием нескольких потоков Numba (см.
parallel=True
иprange
в Numba). Обратите внимание, что два ядра не являются строго эквивалентными, поскольку реализация процессора выполняет вычисления на месте.