How to become an author
Search
Write a publication
Pull to refresh

Изучение OpenCL на примере взлома паролей

Reading time10 min
Views10K
GPGPU*Python*
From sandbox

Вступление


Недавно, почитав различных статей и презентаций про GPGPU, я решил тоже попробовать для себя программирование под видеокарты. Фактически, выбор технологий в этой области не велик — сейчас живы и развиваются только CUDA (проприетарный стандарт nVidia) и OpenCL (свободный стандарт, работает на GPU от ATI, nVidia, а также на центральных процессорах). В связи с тем, что мой ноутбук располагает видеокартой ATI (Mobility Radeon 5650 HD), то выбор и вовсе свёлся к одному варианту — OpenCL. В этой статье речь пойдёт о процессе изучения OpenCL с нуля, а также о том, что из этого получилось.

Обзор OpenCL и PyOpenCl


На первый взгляд мне показалось всё очень запутанным, как управляющий код на Си, так и код так называемых ядер — kernels. В предоставляемом C API даже запуск простейшей программы занимает большое количество строк, особенно с обработкой хоть каких-то ошибок, поэтому захотелось найти что-то более удобное и человечное. Выбор пал на библиотеку PyOpenCL, из названия которой ясно, что управляющий код пишется на Python. В ней всё выглядит уже более понятно, даже для того, кто видит код на OpenCL первый раз (разумеется, это относится только к простым примерам). Однако, код самих ядер пишется по-прежнему на несколько изменённом Си, поэтому его всё-таки придётся изучать. Полную документацию по нему можно получить на сайте разработчика стандарта (Khronos), а информацию о конкретных реалиазациях — на сайтах ATI и nVidia соответственно.
Получить первое впечатление о языке можно по простейшему примеру (сложение двух массивов):

    __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];
    }


А вот полный необходимый код для запуска этого примера и проверки правильности (взято из документации PyOpenCL):

Скрытый текст
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 — это модифицированный Си, можно взять готовую реализацию функции хеширования, например 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)
{
	unsigned int A = a0;
	unsigned int B = b0;
	unsigned int C = c0;
	unsigned int D = d0;
	unsigned int E = e0;
	
	((__local char *)eData)[length] = 0x80;
	for (int i = 0; i <= length / 4; i++)
	{
		REVERSE(eData[i]);
	}
		
	eData[14] = 0;	
	eData[15] = length * 8;		

	subRound( A, B, C, D, E, f1, K1, eData[  0 ] );
	subRound( E, A, B, C, D, f1, K1, eData[  1 ] );
	subRound( D, E, A, B, C, f1, K1, eData[  2 ] );
	subRound( C, D, E, A, B, f1, K1, eData[  3 ] );
	subRound( B, C, D, E, A, f1, K1, eData[  4 ] );
	subRound( A, B, C, D, E, f1, K1, eData[  5 ] );
	subRound( E, A, B, C, D, f1, K1, eData[  6 ] );
	subRound( D, E, A, B, C, f1, K1, eData[  7 ] );
	subRound( C, D, E, A, B, f1, K1, eData[  8 ] );
	subRound( B, C, D, E, A, f1, K1, eData[  9 ] );
	subRound( A, B, C, D, E, f1, K1, eData[ 10 ] );
	subRound( E, A, B, C, D, f1, K1, eData[ 11 ] );
	subRound( D, E, A, B, C, f1, K1, eData[ 12 ] );
	subRound( C, D, E, A, B, f1, K1, eData[ 13 ] );
	subRound( B, C, D, E, A, f1, K1, eData[ 14 ] );
	subRound( A, B, C, D, E, f1, K1, eData[ 15 ] );
	subRound( E, A, B, C, D, f1, K1, expand( eData, 16 ) );
	subRound( D, E, A, B, C, f1, K1, expand( eData, 17 ) );
	subRound( C, D, E, A, B, f1, K1, expand( eData, 18 ) );
	subRound( B, C, D, E, A, f1, K1, expand( eData, 19 ) );

	subRound( A, B, C, D, E, f2, K2, expand( eData, 20 ) );
	subRound( E, A, B, C, D, f2, K2, expand( eData, 21 ) );
	subRound( D, E, A, B, C, f2, K2, expand( eData, 22 ) );
	subRound( C, D, E, A, B, f2, K2, expand( eData, 23 ) );
	subRound( B, C, D, E, A, f2, K2, expand( eData, 24 ) );
	subRound( A, B, C, D, E, f2, K2, expand( eData, 25 ) );
	subRound( E, A, B, C, D, f2, K2, expand( eData, 26 ) );
	subRound( D, E, A, B, C, f2, K2, expand( eData, 27 ) );
	subRound( C, D, E, A, B, f2, K2, expand( eData, 28 ) );
	subRound( B, C, D, E, A, f2, K2, expand( eData, 29 ) );
	subRound( A, B, C, D, E, f2, K2, expand( eData, 30 ) );
	subRound( E, A, B, C, D, f2, K2, expand( eData, 31 ) );
	subRound( D, E, A, B, C, f2, K2, expand( eData, 32 ) );
	subRound( C, D, E, A, B, f2, K2, expand( eData, 33 ) );
	subRound( B, C, D, E, A, f2, K2, expand( eData, 34 ) );
	subRound( A, B, C, D, E, f2, K2, expand( eData, 35 ) );
	subRound( E, A, B, C, D, f2, K2, expand( eData, 36 ) );
	subRound( D, E, A, B, C, f2, K2, expand( eData, 37 ) );
	subRound( C, D, E, A, B, f2, K2, expand( eData, 38 ) );
	subRound( B, C, D, E, A, f2, K2, expand( eData, 39 ) );

	subRound( A, B, C, D, E, f3, K3, expand( eData, 40 ) );
	subRound( E, A, B, C, D, f3, K3, expand( eData, 41 ) );
	subRound( D, E, A, B, C, f3, K3, expand( eData, 42 ) );
	subRound( C, D, E, A, B, f3, K3, expand( eData, 43 ) );
	subRound( B, C, D, E, A, f3, K3, expand( eData, 44 ) );
	subRound( A, B, C, D, E, f3, K3, expand( eData, 45 ) );
	subRound( E, A, B, C, D, f3, K3, expand( eData, 46 ) );
	subRound( D, E, A, B, C, f3, K3, expand( eData, 47 ) );
	subRound( C, D, E, A, B, f3, K3, expand( eData, 48 ) );
	subRound( B, C, D, E, A, f3, K3, expand( eData, 49 ) );
	subRound( A, B, C, D, E, f3, K3, expand( eData, 50 ) );
	subRound( E, A, B, C, D, f3, K3, expand( eData, 51 ) );
	subRound( D, E, A, B, C, f3, K3, expand( eData, 52 ) );
	subRound( C, D, E, A, B, f3, K3, expand( eData, 53 ) );
	subRound( B, C, D, E, A, f3, K3, expand( eData, 54 ) );
	subRound( A, B, C, D, E, f3, K3, expand( eData, 55 ) );
	subRound( E, A, B, C, D, f3, K3, expand( eData, 56 ) );
	subRound( D, E, A, B, C, f3, K3, expand( eData, 57 ) );
	subRound( C, D, E, A, B, f3, K3, expand( eData, 58 ) );
	subRound( B, C, D, E, A, f3, K3, expand( eData, 59 ) );

	subRound( A, B, C, D, E, f4, K4, expand( eData, 60 ) );
	subRound( E, A, B, C, D, f4, K4, expand( eData, 61 ) );
	subRound( D, E, A, B, C, f4, K4, expand( eData, 62 ) );
	subRound( C, D, E, A, B, f4, K4, expand( eData, 63 ) );
	subRound( B, C, D, E, A, f4, K4, expand( eData, 64 ) );
	subRound( A, B, C, D, E, f4, K4, expand( eData, 65 ) );
	subRound( E, A, B, C, D, f4, K4, expand( eData, 66 ) );
	subRound( D, E, A, B, C, f4, K4, expand( eData, 67 ) );
	subRound( C, D, E, A, B, f4, K4, expand( eData, 68 ) );
	subRound( B, C, D, E, A, f4, K4, expand( eData, 69 ) );
	subRound( A, B, C, D, E, f4, K4, expand( eData, 70 ) );
	subRound( E, A, B, C, D, f4, K4, expand( eData, 71 ) );
	subRound( D, E, A, B, C, f4, K4, expand( eData, 72 ) );
	subRound( C, D, E, A, B, f4, K4, expand( eData, 73 ) );
	subRound( B, C, D, E, A, f4, K4, expand( eData, 74 ) );
	subRound( A, B, C, D, E, f4, K4, expand( eData, 75 ) );
	subRound( E, A, B, C, D, f4, K4, expand( eData, 76 ) );
	subRound( D, E, A, B, C, f4, K4, expand( eData, 77 ) );
	subRound( C, D, E, A, B, f4, K4, expand( eData, 78 ) );
	subRound( B, C, D, E, A, f4, K4, expand( eData, 79 ) );
	
	A += a0;
	B += b0;
	C += c0;
	D += d0;
	E += e0;
	
        return as_ulong((uint2)(D, E));
}



