from __future__ import division
from IPython.html.widgets import *
from IPython.display import display
from collections import Counter
from itertools import combinations, chain
from hashlib import sha384
from time import time
from sys import stdout

display(HTMLWidget( value ="""
<script>

CodeMirror.requireMode('clike', function(){
var cells = IPython.notebook.get_cells();
for(var i in cells){
if(cells[i].code_mirror.getValue().indexOf("%%cl_kernel")==0){cells[i].code_mirror.setOption('mode', 'clike');}
}
})

</script>
"""))

from __future__ import division
import numpy as np
import pyopencl as cl
import pyopencl.array
import pyopencl.tools
%load_ext pyopencl.ipython_ext

ctx = cl.create_some_context()
queue = cl.CommandQueue(ctx)
device=ctx.devices[0]
print device.max_mem_alloc_size
print device.max_work_group_size

def python_md5(plaintext, pre_calculate=False):    
    print len(plaintext)
    if pre_calculate:
        assert len(plaintext)>=52
    X = np.zeros(16, dtype=np.uint32)
    DATA =   np.uint32([0x67452301, 0xefcdab89, 0x98badcfe, 0x10325476]) 
    for i,c in enumerate(plaintext):    
        X[i >> 2] |= ord(c) << ((i & 3) << 3)
    X[len(plaintext) >> 2] |=   0x80 << ((len(plaintext) & 3) << 3)
    X[14] = len(plaintext) << 3    
    print map(hex, X)
    S = lambda x,n : ((x << n) | ((x&0xffffffff) >> (32 - n)))& 0xFFFFFFFF    
    
    def P(i0,i1,i2,i3,k,s,t): 
        a,b,c,d = DATA[i0], DATA[i1], DATA[i2], DATA[i3]
        a += F(b,c,d) + X[k] + t
        a = S(a,s) + b
        DATA[i0] = a & 0xFFFFFFFF
        
    A,B,C,D = 0,1,2,3   
    F = lambda x,y,z: (z ^ (x & (y ^ z)))
    P( A, B, C, D,  0,  7, 0xD76AA478 );
    P( D, A, B, C,  1, 12, 0xE8C7B756 );
    P( C, D, A, B,  2, 17, 0x242070DB );
    P( B, C, D, A,  3, 22, 0xC1BDCEEE );
    P( A, B, C, D,  4,  7, 0xF57C0FAF );
    P( D, A, B, C,  5, 12, 0x4787C62A );
    P( C, D, A, B,  6, 17, 0xA8304613 );
    P( B, C, D, A,  7, 22, 0xFD469501 );
    P( A, B, C, D,  8,  7, 0x698098D8 );
    P( D, A, B, C,  9, 12, 0x8B44F7AF );
    P( C, D, A, B, 10, 17, 0xFFFF5BB1 );
    if pre_calculate:
        with open("lowest_md5.h","w") as f:
            for i in range(16):
                if i not in [11, 12]:                
                    print >>f, "#define X%d 0x%xU"%(i, X[i])
                    # print "#define X%d 0x%xU"%(i, X[i])
            for i in range(4):
                    print >>f, "#define DATA%d 0x%xU"%(i, DATA[i])
                    # print "#define DATA%d 0x%xU"%(i, DATA[i])                
    P( B, C, D, A, 11, 22, 0x895CD7BE );
    P( A, B, C, D, 12,  7, 0x6B901122 );
    P( D, A, B, C, 13, 12, 0xFD987193 );
    P( C, D, A, B, 14, 17, 0xA679438E );
    P( B, C, D, A, 15, 22, 0x49B40821 );
    
    F = lambda x,y,z:  (y ^ (z & (x ^ y)))
    P( A, B, C, D,  1,  5, 0xF61E2562 );
    P( D, A, B, C,  6,  9, 0xC040B340 )
    P( C, D, A, B, 11, 14, 0x265E5A51 )
    P( B, C, D, A,  0, 20, 0xE9B6C7AA );
    P( A, B, C, D,  5,  5, 0xD62F105D );
    P( D, A, B, C, 10,  9, 0x02441453 );
    P( C, D, A, B, 15, 14, 0xD8A1E681 );
    P( B, C, D, A,  4, 20, 0xE7D3FBC8 );
    P( A, B, C, D,  9,  5, 0x21E1CDE6 );
    P( D, A, B, C, 14,  9, 0xC33707D6 );
    P( C, D, A, B,  3, 14, 0xF4D50D87 );
    P( B, C, D, A,  8, 20, 0x455A14ED );
    P( A, B, C, D, 13,  5, 0xA9E3E905 );
    P( D, A, B, C,  2,  9, 0xFCEFA3F8 )
    P( C, D, A, B,  7, 14, 0x676F02D9 )
    P( B, C, D, A, 12, 20, 0x8D2A4C8A );
    
    F = lambda x,y,z:  x ^ y ^ z
    P( A, B, C, D,  5,  4, 0xFFFA3942 );
    P( D, A, B, C,  8, 11, 0x8771F681 );
    P( C, D, A, B, 11, 16, 0x6D9D6122 );
    P( B, C, D, A, 14, 23, 0xFDE5380C );
    P( A, B, C, D,  1,  4, 0xA4BEEA44 );
    P( D, A, B, C,  4, 11, 0x4BDECFA9 );
    P( C, D, A, B,  7, 16, 0xF6BB4B60 );
    P( B, C, D, A, 10, 23, 0xBEBFBC70 );
    P( A, B, C, D, 13,  4, 0x289B7EC6 );
    P( D, A, B, C,  0, 11, 0xEAA127FA );
    P( C, D, A, B,  3, 16, 0xD4EF3085 );
    P( B, C, D, A,  6, 23, 0x04881D05 );
    P( A, B, C, D,  9,  4, 0xD9D4D039 );
    P( D, A, B, C, 12, 11, 0xE6DB99E5 );
    P( C, D, A, B, 15, 16, 0x1FA27CF8 );
    P( B, C, D, A,  2, 23, 0xC4AC5665 );

    F = lambda x,y,z:  (y ^ (x | ~z))

    P( A, B, C, D,  0,  6, 0xF4292244 );
    P( D, A, B, C,  7, 10, 0x432AFF97 );
    P( C, D, A, B, 14, 15, 0xAB9423A7 );
    P( B, C, D, A,  5, 21, 0xFC93A039 );
    P( A, B, C, D, 12,  6, 0x655B59C3 );
    P( D, A, B, C,  3, 10, 0x8F0CCC92 );
    P( C, D, A, B, 10, 15, 0xFFEFF47D );
    P( B, C, D, A,  1, 21, 0x85845DD1 );
    P( A, B, C, D,  8,  6, 0x6FA87E4F );
    P( D, A, B, C, 15, 10, 0xFE2CE6E0 );
    P( C, D, A, B,  6, 15, 0xA3014314 );
    P( B, C, D, A, 13, 21, 0x4E0811A1 );
    P( A, B, C, D,  4,  6, 0xF7537E82 );
    P( D, A, B, C, 11, 10, 0xBD3AF235 );
    P( C, D, A, B,  2, 15, 0x2AD7D2BB );
    P( B, C, D, A,  9, 21, 0xEB86D391 );
    DATA[0]=(DATA[0]+0x67452301)&0xffffffff
    DATA[1]=(DATA[1]+0xefcdab89)&0xffffffff
    DATA[2]=(DATA[2]+0x98badcfe)&0xffffffff
    DATA[3]=(DATA[3]+0x10325476)&0xffffffff
    return DATA


s0 = "http://weijr-eng.blogspot.com"
template_str = s0+" "*(52-len(s0)-8)+"AAAAAAAA"
print map(hex, python_md5(template_str, True))
from hashlib import md5
md5(template_str).hexdigest()

%%cl_kernel -o "-I . -cl-strict-aliasing"
#include "lowest_md5.h"
__kernel void kernel_md5(__const ulong base, __global  uint * out_buffer)
{    
	uint id = get_global_id(0);
    ulong code = base + id;
    __private uint X11 = code >> 32;
    __private uint X12 = code &0xffffffff;		
	__private uint A, B, C, D;
#define S(x,n)   rotate(x,n)
#define P(a,b,c,d,k,s,t)	a += F(b,c,d) + X##k + t; a = S(a,s) + b;		

    A = DATA0;
    B = DATA1;
    C = DATA2;
    D = DATA3;
   
#define F(x,y,z) bitselect(z, y,x)

    P( B, C, D, A, 11, 22U, 0x895CD7BE );
    P( A, B, C, D, 12,  7U, 0x6B901122 );
    P( D, A, B, C, 13, 12U, 0xFD987193 );
    P( C, D, A, B, 14, 17U, 0xA679438E );
    P( B, C, D, A, 15, 22U, 0x49B40821 );

#undef F
#define F(x,y,z) bitselect(y,x, z)

    P( A, B, C, D,  1,  5U, 0xF61E2562 );
    P( D, A, B, C,  6,  9U, 0xC040B340 );
	P( C, D, A, B, 11, 14U, 0x265E5A51 );
    P( B, C, D, A,  0, 20U, 0xE9B6C7AA );
    P( A, B, C, D,  5,  5U, 0xD62F105D );
    P( D, A, B, C, 10,  9U, 0x02441453 );
    P( C, D, A, B, 15, 14U, 0xD8A1E681 );
    P( B, C, D, A,  4, 20U, 0xE7D3FBC8 );
    P( A, B, C, D,  9,  5U, 0x21E1CDE6 );
    P( D, A, B, C, 14,  9U, 0xC33707D6 );
    P( C, D, A, B,  3, 14U, 0xF4D50D87 );
    P( B, C, D, A,  8, 20U, 0x455A14ED );
    P( A, B, C, D, 13,  5U, 0xA9E3E905 );
    P( D, A, B, C,  2,  9U, 0xFCEFA3F8 );
    P( C, D, A, B,  7, 14U, 0x676F02D9 );
    P( B, C, D, A, 12, 20U, 0x8D2A4C8A );

#undef F
    
#define F(x,y,z) (x ^ y ^ z)

    P( A, B, C, D,  5,  4U, 0xFFFA3942 );
    P( D, A, B, C,  8, 11U, 0x8771F681 );
    P( C, D, A, B, 11, 16U, 0x6D9D6122 );
    P( B, C, D, A, 14, 23U, 0xFDE5380C );
    P( A, B, C, D,  1,  4U, 0xA4BEEA44 );
    P( D, A, B, C,  4, 11U, 0x4BDECFA9 );
    P( C, D, A, B,  7, 16U, 0xF6BB4B60 );
    P( B, C, D, A, 10, 23U, 0xBEBFBC70 );
    P( A, B, C, D, 13,  4U, 0x289B7EC6 );
    P( D, A, B, C,  0, 11U, 0xEAA127FA );
    P( C, D, A, B,  3, 16U, 0xD4EF3085 );
    P( B, C, D, A,  6, 23U, 0x04881D05 );
    P( A, B, C, D,  9,  4U, 0xD9D4D039 );
    P( D, A, B, C, 12, 11U, 0xE6DB99E5 );
    P( C, D, A, B, 15, 16U, 0x1FA27CF8 );
    P( B, C, D, A,  2, 23U, 0xC4AC5665 );

#undef F

#define F(x,y,z) (y ^ (x | ~z))

    P( A, B, C, D,  0,  6U, 0xF4292244 );
    P( D, A, B, C,  7, 10U, 0x432AFF97 );
    P( C, D, A, B, 14, 15U, 0xAB9423A7 );
    P( B, C, D, A,  5, 21U, 0xFC93A039 );
    P( A, B, C, D, 12,  6U, 0x655B59C3 );
    P( D, A, B, C,  3, 10U, 0x8F0CCC92 );
    P( C, D, A, B, 10, 15U, 0xFFEFF47D );
    P( B, C, D, A,  1, 21U, 0x85845DD1 );
    P( A, B, C, D,  8,  6U, 0x6FA87E4F );
    P( D, A, B, C, 15, 10U, 0xFE2CE6E0 );
    P( C, D, A, B,  6, 15U, 0xA3014314 );
    P( B, C, D, A, 13, 21U, 0x4E0811A1 );
    P( A, B, C, D,  4,  6U, 0xF7537E82 );
    
    if( (A&0xFFFFFFFEU) == 2562383102U ) {        
        P( D, A, B, C, 11, 10U, 0xBD3AF235 );
        P( C, D, A, B,  2, 15U, 0x2AD7D2BB );
        P( B, C, D, A,  9, 21U, 0xEB86D391 );
        A+=0x67452301;
        B += 0xefcdab89;
    if( ( (B^A)&0xff)==0)
   out_buffer[id&0xff]=0x80000000|id; 
   }
    
#undef F

}


import sys
from time import sleep, time
step = 2**28
output=cl.array.zeros(queue, 16, dtype=np.uint32) 
t0=time()
for i in range(0, 2**38, step):
    output.fill(0, queue)
    ret = kernel_md5(queue, (step,), None, np.uint64(i), output.data)
    ret.wait()
    result = output.get()
    for j in np.where(result>0)[0]:
        x =i+(result[j]&0x7fffffff)
        y = x >>32
        x = x&0xffffffff
        s = template_str.replace('AAAAAAAA', "".join(chr((y>>(t*8))&0xff) for t in range(4)) +      "".join(chr((x>>(t*8))&0xff) for t in range(4)))
        print hex(x), md5(s).hexdigest(),repr(s)
        sys.stdout.flush()
    #sleep(1)
print time()-t0

import atitweak as ati
from adl3 import ADLTemperature, ADL_Overdrive5_Temperature_Get, ADL_OK
from adl3 import ADL_DL_FANCTRL_SPEED_TYPE_PERCENT, ADLFanSpeedValue, ADL_Overdrive5_FanSpeed_Get
from adl3 import ADLPMActivity, ADL_Overdrive5_CurrentActivity_Get
from ctypes import sizeof, byref

ati.initialize()
def gpu_activity():
    adapters = ati.get_adapter_info()
    adapter = adapters[0]
    activity = ADLPMActivity()
    activity.iSize = sizeof(activity)            
    if ADL_Overdrive5_CurrentActivity_Get(adapter.iAdapterIndex, byref(activity)) != ADL_OK:
                return None
    return (activity.iEngineClock/100.0, activity.iMemoryClock/100.0, activity.iVddc/1000.0,
                         activity.iCurrentPerformanceLevel, activity.iActivityPercent)
                            
