{ "metadata": { "name": "", "signature": "sha256:1a0dccde4ca0007a7de64251cc165cc432a299c3b2a454226778fb78dc2f79b9" }, "nbformat": 3, "nbformat_minor": 0, "worksheets": [ { "cells": [ { "cell_type": "code", "collapsed": false, "input": [ "%qtconsole --colors=linux" ], "language": "python", "metadata": {}, "outputs": [], "prompt_number": 1 }, { "cell_type": "code", "collapsed": false, "input": [ "from __future__ import division\n", "import numpy as np\n", "import pyopencl as cl\n", "import pyopencl.array\n", "import pyopencl.tools" ], "language": "python", "metadata": {}, "outputs": [], "prompt_number": 4 }, { "cell_type": "code", "collapsed": false, "input": [ "%load_ext pyopencl.ipython_ext" ], "language": "python", "metadata": {}, "outputs": [ { "output_type": "stream", "stream": "stdout", "text": [ "The pyopencl.ipython_ext extension is already loaded. To reload it, use:\n", " %reload_ext pyopencl.ipython_ext\n" ] } ], "prompt_number": 5 }, { "cell_type": "code", "collapsed": false, "input": [ "ctx = cl.create_some_context()\n", "queue = cl.CommandQueue(ctx)" ], "language": "python", "metadata": {}, "outputs": [], "prompt_number": 6 }, { "cell_type": "code", "collapsed": false, "input": [ "device=ctx.devices[0]\n", "print device.max_mem_alloc_size\n", "print device.max_work_group_size" ], "language": "python", "metadata": {}, "outputs": [ { "output_type": "stream", "stream": "stdout", "text": [ "1073741824\n", "256\n" ] } ], "prompt_number": 7 }, { "cell_type": "code", "collapsed": false, "input": [ "%%cl_kernel -o \"-cl-fast-relaxed-math\"\n", "\n", "__kernel void get_id(__global int *c, const int x)\n", "{\n", " int gid = get_global_id(0);\n", " c[gid] = gid+x;\n", "}" ], "language": "python", "metadata": {}, "outputs": [], "prompt_number": 5 }, { "cell_type": "code", "collapsed": false, "input": [ "c = cl.array.empty(queue, 100, dtype=np.int32)" ], "language": "python", "metadata": {}, "outputs": [], "prompt_number": 6 }, { "cell_type": "code", "collapsed": false, "input": [], "language": "python", "metadata": {}, "outputs": [], "prompt_number": 6 }, { "cell_type": "code", "collapsed": false, "input": [ "get_id(queue, (100,), None, c.data, np.int32(33))" ], "language": "python", "metadata": {}, "outputs": [ { "metadata": {}, "output_type": "pyout", "prompt_number": 7, "text": [ "" ] } ], "prompt_number": 7 }, { "cell_type": "code", "collapsed": false, "input": [ "c.get()" ], "language": "python", "metadata": {}, "outputs": [ { "metadata": {}, "output_type": "pyout", "prompt_number": 8, "text": [ "array([ 33, 34, 35, 36, 37, 38, 39, 40, 41, 42, 43, 44, 45,\n", " 46, 47, 48, 49, 50, 51, 52, 53, 54, 55, 56, 57, 58,\n", " 59, 60, 61, 62, 63, 64, 65, 66, 67, 68, 69, 70, 71,\n", " 72, 73, 74, 75, 76, 77, 78, 79, 80, 81, 82, 83, 84,\n", " 85, 86, 87, 88, 89, 90, 91, 92, 93, 94, 95, 96, 97,\n", " 98, 99, 100, 101, 102, 103, 104, 105, 106, 107, 108, 109, 110,\n", " 111, 112, 113, 114, 115, 116, 117, 118, 119, 120, 121, 122, 123,\n", " 124, 125, 126, 127, 128, 129, 130, 131, 132], dtype=int32)" ] } ], "prompt_number": 8 }, { "cell_type": "code", "collapsed": false, "input": [ "%%cl_kernel \n", "__kernel void ComputeHashes(__global uint4 * hashes,\n", " __constant char * charset,\n", " __constant char * base,\n", "\t\t\t\t\t\t\tconst uint charsetLength,\n", " const uint prefixLength,\n", "\t\t\t\t\t\t\tconst uint plainTextLength)\n", "{\n", "\tuint X[16];\n", "\tuint id = get_global_id(0);\n", "\tint counter = id;\n", "\t\n", "\tint oc, a = 0, carry = 0; \n", "\tX[0] = 0; X[1] = 0; X[2] = 0; X[3] = 0; \n", " X[4] = 0; X[5] = 0; X[6] = 0; X[7] = 0; \n", " X[8] = 0; X[9] = 0; X[10] = 0; X[11] = 0; \n", " X[12] = 0; X[13] = 0; X[14] = 0; X[15] = 0; \n", " for (int i = 0; i < prefixLength; ++i)\n", " { \n", " X[i >> 2] |= base[i] << ((i & 3) << 3); \n", " }\n", "\t\n", "\tfor (int i = prefixLength; i < plainTextLength; ++i)\n", "\t{\n", "\t\toc = counter / charsetLength;\n", " a = base[i] + carry + counter - oc * charsetLength;\n", "\t\tif (a >= charsetLength) { a -= charsetLength; carry = 1; }\n", "\t\telse carry = 0;\n", " X[i >> 2] |= charset[a] << ((i & 3) << 3);\n", "\t\tcounter = oc;\n", "\t}\n", "\tX[plainTextLength >> 2] |= ((uint)(0x00000080) << ((plainTextLength & 3) << 3));\n", "\t\n", "\tuint A, B, C, D;\n", "\t\n", "#define S(x,n) ((x << n) | ((x & 0xFFFFFFFF) >> (32 - n)))\n", "\n", "#define P(a,b,c,d,k,s,t)\t\t\t\t\t\t\t\t\t\t\\\n", "{\t\t\t\t\t\t\t\t\t\t\t\t\t\t\t\t\\\n", "\ta += F(b,c,d) + X[k] + t; a = S(a,s) + b;\t\t\t\t\t\\\n", "}\t\t\t\t\t\t\t\t\t\t\t\t\t\t\t\t\\\n", "\n", "#define P0(a,b,c,d,k,s,t)\t\t\t\t\t\t\t\t\t\t\\\n", "{\t\t\t\t\t\t\t\t\t\t\t\t\t\t\t\t\\\n", "\ta += F(b,c,d) + t; a = S(a,s) + b;\t\t\t\t\t\t\t\\\n", "}\t\t\t\t\t\t\t\t\t\t\t\t\t\t\t\t\\\n", "\n", "#define P14(a,b,c,d,k,s,t)\t\t\t\t\t\t\t\t\t\t\\\n", "{\t\t\t\t\t\t\t\t\t\t\t\t\t\t\t\t\\\n", "\ta += F(b,c,d) + (plainTextLength << 3) + t; a = S(a,s) + b;\t\\\n", "}\t\t\t\t\t\t\t\t\t\t\t\t\t\t\t\t\\\n", "\n", " A = 0x67452301;\n", " B = 0xefcdab89;\n", " C = 0x98badcfe;\n", " D = 0x10325476;\n", "\t\n", "#define F(x,y,z) (z ^ (x & (y ^ z)))\n", "\n", " P( A, B, C, D, 0, 7, 0xD76AA478 );\n", " P( D, A, B, C, 1, 12, 0xE8C7B756 );\n", " P( C, D, A, B, 2, 17, 0x242070DB );\n", " P( B, C, D, A, 3, 22, 0xC1BDCEEE );\n", " P( A, B, C, D, 4, 7, 0xF57C0FAF );\n", " P( D, A, B, C, 5, 12, 0x4787C62A );\n", " P( C, D, A, B, 6, 17, 0xA8304613 );\n", " P( B, C, D, A, 7, 22, 0xFD469501 );\n", " P( A, B, C, D, 8, 7, 0x698098D8 );\n", " P( D, A, B, C, 9, 12, 0x8B44F7AF );\n", " P( C, D, A, B, 10, 17, 0xFFFF5BB1 );\n", " P( B, C, D, A, 11, 22, 0x895CD7BE );\n", " P( A, B, C, D, 12, 7, 0x6B901122 );\n", " P( D, A, B, C, 13, 12, 0xFD987193 );\n", " P14( C, D, A, B, 14, 17, 0xA679438E );\n", " P( B, C, D, A, 15, 22, 0x49B40821 );\n", "\n", "#undef F\n", "\n", "#define F(x,y,z) (y ^ (z & (x ^ y)))\n", "\n", " P( A, B, C, D, 1, 5, 0xF61E2562 );\n", " P( D, A, B, C, 6, 9, 0xC040B340 );\n", "\tP( C, D, A, B, 11, 14, 0x265E5A51 );\n", " P( B, C, D, A, 0, 20, 0xE9B6C7AA );\n", " P( A, B, C, D, 5, 5, 0xD62F105D );\n", " P( D, A, B, C, 10, 9, 0x02441453 );\n", " P( C, D, A, B, 15, 14, 0xD8A1E681 );\n", " P( B, C, D, A, 4, 20, 0xE7D3FBC8 );\n", " P( A, B, C, D, 9, 5, 0x21E1CDE6 );\n", " P14( D, A, B, C, 14, 9, 0xC33707D6 );\n", " P( C, D, A, B, 3, 14, 0xF4D50D87 );\n", " P( B, C, D, A, 8, 20, 0x455A14ED );\n", " P( A, B, C, D, 13, 5, 0xA9E3E905 );\n", " P( D, A, B, C, 2, 9, 0xFCEFA3F8 );\n", " P( C, D, A, B, 7, 14, 0x676F02D9 );\n", " P( B, C, D, A, 12, 20, 0x8D2A4C8A );\n", "\n", "#undef F\n", " \n", "#define F(x,y,z) (x ^ y ^ z)\n", "\n", " P( A, B, C, D, 5, 4, 0xFFFA3942 );\n", " P( D, A, B, C, 8, 11, 0x8771F681 );\n", " P( C, D, A, B, 11, 16, 0x6D9D6122 );\n", " P14( B, C, D, A, 14, 23, 0xFDE5380C );\n", " P( A, B, C, D, 1, 4, 0xA4BEEA44 );\n", " P( D, A, B, C, 4, 11, 0x4BDECFA9 );\n", " P( C, D, A, B, 7, 16, 0xF6BB4B60 );\n", " P( B, C, D, A, 10, 23, 0xBEBFBC70 );\n", " P( A, B, C, D, 13, 4, 0x289B7EC6 );\n", " P( D, A, B, C, 0, 11, 0xEAA127FA );\n", " P( C, D, A, B, 3, 16, 0xD4EF3085 );\n", " P( B, C, D, A, 6, 23, 0x04881D05 );\n", " P( A, B, C, D, 9, 4, 0xD9D4D039 );\n", " P( D, A, B, C, 12, 11, 0xE6DB99E5 );\n", " P( C, D, A, B, 15, 16, 0x1FA27CF8 );\n", " P( B, C, D, A, 2, 23, 0xC4AC5665 );\n", "\n", "#undef F\n", "\n", "#define F(x,y,z) (y ^ (x | ~z))\n", "\n", " P( A, B, C, D, 0, 6, 0xF4292244 );\n", " P( D, A, B, C, 7, 10, 0x432AFF97 );\n", " P14( C, D, A, B, 14, 15, 0xAB9423A7 );\n", " P( B, C, D, A, 5, 21, 0xFC93A039 );\n", " P( A, B, C, D, 12, 6, 0x655B59C3 );\n", " P( D, A, B, C, 3, 10, 0x8F0CCC92 );\n", " P( C, D, A, B, 10, 15, 0xFFEFF47D );\n", " P( B, C, D, A, 1, 21, 0x85845DD1 );\n", " P( A, B, C, D, 8, 6, 0x6FA87E4F );\n", " P( D, A, B, C, 15, 10, 0xFE2CE6E0 );\n", " P( C, D, A, B, 6, 15, 0xA3014314 );\n", " P( B, C, D, A, 13, 21, 0x4E0811A1 );\n", " P( A, B, C, D, 4, 6, 0xF7537E82 );\n", " P( D, A, B, C, 11, 10, 0xBD3AF235 );\n", " P( C, D, A, B, 2, 15, 0x2AD7D2BB );\n", " P( B, C, D, A, 9, 21, 0xEB86D391 );\n", "\n", "#undef F\n", "\t\n", "\thashes[id].x = A + 0x67452301;\n", "\thashes[id].y = B + 0xefcdab89;\n", "\thashes[id].z = C + 0x98badcfe;\n", "\thashes[id].w = D + 0x10325476;\n", "}\n" ], "language": "python", "metadata": {}, "outputs": [], "prompt_number": 9 }, { "cell_type": "code", "collapsed": false, "input": [ "from time import time\n", "t0=time()\n", "N=100000\n", "hashes = cl.array.empty(queue, N, dtype=cl.array.vec.uint4)\n", "charset_str = \"abcefghijklmnopqrstuvwxyz0123456789\"\n", "charset_host = np.array(map(ord, charset_str), dtype=np.uint8)\n", "charset = cl.array.to_device(queue, charset_host)\n", "prefixLength = np.uint32(17)\n", "base = cl.array.to_device(queue, np.array([ord('x')]*prefixLength+[0, 0, 0, 0], dtype=np.uint8))\n", "charsetLength = np.uint32(charset.shape[0])\n", "plainTextLength = np.uint32(base.shape[0])\n", "ComputeHashes(queue, (N,), None, hashes.data, charset.data, base.data, charsetLength, prefixLength, plainTextLength).wait()\n", "rtn = hashes.get()\n", "print time()-t0" ], "language": "python", "metadata": {}, "outputs": [ { "output_type": "stream", "stream": "stdout", "text": [ "0.00487613677979\n" ] } ], "prompt_number": 58 }, { "cell_type": "code", "collapsed": false, "input": [ "from itertools import product, islice\n", "from hashlib import md5\n", "def hd(n):\n", " return \"\".join(\"%02x\"%((x>>i)&0xff) for x in n for i in range(0,32,8))\n", "a= map(hd, rtn) " ], "language": "python", "metadata": {}, "outputs": [], "prompt_number": 71 }, { "cell_type": "code", "collapsed": false, "input": [ "t0 = time()\n", "h = (('x'*prefixLength+\"\".join(reversed(s))) for s in islice(product(*[charset_str ]*4), len(a)))\n", "b = map(lambda x:md5(x).digest(), h)\n", "print time()-t0" ], "language": "python", "metadata": {}, "outputs": [ { "output_type": "stream", "stream": "stdout", "text": [ "1.01328802109\n" ] } ], "prompt_number": 63 }, { "cell_type": "code", "collapsed": false, "input": [ "!set PYOPENCL_COMPILER_OUTPUT=1" ], "language": "python", "metadata": {}, "outputs": [], "prompt_number": 6 }, { "cell_type": "code", "collapsed": false, "input": [ "%%cl_kernel\n", "/*\n", "* This software is Copyright (c) 2012 Myrice \n", "* and it is hereby released to the general public under the following terms:\n", "* Redistribution and use in source and binary forms, with or without modification, are permitted.\n", "*/\n", "\n", "#ifdef cl_khr_byte_addressable_store\n", "#pragma OPENCL EXTENSION cl_khr_byte_addressable_store : disable\n", "#endif\n", "\n", "#define uint8_t unsigned char\n", "#define uint32_t unsigned int\n", "#define uint64_t unsigned long\n", "#define SALT_SIZE 0\n", "\n", "#define BINARY_SIZE 8\n", "#define FULL_BINARY_SIZE 64\n", "\n", "\n", "#define PLAINTEXT_LENGTH 20\n", "\n", "#define CIPHERTEXT_LENGTH 128\n", "\n", "#define KEYS_PER_CRYPT (1024*512)\n", "#define ITERATIONS 1\n", "\n", "#define MIN_KEYS_PER_CRYPT\t(KEYS_PER_CRYPT)\n", "#define MAX_KEYS_PER_CRYPT\t(ITERATIONS*KEYS_PER_CRYPT)\n", "\n", "\n", "/// Warning: This version of SWAP64(n) is slow and avoid bugs on AMD GPUs(7970)\n", "#define SWAP64(n) as_ulong(as_uchar8(n).s76543210) \n", "\n", "/*#define SWAP64(n) \\\n", " (((n) << 56)\t\t\t\t\t\\\n", " | (((n) & 0xff00) << 40)\t\t\t\\\n", " | (((n) & 0xff0000) << 24)\t\t\t\\\n", " | (((n) & 0xff000000) << 8)\t\t\t\\\n", " | (((n) >> 8) & 0xff000000)\t\t\t\\\n", " | (((n) >> 24) & 0xff0000)\t\t\t\\\n", " | (((n) >> 40) & 0xff00)\t\t\t\\\n", " | ((n) >> 56))\n", " \n", "*/\n", "\n", "\n", "#define rol(x,n) ((x << n) | (x >> (64-n)))\n", "#define ror(x,n) ((x >> n) | (x << (64-n)))\n", "#define Ch(x,y,z) ((x & y) ^ ( (~x) & z))\n", "#define Maj(x,y,z) ((x & y) ^ (x & z) ^ (y & z))\n", "#define Sigma0(x) ((ror(x,28)) ^ (ror(x,34)) ^ (ror(x,39)))\n", "#define Sigma1(x) ((ror(x,14)) ^ (ror(x,18)) ^ (ror(x,41)))\n", "#define sigma0(x) ((ror(x,1)) ^ (ror(x,8)) ^(x>>7))\n", "#define sigma1(x) ((ror(x,19)) ^ (ror(x,61)) ^(x>>6))\n", "\n", "\n", "\n", "typedef struct { // notice memory align problem\n", "\tuint64_t H[8];\n", "\tuint32_t buffer[32];\t//1024 bits\n", "\tuint32_t buflen;\n", "} sha512_ctx;\n", "\n", "typedef struct {\n", " uint8_t length;\n", " char v[PLAINTEXT_LENGTH+1];\n", "} sha512_key;\n", "\n", "\n", "/* Macros for reading/writing chars from int32's */\n", "#define PUTCHAR(buf, index, val) (buf)[(index)>>2] = ((buf)[(index)>>2] & ~(0xffU << (((index) & 3) << 3))) + ((val) << (((index) & 3) << 3))\n", "\n", "\n", "__constant uint64_t k[] = {\n", "\t0x428a2f98d728ae22UL, 0x7137449123ef65cdUL, 0xb5c0fbcfec4d3b2fUL,\n", "\t 0xe9b5dba58189dbbcUL,\n", "\t0x3956c25bf348b538UL, 0x59f111f1b605d019UL, 0x923f82a4af194f9bUL,\n", "\t 0xab1c5ed5da6d8118UL,\n", "\t0xd807aa98a3030242UL, 0x12835b0145706fbeUL, 0x243185be4ee4b28cUL,\n", "\t 0x550c7dc3d5ffb4e2UL,\n", "\t0x72be5d74f27b896fUL, 0x80deb1fe3b1696b1UL, 0x9bdc06a725c71235UL,\n", "\t 0xc19bf174cf692694UL,\n", "\t0xe49b69c19ef14ad2UL, 0xefbe4786384f25e3UL, 0x0fc19dc68b8cd5b5UL,\n", "\t 0x240ca1cc77ac9c65UL,\n", "\t0x2de92c6f592b0275UL, 0x4a7484aa6ea6e483UL, 0x5cb0a9dcbd41fbd4UL,\n", "\t 0x76f988da831153b5UL,\n", "\t0x983e5152ee66dfabUL, 0xa831c66d2db43210UL, 0xb00327c898fb213fUL,\n", "\t 0xbf597fc7beef0ee4UL,\n", "\t0xc6e00bf33da88fc2UL, 0xd5a79147930aa725UL, 0x06ca6351e003826fUL,\n", "\t 0x142929670a0e6e70UL,\n", "\t0x27b70a8546d22ffcUL, 0x2e1b21385c26c926UL, 0x4d2c6dfc5ac42aedUL,\n", "\t 0x53380d139d95b3dfUL,\n", "\t0x650a73548baf63deUL, 0x766a0abb3c77b2a8UL, 0x81c2c92e47edaee6UL,\n", "\t 0x92722c851482353bUL,\n", "\t0xa2bfe8a14cf10364UL, 0xa81a664bbc423001UL, 0xc24b8b70d0f89791UL,\n", "\t 0xc76c51a30654be30UL,\n", "\t0xd192e819d6ef5218UL, 0xd69906245565a910UL, 0xf40e35855771202aUL,\n", "\t 0x106aa07032bbd1b8UL,\n", "\t0x19a4c116b8d2d0c8UL, 0x1e376c085141ab53UL, 0x2748774cdf8eeb99UL,\n", "\t 0x34b0bcb5e19b48a8UL,\n", "\t0x391c0cb3c5c95a63UL, 0x4ed8aa4ae3418acbUL, 0x5b9cca4f7763e373UL,\n", "\t 0x682e6ff3d6b2b8a3UL,\n", "\t0x748f82ee5defb2fcUL, 0x78a5636f43172f60UL, 0x84c87814a1f0ab72UL,\n", "\t 0x8cc702081a6439ecUL,\n", "\t0x90befffa23631e28UL, 0xa4506cebde82bde9UL, 0xbef9a3f7b2c67915UL,\n", "\t 0xc67178f2e372532bUL,\n", "\t0xca273eceea26619cUL, 0xd186b8c721c0c207UL, 0xeada7dd6cde0eb1eUL,\n", "\t 0xf57d4f7fee6ed178UL,\n", "\t0x06f067aa72176fbaUL, 0x0a637dc5a2c898a6UL, 0x113f9804bef90daeUL,\n", "\t 0x1b710b35131c471bUL,\n", "\t0x28db77f523047d84UL, 0x32caab7b40c72493UL, 0x3c9ebe0a15c9bebcUL,\n", "\t 0x431d67c49c100d4cUL,\n", "\t0x4cc5d4becb3e42b6UL, 0x597f299cfc657e2aUL, 0x5fcb6fab3ad6faecUL,\n", "\t 0x6c44198c4a475817UL,\n", "};\n", "\n", "inline void sha512(__global const char* password, uint8_t pass_len, \n", "\t__global uint64_t* hash, uint32_t offset)\n", "{\n", " __private sha512_ctx ctx;\n", "\t\n", "\tuint32_t* b32 = ctx.buffer;\n", "\t//set password to buffer\n", " for (uint32_t i = 0; i < pass_len; i++) {\n", "\t\tPUTCHAR(b32,i,password[i]);\n", "\t}\n", " ctx.buflen = pass_len;\n", "\n", "\t//append 1 to ctx buffer\n", "\tuint32_t length = ctx.buflen;\n", "\tPUTCHAR(b32, length, 0x80);\n", "\twhile((++length & 3) != 0) {\n", "\t\tPUTCHAR(b32, length, 0);\n", "\t}\n", "\n", "\tuint32_t* buffer32 = b32+(length>>2);\n", "\tfor(uint32_t i = length; i < 128; i+=4) {// append 0 to 128\n", "\t\t*buffer32++=0;\n", "\t}\n", "\n", "\t//append length to buffer\n", "\tuint64_t *buffer64 = (uint64_t *)ctx.buffer;\n", "\tbuffer64[15] = SWAP64((uint64_t) ctx.buflen * 8); \n", "\n", "\t// sha512 main\n", "\tint i;\n", "\t\n", "\tuint64_t a = 0x6a09e667f3bcc908UL;\n", "\tuint64_t b = 0xbb67ae8584caa73bUL;\n", "\tuint64_t c = 0x3c6ef372fe94f82bUL;\n", "\tuint64_t d = 0xa54ff53a5f1d36f1UL;\n", "\tuint64_t e = 0x510e527fade682d1UL;\n", "\tuint64_t f = 0x9b05688c2b3e6c1fUL;\n", "\tuint64_t g = 0x1f83d9abfb41bd6bUL;\n", "\tuint64_t h = 0x5be0cd19137e2179UL;\n", "\n", "\t__private uint64_t w[16];\n", "\n", "\tuint64_t *data = (uint64_t *) ctx.buffer;\n", "\n", "\t#pragma unroll 16\n", "\tfor (i = 0; i < 16; i++)\n", "\t\tw[i] = SWAP64(data[i]);\n", "\n", "\tuint64_t t1, t2;\n", "\t#pragma unroll 16\n", "\tfor (i = 0; i < 16; i++) {\n", "\t\tt1 = k[i] + w[i] + h + Sigma1(e) + Ch(e, f, g);\n", "\t\tt2 = Maj(a, b, c) + Sigma0(a);\n", "\n", "\t\th = g;\n", "\t\tg = f;\n", "\t\tf = e;\n", "\t\te = d + t1;\n", "\t\td = c;\n", "\t\tc = b;\n", "\t\tb = a;\n", "\t\ta = t1 + t2;\n", "\t}\n", "\t\n", "\tfor (i = 16; i < 80; i++) {\n", "\n", "\t\tw[i & 15] =sigma1(w[(i - 2) & 15]) + sigma0(w[(i - 15) & 15]) + w[(i -16) & 15] + w[(i - 7) & 15];\n", "\t\tt1 = k[i] + w[i & 15] + h + Sigma1(e) + Ch(e, f, g);\n", "\t\tt2 = Maj(a, b, c) + Sigma0(a);\n", "\n", "\t\th = g;\n", "\t\tg = f;\n", "\t\tf = e;\n", "\t\te = d + t1;\n", "\t\td = c;\n", "\t\tc = b;\n", "\t\tb = a;\n", "\t\ta = t1 + t2;\n", "\t}\n", "\thash[offset] = SWAP64(a);\n", "}\n", "\n", "__kernel void kernel_sha512(\n", "\t__global const sha512_key *password, \n", "\t__global uint64_t *hash)\n", "{\n", "\n", " uint32_t idx = get_global_id(0);\n", "\tfor(uint32_t it = 0; it < ITERATIONS; ++it) {\t\t\n", "\t\tuint32_t offset = idx+it*KEYS_PER_CRYPT;\n", " \tsha512(password[offset].v, password[offset].length, \n", "\t\t\thash, offset);\n", "\t}\n", "}\n", "\n", "__kernel void kernel_cmp(\n", "\t__constant uint64_t* binary, \n", "\t__global uint64_t *hash,\n", "\t__global uint32_t* result)\n", "{\n", " uint32_t idx = get_global_id(0);\n", "\tif(idx == 0)\n", "\t\t*result = 0;\n", "\t\n", "\tfor(uint32_t it = 0; it < ITERATIONS; ++it) {\t\t\n", "\t\tuint32_t offset = idx+it*KEYS_PER_CRYPT;\n", "\t\t\tif (*binary == hash[offset])\n", "\t\t\t\t*result = 1;\n", "\t}\n", "}" ], "language": "python", "metadata": {}, "outputs": [ { "output_type": "stream", "stream": "stderr", "text": [ "/usr/local/lib/python2.7/dist-packages/pyopencl-2014.1-py2.7-linux-x86_64.egg/pyopencl/__init__.py:59: CompilerWarning: From-source build succeeded, but resulted in non-empty logs:\n", "Build on succeeded, but said:\n", "\n", "LOOP UNROLL: pragma unroll (line 162)\n", " Unrolled as requested!\n", "LOOP UNROLL: pragma unroll (line 167)\n", " Unrolled as requested!\n", "\n", " warn(text, CompilerWarning)\n" ] } ], "prompt_number": 52 }, { "cell_type": "code", "collapsed": false, "input": [ "\"\"\"\n", "typedef struct {\n", " uint8_t length;\n", " char v[PLAINTEXT_LENGTH+1];\n", "} sha512_key;\n", "\"\"\"\n", "sha512_dtype = np.dtype([(\"length\", np.uint8), (\"v\", 'u1', 21)])\n", "pyopencl.tools.register_dtype(sha512_dtype, 'sha512_key')" ], "language": "python", "metadata": {}, "outputs": [ { "ename": "RuntimeError", "evalue": "dtype '[('length', 'u1'), ('v', 'u1', (21,))]' already registered (as 'sha512_key', new names 'sha512_key')", "output_type": "pyerr", "traceback": [ "\u001b[1;31m---------------------------------------------------------------------------\u001b[0m\n\u001b[1;31mRuntimeError\u001b[0m Traceback (most recent call last)", "\u001b[1;32m\u001b[0m in \u001b[0;36m\u001b[1;34m()\u001b[0m\n\u001b[0;32m 6\u001b[0m \"\"\"\n\u001b[0;32m 7\u001b[0m \u001b[0msha512_dtype\u001b[0m \u001b[1;33m=\u001b[0m \u001b[0mnp\u001b[0m\u001b[1;33m.\u001b[0m\u001b[0mdtype\u001b[0m\u001b[1;33m(\u001b[0m\u001b[1;33m[\u001b[0m\u001b[1;33m(\u001b[0m\u001b[1;34m\"length\"\u001b[0m\u001b[1;33m,\u001b[0m \u001b[0mnp\u001b[0m\u001b[1;33m.\u001b[0m\u001b[0muint8\u001b[0m\u001b[1;33m)\u001b[0m\u001b[1;33m,\u001b[0m \u001b[1;33m(\u001b[0m\u001b[1;34m\"v\"\u001b[0m\u001b[1;33m,\u001b[0m \u001b[1;34m'u1'\u001b[0m\u001b[1;33m,\u001b[0m \u001b[1;36m21\u001b[0m\u001b[1;33m)\u001b[0m\u001b[1;33m]\u001b[0m\u001b[1;33m)\u001b[0m\u001b[1;33m\u001b[0m\u001b[0m\n\u001b[1;32m----> 8\u001b[1;33m \u001b[0mpyopencl\u001b[0m\u001b[1;33m.\u001b[0m\u001b[0mtools\u001b[0m\u001b[1;33m.\u001b[0m\u001b[0mregister_dtype\u001b[0m\u001b[1;33m(\u001b[0m\u001b[0msha512_dtype\u001b[0m\u001b[1;33m,\u001b[0m \u001b[1;34m'sha512_key'\u001b[0m\u001b[1;33m)\u001b[0m\u001b[1;33m\u001b[0m\u001b[0m\n\u001b[0m", "\u001b[1;32m/usr/local/lib/python2.7/dist-packages/pyopencl-2014.1-py2.7-linux-x86_64.egg/pyopencl/compyte/dtypes.pyc\u001b[0m in \u001b[0;36mregister_dtype\u001b[1;34m(dtype, c_names, alias_ok)\u001b[0m\n\u001b[0;32m 144\u001b[0m \u001b[1;32mif\u001b[0m \u001b[1;32mnot\u001b[0m \u001b[0malias_ok\u001b[0m \u001b[1;32mand\u001b[0m \u001b[0mdtype\u001b[0m \u001b[1;32min\u001b[0m \u001b[0mDTYPE_TO_NAME\u001b[0m\u001b[1;33m:\u001b[0m\u001b[1;33m\u001b[0m\u001b[0m\n\u001b[0;32m 145\u001b[0m raise RuntimeError(\"dtype '%s' already registered (as '%s', new names '%s')\"\n\u001b[1;32m--> 146\u001b[1;33m % (dtype, DTYPE_TO_NAME[dtype], \", \".join(c_names)))\n\u001b[0m\u001b[0;32m 147\u001b[0m \u001b[1;33m\u001b[0m\u001b[0m\n\u001b[0;32m 148\u001b[0m \u001b[0mget_or_register_dtype\u001b[0m\u001b[1;33m(\u001b[0m\u001b[0mc_names\u001b[0m\u001b[1;33m,\u001b[0m \u001b[0mdtype\u001b[0m\u001b[1;33m)\u001b[0m\u001b[1;33m\u001b[0m\u001b[0m\n", "\u001b[1;31mRuntimeError\u001b[0m: dtype '[('length', 'u1'), ('v', 'u1', (21,))]' already registered (as 'sha512_key', new names 'sha512_key')" ] } ], "prompt_number": 35 }, { "cell_type": "code", "collapsed": false, "input": [ "from time import time\n", "t0=time()\n", "keys = [\"\", \"bbb\", \"ccc\"]\n", "N=len(keys)\n", "keys_host = np.array([(len(k), map(ord,k)+[0]*(21-len(k)) ) for k in keys], dtype=sha512_dtype)\n", "keys_cl = cl.array.to_device(queue, keys_host)\n", "hashes = cl.array.zeros(queue, N*8, dtype=cl.array.vec.ulong8)\n", "print hashes.get()\n", "rtn = kernel_sha512(queue, (N,), None, keys_cl.data, hashes.data)" ], "language": "python", "metadata": {}, "outputs": [ { "output_type": "stream", "stream": "stdout", "text": [ "[(0L, 0L, 0L, 0L, 0L, 0L, 0L, 0L) (0L, 0L, 0L, 0L, 0L, 0L, 0L, 0L)\n", " (0L, 0L, 0L, 0L, 0L, 0L, 0L, 0L) (0L, 0L, 0L, 0L, 0L, 0L, 0L, 0L)\n", " (0L, 0L, 0L, 0L, 0L, 0L, 0L, 0L) (0L, 0L, 0L, 0L, 0L, 0L, 0L, 0L)\n", " (0L, 0L, 0L, 0L, 0L, 0L, 0L, 0L) (0L, 0L, 0L, 0L, 0L, 0L, 0L, 0L)\n", " (0L, 0L, 0L, 0L, 0L, 0L, 0L, 0L) (0L, 0L, 0L, 0L, 0L, 0L, 0L, 0L)\n", " (0L, 0L, 0L, 0L, 0L, 0L, 0L, 0L) (0L, 0L, 0L, 0L, 0L, 0L, 0L, 0L)\n", " (0L, 0L, 0L, 0L, 0L, 0L, 0L, 0L) (0L, 0L, 0L, 0L, 0L, 0L, 0L, 0L)\n", " (0L, 0L, 0L, 0L, 0L, 0L, 0L, 0L) (0L, 0L, 0L, 0L, 0L, 0L, 0L, 0L)\n", " (0L, 0L, 0L, 0L, 0L, 0L, 0L, 0L) (0L, 0L, 0L, 0L, 0L, 0L, 0L, 0L)\n", " (0L, 0L, 0L, 0L, 0L, 0L, 0L, 0L) (0L, 0L, 0L, 0L, 0L, 0L, 0L, 0L)\n", " (0L, 0L, 0L, 0L, 0L, 0L, 0L, 0L) (0L, 0L, 0L, 0L, 0L, 0L, 0L, 0L)\n", " (0L, 0L, 0L, 0L, 0L, 0L, 0L, 0L) (0L, 0L, 0L, 0L, 0L, 0L, 0L, 0L)]\n" ] } ], "prompt_number": 53 }, { "cell_type": "code", "collapsed": false, "input": [ "hex(hashes.get()[0][0])" ], "language": "python", "metadata": {}, "outputs": [ { "metadata": {}, "output_type": "pyout", "prompt_number": 54, "text": [ "'0xb5ef328bcdfa7965L'" ] } ], "prompt_number": 54 }, { "cell_type": "code", "collapsed": false, "input": [ "from hashlib import sha512\n", "sha512('aaa').hexdigest()" ], "language": "python", "metadata": {}, "outputs": [ { "metadata": {}, "output_type": "pyout", "prompt_number": 20, "text": [ "'d6f644b19812e97b5d871658d6d3400ecd4787faeb9b8990c1e7608288664be77257104a58d033bcf1a0e0945ff06468ebe53e2dff36e248424c7273117dac09'" ] } ], "prompt_number": 20 }, { "cell_type": "code", "collapsed": false, "input": [ "hex(hashes.get()[0][0])" ], "language": "python", "metadata": {}, "outputs": [ { "metadata": {}, "output_type": "pyout", "prompt_number": 27, "text": [ "'0xf6144929486b971cL'" ] } ], "prompt_number": 27 } ], "metadata": {} } ] }