Введение Недавно, прочитав различные статьи и презентации о GPGPU, я тоже решил попробовать программировать для видеокарт. На самом деле выбор технологий в этой области невелик — сейчас живы и развиваются только CUDA (собственный стандарт nVidia) и OpenCL (свободный стандарт, работающий на графических процессорах от ATI, nVidia, а также на центральных процессорах).
В связи с тем, что в моем ноутбуке установлена видеокарта ATI (Mobility Radeon 5650 HD), выбор полностью свелся к одному варианту — OpenCL. В этой статье пойдет речь о процессе изучения OpenCL с нуля, а также о том, что из этого получилось.
Обзор OpenCL и PyOpenCl
На первый взгляд мне все показалось очень запутанным, как управляющий код на C, так и код так называемых ядер.В предоставленном C API даже запуск самой простой программы занимает большое количество строк, особенно с обработкой хоть каких-то ошибок, поэтому хотелось найти что-то более удобное и гуманное.
Выбор пал на библиотеку.
PyOpenCL , из названия которого понятно, что управляющий код написан на Python. В нем все выглядит более понятно даже для того, кто впервые видит код OpenCL (конечно, это касается только простых примеров).
Однако код самих ядер по-прежнему написан на слегка модифицированном языке C, поэтому его вам все равно придется изучить.
Полную документацию по нему можно получить на сайте разработчика стандарта ( Хронос ), а информацию о конкретных реализациях можно найти на сайтах ATI и nVidia соответственно.
Первое впечатление о языке можно получить на простом примере (добавление двух массивов):
А вот полный код, необходимый для запуска этого примера и проверки корректности (взято из документации PyOpenCL): Скрытый текст__kernel void sum(__global const float *a, __global const float *b, __global float *c) { int gid = get_global_id(0); c[gid] = a[gid] + b[gid]; }
import pyopencl as cl
import numpy
import numpy.linalg as la
a = numpy.random.rand(50000).
astype(numpy.float32)
b = numpy.random.rand(50000).
astype(numpy.float32)
ctx = cl.create_some_context()
queue = cl.CommandQueue(ctx)
mf = cl.mem_flags
a_buf = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=a)
b_buf = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=b)
dest_buf = cl.Buffer(ctx, mf.WRITE_ONLY, b.nbytes)
prg = cl.Program(ctx, """
__kernel void sum(__global const float *a,
__global const float *b, __global float *c)
{
int gid = get_global_id(0);
c[gid] = a[gid] + b[gid];
}
""").
build()
prg.sum(queue, a.shape, None, a_buf, b_buf, dest_buf)
a_plus_b = numpy.empty_like(a)
cl.enqueue_copy(queue, a_plus_b, dest_buf)
print la.norm(a_plus_b - (a+b))
Сразу видны конкретные строчки: создание контекста, очередей выполнения, создание и копирование буферов на устройство, а также компиляция и запуск самого ядра.
Подробно о контекстах и очередях в OpenCL вы можете прочитать в документации, а для относительно простых программ вам понадобится только одна очередь и один контекст, который будет создан строками, очень похожими на те, что в примере.
В целом структура вычислений в программах OpenCL часто выглядит примерно так:
- создание контекста, очереди, компиляция программы
- копирование данных (буферов) на устройство, не изменяющихся в процессе выполнения
- цикл
-
- копирование данных, специфичных для итерации, на устройство
- выполнение ядра
- копирование вычисленных данных обратно в основную память, возможна некоторая обработка
SHA1-хеширование
Пришло время спуститься на уровень ниже и понять, как работает сам код ядра.Для кого, чтобы функция OpenCL запускалась извне, она должна быть обозначена атрибутом __kernel, иметь тип значения void и ряд аргументов, которые могут быть как прямыми значениями (int, float4,.
) или указатели на области памяти __global, __constant, __local. Также для удобства в программе можно объявлять и другие функции, вызываемые из ядра, и это не влияет на производительность: все функции подставляются автоматически (т. е.
как и в случае с директивой inline).
Это связано с тем, что рекурсия в OpenCL вообще не поддерживается.
Используя тот факт, что язык OpenCL представляет собой модифицированный C, можно взять готовую реализацию хеш-функции, например SHA1, и использовать ее с небольшими модификациями: Скрытый текст #define K1 0x5A827999
#define K2 0x6ED9EBA1
#define K3 0x8F1BBCDC
#define K4 0xCA62C1D6
#define a0 0x67452301;
#define b0 0xEFCDAB89;
#define c0 0x98BADCFE;
#define d0 0x10325476;
#define e0 0xC3D2E1F0;
#define f1(x,y,z) ( z ^ ( x & ( y ^ z ) ) ) /* Rounds 0-19 */
#define f2(x,y,z) ( x ^ y ^ z ) /* Rounds 20-39 */
#define f3(x,y,z) ( ( x & y ) | ( z & ( x | y ) ) ) /* Rounds 40-59 */
#define f4(x,y,z) ( x ^ y ^ z ) /* Rounds 60-79 */
#define ROTL(n,X) ( ( ( X ) << n ) | ( ( X ) >> ( 32 - n ) ) )
#define expand(W,i) ( W[ i & 15 ] = ROTL( 1, ( W[ i & 15 ] ^ W[ (i - 14) & 15 ] ^ \
W[ (i - 8) & 15 ] ^ W[ (i - 3) & 15 ] ) ) )
#define subRound(a, b, c, d, e, f, k, data) \
( e += ROTL( 5, a ) + f( b, c, d ) + k + data, b = ROTL( 30, b ) )
#define REVERSE(value) value = ((value & 0xFF000000) >> 24) | ((value & 0x00FF0000) >> 8) | ((value & 0x0000FF00) << 8) | ((value & 0x000000FF) << 24)
long sha1(uint *eData, const int length)
{
Теги: #opencl #взлом #пароли #python #python #GPGPU
-
Торий
19 Oct, 24 -
Полевая Радиолюбительская Аппаратура
19 Oct, 24 -
Технологии Меняют Мир Прямо Сейчас
19 Oct, 24 -
Бэкэнд-Истории Meetup
19 Oct, 24