def gpu_temperature():
    adapters = ati.get_adapter_info()
    adapter = adapters[0]
    temperature = ADLTemperature()
    temperature.iSize = sizeof(temperature)
    if ADL_Overdrive5_Temperature_Get(adapter.iAdapterIndex, 0, byref(temperature)) != ADL_OK:
                return None    
    return  temperature.iTemperature/1000    
def gpu_fan_speed():
    adapters = ati.get_adapter_info()
    adapter = adapters[0]
    fan_speed_value = ADLFanSpeedValue()
    fan_speed_value.iSize = sizeof(fan_speed_value)
    fan_speed_value.iSpeedType = ADL_DL_FANCTRL_SPEED_TYPE_PERCENT
    if ADL_Overdrive5_FanSpeed_Get(adapter.iAdapterIndex, 0, byref(fan_speed_value)) != ADL_OK:
        return None
    return fan_speed_value.iFanSpeed
print gpu_temperature(), gpu_fan_speed()
# ati.shutdown()

from IPython.html.widgets import *
from IPython.display import display
from time import sleep, time
from sys import stdout
from hashlib import md5
import cgi
prefix_mask = 0xffffffffff000000
class ListTable(list):
    """ Overridden list class which takes a 2-dimensional list of 
        the form [[1,2,3],[4,5,6]], and renders an HTML Table in 
        IPython Notebook. """
    
    def _repr_html_(self):
        html = ["<table>"]
        for row in self:
            html.append("<tr>")            
            html.extend("<td><pre>{}</pre></td>".format(cgi.escape(str(col))) for col in row)            
            html.append("</tr>")
        html.append("</table>")
        return ''.join(html)
    
