Главная
Параллельные функции шума и случайных чисел для ядер OpenCL

Параллельные функции шума и случайных чисел для ядер OpenCL +5

18.03.2016 06:15

saul 0 2400 Источник

Образец кода Noise, прилагаемый к этой статье, включает реализацию алгоритма создания шума Перлина, полезного для формирования текстур естественного вида, таких как мрамор или облака, для трехмерной графики. В состав статьи входит тест, в котором используется алгоритм шума Перлина для создания изображения «облака». (Дополнительные сведения об алгоритме шума Перлина см. в разделе «Справочные материалы».) Включены двухмерная и трехмерная версии алгоритма. Это означает, что функции принимают на вход два или три набора данных, чтобы создать одно выходное значение шума Перлина.
Пример Noise также включает функции генератора псевдослучайных чисел (RNG), выдающих сравнительно неплохие результаты, достаточные для того, чтобы полученное изображение действительно выглядело случайным. Включена одномерная, двухмерная и трехмерная версии: количество измерений и в этом случае равно количеству наборов входных данных, на основе которых формируется одно псевдослучайное выходное значение.

Введение и мотивация

Во многих областях применения требуется некоторая степень «случайности» или, точнее, «псевдослучайности». Речь идет о последовательностях значений, которые покажутся человеку произвольными, беспорядочными, то есть «шумом». При этом для целей повторяемости приложениям часто требуется, чтобы генератор случайных чисел мог с достаточной надежностью выдавать в точности одинаковые последовательности при получении одинаковых входных значений (или наборов значений).

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

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

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

В функциях уровня ядра OpenCL в образце кода Noise применяется подход, более подходящий к задействованному в OpenCL алгоритму разделения работы на параллельные операции.

Создание шума и случайных чисел для OpenCL

В OpenCL глобальное рабочее пространство (массив рабочих элементов) определяется с помощью одного, двух или трех измерений. Каждый рабочий элемент в этом глобальном пространстве обладает уникальным набором идентифицирующих целочисленных значений, соответствующих координатам по осям X, Y и Z в глобальном пространстве.

Функции шума Перлина и генератора случайных чисел в примере Noise создают случайное число (шум) на основе до трех входных значений, которые могут быть глобальными идентификаторами каждого рабочего элемента. Есть и другой алгоритм: одно или несколько значений могут быть созданы путем сочетания глобальных идентификаторов и каких-либо значений данных, полученных или созданных ядром.

Например, в следующем фрагменте кода ядра OpenCL показано создание случайных чисел на основе двухмерного глобального идентификатора рабочего элемента.

