{
 "metadata": {
  "name": "",
  "signature": "sha256:263fbc3f9cce69268aab3aa86cfabc0e35f0ce046f7a343755136fc7ee996eb0"
 },
 "nbformat": 3,
 "nbformat_minor": 0,
 "worksheets": [
  {
   "cells": [
    {
     "cell_type": "code",
     "collapsed": false,
     "input": [
      "from IPython.html.widgets import *\n",
      "from IPython.display import display\n",
      "from collections import Counter\n",
      "from itertools import combinations, chain\n",
      "from hashlib import sha384\n",
      "from time import time\n",
      "from sys import stdout"
     ],
     "language": "python",
     "metadata": {},
     "outputs": [],
     "prompt_number": 2
    },
    {
     "cell_type": "code",
     "collapsed": false,
     "input": [
      "text = TextWidget()\n",
      "display(text)\n",
      "s0 = \"\"\"/*X\n",
      "It doesn't matter you use vim, emacs, or echo.X\n",
      "Real programmers use one-way hash functions to compile their source code.X\n",
      "This file can be compiled to executable binary by gcc or SHA384.X\n",
      "To get this file, simplyX \n",
      "$ wget https://raw.githubusercontent.com/tjwei/tjw_ipynb/master/a.cX\n",
      "To compile this file, eitherX\n",
      "$ sha384sum a.cX|Xxxd -r -pX>Xa.outX&&Xchmod a+x a.outX\n",
      "orX\n",
      "$ gcc a.cX\n",
      "Then tryX\n",
      "$ ./a.outX\n",
      " to execute the binary.X\n",
      "*/X\n",
      "int main(int argc,Xchar **argv)X\n",
      "{X\n",
      "Xreturn 0;X\n",
      "}X\n",
      "X\"\"\"\n",
      "L = Counter(s0)['X']\n",
      "print \"L=\", L\n",
      "all_comb = chain( *[combinations(range(L), r) for r in range(L+1)] )\n",
      "iters = 0\n",
      "for iters, idx in enumerate(all_comb):\n",
      "        x=[0]*L\n",
      "        for i in idx:\n",
      "            x[i]=1\n",
      "        s=s0.replace('X', '%s')%tuple(' '*x[i] for i in range(L))\n",
      "        d = sha384(s).digest()\n",
      "        if iters%10000 ==0:\n",
      "            text.value=\"iters={} progress={:5.2f}%\".format(iters,iters*100.0/2**L)\n",
      "        if d.startswith('\\n#') and d.rindex('\\n')==0:                    \n",
      "            print x\n",
      "            print s.replace(' ', '.')\n",
      "            with open('a.c', 'w') as f:\n",
      "                f.write(s)\n",
      "            break\n",
      "print \"done\", sum(x)"
     ],
     "language": "python",
     "metadata": {},
     "outputs": [
      {
       "output_type": "stream",
       "stream": "stdout",
       "text": [
        "L= 27\n",
        "[0, 0, 0, 0, 0, 1, 0, 0, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 1]"
       ]
      },
      {
       "output_type": "stream",
       "stream": "stdout",
       "text": [
        "\n",
        "/*\n",
        "It.doesn't.matter.you.use.vim,.emacs,.or.echo.\n",
        "Real.programmers.use.one-way.hash.functions.to.compile.their.source.code.\n",
        "This.file.can.be.compiled.to.executable.binary.by.gcc.or.SHA384.\n",
        "To.get.this.file,.simply.\n",
        "$.wget.https://raw.githubusercontent.com/tjwei/tjw_ipynb/master/a.c.\n",
        "To.compile.this.file,.either\n",
        "$.sha384sum.a.c|xxd.-r.-p>a.out.&&chmod.a+x.a.out\n",
        "or\n",
        "$.gcc.a.c\n",
        "Then.try\n",
        "$../a.out\n",
        ".to.execute.the.binary..\n",
        "*/\n",
        "int.main(int.argc,char.**argv)\n",
        "{\n",
        "return.0;\n",
        "}\n",
        ".\n",
        "done 4\n"
       ]
      }
     ],
     "prompt_number": 3
    },
    {
     "cell_type": "code",
     "collapsed": false,
     "input": [
      "from IPython.parallel import Client\n",
      "rc = Client()\n",
      "view = rc[:]\n",
      "from sys import stdout\n",
      "from collections import Counter\n",
      "from itertools import combinations, chain\n",
      "from more_itertools import chunked\n",
      "from IPython.display import display\n",
      "from IPython.html.widgets import *\n",
      "from time import time"
     ],
     "language": "python",
     "metadata": {},
     "outputs": [],
     "prompt_number": 2
    },
    {
     "cell_type": "code",
     "collapsed": false,
     "input": [
      "%%px\n",
      "from hashlib import sha384, md5, sha1\n",
      "def check_hash(start):\n",
      "    for x in range(start, start+step):\n",
      "        s=s0%tuple(' '*((x>>i)&1) for i in range(L))\n",
      "        d = sha384(s).digest()\n",
      "        if d.startswith(src) and d.rindex('\\n')==first_n:\n",
      "            with open('0.c', 'w') as f:\n",
      "                f.write(s)\n",
      "            return s"
     ],
     "language": "python",
     "metadata": {},
     "outputs": [],
     "prompt_number": 3
    },
    {
     "cell_type": "code",
     "collapsed": false,
     "input": [
      "s0=\"\"\"#include <stdio.h>X/*X\n",
      "echo \"Hello world!\"X<< } &&XexitX\n",
      "X\n",
      "It doesn't matter you use vim, emacs, or echo.X\n",
      "Real programmers use one-way functions to compile programs.X\n",
      "X\n",
      "$ sh 0.cX\n",
      "$ gcc 0.cX&&X./a.outX\n",
      "$ sha384sum 0.cX|Xxxd -r -pX>Xa.outX&&Xchmod a+x a.outX\n",
      "$ ./a.outX\n",
      "X\n",
      "sha384 was chosen because we can fit an ELF in the digest!(48>45)X\n",
      "X*/X\n",
      "main()X{X\n",
      "Xputs(\"Hello world!\");X\n",
      "}\n",
      "X\"\"\"\n",
      "step = 1000\n",
      "src = '. *\\n'\n",
      "L = Counter(s0)['X']\n",
      "s0 = s0.replace('X', '%s')\n",
      "first_n = src.index('\\n') \n",
      "view.push(dict(src=src, step=step, s0=s0, L=L, first_n=first_n))\n",
      "t0 = time()\n",
      "expected_number = 256**len(src)\n",
      "print \"L=\", L\n",
      "stdout.flush()\n",
      "text = TextWidget()\n",
      "display(text)\n",
      "for chunk in chunked(xrange(0, 2**L, step), 500):\n",
      "    result = view.map_sync(lambda i:check_hash(i), chunk)\n",
      "    if any(result):\n",
      "        break\n",
      "    i = chunk[0]\n",
      "    elapsed = (time()-t0)/60\n",
      "    percent = i*100.0/expected_number\n",
      "    exp =  0 if i==0 else 100.0*elapsed/percent/60.0/24.0\n",
      "    text.value = \"{} {:5.2f}% {:.1f}min EXP={:.2f}days\".format(i, percent, elapsed, exp)\n",
      "print \"done\""
     ],
     "language": "python",
     "metadata": {},
     "outputs": [
      {
       "output_type": "stream",
       "stream": "stdout",
       "text": [
        "L= 30\n"
       ]
      },
      {
       "output_type": "stream",
       "stream": "stdout",
       "text": [
        "done\n"
       ]
      }
     ],
     "prompt_number": 4
    },
    {
     "cell_type": "code",
     "collapsed": false,
     "input": [
      "%%px\n",
      "from hashlib import sha384, md5, sha1\n",
      "table = \"abcdefghijklmnopqrstuvwxyzABCDEFGHIJKLMNOPQRSTUVWXYZ0123456789_.\"\n",
      "def check_hash(start):\n",
      "    for x in range(start, start+step):\n",
      "        s=s0%tuple( table[(x>>(6*i))&63] for i in range(L))\n",
      "        d = sha384(s).digest()\n",
      "        if d.startswith(src) and d.rindex('\\n')==first_n:\n",
      "            with open('0.c', 'w') as f:\n",
      "                f.write(s)\n",
      "            return s"
     ],
     "language": "python",
     "metadata": {},
     "outputs": [],
     "prompt_number": 2
    },
    {
     "cell_type": "code",
     "collapsed": false,
     "input": [
      "s0 = \"\"\"#define echo main(){puts(\n",
      "#define exit );}\n",
      "echo \"Hello world!\"\n",
      "exit\n",
      "\"\"\"\n",
      "step = 1000\n",
      "src = 'sh *\\n#'\n",
      "L = Counter(s0)['X']\n",
      "s0 = s0.replace('X', '%s')\n",
      "first_n = src.index('\\n') \n",
      "view.push(dict(src=src, step=step, s0=s0, L=L, first_n=first_n))\n",
      "t0 = time()\n",
      "expected_number = 256**len(src)\n",
      "print \"L=\", L\n",
      "stdout.flush()\n",
      "text = TextWidget()\n",
      "display(text)\n",
      "for chunk in chunked(xrange(0, 2**(6*L), step), 500):\n",
      "    result = view.map_sync(lambda i:check_hash(i), chunk)\n",
      "    if any(result):\n",
      "        break\n",
      "    i = chunk[0]\n",
      "    elapsed = (time()-t0)/60\n",
      "    percent = i*100.0/expected_number\n",
      "    exp =  0 if i==0 else 100.0*elapsed/percent/60.0/24.0\n",
      "    text.value = \"{} {:5.2f}% {:.1f}min EXP={:.2f}days\".format(i, percent, elapsed, exp)\n",
      "print \"done\""
     ],
     "language": "python",
     "metadata": {},
     "outputs": [
      {
       "output_type": "stream",
       "stream": "stdout",
       "text": [
        "L= 7\n"
       ]
      }
     ]
    },
    {
     "cell_type": "code",
     "collapsed": false,
     "input": [
      "# OPENCL VERSION STARTS HERE\n",
      "from collections import Counter\n",
      "s0=\"\"\"#include <stdio.h>X/*X\n",
      "echo \"Hello world!\"X<< '}'X&&XexitX\n",
      "X\n",
      "It doesn't matter you use vim, emacs, or echo.X\n",
      "Real programmers use one-way functions to compile programs.X\n",
      "X\n",
      "Usage:X\n",
      "$ mkdir testX\n",
      "$ cd testX\n",
      "$ wget https://raw.githubusercontent.com/tjwei/tjw_ipynb/master/0.cX\n",
      "$ sh 0.cX\n",
      "Hello World!X\n",
      "$ gcc 0.cX&&X./a.outX\n",
      "Hello World!X\n",
      "$ shasum -a 384 0.cX|Xxxd -r -pX>Xa.outX&&Xchmod a+x a.outX\n",
      "$ ./a.outX\n",
      "Hello World!X\n",
      "X\n",
      "sha384 was chosenX\n",
      "because we can fit an ELF in the digest!(48>45)X\n",
      "X*/X\n",
      "main()X{X\n",
      "Xputs(\"Hello world!\");X\n",
      "}\n",
      "X\"\"\"\n",
      "len(s0), Counter(s0)['X']"
     ],
     "language": "python",
     "metadata": {},
     "outputs": [
      {
       "metadata": {},
       "output_type": "pyout",
       "prompt_number": 19,
       "text": [
        "(531, 39)"
       ]
      }
     ],
     "prompt_number": 19
    },
    {
     "cell_type": "code",
     "collapsed": false,
     "input": [
      "from more_itertools import chunked\n",
      "sp = s0.split('X')[:-1]\n",
      "prefix = '. *\\n'\n",
      "s0_ord = map(ord, s0.replace('X', ' '))\n",
      "with open(\"opencl_realman.h\" ,\"w\") as f:    \n",
      "    print >>f, \"__constant uint8_t SRC[%d]={\"%len(s0)+ \",\\n\".join(str(x)[1:-1] for x in chunked(s0_ord, 20))+\"};\"\n",
      "    pos = 0\n",
      "    print >>f, \"#define UPDATE \\\\\"\n",
      "    xx = 1\n",
      "    for l in sp:\n",
      "        L=len(l)+1\n",
      "        print >>f, \"ctx_update(ctx, SRC+%d, (x&0x%xUL) ? %d : %d);\\\\\"%(pos, xx, L, L-1)\n",
      "        pos=pos+L\n",
      "        xx<<=1\n",
      "    print >>f\n",
      "    prefix_hex=prefix_mask=0\n",
      "    for i, c in enumerate(prefix):\n",
      "        offset = 56-8*i\n",
      "        prefix_hex |= ord(c)<<offset\n",
      "        prefix_mask |= 0xff<<offset\n",
      "    print >>f, \"#define CONDITION (ctx->H[0]&0x%xUL) == 0x%xUL\"%(prefix_mask, prefix_hex)\n",
      "with open(\"opencl_realman.h\" ,\"r\") as f:\n",
      "    print f.read()"
     ],
     "language": "python",
     "metadata": {},
     "outputs": [
      {
       "output_type": "stream",
       "stream": "stdout",
       "text": [
        "__constant uint8_t SRC[531]={35, 105, 110, 99, 108, 117, 100, 101, 32, 60, 115, 116, 100, 105, 111, 46, 104, 62, 32, 47,\n",
        "42, 32, 10, 101, 99, 104, 111, 32, 34, 72, 101, 108, 108, 111, 32, 119, 111, 114, 108, 100,\n",
        "33, 34, 32, 60, 60, 32, 39, 125, 39, 32, 38, 38, 32, 101, 120, 105, 116, 32, 10, 32,\n",
        "10, 73, 116, 32, 100, 111, 101, 115, 110, 39, 116, 32, 109, 97, 116, 116, 101, 114, 32, 121,\n",
        "111, 117, 32, 117, 115, 101, 32, 118, 105, 109, 44, 32, 101, 109, 97, 99, 115, 44, 32, 111,\n",
        "114, 32, 101, 99, 104, 111, 46, 32, 10, 82, 101, 97, 108, 32, 112, 114, 111, 103, 114, 97,\n",
        "109, 109, 101, 114, 115, 32, 117, 115, 101, 32, 111, 110, 101, 45, 119, 97, 121, 32, 102, 117,\n",
        "110, 99, 116, 105, 111, 110, 115, 32, 116, 111, 32, 99, 111, 109, 112, 105, 108, 101, 32, 112,\n",
        "114, 111, 103, 114, 97, 109, 115, 46, 32, 10, 32, 10, 85, 115, 97, 103, 101, 58, 32, 10,\n",
        "36, 32, 109, 107, 100, 105, 114, 32, 116, 101, 115, 116, 32, 10, 36, 32, 99, 100, 32, 116,\n",
        "101, 115, 116, 32, 10, 36, 32, 119, 103, 101, 116, 32, 104, 116, 116, 112, 115, 58, 47, 47,\n",
        "114, 97, 119, 46, 103, 105, 116, 104, 117, 98, 117, 115, 101, 114, 99, 111, 110, 116, 101, 110,\n",
        "116, 46, 99, 111, 109, 47, 116, 106, 119, 101, 105, 47, 116, 106, 119, 95, 105, 112, 121, 110,\n",
        "98, 47, 109, 97, 115, 116, 101, 114, 47, 48, 46, 99, 32, 10, 36, 32, 115, 104, 32, 48,\n",
        "46, 99, 32, 10, 72, 101, 108, 108, 111, 32, 87, 111, 114, 108, 100, 33, 32, 10, 36, 32,\n",
        "103, 99, 99, 32, 48, 46, 99, 32, 38, 38, 32, 46, 47, 97, 46, 111, 117, 116, 32, 10,\n",
        "72, 101, 108, 108, 111, 32, 87, 111, 114, 108, 100, 33, 32, 10, 36, 32, 115, 104, 97, 115,\n",
        "117, 109, 32, 45, 97, 32, 51, 56, 52, 32, 48, 46, 99, 32, 124, 32, 120, 120, 100, 32,\n",
        "45, 114, 32, 45, 112, 32, 62, 32, 97, 46, 111, 117, 116, 32, 38, 38, 32, 99, 104, 109,\n",
        "111, 100, 32, 97, 43, 120, 32, 97, 46, 111, 117, 116, 32, 10, 36, 32, 46, 47, 97, 46,\n",
        "111, 117, 116, 32, 10, 72, 101, 108, 108, 111, 32, 87, 111, 114, 108, 100, 33, 32, 10, 32,\n",
        "10, 115, 104, 97, 51, 56, 52, 32, 119, 97, 115, 32, 99, 104, 111, 115, 101, 110, 32, 10,\n",
        "98, 101, 99, 97, 117, 115, 101, 32, 119, 101, 32, 99, 97, 110, 32, 102, 105, 116, 32, 97,\n",
        "110, 32, 69, 76, 70, 32, 105, 110, 32, 116, 104, 101, 32, 100, 105, 103, 101, 115, 116, 33,\n",
        "40, 52, 56, 62, 52, 53, 41, 32, 10, 32, 42, 47, 32, 10, 109, 97, 105, 110, 40, 41,\n",
        "32, 123, 32, 10, 32, 112, 117, 116, 115, 40, 34, 72, 101, 108, 108, 111, 32, 119, 111, 114,\n",
        "108, 100, 33, 34, 41, 59, 32, 10, 125, 10, 32};\n",
        "#define UPDATE \\\n",
        "ctx_update(ctx, SRC+0, (x&0x1UL) ? 19 : 18);\\\n",
        "ctx_update(ctx, SRC+19, (x&0x2UL) ? 3 : 2);\\\n",
        "ctx_update(ctx, SRC+22, (x&0x4UL) ? 21 : 20);\\\n",
        "ctx_update(ctx, SRC+43, (x&0x8UL) ? 7 : 6);\\\n",
        "ctx_update(ctx, SRC+50, (x&0x10UL) ? 3 : 2);\\\n",
        "ctx_update(ctx, SRC+53, (x&0x20UL) ? 5 : 4);\\\n",
        "ctx_update(ctx, SRC+58, (x&0x40UL) ? 2 : 1);\\\n",
        "ctx_update(ctx, SRC+60, (x&0x80UL) ? 48 : 47);\\\n",
        "ctx_update(ctx, SRC+108, (x&0x100UL) ? 61 : 60);\\\n",
        "ctx_update(ctx, SRC+169, (x&0x200UL) ? 2 : 1);\\\n",
        "ctx_update(ctx, SRC+171, (x&0x400UL) ? 8 : 7);\\\n",
        "ctx_update(ctx, SRC+179, (x&0x800UL) ? 14 : 13);\\\n",
        "ctx_update(ctx, SRC+193, (x&0x1000UL) ? 11 : 10);\\\n",
        "ctx_update(ctx, SRC+204, (x&0x2000UL) ? 69 : 68);\\\n",
        "ctx_update(ctx, SRC+273, (x&0x4000UL) ? 10 : 9);\\\n",
        "ctx_update(ctx, SRC+283, (x&0x8000UL) ? 14 : 13);\\\n",
        "ctx_update(ctx, SRC+297, (x&0x10000UL) ? 11 : 10);\\\n",
        "ctx_update(ctx, SRC+308, (x&0x20000UL) ? 3 : 2);\\\n",
        "ctx_update(ctx, SRC+311, (x&0x40000UL) ? 8 : 7);\\\n",
        "ctx_update(ctx, SRC+319, (x&0x80000UL) ? 14 : 13);\\\n",
        "ctx_update(ctx, SRC+333, (x&0x100000UL) ? 21 : 20);\\\n",
        "ctx_update(ctx, SRC+354, (x&0x200000UL) ? 2 : 1);\\\n",
        "ctx_update(ctx, SRC+356, (x&0x400000UL) ? 10 : 9);\\\n",
        "ctx_update(ctx, SRC+366, (x&0x800000UL) ? 2 : 1);\\\n",
        "ctx_update(ctx, SRC+368, (x&0x1000000UL) ? 6 : 5);\\\n",
        "ctx_update(ctx, SRC+374, (x&0x2000000UL) ? 3 : 2);\\\n",
        "ctx_update(ctx, SRC+377, (x&0x4000000UL) ? 16 : 15);\\\n",
        "ctx_update(ctx, SRC+393, (x&0x8000000UL) ? 11 : 10);\\\n",
        "ctx_update(ctx, SRC+404, (x&0x10000000UL) ? 14 : 13);\\\n",
        "ctx_update(ctx, SRC+418, (x&0x20000000UL) ? 2 : 1);\\\n",
        "ctx_update(ctx, SRC+420, (x&0x40000000UL) ? 19 : 18);\\\n",
        "ctx_update(ctx, SRC+439, (x&0x80000000UL) ? 49 : 48);\\\n",
        "ctx_update(ctx, SRC+488, (x&0x100000000UL) ? 2 : 1);\\\n",
        "ctx_update(ctx, SRC+490, (x&0x200000000UL) ? 3 : 2);\\\n",
        "ctx_update(ctx, SRC+493, (x&0x400000000UL) ? 8 : 7);\\\n",
        "ctx_update(ctx, SRC+501, (x&0x800000000UL) ? 2 : 1);\\\n",
        "ctx_update(ctx, SRC+503, (x&0x1000000000UL) ? 2 : 1);\\\n",
        "ctx_update(ctx, SRC+505, (x&0x2000000000UL) ? 22 : 21);\\\n",
        "ctx_update(ctx, SRC+527, (x&0x4000000000UL) ? 4 : 3);\\\n",
        "\n",
        "#define CONDITION (ctx->H[0]&0xffffffff00000000UL) == 0x2e202a0a00000000UL\n",
        "\n"
       ]
      }
     ],
     "prompt_number": 20
    },
    {
     "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\n",
      "%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": 21
    },