Готовый перевод PyCUDA documentation / Документация PyCUDA: Metaprogramming with PyCUDA http://documen.tician.de/pycuda/metaprog.html#metaprog

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

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

Выглядит довольно сложным, поэтому для начала рассмотрим, почему эта идея, тем не менее, может оказаться разумной.

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

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

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

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

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

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

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

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

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

Типы данных

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

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

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

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

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

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

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

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

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

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

from jinja2 import Template

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;

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

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

mod = SourceModule(rendered_tpl)

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

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

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

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

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

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

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)]))])

mod = SourceModule(mod)

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

Переведено на Нотабеноиде
http://notabenoid.org/book/20422/66701

Переводчики: wuffur