PyCUDA documentation / Документация PyCUDA: Tutorial Introduction http://documen.tician.de/pycuda/tutorial.html

До того, как вы сможете использовать PyCUDA, вам нужно будет импортировать и инициализировать его:

wuffur 03.08.11 в 11:35

import pycuda.driver as cuda
import pycuda.autoinit
from pycuda.compiler import SourceModule

wuffur 03.08.11 в 11:35

Отметим, что вам не нужно использовать инициализацию pycuda.autoinit, контекстное создание, и очистку тоже можно произвести вручную, если вы этого желаете.

wuffur 03.08.11 в 11:37

Передача данных

wuffur 03.08.11 в 11:37

Следдующим шагом в большинстве программ является передача данных в прибор. В PyCUDA, вы будете в основном передавать данный через массивы numpy на хост.(но в действительности, все, что заставляет работать интерфейс буфера Pythona будет работать, даже строка). Давайте сделаем массив 4x4 случайных чисел:

wuffur 03.08.11 в 11:41

import numpy
a = numpy.random.randn(4,4)

wuffur 03.08.11 в 12:04

Но подождите - это совокупность чисел двойной точности, но большинство приборов nVidia поддерживают только числа одинарной точности:

wuffur 03.08.11 в 12:06

a = a.astype(numpy.float32)

wuffur 03.08.11 в 12:06

Наконец, нам нужно куда-то передавать данные, поэтому нам нужно выделить память на устройстве:

wuffur 03.08.11 в 12:06

a_gpu = cuda.mem_alloc(a.nbytes)

wuffur 03.08.11 в 12:07

В качестве последнего шага, нам нужно передать данные в GPU:

wuffur 03.08.11 в 12:07

cuda.memcpy_htod(a_gpu, a)

wuffur 03.08.11 в 12:08

Выполнение Ядра

wuffur 03.08.11 в 12:08

В этом руководстве мы будем придерживаться чего-то простого: будем писать код в объекты из переменных двойной точности в a_gpu. К концу этой документации, мы напишем соответствующий им код Си на CUDA, и передадим этот код конструктору pycuda.compiler.SourceModule:

wuffur 03.08.11 в 12:10

mod = SourceModule("""
__global__ void doublify(float *a)
{
int idx = threadIdx.x + threadIdx.y*4;
a[idx] *= 2;
}
""")

wuffur 03.08.11 в 12:11

Если здесь нет никаких ошибок, то код теперь откомпилируется и загрузится в прибор. Мы нашли ссылку на нашу pycuda.driver.Function и вызвали ее, специфицируя a_gpu как аргумент и установив размер блока 4x4:

wuffur 03.08.11 в 20:33

func = mod.get_function("doublify")
func(a_gpu, block=(4,4,1))

wuffur 03.08.11 в 20:33

И наконец, мы произведем выборку данных из GPU и выведим их на экран вместе с оригинальным a:

wuffur 04.08.11 в 04:37

a_doubled = numpy.empty_like(a)
cuda.memcpy_dtoh(a_doubled, a_gpu)
print a_doubled
print a

wuffur 04.08.11 в 04:37

Этот кусок кода выведет, что-то похожее на это:

wuffur 04.08.11 в 04:38

[[ 0.51360393 1.40589952 2.25009012 3.02563429]
[-0.75841576 -1.18757617 2.72269917 3.12156057]
[ 0.28826082 -2.92448163 1.21624792 2.86353827]
[ 1.57651746 0.63500965 2.21570683 -0.44537592]]
[[ 0.25680196 0.70294976 1.12504506 1.51281714]
[-0.37920788 -0.59378809 1.36134958 1.56078029]
[ 0.14413041 -1.46224082 0.60812396 1.43176913]
[ 0.78825873 0.31750482 1.10785341 -0.22268796]]

wuffur 04.08.11 в 04:38

Она работает! На этом закончим наше пошаговое руководство. К счастью, PyCuda наследует(запоминает эти данные) это и все вам подчищает, поэтому вы закончили. Однако не пропустите некоторый бонусный материал в следующем разделе.

wuffur 04.08.11 в 04:39

(Вы можете найти код для этого демо, по адресу examples/demo.py, в дистрибутиве исходников PyCUDA.)

wuffur 04.08.11 в 05:02

Сочетания клавиш для явных копий памяти

wuffur 04.08.11 в 05:07

Обработчик аргумента pycuda.driver.In, pycuda.driver.Out и pycuda.driver.InOut может упростить некоторые процессы передачи памяти. Например, вместо создания a_gpu, если замещение a прошло хорошо, можно использовать следующий код:

wuffur 04.08.11 в 07:39

func(cuda.InOut(a), block=(4, 4, 1))

wuffur 04.08.11 в 07:39

Подготовленные вызовы.

wuffur 04.08.11 в 07:39

