{ "metadata": { "name": "", "signature": "sha256:7e9e35823d13aa87d70a15a8a8ad888d4b8898ad4de39fc3f8cfd8ab70b26182" }, "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\n", "from __future__ import division" ], "language": "python", "metadata": {}, "outputs": [], "prompt_number": 1 }, { "cell_type": "code", "collapsed": false, "input": [ "# OPENCL VERSION STARTS HERE\n", "from collections import Counter\n", "s0=\"\"\"#include/*\n", "#usage: http://x.co/AAAAAAA\n", "echo \"Hello world!\"<<'}'\n", "*/\n", "main(){puts(\"Hello world!\");\n", "}\"\"\"\n", "len(s0), s0.find('A')" ], "language": "python", "metadata": {}, "outputs": [ { "metadata": {}, "output_type": "pyout", "prompt_number": 2, "text": [ "(106, 40)" ] } ], "prompt_number": 2 }, { "cell_type": "code", "collapsed": false, "input": [ "from more_itertools import chunked\n", "prefix = 'sh * '\n", "s0_ord = map(ord, s0)+[0x80]+[0]*(128-len(s0)-3)+[len(s0)>>5, (len(s0)<<3)&0xff]\n", "def hex8(l):\n", " return \"0x\"+\"\".join(\"%02x\"%x for x in l)\n", "with open(\"opencl_realman2.h\" ,\"w\") as f: \n", " print >>f, \"__constant uint64_t SRC[16]={\"+ \",\\n\".join(hex8(x)+\"UL\" for x in chunked(s0_ord, 8))+\"};\"\n", " #for i, x in enumerate(chunked(s0_ord, 8)):\n", " # print >>f, \"#define SRC%x %sUL\"%(i, hex8(x))\n", " print >>f, \"#define A %d\"%s0.find('A')\n", " print >>f, \"#define A8 %x\"%(s0.find('A')//8)\n", " print >>f, \"#define B %d\"%(ord(s0[s0.find('A')+7]))\n", " print >>f, \"#define LEN %d\"%len(s0)\n", " prefix_hex=prefix_mask=0\n", " charset = \"abcdefghijklmnopqrstuvwxyzABCDEFGHIJKLMNOPQRSTUVWXYZ0123456789?#\"\n", " _charset = map(ord, charset)\n", " print >>f, \"__constant uint64_t charset[64]={\" +str(_charset)[1:-1]+ \"};\"\n", " for i, c in enumerate(prefix):\n", " offset = 56-8*i\n", " prefix_hex |= ord(c)<>f, \"#define CONDITION (v.s0&0x%xUL) == 0x%xUL\"%(prefix_mask, prefix_hex)\n", "with open(\"opencl_realman2.h\" ,\"r\") as f:\n", " print f.read()" ], "language": "python", "metadata": {}, "outputs": [ { "output_type": "stream", "stream": "stdout", "text": [ "__constant uint64_t SRC[16]={0x23696e636c756465UL,\n", "0x3c737464696f2e68UL,\n", "0x3e2f2a0a23757361UL,\n", "0x67653a2068747470UL,\n", "0x3a2f2f782e636f2fUL,\n", "0x414141414141410aUL,\n", "0x6563686f20224865UL,\n", "0x6c6c6f20776f726cUL,\n", "0x6421223c3c277d27UL,\n", "0x0a2a2f0a6d61696eUL,\n", "0x28297b7075747328UL,\n", "0x2248656c6c6f2077UL,\n", "0x6f726c642122293bUL,\n", "0x0a7d800000000000UL,\n", "0x0000000000000000UL,\n", "0x0000000000000350UL};\n", "#define A 40\n", "#define A8 5\n", "#define B 10\n", "#define LEN 106\n", "__constant uint64_t charset[64]={97, 98, 99, 100, 101, 102, 103, 104, 105, 106, 107, 108, 109, 110, 111, 112, 113, 114, 115, 116, 117, 118, 119, 120, 121, 122, 65, 66, 67, 68, 69, 70, 71, 72, 73, 74, 75, 76, 77, 78, 79, 80, 81, 82, 83, 84, 85, 86, 87, 88, 89, 90, 48, 49, 50, 51, 52, 53, 54, 55, 56, 57, 63, 35};\n", "#define CONDITION (v.s0&0xffffffffff000000UL) == 0x7368202a20000000UL\n", "\n" ] } ], "prompt_number": 3 }, { "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": [], "prompt_number": 4 }, { "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": 5 }, { "cell_type": "code", "collapsed": false, "input": [ "%%cl_kernel -o \"-I . -cl-strict-aliasing\"\n", "//Type names definition.\n", "#define uint8_t unsigned char\n", "#define uint16_t unsigned short\n", "#define uint32_t unsigned int\n", "#define uint64_t unsigned long //Tip: unsigned long long int failed on compile (AMD).\n", "\n", "#define _OPENCL_COMPILER\n", "//Macros.\n", "#pragma OPENCL EXTENSION cl_amd_media_ops : enable\n", "#define Ch(x,y,z) bitselect(z, y, x)\n", "#define Maj(x,y,z) bitselect(x, y, z ^ x)\n", "#define ror(x, n) rotate(x, (uint64_t) 64-n)\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", "#include \"opencl_realman2.h\"\n", "/*\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", "#define k0 0x428a2f98d728ae22U\n", "#define k1 0x7137449123ef65cdUL\n", "#define k2 0xb5c0fbcfec4d3b2fUL\n", "#define k3 0xe9b5dba58189dbbcUL\n", "#define k4 0x3956c25bf348b538UL\n", "#define k5 0x59f111f1b605d019UL\n", "#define k6 0x923f82a4af194f9bUL\n", "#define k7 0xab1c5ed5da6d8118UL\n", "#define k8 0xd807aa98a3030242UL\n", "#define k9 0x12835b0145706fbeUL\n", "#define k10 0x243185be4ee4b28cUL\n", "#define k11 0x550c7dc3d5ffb4e2UL\n", "#define k12 0x72be5d74f27b896fUL\n", "#define k13 0x80deb1fe3b1696b1UL\n", "#define k14 0x9bdc06a725c71235UL\n", "#define k15 0xc19bf174cf692694UL\n", "#define k16 0xe49b69c19ef14ad2UL\n", "#define k17 0xefbe4786384f25e3UL\n", "#define k18 0x0fc19dc68b8cd5b5UL\n", "#define k19 0x240ca1cc77ac9c65UL\n", "#define k20 0x2de92c6f592b0275UL\n", "#define k21 0x4a7484aa6ea6e483UL\n", "#define k22 0x5cb0a9dcbd41fbd4UL\n", "#define k23 0x76f988da831153b5UL\n", "#define k24 0x983e5152ee66dfabUL\n", "#define k25 0xa831c66d2db43210UL\n", "#define k26 0xb00327c898fb213fUL\n", "#define k27 0xbf597fc7beef0ee4UL\n", "#define k28 0xc6e00bf33da88fc2UL\n", "#define k29 0xd5a79147930aa725UL\n", "#define k30 0x06ca6351e003826fUL\n", "#define k31 0x142929670a0e6e70UL\n", "#define k32 0x27b70a8546d22ffcUL\n", "#define k33 0x2e1b21385c26c926UL\n", "#define k34 0x4d2c6dfc5ac42aedUL\n", "#define k35 0x53380d139d95b3dfUL\n", "#define k36 0x650a73548baf63deUL\n", "#define k37 0x766a0abb3c77b2a8UL\n", "#define k38 0x81c2c92e47edaee6UL\n", "#define k39 0x92722c851482353bUL\n", "#define k40 0xa2bfe8a14cf10364UL\n", "#define k41 0xa81a664bbc423001UL\n", "#define k42 0xc24b8b70d0f89791UL\n", "#define k43 0xc76c51a30654be30UL\n", "#define k44 0xd192e819d6ef5218UL\n", "#define k45 0xd69906245565a910UL\n", "#define k46 0xf40e35855771202aUL\n", "#define k47 0x106aa07032bbd1b8UL\n", "#define k48 0x19a4c116b8d2d0c8UL\n", "#define k49 0x1e376c085141ab53UL\n", "#define k50 0x2748774cdf8eeb99UL\n", "#define k51 0x34b0bcb5e19b48a8UL\n", "#define k52 0x391c0cb3c5c95a63UL\n", "#define k53 0x4ed8aa4ae3418acbUL\n", "#define k54 0x5b9cca4f7763e373UL\n", "#define k55 0x682e6ff3d6b2b8a3UL\n", "#define k56 0x748f82ee5defb2fcUL\n", "#define k57 0x78a5636f43172f60UL\n", "#define k58 0x84c87814a1f0ab72UL\n", "#define k59 0x8cc702081a6439ecUL\n", "#define k60 0x90befffa23631e28UL\n", "#define k61 0xa4506cebde82bde9UL\n", "#define k62 0xbef9a3f7b2c67915UL\n", "#define k63 0xc67178f2e372532bUL\n", "#define k64 0xca273eceea26619cUL\n", "#define k65 0xd186b8c721c0c207UL\n", "#define k66 0xeada7dd6cde0eb1eUL\n", "#define k67 0xf57d4f7fee6ed178UL\n", "#define k68 0x06f067aa72176fbaUL\n", "#define k69 0x0a637dc5a2c898a6UL\n", "#define k70 0x113f9804bef90daeUL\n", "#define k71 0x1b710b35131c471bUL\n", "#define k72 0x28db77f523047d84UL\n", "#define k73 0x32caab7b40c72493UL\n", "#define k74 0x3c9ebe0a15c9bebcUL\n", "#define k75 0x431d67c49c100d4cUL\n", "#define k76 0x4cc5d4becb3e42b6UL\n", "#define k77 0x597f299cfc657e2aUL\n", "#define k78 0x5fcb6fab3ad6faecUL\n", "#define k79 0x6c44198c4a475817UL\n", "\n", "\n", "#define ALPHADIGITS\n", "#define H0 0xcbbb9d5dc1059ed8UL\n", "#define H1 0x629a292a367cd507UL\n", "#define H2 0x9159015a3070dd17UL\n", "#define H3 0x152fecd8f70e5939UL\n", "#define H4 0x67332667ffc00b31UL\n", "#define H5 0x8eb44a8768581511UL\n", "#define H6 0xdb0c2e0d64f98fa7UL\n", "#define H7 0x47b5481dbefa4fa4UL\n", "\n", "\n", "inline bool sha384_block(__const uint64_t x) {\n", " __private ulong8 v = {H0, H1, H2, H3, H4,H5,H6,H7};\n", " __private uint64_t w[16]; //={SRC0, SRC1, SRC2, SRC3, SRC4, SRC5, SRC6, SRC7, SRC8, SRC9, SRCa, SRCb,SRCc,SRCd, SRCe, SRCf}; \n", " // __private uint64_t t0, t1;\n", " __private ulong2 _t;\n", " #define t0 _t.s0\n", " #define t1 _t.s1\n", " \n", "/*\n", "#ifdef ALPHADIGITS\n", " t.s1=x;\n", " #define SET_W_A8 \\\n", " w[A8] = (charset[t.s1&63]<<56)^(charset[(t.s1>>6)&63]<<48)^(charset[(t.s1>>12)&63]<<40) \\\n", " ^(charset[(t.s1>>18)&63]<<32)^(charset[(t.s1>>24)&63]<<24) \\\n", " ^(charset[(t.s1>>30)&63]<<16)^(charset[(t.s1>>36)&63]<<8)^B; \n", " \n", " SET_W_A8\n", "#else\n", " w[A8]= (((x<<14)^x)&0x3f3f3f3f3f3f3f3f)+0x3f3f3f3f3f3f3f3f;\n", "#endif \n", " */\n", " #pragma unroll\n", " for (int i = 0; i < A8; i++) w[i] = SRC[i];\n", " \n", "#ifdef ALPHADIGITS \n", " w[A8] = (charset[x&63]<<56)^(charset[(x>>6)&63]<<48)^(charset[(x>>12)&63]<<40) \\\n", " ^(charset[(x>>18)&63]<<32)^(charset[(x>>24)&63]<<24) \\\n", " ^(charset[(x>>30)&63]<<16)^(charset[(x>>36)&63]<<8)^B; \n", "#else\n", "\n", " w[A8]= (((x<<14)^x)&0x3f3f3f3f3f3f3f3f)+0x3f3f3f3f3f3f3f3f;\n", "#endif \n", " #pragma unroll\n", " for (int i = A8+1; i < 16; i++) w[i] = SRC[i];\n", " \n", "\n", " \n", " /* i, i-2, i-15, i, i-7*/\n", " #define T1(a,b,c,d,e,f,g,h,i,j1) k##i + w[0x##j1] + v.s##h + Sigma1(v.s##e) + Ch(v.s##e, v.s##f, v.s##g)\n", " #define T2(a,b,c) Maj(v.s##a, v.s##b, v.s##c) + Sigma0(v.s##a)\n", " #define R0(a,b,c,d,e,f,g,h, i, j1, j2, j3, j4) \\\n", " t0 = T1(a,b,c,d,e,f,g,h, i, j1); \\\n", " t1 = T2(a,b,c);\\\n", " v.s##h= v.s##d + t0; v.s##d = t0 + t1;\n", " \n", " #define R(a,b,c,d,e,f,g,h, i, j1, j2, j3, j4)\\\n", " w[0x##j1] = sigma1(w[0x##j2]) + sigma0(w[0x##j3]) + w[0x##j1] + w[0x##j4]; \\\n", " R0(a,b,c,d,e,f,g,h, i, j1, j2, j3, j4)\n", " #define Rlast(a,b,c,d,e,f,g,h, i, j1, j2, j3, j4)\\\n", " v.s##d = k##i + sigma1(w[0x##j2]) + sigma0(w[0x##j3]) + w[0x##j1] + w[0x##j4] + v.s##h + Sigma1(v.s##e) + Ch(v.s##e, v.s##f, v.s##g) + T2(a,b,c) +H0;\n", " \n", "\n", "R0(0, 1, 2, 3, 4, 5, 6, 7, 0, 0, e, 1, 9);\n", "R0(3, 0, 1, 2, 7, 4, 5, 6, 1, 1, f, 2, a);\n", "R0(2, 3, 0, 1, 6, 7, 4, 5, 2, 2, 0, 3, b);\n", "R0(1, 2, 3, 0, 5, 6, 7, 4, 3, 3, 1, 4, c);\n", "R0(0, 1, 2, 3, 4, 5, 6, 7, 4, 4, 2, 5, d);\n", "R0(3, 0, 1, 2, 7, 4, 5, 6, 5, 5, 3, 6, e);\n", "R0(2, 3, 0, 1, 6, 7, 4, 5, 6, 6, 4, 7, f);\n", "R0(1, 2, 3, 0, 5, 6, 7, 4, 7, 7, 5, 8, 0);\n", "R0(0, 1, 2, 3, 4, 5, 6, 7, 8, 8, 6, 9, 1);\n", "R0(3, 0, 1, 2, 7, 4, 5, 6, 9, 9, 7, a, 2);\n", "R0(2, 3, 0, 1, 6, 7, 4, 5, 10, a, 8, b, 3);\n", "R0(1, 2, 3, 0, 5, 6, 7, 4, 11, b, 9, c, 4);\n", "R0(0, 1, 2, 3, 4, 5, 6, 7, 12, c, a, d, 5);\n", "R0(3, 0, 1, 2, 7, 4, 5, 6, 13, d, b, e, 6);\n", "R0(2, 3, 0, 1, 6, 7, 4, 5, 14, e, c, f, 7);\n", "R0(1, 2, 3, 0, 5, 6, 7, 4, 15, f, d, 0, 8);\n", "R(0, 1, 2, 3, 4, 5, 6, 7, 16, 0, e, 1, 9);\n", "R(3, 0, 1, 2, 7, 4, 5, 6, 17, 1, f, 2, a);\n", "R(2, 3, 0, 1, 6, 7, 4, 5, 18, 2, 0, 3, b);\n", "R(1, 2, 3, 0, 5, 6, 7, 4, 19, 3, 1, 4, c);\n", "R(0, 1, 2, 3, 4, 5, 6, 7, 20, 4, 2, 5, d);\n", "R(3, 0, 1, 2, 7, 4, 5, 6, 21, 5, 3, 6, e);\n", "R(2, 3, 0, 1, 6, 7, 4, 5, 22, 6, 4, 7, f);\n", "R(1, 2, 3, 0, 5, 6, 7, 4, 23, 7, 5, 8, 0);\n", "R(0, 1, 2, 3, 4, 5, 6, 7, 24, 8, 6, 9, 1);\n", "R(3, 0, 1, 2, 7, 4, 5, 6, 25, 9, 7, a, 2);\n", "R(2, 3, 0, 1, 6, 7, 4, 5, 26, a, 8, b, 3);\n", "R(1, 2, 3, 0, 5, 6, 7, 4, 27, b, 9, c, 4);\n", "R(0, 1, 2, 3, 4, 5, 6, 7, 28, c, a, d, 5);\n", "R(3, 0, 1, 2, 7, 4, 5, 6, 29, d, b, e, 6);\n", "R(2, 3, 0, 1, 6, 7, 4, 5, 30, e, c, f, 7);\n", "R(1, 2, 3, 0, 5, 6, 7, 4, 31, f, d, 0, 8);\n", "R(0, 1, 2, 3, 4, 5, 6, 7, 32, 0, e, 1, 9);\n", "R(3, 0, 1, 2, 7, 4, 5, 6, 33, 1, f, 2, a);\n", "R(2, 3, 0, 1, 6, 7, 4, 5, 34, 2, 0, 3, b);\n", "R(1, 2, 3, 0, 5, 6, 7, 4, 35, 3, 1, 4, c);\n", "R(0, 1, 2, 3, 4, 5, 6, 7, 36, 4, 2, 5, d);\n", "R(3, 0, 1, 2, 7, 4, 5, 6, 37, 5, 3, 6, e);\n", "R(2, 3, 0, 1, 6, 7, 4, 5, 38, 6, 4, 7, f);\n", "R(1, 2, 3, 0, 5, 6, 7, 4, 39, 7, 5, 8, 0);\n", "R(0, 1, 2, 3, 4, 5, 6, 7, 40, 8, 6, 9, 1);\n", "R(3, 0, 1, 2, 7, 4, 5, 6, 41, 9, 7, a, 2);\n", "R(2, 3, 0, 1, 6, 7, 4, 5, 42, a, 8, b, 3);\n", "R(1, 2, 3, 0, 5, 6, 7, 4, 43, b, 9, c, 4);\n", "R(0, 1, 2, 3, 4, 5, 6, 7, 44, c, a, d, 5);\n", "R(3, 0, 1, 2, 7, 4, 5, 6, 45, d, b, e, 6);\n", "R(2, 3, 0, 1, 6, 7, 4, 5, 46, e, c, f, 7);\n", "R(1, 2, 3, 0, 5, 6, 7, 4, 47, f, d, 0, 8);\n", "R(0, 1, 2, 3, 4, 5, 6, 7, 48, 0, e, 1, 9);\n", "R(3, 0, 1, 2, 7, 4, 5, 6, 49, 1, f, 2, a);\n", "R(2, 3, 0, 1, 6, 7, 4, 5, 50, 2, 0, 3, b);\n", "R(1, 2, 3, 0, 5, 6, 7, 4, 51, 3, 1, 4, c);\n", "R(0, 1, 2, 3, 4, 5, 6, 7, 52, 4, 2, 5, d);\n", "R(3, 0, 1, 2, 7, 4, 5, 6, 53, 5, 3, 6, e);\n", "R(2, 3, 0, 1, 6, 7, 4, 5, 54, 6, 4, 7, f);\n", "R(1, 2, 3, 0, 5, 6, 7, 4, 55, 7, 5, 8, 0);\n", "R(0, 1, 2, 3, 4, 5, 6, 7, 56, 8, 6, 9, 1);\n", "R(3, 0, 1, 2, 7, 4, 5, 6, 57, 9, 7, a, 2);\n", "R(2, 3, 0, 1, 6, 7, 4, 5, 58, a, 8, b, 3);\n", "R(1, 2, 3, 0, 5, 6, 7, 4, 59, b, 9, c, 4);\n", "R(0, 1, 2, 3, 4, 5, 6, 7, 60, c, a, d, 5);\n", "R(3, 0, 1, 2, 7, 4, 5, 6, 61, d, b, e, 6);\n", "R(2, 3, 0, 1, 6, 7, 4, 5, 62, e, c, f, 7);\n", "R(1, 2, 3, 0, 5, 6, 7, 4, 63, f, d, 0, 8);\n", "R(0, 1, 2, 3, 4, 5, 6, 7, 64, 0, e, 1, 9);\n", "R(3, 0, 1, 2, 7, 4, 5, 6, 65, 1, f, 2, a);\n", "R(2, 3, 0, 1, 6, 7, 4, 5, 66, 2, 0, 3, b);\n", "R(1, 2, 3, 0, 5, 6, 7, 4, 67, 3, 1, 4, c);\n", "R(0, 1, 2, 3, 4, 5, 6, 7, 68, 4, 2, 5, d);\n", "R(3, 0, 1, 2, 7, 4, 5, 6, 69, 5, 3, 6, e);\n", "R(2, 3, 0, 1, 6, 7, 4, 5, 70, 6, 4, 7, f);\n", "R(1, 2, 3, 0, 5, 6, 7, 4, 71, 7, 5, 8, 0);\n", "R(0, 1, 2, 3, 4, 5, 6, 7, 72, 8, 6, 9, 1);\n", "R(3, 0, 1, 2, 7, 4, 5, 6, 73, 9, 7, a, 2);\n", "R(2, 3, 0, 1, 6, 7, 4, 5, 74, a, 8, b, 3);\n", "R(1, 2, 3, 0, 5, 6, 7, 4, 75, b, 9, c, 4);\n", "R(0, 1, 2, 3, 4, 5, 6, 7, 76, c, a, d, 5);\n", "R(3, 0, 1, 2, 7, 4, 5, 6, 77, d, b, e, 6);\n", "R(2, 3, 0, 1, 6, 7, 4, 5, 78, e, c, f, 7);\n", "Rlast(1, 2, 3, 0, 5, 6, 7, 4, 79, f, d, 0, 8);\n", "\n", "\n", "/* R(1, 2, 3, 0, 5, 6, 7, 4, 79);*/\n", "//v.s0+=H0;\n", "//v.s0 = sigma1(w[13]) + sigma0(w[0]) + w[15] + w[8] + k79 + v.s4 + Sigma1(v.s5) + Ch(v.s5, v.s6, v.s7) + Maj(v.s1, v.s2, v.s3) + Sigma0(v.s1) +H0;\n", " return CONDITION;\n", "}\n", "\n", "__kernel\n", "void kernel_sha384(__const uint64_t base, __global uint32_t * out_buffer) {\n", "if(sha384_block(base + get_global_id(0)))\n", " out_buffer[get_global_id(0)>>24] = get_global_id(0)&0xffffff;\n", "\n", "}\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:63: CompilerWarning: Non-empty compiler output encountered. Set the environment variable PYOPENCL_COMPILER_OUTPUT=1 to see more.\n", " \"to see more.\", CompilerWarning)\n" ] } ], "prompt_number": 6 }, { "cell_type": "code", "collapsed": false, "input": [ "import atitweak as ati\n", "from adl3 import ADLTemperature, ADL_Overdrive5_Temperature_Get, ADL_OK\n", "from adl3 import ADL_DL_FANCTRL_SPEED_TYPE_PERCENT, ADLFanSpeedValue, ADL_Overdrive5_FanSpeed_Get\n", "from adl3 import ADLPMActivity, ADL_Overdrive5_CurrentActivity_Get\n", "from ctypes import sizeof, byref\n", "\n", "ati.initialize()\n", "def gpu_activity():\n", " adapters = ati.get_adapter_info()\n", " adapter = adapters[0]\n", " activity = ADLPMActivity()\n", " activity.iSize = sizeof(activity) \n", " if ADL_Overdrive5_CurrentActivity_Get(adapter.iAdapterIndex, byref(activity)) != ADL_OK:\n", " return None\n", " return (activity.iEngineClock/100.0, activity.iMemoryClock/100.0, activity.iVddc/1000.0,\n", " activity.iCurrentPerformanceLevel, activity.iActivityPercent)\n", " \n", " # engine clock %gMHz, memory clock %gMHz, core voltage %gVDC, performance level %d, utilization %d%%\" % \n", " \n", "def gpu_temperature():\n", " adapters = ati.get_adapter_info()\n", " adapter = adapters[0]\n", " temperature = ADLTemperature()\n", " temperature.iSize = sizeof(temperature)\n", " if ADL_Overdrive5_Temperature_Get(adapter.iAdapterIndex, 0, byref(temperature)) != ADL_OK:\n", " return None \n", " return temperature.iTemperature/1000 \n", "def gpu_fan_speed():\n", " adapters = ati.get_adapter_info()\n", " adapter = adapters[0]\n", " fan_speed_value = ADLFanSpeedValue()\n", " fan_speed_value.iSize = sizeof(fan_speed_value)\n", " fan_speed_value.iSpeedType = ADL_DL_FANCTRL_SPEED_TYPE_PERCENT\n", " if ADL_Overdrive5_FanSpeed_Get(adapter.iAdapterIndex, 0, byref(fan_speed_value)) != ADL_OK:\n", " return None\n", " return fan_speed_value.iFanSpeed\n", "print gpu_temperature(), gpu_fan_speed()\n", "# ati.shutdown()" ], "language": "python", "metadata": {}, "outputs": [ { "output_type": "stream", "stream": "stdout", "text": [ "47.0 35\n" ] } ], "prompt_number": 7 }, { "cell_type": "code", "collapsed": false, "input": [ "from IPython.html.widgets import *\n", "from IPython.display import display\n", "from time import sleep, time\n", "from sys import stdout\n", "from hashlib import sha384\n", "step = 2**28\n", "t1 = t0 = time()\n", "expected_number = 256**len(prefix)\n", "text = TextWidget()\n", "tempw = FloatProgressWidget(min = 50, max = 82, description=\"Temperature:\")\n", "fanw = FloatProgressWidget(min=0, max = 100, description = \"Fan Speed:\")\n", "keyw = TextWidget(description = \"Current Key:\")\n", "ratew = TextWidget(description=\"Hashes/S:\")\n", "elapsedw = TextWidget(description=\"Time Elapsed:\")\n", "actw = TextWidget()\n", "actw.set_css('width', 800)\n", "gpu_clockw = FloatProgressWidget(min=0, max=1050, description=\"GPU Clock\")\n", "gpu_loadw = FloatProgressWidget(min=0, max=100, description=\"GPU Load\")\n", "for w in [elapsedw, text, tempw, fanw, keyw, ratew, actw, gpu_clockw, gpu_loadw]:\n", " display(w)\n", "text.set_css('width', 800)\n", "sleep_time = 0.0\n", "START = 2168053219804\n", "START = 0x8ba0000000\n", "output=cl.array.zeros(queue, step>>24, dtype=np.uint32) \n", "output0=cl.array.zeros(queue, step>>24, dtype=np.uint32) \n", "ret = None\n", "found = 0\n", "for i0 in xrange(START, 64**7, step): \n", " output0.fill(0, queue)\n", " ret0 = kernel_sha384(queue, (step,), None, np.uint64(i0), output0.data)\n", " \n", " if ret is None:\n", " ret=ret0\n", " output, output0 = output0, output\n", " i = i0\n", " continue\n", " \n", " #ret,i,output = ret0, i0, output0 \n", " ret.wait()\n", " result = output.get()\n", " for j in np.where(result>0)[0]: \n", " found+=1\n", " x =i+(j<<24)+result[j]\n", " key = \"\".join(charset[(x>>(6*h))&63]for h in range(7))\n", " s=s0.replace('A'*7, key)\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", " exp_found = (i-START)/expected_number\n", " if elapsed >0:\n", " rate = (i-START)/60/elapsed\n", " else:\n", " rate=0\n", " exp = 0 if i==START else elapsed/exp_found/60.0\n", " expday = int(exp/24)\n", " exphr = exp-expday*24\n", " key = \"\".join(charset[(i>>(6*h))&63]for h in range(7))\n", " temp = gpu_temperature()\n", " fan = gpu_fan_speed()\n", " activity = gpu_activity()\n", " elapsedw.value = \" {:.2f}min\".format(elapsed)\n", " text.value = \"{} exp_found={:.4f} found={} EXP={}days {:.2f}hrs sleep_time={} temp={}\".format(hex(i), exp_found, found, expday, exphr, sleep_time, temp)\n", " current_rate = step/(time()-t1)\n", " t1 = time()\n", " ratew.value = \"avg={:.2f}M/s current={:.2f}M/s\".format(rate/10**6, current_rate/10**6)\n", " keyw.value=key\n", " fanw.value = fan\n", " tempw.value = temp\n", " actw.value = \"engine clock %gMHz, memory clock %gMHz, core voltage %gVDC, performance level %d, utilization %d%%\" % activity\n", " gpu_clockw.value = activity[0]\n", " gpu_loadw.value = activity[4]\n", " \n", " if temp > 80:\n", " sleep_time +=0.05\n", " if temp is not None and temp < 78.5:\n", " sleep_time -=0.05 \n", " if sleep_time > 0:\n", " sleep(sleep_time)\n", " else:\n", " sleep_time = 0.0\n", " ret, i = ret0, i0\n", " output, output0 = output0, output\n", " i = i0\n", "print \"done\"" ], "language": "python", "metadata": {}, "outputs": [ { "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\u001b[0m in \u001b[0;36m\u001b[1;34m()\u001b[0m\n\u001b[0;32m 39\u001b[0m \u001b[1;31m#ret,i,output = ret0, i0, output0\u001b[0m\u001b[1;33m\u001b[0m\u001b[1;33m\u001b[0m\u001b[0m\n\u001b[0;32m 40\u001b[0m \u001b[0mret\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[1;32m---> 41\u001b[1;33m \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[0m\u001b[0;32m 42\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;36m0\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[0;32m 43\u001b[0m \u001b[0mfound\u001b[0m\u001b[1;33m+=\u001b[0m\u001b[1;36m1\u001b[0m\u001b[1;33m\u001b[0m\u001b[0m\n", "\u001b[1;32m/usr/local/lib/python2.7/dist-packages/pyopencl-2014.1-py2.7-linux-x86_64.egg/pyopencl/array.pyc\u001b[0m in \u001b[0;36mget\u001b[1;34m(self, queue, ary, async)\u001b[0m\n\u001b[0;32m 690\u001b[0m cl.enqueue_copy(queue or self.queue, ary, self.base_data,\n\u001b[0;32m 691\u001b[0m \u001b[0mdevice_offset\u001b[0m\u001b[1;33m=\u001b[0m\u001b[0mself\u001b[0m\u001b[1;33m.\u001b[0m\u001b[0moffset\u001b[0m\u001b[1;33m,\u001b[0m\u001b[1;33m\u001b[0m\u001b[0m\n\u001b[1;32m--> 692\u001b[1;33m is_blocking=not async)\n\u001b[0m\u001b[0;32m 693\u001b[0m \u001b[1;33m\u001b[0m\u001b[0m\n\u001b[0;32m 694\u001b[0m \u001b[1;32mreturn\u001b[0m \u001b[0mary\u001b[0m\u001b[1;33m\u001b[0m\u001b[0m\n", "\u001b[1;32m/usr/local/lib/python2.7/dist-packages/pyopencl-2014.1-py2.7-linux-x86_64.egg/pyopencl/__init__.pyc\u001b[0m in \u001b[0;36menqueue_copy\u001b[1;34m(queue, dest, src, **kwargs)\u001b[0m\n\u001b[0;32m 1088\u001b[0m \u001b[1;32mreturn\u001b[0m \u001b[0m_cl\u001b[0m\u001b[1;33m.\u001b[0m\u001b[0m_enqueue_read_buffer_rect\u001b[0m\u001b[1;33m(\u001b[0m\u001b[0mqueue\u001b[0m\u001b[1;33m,\u001b[0m \u001b[0msrc\u001b[0m\u001b[1;33m,\u001b[0m \u001b[0mdest\u001b[0m\u001b[1;33m,\u001b[0m \u001b[1;33m**\u001b[0m\u001b[0mkwargs\u001b[0m\u001b[1;33m)\u001b[0m\u001b[1;33m\u001b[0m\u001b[0m\n\u001b[0;32m 1089\u001b[0m \u001b[1;32melse\u001b[0m\u001b[1;33m:\u001b[0m\u001b[1;33m\u001b[0m\u001b[0m\n\u001b[1;32m-> 1090\u001b[1;33m \u001b[1;32mreturn\u001b[0m \u001b[0m_cl\u001b[0m\u001b[1;33m.\u001b[0m\u001b[0m_enqueue_read_buffer\u001b[0m\u001b[1;33m(\u001b[0m\u001b[0mqueue\u001b[0m\u001b[1;33m,\u001b[0m \u001b[0msrc\u001b[0m\u001b[1;33m,\u001b[0m \u001b[0mdest\u001b[0m\u001b[1;33m,\u001b[0m \u001b[1;33m**\u001b[0m\u001b[0mkwargs\u001b[0m\u001b[1;33m)\u001b[0m\u001b[1;33m\u001b[0m\u001b[0m\n\u001b[0m\u001b[0;32m 1091\u001b[0m \u001b[1;32melif\u001b[0m \u001b[0msrc\u001b[0m\u001b[1;33m.\u001b[0m\u001b[0mtype\u001b[0m \u001b[1;32min\u001b[0m \u001b[1;33m[\u001b[0m\u001b[0mmem_object_type\u001b[0m\u001b[1;33m.\u001b[0m\u001b[0mIMAGE2D\u001b[0m\u001b[1;33m,\u001b[0m \u001b[0mmem_object_type\u001b[0m\u001b[1;33m.\u001b[0m\u001b[0mIMAGE3D\u001b[0m\u001b[1;33m]\u001b[0m\u001b[1;33m:\u001b[0m\u001b[1;33m\u001b[0m\u001b[0m\n\u001b[0;32m 1092\u001b[0m \u001b[0morigin\u001b[0m \u001b[1;33m=\u001b[0m \u001b[0mkwargs\u001b[0m\u001b[1;33m.\u001b[0m\u001b[0mpop\u001b[0m\u001b[1;33m(\u001b[0m\u001b[1;34m\"origin\"\u001b[0m\u001b[1;33m)\u001b[0m\u001b[1;33m\u001b[0m\u001b[0m\n", "\u001b[1;31mKeyboardInterrupt\u001b[0m: " ] } ], "prompt_number": 8 }, { "cell_type": "code", "collapsed": false, "input": [ "display(HTMLWidget( value =\"\"\"\n", "\n", "\"\"\"))" ], "language": "python", "metadata": {}, "outputs": [] }, { "cell_type": "code", "collapsed": false, "input": [ "%qtconsole --colors=linux" ], "language": "python", "metadata": {}, "outputs": [], "prompt_number": 9 }, { "cell_type": "code", "collapsed": false, "input": [ "for i in range(0,80):\n", " print \"R%s(%d, %d, %d, %d, %d, %d, %d, %d, %d, %x, %x, %x, %x);\"%('0' if i<16 else '', (-i)%4, (1-i)%4, (2-i)%4, (3-i)%4,\n", " (-i)%4+4, (1-i)%4+4, (2-i)%4+4, (3-i)%4+4, i, i%16, (i-2)%16, (i-15)%16, (i-7)%16)\n", " \n" ], "language": "python", "metadata": {}, "outputs": [] }, { "cell_type": "code", "collapsed": false, "input": [], "language": "python", "metadata": {}, "outputs": [] } ], "metadata": {} } ] }