PyCUDA documentation / Документация PyCUDA: Metaprogramming with PyCUDA http://documen.tician.de/pycuda/metaprog.html#metaprog

Метапрограммирование с PyCUDA

wuffur 02.08.11 в 07:47

При "Обычном" программировании, можно написать программу, которая выполнит задачу. В метопрограммировании, можно написать программу, которая напишет программу, которая выполнит задачу.

wuffur 02.08.11 в 07:49

Это звучит довольно запутанно(сложно), поэтому, во-первых, мы посмотрим почему может быть хорошей идеей несмотря ни на что.

wuffur 02.08.11 в 07:51

Почему мы используем Метапрограммирование?

wuffur 02.08.11 в 09:47

Автоматизированная настройка

wuffur 02.08.11 в 09:47

Значительная часть времени работы программиста на CUDA обычно тратиться на настройку кода. Настройка отвечает на такие вопросы, как:

wuffur 02.08.11 в 09:49

Какое оптимальное количество потоков идет на каждый блок?

wuffur 02.08.11 в 09:51

С насколько большим объемом данных я должен работать за раз?

wuffur 02.08.11 в 09:51

Какие данные должны быть загружены в разделяемую память и насколько большими должны быть соответствующие блоки?

wuffur 02.08.11 в 09:52

И если вы будете достаточно удачливыми, вы сможете найти шаблон(паттерн) во времени выполнения и придумать эвристику(эвристический алгоритм), которая позволит надежно подбирать самую быструю версию. К сожалению, эта эвристика может стать ненадежной или даже совсем не сработать на новом поколении оборудования(графических карт). Решение проблемы, которое PyCUDA пытается нам пообещать такое:

wuffur 02.08.11 в 09:57

Забудьте эвристику. Тесты производительности происходят во время выполнения и используется то, что быстрее работает.

wuffur 02.08.11 в 09:59

Это важное преимущество PyCUDA над временем выполнения CUDA API: Это позволяет вам принять эти решения когда код запущен. Ряд известных вычислительных пакетов используют схожие методы, вроде ATLAS и FFTW. И хотя это требует довольно сложную технику оптимизации программ драйверов, вы можете управлять PyCUDA из комфортного Python.

wuffur 02.08.11 в 10:05

Типы данных

wuffur 02.08.11 в 10:07

Ваш код может работать с различными типами данных во время выполнения. Это может, например, работать на обоих сторонах с числами с плавающей запятой одинарной и двойной точности. Вы можете просто передкомпилировать версии для обоих, поэтому почему? Просто сгенерируйте тот код, который необходимо прямо тогда, когда это нужно.

wuffur 02.08.11 в 10:19

Специализация кода для поставленной проблемы.

wuffur 02.08.11 в 10:26

Если вы пишите библиотеку, то тогда ваши пользователи будут спрашивать вашу библиотеку для того, чтобы выполнить ряд задач. Представьте себе как свободно было бы если бы вы могли генерировать код специально для той проблемы, которую вы просите разрешить, вместо того, чтобы писать код, который имеет излишне общий характер и тем самым был медленным. PyCUDA делает это реальностью.

wuffur 02.08.11 в 10:27

Константы быстрее, чем переменные.

wuffur 02.08.11 в 10:45

Если ваша проблема распространяется на большое расстояние между запусками, но вы выполняете более большое количество вызовов ядра на данных идентичного размера, вам можете захотеть считать, что размер компилируемых данных в вашем коде константа. Это может иметь значительную преимущество по производительности, приводя в основном к тому, что уменьшается количество раз выборки и регистрируется меньшее давление. В особенности, произведение констант намного более эффективно проводит эти общие произведения переменная на переменную.

wuffur 02.08.11 в 10:56

Развертывание цикла.

wuffur 02.08.11 в 11:01

Руководство по программированию CUDA хорошо говорит о nvcc и как он будет развертывать циклы для вас. Поскольку ваша версия 2.1, это просто неправда, и директива #pragma развертывается просто как холостая команда, по крайней мере, опираясь на мой опыт. С использованием метапрограммирования, вы можете динамически развертывать ваши циклы до ниобходимого Python размера.

wuffur 02.08.11 в 11:06

Метапрограммирование с использование Шаблонизатора

wuffur 02.08.11 в 11:06

Если метапрограммированию нужен довольно простой и, возможно, самый легкий способ генерации кода во время выполнения это пропускать его через шаблонизатор. Существует много шаблонизаторов у Python, два наиболее известных из них - это Jinja 2 и Cheetah.

wuffur 02.08.11 в 11:09

Следующее - это простая метапрограмма, которая позволяет добавлять вектор на конфирируемые размеры блоков. Она иллюстрирует основанные на шаблонизирование метод метопрограммирования:

wuffur 02.08.11 в 11:16

from jinja2 import Template

wuffur 02.08.11 в 11:16

tpl = Template("""
__global__ void add(
{{ type_name }} *tgt,
{{ type_name }} *op1,
{{ type_name }} *op2)
{
int idx = threadIdx.x +
{{ thread_block_size }} * {{block_size}}
* blockIdx.x;

wuffur 02.08.11 в 11:16

{% for i in range(block_size) %}
{% set offset = i*thread_block_size %}
tgt[idx + {{ offset }}] =
op1[idx + {{ offset }}]
+ op2[idx + {{ offset }}];
{% endfor %}
}""")

wuffur 02.08.11 в 11:16

rendered_tpl = tpl.render(
type_name="float", block_size=block_size,
thread_block_size=thread_block_size)

wuffur 02.08.11 в 11:16

mod = SourceModule(rendered_tpl)

wuffur 02.08.11 в 11:16

Следующий фрагмент кода в рабочей окружении можно найти в
examples/demo_meta_template.py.

wuffur 03.08.11 в 07:12

Вы также можете найти пример оптимизации произведения матриц используя метапрограммирование шаблонов с Cheetah в
demo_meta_matrixmul_cheetah.py и demo_meta_matrixmul_cheetah.template.cu.

wuffur 03.08.11 в 07:14

Метапрограммирование используя codepy

wuffur 03.08.11 в 07:14

Для более сложных метапрограмм, желательно иметь больше программного контроля надл сборкой исходной кода чем шаблонизатор может предоставить. Пакет codepy предоставляет средство генерации исходного кода CUDA из структуры данных Python.

wuffur 03.08.11 в 07:16

Следующий пример демонстрирует использование codepy для метапрограммирования. Он завершает в точности ту же, что и описана выше программу.

wuffur 03.08.11 в 07:17

from codepy.cgen import FunctionBody, \
FunctionDeclaration, Typedef, POD, Value, \
Pointer, Module, Block, Initializer, Assign
from codepy.cgen.cuda import CudaGlobal
35 mod = Module([

wuffur 03.08.11 в 07:17

mod = Module([
FunctionBody(
CudaGlobal(FunctionDeclaration(
Value("void", "add"),
arg_decls=[Pointer(POD(dtype, name))
for name in ["tgt", "op1", "op2"]])),
Block([
Initializer(
POD(numpy.int32, "idx"),
"threadIdx.x + %d*blockIdx.x"
% (thread_block_size*block_size)),
]+[
Assign(
"tgt[idx+%d]" % (o*thread_block_size),
"op1[idx+%d] + op2[idx+%d]" % (
o*thread_block_size,
o*thread_block_size))
for o in range(block_size)]))])

wuffur 03.08.11 в 07:18

mod = SourceModule(mod)

wuffur 03.08.11 в 07:18

Этот фрагмент кода в рабочем окружении может быть найден в
examples/demo_meta_codepy.py.

wuffur 03.08.11 в 07:18