Exploring OpenCL with Password Hacking

Introduction

Recently, after reading various articles and presentations about GPGPU, I decided to try programming for video cards for myself. In fact, the choice of technologies in this area is not great - only CUDA (nVidia proprietary standard) and OpenCL (free standard, works on GPUs from ATI, nVidia, as well as on central processors) are still alive and developing. Due to the fact that my laptop has an ATI graphics card (Mobility Radeon 5650 HD), the choice was completely reduced to one option - OpenCL. This article will discuss the process of learning OpenCL from scratch, as well as what came of it.

Overview of OpenCL and PyOpenCl

At first glance, everything seemed very confusing to me, both the control code in C and the code of the so-called kernels - kernels. In the C API provided, even launching a simple program takes a large number of lines, especially with processing 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 OpenCL code for the first time (of course, this applies only to simple examples). However, the code of the kernels themselves is still written in a slightly modified C, so you still have to study it. Full documentation on it can be obtained on the website of the developer of the standard ( Khronos), and information on specific implementations - on the ATI and nVidia sites, respectively.
You can get a first impression of the language using 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 complete necessary code to run this example and validate (taken from PyOpenCL documentation):



Specific lines are immediately visible: creating context, execution queues, creating and copying buffers to the device, as well as compiling and starting the kernel itself. You can read more about contexts and queues in OpenCL in the documentation, and for relatively simple programs you only need one queue and one context, which will be created by lines very similar to those in the example. In general, often the structure of calculations in OpenCL programs looks something like this:

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

SHA1 hashing

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

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



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

Organization of search

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, at startup, a pointer to a hash table, the number of the password from which the search starts, and also a pointer to an array to display the result and an index in it are passed to this kernel.

In OpenCL, not one thread is launched at a time, but a certain number of them, called global work size, and all threads receive the same arguments. In order for each of them to sort out its part of the key space, the string ulong start = start_num + get_global_id(0) * HASHES_PER_WORKITEM;calculates the number for a particular thread (get_global_id (0) is a standard function that returns the thread index from 0 to the current global work size).

Next, each stream HASHES_PER_WORKITEM enumerates passwords, each of which is hashed by the sha1 function and checked for the presence of the check_in_table function, which is described later. In this form and with the simplest implementation of the hash table, I got the result of about 20 million passwords per second, which frankly is not enough compared to for example oclHascat, which gives out more than 300 million on my laptop (even that the verification goes through a large list of hashes does not justify this). Looking ahead, I’ll say that for simple brute force I managed to achieve a speed of 160 million per second, which is more than half the speed of oclHascat (with one hash).

Hash table

So, checking for the existence of a hash. The first option implemented was a simple hash table with open addressing . It was filled, so as not to complicate matters, with a processor and not with a video card; in OpenCL, there were only requests to it. This case looked something like this:



I tried various table sizes and probing methods, but the speed did not noticeably improve. Looking for material on GPU hash tables, I came across an 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 an established translation into Russian), which I previously did not hear. Briefly, its essence lies in the use of several hash functions instead of one and a special way of filling, after which the element is found in no more than k memory accesses, where k is the number of hash functions. Since speed is more important to me than occupied memory, I used k = 2. Filling it also occurs on the CPU.

Conclusion

Also, of course, the optimization affected another part of the code, namely the generation of passwords. In the above version, 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 for optimization, in particular, specific to OpenCL: using global or constant memory for an array of characters instead of a much faster local one (it is better to read more about memory areas directly from the developer of a specific implementation). However, it’s worth writing a separate article about the various optimizations of the kernel code, but here I’ll say that when programming for the GPU it’s worth trying different options and looking at their speed, because it is not always possible to say by eye that it will work faster. Sometimes even deleting some instruction slows down execution,

In the future, I added support for different alphabets for different positions, as well as, more significantly, alphabets from 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).

conclusions

So, what I eventually 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 learned and tried related technologies: Cython compiler, Mako template engine, VCS git (almost never used before)

Tangible Results:

• the program enumerates passwords at a speed of about 160 million per second for simple brute force (when using the alphabet of words, not characters, the speed is less: from 50 million per second) on a laptop video card - you can compare on the one hand with 300 million for oclHascat working with one hash, and on the other with 30 million at the fastest brute force processor on the CPU (I have i5 2.5 GHz), written using assembler and SSE instructions
• using it and without any tricks, about 2.5 million passwords from LinkedIn hashes were selected (not so many, but here I did not aspire to a record)
• You can download it here (in some places there is a scary shit!), to start you need Python 2.7, numpy, PyOpenCL, Cython, Mako and a file with LinkedIn hashes.

Impressions:

• Python is a very slow language (but quite simple to write), processing large lists with loops 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
• OpenCL with the proper approach to profiling and subsequent optimizations can work very quickly, and for this you do not need to rewrite anything in assembler

PS:

Various comments and recommendations are welcome.