Вызов функции использует встроенный метод pycuda.driver.Function.__call__(), берущий на себя накладные расходы на определение типа(смотрите "Справочную документацию по интерфейсу устройств"). Чтобы достичь того же эффекта что и выше без этих накладных расходов, функций связывается с типами аргументов(как спроектировано в структурном модуле стандартной библиотеки Python), и затем вызывается. Это также помогает избежать объявления неявных размеров аргументов, используя классы numpy.number:

wuffur 04.08.11 в 07:46

func.prepare("P", block=(4,4,1))
func.prepared_call((1, 1), a_gpu)

wuffur 04.08.11 в 07:48

Бонус: Абстрагирование от осложнений.

wuffur 04.08.11 в 07:49

Использование pycuda.gpuarray.GPUArray,тот же эффект можно достичь написанием значительно меньшего количества кода:

wuffur 04.08.11 в 07:50

import pycuda.gpuarray as gpuarray
import pycuda.driver as cuda
import pycuda.autoinit
import numpy

wuffur 04.08.11 в 07:50

a_gpu = gpuarray.to_gpu(numpy.random.randn(4,4).astype(numpy.float32))
a_doubled = (2*a_gpu).get()
print a_doubled
print a_gpu

wuffur 04.08.11 в 07:50

Продвинутые темы
Структуры

wuffur 04.08.11 в 07:51

(добавлено Nicholas Tung, найдите код в examples/demo_struct.py)

wuffur 04.08.11 в 07:52

Предполагается, что у вас есть следующая структура, для удвоения числа переменных длин массивов:

wuffur 04.08.11 в 07:53

mod = SourceModule("""
struct DoubleOperation {
int datalen, __padding; // поэтому 64-битные указатели могут быть выравнены
float *ptr;
};

wuffur 04.08.11 в 07:53

__global__ void double_array(DoubleOperation *a) {
a = &a[blockIdx.x];
for (int idx = threadIdx.x; idx < a->datalen; idx += blockDim.x) {
a->ptr[idx] *= 2;
}
}
""")

wuffur 04.08.11 в 07:53

Каждый блок в грид(смотрите документацию CUDA) удвоит один из массивов. Цикл for позволяет удвоить большее количиство элементов данных, чем потоки, хотя и не эффективен, если нельзя гарантировать, что будет необходимое количество потоков. Затем создается класс обертки для струкутры, и два массива инициализируются:

wuffur 04.08.11 в 07:57

class DoubleOpStruct:
mem_size = 8 + numpy.intp(0).nbytes
def __init__(self, array, struct_arr_ptr):
self.data = cuda.to_device(array)
self.shape, self.dtype = array.shape, array.dtype
cuda.memcpy_htod(int(struct_arr_ptr), numpy.int32(array.size))
cuda.memcpy_htod(int(struct_arr_ptr) + 8, numpy.intp(int(self.data)))
def __str__(self):
return str(cuda.from_device(self.data, self.shape, self.dtype))

wuffur 04.08.11 в 08:03

struct_arr = cuda.mem_alloc(2 * DoubleOpStruct.mem_size)
do2_ptr = int(struct_arr) + DoubleOpStruct.mem_size

wuffur 04.08.11 в 08:03

array1 = DoubleOpStruct(numpy.array([1, 2, 3], dtype=numpy.float32), struct_arr)
array2 = DoubleOpStruct(numpy.array([0, 4], dtype=numpy.float32), do2_ptr)
print("original arrays", array1, array2)

wuffur 04.08.11 в 08:03

Этот код использует функции pycuda.driver.to_device() и pycuda.driver.from_device(), чтобы выделить(локализировать) и скопировать значения и демонстрирует какие отступы выделенных блоков памяти могут использоваться. Наконец, код можно выполнить. Следующий код сначала демонстрирует удвоение обоих массивов, а затем только второй.

wuffur 04.08.11 в 08:05

func = mod.get_function("double_array")
func(struct_arr, block = (32, 1, 1), grid=(2, 1))
print("doubled arrays", array1, array2)

wuffur 04.08.11 в 08:07

func(numpy.intp(do2_ptr), block = (32, 1, 1), grid=(1, 1))
print("doubled second only", array1, array2, "\n")

wuffur 04.08.11 в 08:07

Куда теперь направится

wuffur 04.08.11 в 08:07

Когда вы почувствуете себя достаточно знакомыми с основами, можете свободно смотреть следующий раздел (Device Interface Reference Documentation). Для получения большего количества примеров, проверьте поддиректорию examples/ дистрибутива. Этот ярлык может также содержать некоторые тесты производительности, которые показывают разницу между вычислениями на GPU и CPU. В качестве руководства по тому, как это было сделано также может помочь PyCuda’s test suite, расположенный в поддиректории test/ дистрибутива.

wuffur 04.08.11 в 08:12