step = 2**31
t1 = t0 = time()
expected_number = 1
leading_zeros = 0
x = prefix_mask
while x:
    if x&1:
        leading_zeros +=1
        expected_number *=2
    x >>=1
leading_zeros= leading_zeros//4
leading_zeros = max(leading_zeros-1,0)
print "leading_zeros=%d"% leading_zeros
stdout.flush()
text = TextWidget()
tempw = FloatProgressWidget(min = 50, max = 82, description="Temperature:")
fanw = FloatProgressWidget(min=0, max = 100, description = "Fan Speed:")
ratew = TextWidget(description="Hashes/S:")
elapsedw = TextWidget(description="Time Elapsed:")
actw = TextWidget()
actw.set_css('width', 800)
gpu_clockw = FloatProgressWidget(min=0, max=1050, description="GPU Clock")
gpu_loadw  = FloatProgressWidget(min=0, max=100, description="GPU Load")

for w in [elapsedw, text, tempw, fanw, ratew, actw, gpu_clockw, gpu_loadw]:
    display(w)
text.set_css('width', 800)
sleep_time = 0.0
START = 0x1574500000000
output=cl.array.zeros(queue, 256, dtype=np.uint32) 
output0=cl.array.zeros(queue, 256, dtype=np.uint32) 
ret = None
found = [0]*4
leader_board = []
leaderw = HTMLWidget()
display(leaderw)
for i0 in xrange(START, 2**63-2, step):       
    output0.fill(0, queue)
    ret0 = kernel_md5(queue, (step,), None, np.uint64(i0), output0.data)
    
    if ret is None:
        ret=ret0
        output, output0 = output0, output
        i = i0
        continue
    
    #ret,i,output = ret0, i0, output0    
    ret.wait()
    result = output.get()
    for j in np.where(result>0)[0]:            
        x =i+(result[j]&0x7fffffff)
        y = x >>32
        x = x&0xffffffff
        s = template_str.replace('AAAAAAAA', "".join(chr((y>>(t*8))&0xff) for t in range(4)) +      "".join(chr((x>>(t*8))&0xff) for t in range(4)))                        
        k = md5(s).hexdigest()
        leader_board.append((k, repr(s)))
        for zi in range(4):
            if k [leading_zeros+zi]==k[leading_zeros]:
                found[zi]+=1                
            else:
                break
        print i
        print repr(s), k
        stdout.flush()
        leader_board = sorted(leader_board)
        if len(leader_board) > 10:
            leader_board=leader_board[:5]+leader_board[-5:]
        leaderw.value = ListTable((i,s,k[leading_zeros:20]) for i,(k,s) in enumerate(leader_board))._repr_html_()
    elapsed = (time()-t0)/60
    exp_found = (i-START)/expected_number
    if elapsed >0:
        rate = (i-START)/60/elapsed
    else:
        rate=0
    exp =  0 if i==START else elapsed/exp_found/60.0
    expday = int(exp/24)
    exphr = exp-expday*24
    temp = gpu_temperature()
    fan = gpu_fan_speed()
    activity = gpu_activity()
    elapsedw.value = " {:.2f}min".format(elapsed)
    text.value = "{} exp_found={:.1f} {:.1f} {:.1f} {:.1f} found={} EXP={}days {:.2f}hrs sleep_time={} temp={}".format(hex(i), 
                                                                                                                exp_found, exp_found/16, exp_found/256, exp_found/256/16, found, 
                                                                                                                expday, exphr, sleep_time, temp)
    current_rate = step/(time()-t1)
    t1 = time()
    ratew.value = "avg={:.2f}M/s current={:.2f}M/s".format(rate/10**6, current_rate/10**6)
    fanw.value = fan
    tempw.value = temp
    actw.value = "engine clock %gMHz, memory clock %gMHz, core voltage %gVDC, performance level %d, utilization %d%%" % activity
    gpu_clockw.value = activity[0]
    gpu_loadw.value = activity[4]
    
    if temp > 80:
        sleep_time +=0.05
    if temp > 82:            
            print "GPU temperature too high"
            break
    if temp is not None and temp < 78.5:
        sleep_time -=0.05 
    if sleep_time > 0:
        sleep(sleep_time)
    else:
        sleep_time = 0.0
    ret, i = ret0, i0
    output, output0 = output0, output
    i = i0
