In [ ]:
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
In [ ]:
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>
"""))
In [ ]:
from __future__ import division
import numpy as np
import pyopencl as cl
import pyopencl.array
import pyopencl.tools
%load_ext pyopencl.ipython_ext
In [ ]:
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
In [ ]:
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
In [ ]:
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()
In [ ]:
%%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
}
In [ ]:
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
In [ ]:
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()
In [197]:
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"
In [ ]:
%qtconsole --colors=linux
In [ ]:
%%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
}