Здесь нужно сделать некоторые пояснения. В качестве «подопытых» хешей паролей для взлома я взял утёкшие хеши LinkedIn, которых почти 6 миллионов (уникальных). Есть несколько вариантов достаточно быстрой проверки наличия в списке, я использовал хэш-таблицы (подробнее дальше). Для уменьшения расхода памяти и ускорения работы появилась мысль хранить не полные 20 байт SHA1, а последние 8 байт, т.е. одно значение long/ulong. Конечно, это повышает вероятность ложного совпадения, но она остаётся весьма небольшой: из всех перебраных паролей у меня было только 6 таких случаев, что совсем не критично. Поэтому из вышеприведённой функции возвращается сразу обрезанное значение (последние 8 байт). В остальном всё стандартно, алгоритм SHA1 реализован напрямую по спецификации для случая строк менее 56 байт.

Организация перебора


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

Скрытый текст
__kernel void do_brute(
	__global const long *table,
	const ulong start_num,
	__global ulong *result,
	__global uint *res_ind)
{
	char s[64];
	uint *s_l = (__local uint *)s;
	int i, j;
	ulong _n, n;
	
	ulong start = start_num + get_global_id(0) * HASHES_PER_WORKITEM;
	for (j = 0; j < HASHES_PER_WORKITEM; j++)
	{
		n = _n = j + start;
		
		for (i = 15; i >= 0; i--)
		{
			s_l[i] = 0;
		}
		
		for (i = COMB_LEN - 1; i >= 0; i--)
		{
			s[i] = charset[n % CHARS_CNT];
			n /= CHARS_CNT;
		}
		
		if (check_in_table(table, sha1(s_l, COMB_LEN)))
		{
			result[atomic_inc(res_ind)] = _n;
		}
	}
}


Здесь HASHES_PER_WORKITEM — количество хэшей, обрабатываемых за один запуск одним work item (потоком), COMB_LEN — длина комбинации, charset — массив символов, CHARS_CNT — количество символов в массиве. Как можно увидеть, этому ядру при запуске передаются указатель на хэш-таблицу, номер пароля, с которого начинается перебор, а также указатель на массив для вывода результата и индекс в нём.

В OpenCL за раз запускается не один поток, а некоторое их количество, называемое global work size, причём все потоки получают одни и те же аргументы. Для того, чтобы каждый из них перебирал свою часть пространства ключей, строка ulong start = start_num + get_global_id(0) * HASHES_PER_WORKITEM; вычисляет номер для конкретного потока (get_global_id(0) — стандартная функция, возвращающая индекс потока от 0 до текущего global work size).

Далее происходит перебор каждым потоком HASHES_PER_WORKITEM паролей, каждый из которых хэшируется функцией sha1 и проверяется на наличие функцией check_in_table, о которой далее. В таком виде и с простейшей реализацией хэш-таблицы у меня получился результат около 20 миллионов паролей в секунду, что откровенно говоря мало по сравнению например с oclHascat, который выдаёт на моём ноутбуке более 300 миллионов (даже то, что проверка идёт по большому списку хэшей, не оправдывает этого). Забегая вперёд скажу, что для простого брутфорса у меня получилось достичь скорости в 160 миллионов в секунду, что составляет более половины от скорости oclHascat (с одним хэшем).