print "done"

%qtconsole --colors=linux

%%cl_kernel 
__kernel void kernel_md5(__const ulong base, __global  uint * out_buffer)
{
    __private uint X[16]={0x77777777,0,0,0x80,0,0,0,0,0,0,0,0,0,0,0,0};
	uint id = get_global_id(0);
    ulong code = base + id;
    X[1] = code >> 32;
    X[2] = code &0xffffffff;		
	__private uint A, B, C, D;
#define S(x,n)   rotate(x,n)
#define P(a,b,c,d,k,s,t)	a += F(b,c,d) + X[k] + t; a = S(a,s) + b;		
#define P0(a,b,c,d,k,s,t)	a += F(b,c,d) + t; a = S(a,s) + b;			
#define P14(a,b,c,d,k,s,t) 	a += F(b,c,d) + 96 + t; a = S(a,s) + b;	

    A = 0x67452301;
    B = 0xefcdab89;
    C = 0x98badcfe;
    D = 0x10325476;
   
#define F(x,y,z) bitselect(z, y,x)

    P( A, B, C, D,  0,  7U, 0xD76AA478 );
    P( D, A, B, C,  1, 12U, 0xE8C7B756 );
    P( C, D, A, B,  2, 17U, 0x242070DB );
    P( B, C, D, A,  3, 22U, 0xC1BDCEEE );
    P0( A, B, C, D,  4,  7U, 0xF57C0FAF );
    P0( D, A, B, C,  5, 12U, 0x4787C62A );
    P0( C, D, A, B,  6, 17U, 0xA8304613 );
    P0( B, C, D, A,  7, 22U, 0xFD469501 );
    P0( A, B, C, D,  8,  7U, 0x698098D8 );
    P0( D, A, B, C,  9, 12U, 0x8B44F7AF );
    P0( C, D, A, B, 10, 17U, 0xFFFF5BB1 );
    P0( B, C, D, A, 11, 22U, 0x895CD7BE );
    P0( A, B, C, D, 12,  7U, 0x6B901122 );
    P0( D, A, B, C, 13, 12U, 0xFD987193 );
    P14( C, D, A, B, 14, 17U, 0xA679438E );
    P0( B, C, D, A, 15, 22U, 0x49B40821 );

#undef F
#define F(x,y,z) bitselect(y,x, z)

    P( A, B, C, D,  1,  5U, 0xF61E2562 );
    P0( D, A, B, C,  6,  9U, 0xC040B340 );
	P0( C, D, A, B, 11, 14U, 0x265E5A51 );
    P( B, C, D, A,  0, 20U, 0xE9B6C7AA );
    P0( A, B, C, D,  5,  5U, 0xD62F105D );
    P0( D, A, B, C, 10,  9U, 0x02441453 );
    P0( C, D, A, B, 15, 14U, 0xD8A1E681 );
    P0( B, C, D, A,  4, 20U, 0xE7D3FBC8 );
    P0( A, B, C, D,  9,  5U, 0x21E1CDE6 );
    P14( D, A, B, C, 14,  9U, 0xC33707D6 );
    P( C, D, A, B,  3, 14U, 0xF4D50D87 );
    P0( B, C, D, A,  8, 20U, 0x455A14ED );
    P0( A, B, C, D, 13,  5U, 0xA9E3E905 );
    P( D, A, B, C,  2,  9U, 0xFCEFA3F8 );
    P0( C, D, A, B,  7, 14U, 0x676F02D9 );
    P0( B, C, D, A, 12, 20U, 0x8D2A4C8A );

#undef F
    
#define F(x,y,z) (x ^ y ^ z)

    P( A, B, C, D,  5,  4U, 0xFFFA3942 );
    P( D, A, B, C,  8, 11U, 0x8771F681 );
    P( C, D, A, B, 11, 16U, 0x6D9D6122 );
    P14( B, C, D, A, 14, 23U, 0xFDE5380C );
    P( A, B, C, D,  1,  4U, 0xA4BEEA44 );
    P( D, A, B, C,  4, 11U, 0x4BDECFA9 );
    P( C, D, A, B,  7, 16U, 0xF6BB4B60 );
    P( B, C, D, A, 10, 23U, 0xBEBFBC70 );
    P( A, B, C, D, 13,  4U, 0x289B7EC6 );
    P( D, A, B, C,  0, 11U, 0xEAA127FA );
    P( C, D, A, B,  3, 16U, 0xD4EF3085 );
    P( B, C, D, A,  6, 23U, 0x04881D05 );
    P( A, B, C, D,  9,  4U, 0xD9D4D039 );
    P( D, A, B, C, 12, 11U, 0xE6DB99E5 );
    P( C, D, A, B, 15, 16U, 0x1FA27CF8 );
    P( B, C, D, A,  2, 23U, 0xC4AC5665 );

#undef F

#define F(x,y,z) (y ^ (x | ~z))

    P( A, B, C, D,  0,  6U, 0xF4292244 );
    P( D, A, B, C,  7, 10U, 0x432AFF97 );
    P14( C, D, A, B, 14, 15U, 0xAB9423A7 );
    P( B, C, D, A,  5, 21U, 0xFC93A039 );
    P( A, B, C, D, 12,  6U, 0x655B59C3 );
    P( D, A, B, C,  3, 10U, 0x8F0CCC92 );
    P( C, D, A, B, 10, 15U, 0xFFEFF47D );
    P( B, C, D, A,  1, 21U, 0x85845DD1 );
    P( A, B, C, D,  8,  6U, 0x6FA87E4F );
    P( D, A, B, C, 15, 10U, 0xFE2CE6E0 );
    P( C, D, A, B,  6, 15U, 0xA3014314 );
    P( B, C, D, A, 13, 21U, 0x4E0811A1 );
    P( A, B, C, D,  4,  6U, 0xF7537E82 );
    
    if( (A&0xFFFFFFFEU) == 2562383102U ) {        
        P( D, A, B, C, 11, 10U, 0xBD3AF235 );
        P( C, D, A, B,  2, 15U, 0x2AD7D2BB );
        P( B, C, D, A,  9, 21U, 0xEB86D391 );
        A+=0x67452301;
        B += 0xefcdab89;
    if( ( (B^A)&0xf0)==0)
   out_buffer[id&0xf]=0x80000000|id; 
   }
    
#undef F

}