
    {
     "cell_type": "code",
     "collapsed": false,
     "input": [
      "ctx = cl.create_some_context()\n",
      "queue = cl.CommandQueue(ctx)\n",
      "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": 22
    },
    {
     "cell_type": "code",
     "collapsed": false,
     "input": [
      "%%cl_kernel -o \"-I .\"\n",
      "#define _OPENCL_COMPILER\n",
      "#define DEVICE_INFO 2|64|1024\n",
      "#include \"opencl_cryptsha512.h\"\n",
      "#include \"opencl_realman.h\"\n",
      "__constant uint64_t k[] = {\n",
      "    0x428a2f98d728ae22UL, 0x7137449123ef65cdUL, 0xb5c0fbcfec4d3b2fUL, 0xe9b5dba58189dbbcUL,\n",
      "    0x3956c25bf348b538UL, 0x59f111f1b605d019UL, 0x923f82a4af194f9bUL, 0xab1c5ed5da6d8118UL,\n",
      "    0xd807aa98a3030242UL, 0x12835b0145706fbeUL, 0x243185be4ee4b28cUL, 0x550c7dc3d5ffb4e2UL,\n",
      "    0x72be5d74f27b896fUL, 0x80deb1fe3b1696b1UL, 0x9bdc06a725c71235UL, 0xc19bf174cf692694UL,\n",
      "    0xe49b69c19ef14ad2UL, 0xefbe4786384f25e3UL, 0x0fc19dc68b8cd5b5UL, 0x240ca1cc77ac9c65UL,\n",
      "    0x2de92c6f592b0275UL, 0x4a7484aa6ea6e483UL, 0x5cb0a9dcbd41fbd4UL, 0x76f988da831153b5UL,\n",
      "    0x983e5152ee66dfabUL, 0xa831c66d2db43210UL, 0xb00327c898fb213fUL, 0xbf597fc7beef0ee4UL,\n",
      "    0xc6e00bf33da88fc2UL, 0xd5a79147930aa725UL, 0x06ca6351e003826fUL, 0x142929670a0e6e70UL,\n",
      "    0x27b70a8546d22ffcUL, 0x2e1b21385c26c926UL, 0x4d2c6dfc5ac42aedUL, 0x53380d139d95b3dfUL,\n",
      "    0x650a73548baf63deUL, 0x766a0abb3c77b2a8UL, 0x81c2c92e47edaee6UL, 0x92722c851482353bUL,\n",
      "    0xa2bfe8a14cf10364UL, 0xa81a664bbc423001UL, 0xc24b8b70d0f89791UL, 0xc76c51a30654be30UL,\n",
      "    0xd192e819d6ef5218UL, 0xd69906245565a910UL, 0xf40e35855771202aUL, 0x106aa07032bbd1b8UL,\n",
      "    0x19a4c116b8d2d0c8UL, 0x1e376c085141ab53UL, 0x2748774cdf8eeb99UL, 0x34b0bcb5e19b48a8UL,\n",
      "    0x391c0cb3c5c95a63UL, 0x4ed8aa4ae3418acbUL, 0x5b9cca4f7763e373UL, 0x682e6ff3d6b2b8a3UL,\n",
      "    0x748f82ee5defb2fcUL, 0x78a5636f43172f60UL, 0x84c87814a1f0ab72UL, 0x8cc702081a6439ecUL,\n",
      "    0x90befffa23631e28UL, 0xa4506cebde82bde9UL, 0xbef9a3f7b2c67915UL, 0xc67178f2e372532bUL,\n",
      "    0xca273eceea26619cUL, 0xd186b8c721c0c207UL, 0xeada7dd6cde0eb1eUL, 0xf57d4f7fee6ed178UL,\n",
      "    0x06f067aa72176fbaUL, 0x0a637dc5a2c898a6UL, 0x113f9804bef90daeUL, 0x1b710b35131c471bUL,\n",
      "    0x28db77f523047d84UL, 0x32caab7b40c72493UL, 0x3c9ebe0a15c9bebcUL, 0x431d67c49c100d4cUL,\n",
      "    0x4cc5d4becb3e42b6UL, 0x597f299cfc657e2aUL, 0x5fcb6fab3ad6faecUL, 0x6c44198c4a475817UL,\n",
      "};\n",
      "\n",
      "\n",
      "inline void init_ctx384(__local sha512_ctx * ctx) {\n",
      "    ctx->H[0] = 0xcbbb9d5dc1059ed8UL;\n",
      "    ctx->H[1] = 0x629a292a367cd507UL;\n",
      "    ctx->H[2] = 0x9159015a3070dd17UL;\n",
      "    ctx->H[3] = 0x152fecd8f70e5939UL;\n",
      "    ctx->H[4] = 0x67332667ffc00b31UL;\n",
      "    ctx->H[5] = 0x8eb44a8768581511UL;\n",
      "    ctx->H[6] = 0xdb0c2e0d64f98fa7UL;\n",
      "    ctx->H[7] = 0x47b5481dbefa4fa4UL;\n",
      "    ctx->total = 0;\n",
      "    ctx->buflen = 0;\n",
      "}\n",
      "\n",
      "inline void sha512_block(__local sha512_ctx * ctx) {\n",
      "    uint64_t a = ctx->H[0];\n",
      "    uint64_t b = ctx->H[1];\n",
      "    uint64_t c = ctx->H[2];\n",
      "    uint64_t d = ctx->H[3];\n",
      "    uint64_t e = ctx->H[4];\n",
      "    uint64_t f = ctx->H[5];\n",
      "    uint64_t g = ctx->H[6];\n",
      "    uint64_t h = ctx->H[7];\n",
      "    uint64_t t1, t2;\n",
      "    uint64_t w[16];\n",
      "\n",
      "#ifdef VECTOR_USAGE\n",
      "    ulong16  w_vector;\n",
      "    w_vector = vload16(0, ctx->buffer->mem_64);\n",
      "    w_vector = SWAP64_V(w_vector);\n",
      "    vstore16(w_vector, 0, w);\n",
      "#else\n",
      "    #pragma unroll\n",
      "    for (int i = 0; i < 16; i++)\n",
      "        w[i] = SWAP64(ctx->buffer->mem_64[i]);\n",
      "#endif\n",
      "\n",
      "    #pragma unroll\n",
      "    for (int i = 0; i < 80; i++) {\n",
      "\n",
      "        if (i > 15) {\n",
      "            w[i & 15] = sigma1(w[(i - 2) & 15]) + sigma0(w[(i - 15) & 15]) + w[(i - 16) & 15] + w[(i - 7) & 15];\n",
      "        }\n",
      "        t1 = k[i] + w[i & 15] + h + Sigma1(e) + Ch(e, f, g);\n",
      "        t2 = Maj(a, b, c) + Sigma0(a);\n",
      "        h = g;\n",
      "        g = f;\n",
      "        f = e;\n",
      "        e = d + t1;\n",
      "        d = c;\n",
      "        c = b;\n",
      "        b = a;\n",
      "        a = t1 + t2;\n",
      "    }\n",
      "    /* Put checksum in context given as argument. */\n",
      "    ctx->H[0] += a;\n",
      "    ctx->H[1] += b;\n",
      "    ctx->H[2] += c;\n",
      "    ctx->H[3] += d;\n",
      "    ctx->H[4] += e;\n",
      "    ctx->H[5] += f;\n",
      "    ctx->H[6] += g;\n",
      "    ctx->H[7] += h;\n",
      "}\n",
      "\n",
      "inline void insert_to_buffer(__local sha512_ctx    * ctx,\n",
      "                      __constant  uint8_t * string,\n",
      "                      const uint32_t len) {\n",
      "    __local uint8_t * dest;\n",
      "    dest = ctx->buffer->mem_08 + ctx->buflen;\n",
      "\n",
      "    for (int i = 0; i < len; i++)\n",
      "        PUTCHAR(dest, i, GETCHAR(string, i));\n",
      "    ctx->buflen += len;    \n",
      "}\n",
      "\n",
      "inline void ctx_update(__local sha512_ctx * ctx,\n",
      "                __constant uint8_t  * string, uint32_t len) {\n",
      "\n",
      "    ctx->total += len;\n",
      "    uint32_t startpos = ctx->buflen;\n",
      "    uint32_t offset = 0;\n",
      "\n",
      "    while(1){\n",
      "            insert_to_buffer(ctx, string+offset, (len -offset <= 128 - startpos ? len-offset  : 128 - startpos));\n",
      "            if(ctx->buflen<128) break;\n",
      "            mem_fence(CLK_LOCAL_MEM_FENCE);\n",
      "            sha512_block(ctx);            \n",
      "            offset = offset +128 - startpos;\n",
      "            ctx->buflen = 0;\n",
      "            startpos=0;\n",
      "    }\n",
      "\n",
      "}\n",
      "\n",
      "inline void ctx_append_1(__local sha512_ctx * ctx) {\n",
      "\n",
      "    uint32_t length = ctx->buflen;\n",
      "    PUTCHAR(ctx->buffer->mem_08, length, 0x80);\n",
      "\n",
      "    while (++length & 3)\n",
      "        PUTCHAR(ctx->buffer->mem_08, length, 0);\n",
      "\n",
      "    if (length & 7) {\n",
      "        __local uint32_t * l = (__local uint32_t *) (ctx->buffer->mem_08 + length);\n",
      "        *l = 0;\n",
      "        length += 4;\n",
      "    }\n",
      "    __local uint64_t * l = (__local uint64_t *) (ctx->buffer->mem_08 + length);\n",
      "\n",
      "    while (length < 128) {\n",
      "        *l++ = 0;\n",
      "        length += 8;\n",
      "    }\n",
      "}\n",
      "\n",
      "inline void ctx_add_length(__local sha512_ctx * ctx) {\n",
      "\n",
      "    ctx->buffer->mem_64[15] = SWAP64((uint64_t) (ctx->total * 8));\n",
      "}\n",
      "\n",
      "inline void finish_ctx(__local sha512_ctx * ctx) {\n",
      "    ctx_append_1(ctx);\n",
      "    ctx_add_length(ctx);\n",
      "    ctx->buflen = 0;\n",
      "}\n",
      "\n",
      "inline void clear_ctx_buffer(__local sha512_ctx * ctx) {\n",
      "\n",
      "#ifdef VECTOR_USAGE\n",
      "    ulong16  w_vector = 0;\n",
      "    vstore16(w_vector, 0, ctx->buffer->mem_64);\n",
      "#else\n",
      "    #pragma unroll\n",
      "    for (int i = 0; i < 16; i++)\n",
      "        ctx->buffer->mem_64[i] = 0;\n",
      "#endif\n",
      "\n",
      "    ctx->buflen = 0;\n",
      "}\n",
      "\n",
      "inline void sha384_digest(__local sha512_ctx * ctx,\n",
      "                   __global uint8_t   *result) {\n",
      "\n",
      "    if (ctx->buflen <= 111) { //data+0x80+datasize fits in one 1024bit block\n",
      "        finish_ctx(ctx);\n",
      "\n",
      "    } else {\n",
      "        bool moved = true;\n",
      "\n",
      "        if (ctx->buflen < 128) { //data and 0x80 fits in one block\n",
      "            ctx_append_1(ctx);\n",
      "            moved = false;\n",
      "        }\n",
      "        sha512_block(ctx);\n",
      "        clear_ctx_buffer(ctx);\n",
      "\n",
      "        if (moved) //append 1,the rest is already clean\n",
      "            PUTCHAR(ctx->buffer->mem_08, 0, 0x80);\n",
      "        ctx_add_length(ctx);\n",
      "    }\n",
      "    sha512_block(ctx);\n",
      "    \n",
      "   if( CONDITION )\n",
      "           *result=1;\n",
      "else *result=0;\n",
      "/*\n",
      "    #pragma unroll\n",
      "    for (int i = 0; i < 6; i++)\n",
      "        result[i] = SWAP64(ctx->H[i]);\n",
      "*/\n",
      "}\n",
      "\n",
      "inline void sha384crypt(__local  sha512_ctx * ctx,\n",
      "                 __global uint8_t *output, uint64_t x) {    \n",
      "    init_ctx384(ctx);\n",
      "    UPDATE; \n",
      "    sha384_digest(ctx, output);\n",
      "}\n",
      "\n",
      "__kernel\n",
      "void kernel_sha384(uint64_t base, __global   uint8_t * out_buffer,\n",
      "                  __local sha512_ctx  * ctx_memory) {\n",
      "    size_t gid = get_global_id(0);\n",
      "    size_t lid = get_local_id(0);\n",
      "    uint64_t x = base + gid;\n",
      "    sha384crypt(&ctx_memory[lid], &out_buffer[gid], x);\n",
      "}\n"
     ],
     "language": "python",
     "metadata": {},
     "outputs": [],
     "prompt_number": 23
    },
    {
     "cell_type": "code",
     "collapsed": false,
     "input": [
      "from time import sleep\n",
      "step = 10**7\n",
      "t0 = time()\n",
      "expected_number = 256**4\n",
      "L=len(sp)\n",
      "print \"L=\", L\n",
      "stdout.flush()\n",
      "text = TextWidget()\n",
      "display(text)\n",
      "\n",
      "for i in xrange(0,2**L, step):\n",
      "    output=cl.array.zeros(queue, step, dtype=np.uint8)\n",
      "    working = cl.LocalMemory(200*128)\n",
      "    kernel_sha384(queue, (step,), (128,), np.uint64(i), output.data, working).wait()\n",
      "    result = output.get()\n",
      "    for j in np.where(result==1)[0]:        \n",
      "        x =i+j\n",
      "        s=s0.replace('X', '%s')%tuple(' '*((x>>i)&1) for i in range(L))\n",
      "        print x,repr(sha384(s).digest())\n",
      "        with open(\"good/%d.c\"%x, \"w\") as f:\n",
      "            f.write(s)\n",
      "        stdout.flush()    \n",
      "    elapsed = (time()-t0)/60\n",
      "    percent = i*100.0/expected_number\n",
      "    exp =  0 if i==0 else 100.0*elapsed/percent/60.0\n",
      "    text.value = \"{} {:5.2f}% {:.2f}min EXP={:.2f}hrs\".format(i, percent, elapsed, exp)\n",
      "    sleep(2)\n",
      "print \"done\""
     ],
     "language": "python",
     "metadata": {},
     "outputs": [
      {
       "output_type": "stream",
       "stream": "stdout",
       "text": [
        "L= 39\n"
       ]
      },
      {
       "output_type": "stream",
       "stream": "stdout",
       "text": [
        "1580306764 '. *\\n\"@\\xc1\\xe4\\xa0X\\xc4\"\\x13\\xd08i\\xb9\\x88\\xcd\\xbb\\x15\\xbcl\\xf2\\xe7\\xbdX\\xedX\\xc6\\x1e\\xf0>B\\xdf\\xf4\\xb7\\x00\\xf9\\xb0\\x9f\\xcbLr\\xdd}\\xfc\\x81'\n"
       ]
      },
      {
       "output_type": "stream",
       "stream": "stdout",
       "text": [
        "7385759359 \". *\\n8\\xff3`\\x06'd\\xc0\\xee\\xf4\\x9b\\xd3\\xbaA\\xa83\\xb8\\xb9\\x96\\xc9\\x8c\\x1e\\xd4\\xc1\\x13\\xfd\\xfdO\\x00\\x1a9('\\x90\\x04\\x04(<b*^*xe\"\n"
       ]
      },
      {
       "output_type": "stream",
       "stream": "stdout",
       "text": [
        "10346520310 '. *\\ni\\xab\\x17\\x1dQn\\xe8t\\xdbr\\xb7,\\xff\\x9f\\xbf\\x9e\\x9e\\x8fN\\xb8d\\xe3\\xec\\x06^^\\xf6V\\xf6i\\x98\\xe8\\xe0\\xc8\\x933\\x98\\xe7\\xb9\\xcf\\xf4\\t\\x04\\x82'\n"
       ]
      },
      {
       "output_type": "stream",
       "stream": "stdout",
       "text": [
        "15535172170 '. *\\n\\x02\\xd5[\\xab\\x0f-\\xfa4\\x14+\\xfeL\\xfdm\\xaee>ExB\\xae\\xc4\\x99\\xd53\\r\\x80{.\\xb3\\xc3\\x95+F\\xb7\\tJ\\x0c\\xdf\\xaaL\\ns\\xb2'\n"
       ]
      },
      {
       "ename": "KeyboardInterrupt",
       "evalue": "",
       "output_type": "pyerr",
       "traceback": [
        "\u001b[1;31m---------------------------------------------------------------------------\u001b[0m\n\u001b[1;31mKeyboardInterrupt\u001b[0m                         Traceback (most recent call last)",
        "\u001b[1;32m<ipython-input-26-73032072ccc5>\u001b[0m in \u001b[0;36m<module>\u001b[1;34m()\u001b[0m\n\u001b[0;32m     12\u001b[0m     \u001b[0moutput\u001b[0m\u001b[1;33m=\u001b[0m\u001b[0mcl\u001b[0m\u001b[1;33m.\u001b[0m\u001b[0marray\u001b[0m\u001b[1;33m.\u001b[0m\u001b[0mzeros\u001b[0m\u001b[1;33m(\u001b[0m\u001b[0mqueue\u001b[0m\u001b[1;33m,\u001b[0m \u001b[0mstep\u001b[0m\u001b[1;33m,\u001b[0m \u001b[0mdtype\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[0m\n\u001b[0;32m     13\u001b[0m     \u001b[0mworking\u001b[0m \u001b[1;33m=\u001b[0m \u001b[0mcl\u001b[0m\u001b[1;33m.\u001b[0m\u001b[0mLocalMemory\u001b[0m\u001b[1;33m(\u001b[0m\u001b[1;36m200\u001b[0m\u001b[1;33m*\u001b[0m\u001b[1;36m128\u001b[0m\u001b[1;33m)\u001b[0m\u001b[1;33m\u001b[0m\u001b[0m\n\u001b[1;32m---> 14\u001b[1;33m     \u001b[0mkernel_sha384\u001b[0m\u001b[1;33m(\u001b[0m\u001b[0mqueue\u001b[0m\u001b[1;33m,\u001b[0m \u001b[1;33m(\u001b[0m\u001b[0mstep\u001b[0m\u001b[1;33m,\u001b[0m\u001b[1;33m)\u001b[0m\u001b[1;33m,\u001b[0m \u001b[1;33m(\u001b[0m\u001b[1;36m128\u001b[0m\u001b[1;33m,\u001b[0m\u001b[1;33m)\u001b[0m\u001b[1;33m,\u001b[0m \u001b[0mnp\u001b[0m\u001b[1;33m.\u001b[0m\u001b[0muint64\u001b[0m\u001b[1;33m(\u001b[0m\u001b[0mi\u001b[0m\u001b[1;33m)\u001b[0m\u001b[1;33m,\u001b[0m \u001b[0moutput\u001b[0m\u001b[1;33m.\u001b[0m\u001b[0mdata\u001b[0m\u001b[1;33m,\u001b[0m \u001b[0mworking\u001b[0m\u001b[1;33m)\u001b[0m\u001b[1;33m.\u001b[0m\u001b[0mwait\u001b[0m\u001b[1;33m(\u001b[0m\u001b[1;33m)\u001b[0m\u001b[1;33m\u001b[0m\u001b[0m\n\u001b[0m\u001b[0;32m     15\u001b[0m     \u001b[0mresult\u001b[0m \u001b[1;33m=\u001b[0m \u001b[0moutput\u001b[0m\u001b[1;33m.\u001b[0m\u001b[0mget\u001b[0m\u001b[1;33m(\u001b[0m\u001b[1;33m)\u001b[0m\u001b[1;33m\u001b[0m\u001b[0m\n\u001b[0;32m     16\u001b[0m     \u001b[1;32mfor\u001b[0m \u001b[0mj\u001b[0m \u001b[1;32min\u001b[0m \u001b[0mnp\u001b[0m\u001b[1;33m.\u001b[0m\u001b[0mwhere\u001b[0m\u001b[1;33m(\u001b[0m\u001b[0mresult\u001b[0m\u001b[1;33m==\u001b[0m\u001b[1;36m1\u001b[0m\u001b[1;33m)\u001b[0m\u001b[1;33m[\u001b[0m\u001b[1;36m0\u001b[0m\u001b[1;33m]\u001b[0m\u001b[1;33m:\u001b[0m\u001b[1;33m\u001b[0m\u001b[0m\n",
        "\u001b[1;31mKeyboardInterrupt\u001b[0m: "
       ]
      }
     ],
     "prompt_number": 26
    },
    {
     "cell_type": "code",
     "collapsed": false,
     "input": [],
     "language": "python",
     "metadata": {},
     "outputs": []
    }
   ],
   "metadata": {}
  }
 ]
}