Exploring OpenCL through password cracking

Introduction

Recently, after reading various articles and presentations about GPGPU, I decided to also try programming for a video card. In fact, the choice of technologies in this area is not great - only CUDA (proprietary standard nVidia) and OpenCL (free standard, running on ATI's GPU, nVidia, and also on central processors) are alive and well. Due to the fact that my laptop has an ATI video card (Mobility Radeon 5650 HD), the choice has come down to only one option - OpenCL. This article focuses on the process of learning OpenCL from scratch, as well as what came of it.

OpenCL and PyOpenCl Overview

At first glance, it seemed to me very confusing, both the control code in C and the code of the so-called kernels - kernels. In the provided C API, even launching the simplest program takes a large number of lines, especially with handling at least some errors, so I wanted to find something more convenient and humane. The choice fell on the PyOpenCL library, from the name of which it is clear that the control code is written in Python. Everything looks more understandable in it, even for someone who sees the code on OpenCL for the first time (of course, this applies only to simple examples). However, the code of the cores themselves is still written in slightly modified C, so you will still have to study it. Full documentation on it can be obtained on the standard developer’s website ( Khronos ), and information on specific implementations can be obtained on the ATI and nVidia sites, respectively.
You can get the first impression about the language by the simplest example (adding two arrays):

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

And here is the full necessary code to run this example and validate (taken from the PyOpenCL documentation):
')


Specific lines are immediately visible: creating a context, execution queues, creating and copying buffers to the device, as well as compiling and running the actual kernel. About contexts and queues in OpenCL, you can read in detail in the documentation, and for relatively simple programs you only need one queue and one context, which will be created in lines that are very similar to those in the example. In general, often the structure of calculations in programs on OpenCL look like this:

• creating context, queue, program compilation
• copying data (buffers) to the device that do not change during execution
• cycle
• • copying data specific to this iteration to a device
  • kernel execution
  • copying the calculated data back to main memory, maybe some processing
  
  

Hashing SHA1

It is time to go down a level and understand how the code of the kernel itself is arranged. For whom, in order for the OpenCL function to run from the outside, it must be indicated by the __kernel attribute, have the value type void and a number of arguments, which can be either the values ​​itself (int, float4, ...), or pointers to areas memory __global, __constant, __local. Also, for convenience, other functions that are called from the kernel can be declared for convenience, and this does not affect performance: all functions are automatically substituted (that is, as with the inline directive). Related to this is that recursion in OpenCL is not supported at all.

Using the fact that the OpenCL language is a modified C, you can take a ready-made implementation of the hashing function, for example SHA1, and with minor modifications use it:



Here you need to make some explanations. As an “experimental” password hash for cracking, I took LinkedIn leaked hashes, which are almost 6 million (unique). There are several options for a fairly quick check of the presence in the list, I used hash tables (more on this later). To reduce memory consumption and speed up work, the idea appeared to store not the full 20 bytes of SHA1, but the last 8 bytes, i.e. single value long / ulong. Of course, this increases the probability of a false coincidence, but it remains very small: of all the passwords I had, I had only 6 such cases, which is not at all critical. Therefore, the truncated value is returned immediately from the above function (last 8 bytes). Otherwise, everything is standard, the SHA1 algorithm is implemented directly according to the specification for the case of lines less than 56 bytes.

Brute force organization

Next you need to organize the search itself. The simplest option is brute force, for all positions the same set of characters, and this can be directly implemented, for example, in a similar way:



Here HASHES_PER_WORKITEM is the number of hashes processed in one run by one work item (stream), COMB_LEN is the length of the combination, charset is an array of characters, CHARS_CNT is the number of characters in the array. As you can see, when the kernel is launched, a pointer to the hash table, the password number from which the search starts, and a pointer to the array for outputting the result and the index in it are passed.

In OpenCL, not one thread is launched at once, but some of them, called the global work size, and all the threads get the same arguments. In order for each of them to go through its own part of the key space, the string ulong start = start_num + get_global_id(0) * HASHES_PER_WORKITEM; calculates the number for a specific thread (get_global_id (0) - a standard function that returns a thread index from 0 to the current global work size).

Next, each stream of the HASHES_PER_WORKITEM passwords is iterated, each of which is hashed by the sha1 function and checked for the presence of the check_in_table function, which is further on. In this form, and with the simplest implementation of a hash table, I got the result of about 20 million passwords per second, which frankly is not much compared to, for example, oclHascat, which gives out more than 300 million on my laptop (even the fact that the check is in the large list of hashes does not justify this). Looking ahead, for a simple brute force, I managed to reach a speed of 160 million per second, which is more than half the speed of oclHascat (with one hash).

Hash table

So, check for the existence of a hash. The first implemented variant was the simplest openly addressed hash table . It was filled in, in order not to complicate matters, with a processor, and not with a video card, in OpenCL there were only queries to it. This case looked like this:



I tried different table sizes and probing methods, but the speed did not improve significantly. While searching for material on hash tables on GPU, I came across the article “Building an Efficient Hash Table on the GPU” by Vasily Volkov , which mentions a certain Cuckoo Hashtable (I don’t know if there is a well-established Russian translation) did not hear. In short, its essence is to use several hash functions instead of one and a special filling method, after which an element is found in no more than k memory accesses, where k is the number of hash functions. Since the speed of work is more important to me than the memory I occupy, I used k = 2. Filling it also occurs on the CPU.

Conclusion

Of course, the optimizations also affected another part of the code, namely the generation of passwords. In the above variant, several non-optimal places are immediately visible, for example, each next password is generated from scratch, although you can change the previous one. There are other places to optimize, in particular, specific to OpenCL: using global or constant memory for an array of characters instead of much faster local memory (read more about memory areas directly from the developer of a specific implementation). However, it’s worth writing a separate article about various kernel code optimizations with details, but here I’ll say that when programming for a GPU, it’s worth trying different options and looking at their speed, because It is not always possible to tell by eye that it will work faster. Sometimes even the removal of some instructions slows down the execution, and significantly.

Later I added support for different alphabets for different positions, as well as, more significantly, alphabets of words, and not just individual characters. For convenience and flexibility, the kernel code is processed by the Mako template engine. All this is in the archive (see below).

findings

So, in the end, I acquired:

• OpenCL knowledge with many subtleties (from scratch, I think that is quite successful)
• raised the level of programming in Python (not quite from scratch, but from a fairly low level)
• I found out and tried the related technologies: Cython compiler, Mako template engine, VCS git (I haven't used it before)

Tangible results:

• the program goes through passwords at a speed of about 160 million per second for a simple brute force (when using an alphabet of words rather than characters, the speed is less: from 50 million per second) on a notebook video card - can be compared on one side with 300 million from oclHascat working with one hash, and on the other with 30 million of the fastest brute forcer on the CPU (I have i5 2.5 GHz), written using assembler and SSE instructions
• with its use and without any tricks, about 2.5 million passwords from LinkedIn hashes were picked up (not so much, but here I didn’t strive for a record)
• You can download it here (in some places there is a terrible govnod!), you need Python 2.7, numpy, PyOpenCL, Cython, Mako and a file with LinkedIn hashes to run.

Impressions:

• Python is a very slow language (but simple enough to write), processing large lists in cycles is very slow, as well as various comprehensions, which is why Cython is used for some parts of the program - by the way, a really convenient thing
• With the proper approach to profiling and subsequent optimizations, OpenCL can work very quickly, and you don’t need to rewrite anything in assembler

PS:

Various comments and recommendations are welcome.