kernel void genRand()
{
	uint	x = get_global_id(0);
	uint	y = get_global_id(1);

	uint	rand_num = ParallelRNG2( x, y );

	...

Рисунок 1. Пример использования случайных чисел (два измерения)

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

kernel void multi2dNoise( float fScale, float offset )
{
float	fX = fScale * get_global_id(0);
float	fY = fScale * get_global_id(1);
float	fZ = offset;

float	randResult = Noise_3d(  fX,  fY,  fZ );
...

Рисунок 2. Пример использования алгоритма шума Перлина (три измерения)

Ограничения

В функциях Noise_2d и Noise_3d используется один и тот же базовый алгоритм шума Перлина, но различается реализация на основе рекомендаций Перлина. (См. первую ссылку в списке справочных материалов.) В образце Noise только Noise_3d используется в примере noise, а тестовое ядро Noise_2d включено в файл Noise.cl для читателей, желающих изменить этот пример и поэкспериментировать с ним.

Функции Noise_2d и Noise_3d следует вызывать со входными значениями с плавающей запятой. Значения должны охватывать какой-либо диапазон, например (0.0, 128.0), чтобы задать размер «таблицы» (см. рисунок 3) случайных значений. Читателям следует рассмотреть пример с облаками, чтобы понять, как можно преобразовать шум Перлина в разнообразные изображения «естественного вида».

Функция ParallelRNG по умолчанию, используемая в тесте random, предоставляет случайные (и выглядящие таковыми) результаты, но не является самым быстрым алгоритмом генератора случайных чисел. Эта функция построена на основе «хеша Вонга», который изначально не предназначался для использования в качестве генератора случайных чисел. Тем не менее некоторые распространенные функции генераторов случайных чисел (закомментированный пример в файле Noise.cl) показывали видимую повторяемость при заполнении двухмерных изображений, особенно в битах младшего разряда результатов. Читатели могут поэкспериментировать и с другими (более быстрыми) функциями генераторов случайных чисел.

Используемая по умолчанию функция ParallelRNG создает в качестве результатов только 32-битные целые числа без знака. Если требуются значения с плавающей запятой для диапазона, такого как (0.0, 1.0), то приложение должно применить сопоставление для этого диапазона. Образец random сопоставляет случайный целочисленный результат без знака с диапазоном (0, 255) для создания значений пикселей в оттенках серого. Для этого просто применяется двоичная операция AND, чтобы выбрать 8 бит.

Используемая по умолчанию функция ParallelRNG не будет создавать все 4 294 967 296 (2³²) целочисленных значений без знака для последовательных вызовов на основе созданного ранее значения. Для каждого отдельного начального значения величина псевдослучайных последовательностей (циклов) может составлять от всего 7000 уникальных значений до приблизительно 2 миллиардов значений. Функция ParallelRNG создает около 20 различных циклов. Автор считает маловероятным, что какому-либо рабочему элементу ядра OpenCL может потребоваться больше последовательно созданных случайных чисел, чем образуется в самом маленьком цикле.

Двухмерные и трехмерные версии этой функции —ParallelRNG2 и ParallelRNG3 — используют «смешение» циклов путем применения двоичной операции XOR между результатом предыдущего вызова ParallelRNG и следующим входным значением, из-за чего изменяются длины циклов. Тем не менее такое измененное поведение не подвергалось подробному анализу, поэтому читателю рекомендуется внимательно убедиться в том, что функции ParallelRNG отвечают потребностям приложения.

Структура проекта

В этом разделе приводятся только основные элементы исходного кода образца приложения.

NoiseMain.cpp:

main()

основная входная функция. После разбора параметров командной строки она инициализирует OpenCL, собирает программу OpenCL из файла Noise.cl, подготавливает одно из ядер к запуску и вызывает ExecuteNoiseKernel(), а затем ExecuteNoiseReference(). Проверив, что эти две реализации выдают одинаковые результаты, main() выдает информацию о времени работы каждой из этих функций и сохраняет изображения, получившиеся в результате их работы.

ExecuteNoiseKernel()

Настройка и запуск выбранного ядра Noise в OpenCL.

ExecuteNoiseReference()

Настройка и запуск выбранного эталонного кода Noise на языке C.

Noise.cl:

defaut_perm[256]

Таблица случайных значений 0–255 для ядра трехмерного шума Перлина. Для дополнительного повышения случайности эту таблицу можно сформировать и передать в ядро шума Перлина.

grads2d[16]

16 однородно распределенных единичных векторов, градиенты для ядра двухмерного шума Перлина.

grads3d[16]

16 векторных градиентов для ядра трехмерного шума Перлина.

ParallelRNG()

Генератор псевдослучайных чисел, один проход для одного входного значения. Альтернативная функция генератора случайных чисел закомментирована, но добавлена в файл на случай, если у читателей возникнет желание протестировать функцию, которая работает быстрее, но выдает худшие результаты.

ParallelRNG2()

Генератор случайных чисел, делающий 2 прохода для 2 входных значений.

ParallelRNG3()

Генератор случайных чисел, делающий 3 прохода для 3 входных значений.

weight_poly3(), weight_poly5() и WEIGHT()

Это альтернативные функции веса, используемые алгоритмом шума Перлина для обеспечения непрерывности градиентов. Вторая (предпочитаемая) функция также обеспечивает непрерывность второй производной. Макрос WEIGHT выбирает, какая функция используется.

NORM256()

Макрос, преобразующий диапазон (0, 255) в (-1.0, 1.0).

interp()

Билинейная интерполяция, использующая OpenCL.

hash_grad_dot2()

Выбирает градиент и вычисляет скалярное произведение со входными значениями XY в составе функции шума Перлина Noise_2d.

Noise_2d()

Генератор шума Перлина с двумя входными значениями.

hash_grad_dot3()

Выбирает градиент и вычисляет скалярное произведение со входными значениями XYZ в составе функции шума Перлина Noise_3d.

Noise_3d()

Генератор шума Перлина с тремя входными значениями.

cloud()

Создает один пиксель «облачного» выходного изображения для CloudTest с помощью Noise_3d.

map256()

Преобразует выходной диапазон шума Перлина (-1.0, 1.0) в диапазон (0, 255), необходимый для получения пикселей в оттенках серого.

CloudTest()

Тест с созданием изображения облака. Параметр slice передается в функцию cloud, чтобы код системы мог создавать другие изображения облаков.

Noise2dTest()

Тест Noise_2d не используется по умолчанию.

Noise3dTest()

Тест Noise_3d — функции шума Перлина по умолчанию. Использует map256 для получения значений пикселей изображения в оттенках серого.

RandomTest()

Тест ParallelRNG3 использует байт младшего разряда целочисленного результата без знака для получения изображения в оттенках серого.

Предоставляются два файла решений Microsoft Visual Studio для Visual Studio версий 2012 и 2013. Это файлы Noise_2012.sln и Noise_2013.sln. Если читатель использует более новую версию Visual Studio, должно быть возможно использовать функцию обновления решений и проектов Visual Studio для создания нового решения на основе этих файлов.
В обоих решениях предполагается, что в системе установлен компонент Intel® OpenCL™ Code Builder.

Управление образцом

Образец можно запустить из командной строки Microsoft Windows* из папки, в которой находится EXE-файл.

Noise.exe <параметры>

Параметры

-h или --help

Отображение справки в командной строке. Демонстрации при этом не запускаются.

-t или --type [ all | cpu | gpu | acc | default | <константа типа устройства OpenCL>

Выбор типа устройства, для которого запускается ядро OpenCL. Значение по умолчанию: all.

CL_DEVICE_TYPE_ALL | CL_DEVICE_TYPE_CPU | CL_DEVICE_TYPE_GPU |

CL_DEVICE_TYPE_ACCELERATOR | CL_DEVICE_TYPE_DEFAULT

<константа типа устройства OpenCL>

-p или --platform <число или строка>

Выбор используемой платформы. При запуске демонстрации выводится список всех номеров и названий платформ. Справа от используемой платформы будет указано [Selected]. При использовании строки укажите достаточно букв для уникального распознавания названия платформы. Значение по умолчанию: Intel.

-d или --device <число или строка>

Выбор устройства, на котором выполняются ядра OpenCL, по номеру или имени. При запуске демонстрации выводится список всех номеров и названий устройств на используемой платформе. Справа от текущего устройства будет указано [Selected]. Значение по умолчанию: 0.

-r или --run [ random | perlin | clouds ]

Выбор демонстрации функции для запуска. У генератора случайных чисел, алгоритма шума Перлина и генератора изображения облака есть демонстрационные ядра. Значение по умолчанию: random.

-s или --seed <целое число>

Целочисленное входное значение, в зависимости от которого изменяется результат работы алгоритма. Значение по умолчанию: 1.

Noise.exe выводит время работы ядра OpenCL и эталонного эквивалентного кода C, а также имена выходных файлов обоих алгоритмов. По завершении вывода информации программа ожидает нажатия пользователем клавиши ВВОД перед выходом. Обратите внимание, что функции эталонного кода на C не были оптимизированы, их цель состоит лишь в проверке правильности кода ядра OpenCL.

Анализ результатов

После завершения работы Noise.exe просмотрите файлы изображений в формате BMP OutputOpenCL.bmp и OutputReference.bmp в рабочей папке, чтобы сравнить результаты кода OpenCL и C++. Эти два изображения должны быть одинаковыми, хотя возможны весьма незначительные различия между двумя изображениями шума Перлина или между двумя изображениями облаков.
Результат работы алгоритма noise (шум Перлина) должен выглядеть, как показано на рисунке 3.

Рисунок 3. Результат работы алгоритма шума Перлина

Результат работы алгоритма random (генератор случайных чисел) должен выглядеть, как показано на рисунке 4.

Рисунок 4. Результат работы генератора случайных чисел

Результат работы алгоритма cloud (облака) должен выглядеть, как показано на рисунке 5.

Рисунок 5. Результат работы алгоритма создания облаков