Хэш-таблица


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

Скрытый текст
bool check_in_table(
	__global const long *table,
	const long value)
{
	uint index = calc_index(value);
	uint step = calc_step(value);
	for (uint a = 1; ; a++)
	{
		index %= TABLE_SIZE;
		if (table[index] == 0)
		{
			return false;
		}
		if (table[index] == value)
		{
			return true;
		}
		index += a * step + 1;
	}
}


Я пробовал различные размеры таблицы и методы пробирования, но скорость заметно не улучшалась. В поисках материала о хэш-таблицах на GPU я наткнулся на статью «Building an Efficient Hash Table on the GPU» от Vasily Volkov, в которой упоминается некая Cuckoo Hashtable (не знаю, есть ли устоявшийся перевод на русский язык), о которой ранее я не слышал. Вкратце суть её заключается в использовании нескольких хэш-функций вместо одной и особом способе заполнения, после чего нахождение элемента происходит за не более k доступов к памяти, где k — количество хэш-функций. Так как мне важнее скорость работы, чем занимаемая память, я использовал k=2. Заполнение её также происходит на CPU.

Заключение


Также, разумеется, оптимизации затронули и другую часть кода, а именно генерацию паролей. В приведённом выше варианте сразу видны несколько неоптимальных мест, например каждый следующий пароль генерируется с нуля, хотя можно изменять предыдущий. Есть и другие места для оптимизации, в частности специфичные для OpenCL: использование глобальной или константной памяти для массива символов вместо значительно более быстрой локальной (подробнее про области памяти лучше читать прямо у разработчика конкретной реализации). Однако, про различные оптимизации кода ядра с подробностями стоит писать отдельную статью, а здесь скажу, что при программировании для GPU стоит пробовать различные варианты и смотреть на их скорость, т.к. не всегда можно сказать на глаз, что будет работать быстрее. Иногда даже удаление некоторой инструкции замедляет выполнение, причём значительно.

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

Выводы


Итак, что в итоге я приобрёл:
  • знания OpenCL со многими тонкостями (совсем с нуля, считаю, что достаточно успешно)
  • поднял уровень программирования на Python (не совсем с нуля, но с достаточно низкого уровня)
  • узнал и попробовал сопутствующие технологии: компилятор Cython, шаблонизатор Mako, VCS git (раньше почти не использовал)


Осязаемые результаты:
  • программа перебирает пароли со скоростью около 160 миллионов в секунду для простого брутфорса (при использовании алфавита из слов, а не символов, скорость меньше: от 50 миллионов в секунду) на ноутбучной видеокарте — можно сравнить с одной стороны с 300 миллионами у oclHascat, работающего с одним хэшем, а с другой с 30 миллионами у самого быстрого брутфорсера на CPU (у меня i5 2.5 GHz), написанного и использованием ассемблера и инструкций SSE
  • с её использованием и без всяких хитростей подобрано около 2.5 миллионов паролей из хэшей LinkedIn (не так много, но здесь я и не стремился к рекорду)
  • скачать можно здесь (в некоторых местах встречается страшный говнокод!), для запуска необходим Python 2.7, numpy, PyOpenCL, Cython, Mako и файл с хешами LinkedIn.


Впечатления:
  • Python — очень медленный язык (зато достаточно прост в написании), обработка больших списков циклами сильно тормозит, также как и различные comprehensions, именно поэтому для некоторых частей программы используется Cython — кстати, действительно удобная вещь
  • OpenCL при должном подходе к профилированию и последующим оптимизациям может работать очень быстро, причём для этого не требуется переписывать ничего на ассемблере

P.S.:


Различные комментарии и рекомендации приветствуются.
Tags:
Hubs:
Total votes 33: ↑32 and ↓1+31
Comments2
64
Karma
0
Rating
Александр Плавин @chersanya

Исследователь

Send message

Articles

Your account

Sections

Information

Services

Support
© 2006–2024, Habr